compiler error at goto label

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I get the compiler error

with the kernel code

__global__ void kernel53(){

int xx = 0;

int psize=10;

goto L109;

L80:

if (xx>15) goto exception;

if (xx>20) goto exception;

xx+=1;

L109:

if (xx<psize) goto L80;

return;

exception:

return;

}

If the line if (xx>20) goto exception; is changed to if (xx>20) xx+=1; the compilation works.

The same happens in my real program. If I change the last goto exception to some expression that isn’t a goto or a return statement it works just fine and the program returns the correct results.

OS: Kubuntu 8.04

Kernel: 2.6.28.7

Cuda Toolkit: 2.1

Cuda SDK: 2.1

Host Compiler: gcc 4.2.4

Compiler Command: nvcc -g --verbose --shared --ptxas-options=-v -o libtest.so test.cu -L/usr/local/cuda/lib/ -lcudart

Graphic Card: GeForce 8600M GS

2

wow, that remind me when I first start programming basic. There are some valid uses for goto. But this doesn’t seem to fit.

You should try something like this, instead:

__global__ void kernel53(){

   int xx = 0;

   int psize=10;

bool exception = false;

while (xx<psize) {

	  if (xx>15) {

		  exception = true;

		  break;

	  }

	  if (xx>20) { // you do realize this is not going to happen, but since it was on the original code		  

		  exception = true;

		  break;		

	  }

	 xx+=1;

   }

if (exception) {

	 //do what you have to do

   }

}
3

I’m not aware that there are any limitations on the use of goto in c or cuda. Basically your while loop would be translated into the goto statements ;)

As I’m generating the cuda code out of Java Byte Code it is easier to use the already present goto statements. Looking for loop structures and transforming the gotos into loops would be more complicated and I’m not sure it would work in all occasions.

The second condition does not make sense but it is necessary to generate the compiler error ;)