Problems when I use -O3

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello.

I have some problems when I use -O3 in my cuda program. I have simplified the problem to this code:

This is the main file:

[codebox]

#include <stdio.h>

#include “prueba_kernel.cu”

int main(){

dim3 grid (10);

dim3 block (128);

double * array_d, * array_h;

int i;

array_h = (double ) malloc (10sizeof(double));

cudaMalloc((void **)&array_d, 10*sizeof(double));

for (i=0; i<100; i++){

  prueba_kernel<<<grid, block>>>(array_d, i);

  cudaMemcpy(array_h, array_d, 10*sizeof(double), cudaMemcpyDeviceToHost);

  printf("IT: %d, %lf\n", i, array_h[0]);

}

cudaFree((void *)array_d);

free((void *) array_h);

}

[/codebox]

This is the _kernel.cu file:

[codebox]

global void prueba_kernel(double * array_d, int i){

unsigned int gid = (blockIdx.x * blockDim.x + threadIdx.x);

if (gid < 10){

  array_d[gid] = i;

}

}

[/codebox]

And I use this to compile:

[codebox]

nvcc -O3 -arch sm_13 prueba.cu -o prueba

[/codebox]

The result is (when I execute the kernel several times):

[codebox]

IT: 0, 99.000000

IT: 1, 1.000000

IT: 2, 2.000000

IT: 3, 3.000000

…

IT: 98, 98.000000

IT: 99, 99.000000

[/codebox]

gcc (Debian 4.3.4-6) 4.3.4 and cuda toolkit 3.0beta

Any help?

Thank you so much!

2

I can confirm what you are seeing but I can’t explain why. The optimization flags are making no difference to the PTX the compiler is generating, which is really weird…

3

Thank you!

Yes, it’s really strange… Maybe it’s due to gcc…
I’m trying to unroll the first iteration, and then I don’t have the problem, but if the loop is more complicated this solution doesn’t work :(
Nobody can explain us what’s happening?

4

It works fine for me on both CUDA 2.3 and 3.0 on RHEL4.

5

What gcc version do you use?

6

Must be a Ubuntu specific gcc thing…

7

On RHEL4, gcc 3.4.6.
I also tried RHEL5, gcc 4.1.2, it works fine also there.

You could try this:

  1. remove the include of the kernel
  2. nvcc -O3 -arch sm_13 -c prueba_kernel.cu
  3. gcc -O3 prueba.c prueba_kernel.o -L/usr/local/cuda/lib64 -lcudart
8

Thank you! But gcc can’t find some things:

prueba.c:6: error: ‘dim3’ undeclared (first use in this function)

…

prueba.c:18: error: ‘cudaMemcpyDeviceToHost’ undeclared (first use in this function)

…

9

You will need to include cuda_runtime.h.
You cannot use the <<< >>> syntax in .c, so you will need to move all the kernel launch to the .cu file or use the driver API.

10

Ok, now it works, but… I can’t do this kind of changes in my real program!!
If it works now, is the problem in nvcc? Hmmm, the program now is quite different. Maybe the compiler does a different optimization… I’m using a function with the kernel lauch in the .cu file.

What do you thing?

11

I confirm the problem on Debian x64 with gcc 4.3.4 and CUDA Toolkit 3.0 beta.
When the kernel is launched for the first time, cudaSetupArgument ends up being called with some ridiculous value in the ‘offset’ argument (0x7fffffffe1d8), and the launch fails with the error “invalid argument” (please check the return values!).

Looks like a problem between some recent gcc optimizations and the CUDA Runtime. Maybe some ABI mismatch, or just a gcc bug?..

12

Syl,

Does “volatile dim3 …” help?

13

Interesting…
If I install another gcc version, how can use it with nvcc? Is there any enviroment variable?

Thank you.