CUDA bug: use of ""

Hello,

I seem to have found a CUDA bug.

Any time I try to use a string constant in a device or global function I get an internal compiler error.

Code similar to this:

__device__ void somefunction(char*f) {

/* code goes here */

}

__global__ void cuda_function() {

    somefunction("");

}

Produces this error:

### Assertion failure at line 793 of ../../be/cg/NVISA/exp_loadstore.cxx:

### Compiler Error in file /tmp/tmpxft_00002c1d_00000000-1.i during Code_Expansion phase:

### variable not in memory space

nvopencc INTERNAL ERROR: /usr/local/cuda/open64/lib//be returned non-zero status 1

If I remove the string and replace it this, then there is no error:

__device__ void somefunction(char*f) {

/* code goes here */

}

__global__ void cuda_function() {

 char f = (char)0;

 somefunction(&f);

}

Sorry, I can’t post the real code. :(

The compiler should not assert, so this looks likely to be a bug. However, I’d need a means of reproducing it to investigate further. If you can provide the cude/app privately, please send me a PM.

I spent a bit of time and made it into something I can put out into public:

Also attached Linux makefile.txt (rename to Makefile) that assumes file is named bugbug.cu, and that cuda is in the default install location.

The system is RHEL 5 x86_64, Dual 8800GTX.

#include<stdio.h>

__global__ void cuda_function(char* foo) {

   foo[0] = 'a';

   foo[1] = (char) 0;

   char * bar = "";  

  

   // are the two pointers the same?

   // should NEVER be true.

   // output should always be 'a'

  // BUGBUG

   if (foo == bar) { foo[0] = 'b'; }

}

int main () {

   char bar[10];

   char* foo;

  cudaMalloc((void**)&foo,sizeof(char)*10);

   dim3 block_size (1);

   dim3 num_blocks(1);

  cuda_function<<<num_blocks,block_size>>>(foo);

   cudaMemcpy(bar,foo,sizeof(char)*10,cudaMemcpyDeviceToHost);

   printf("%s\n",bar); 

   cudaFree(foo);

}

Exact error for this code is:

### Assertion failure at line 793 of ../../be/cg/NVISA/exp_loadstore.cxx:

### Compiler Error in file /tmp/tmpxft_00002ed0_00000000-1.i during Code_Expansion phase:

### variable not in memory space

nvopencc INTERNAL ERROR: /usr/local/cuda/open64/lib//be returned non-zero status 1

make: *** [default] Error 255

Makefile.txt (1.84 KB)

Well, “” is a pointer to the data segment, which is on CPU.
Unless the compiler automatically detect and copy the thing to const space, the kernel won’t run anyway (if the content of the char* is used).
From the look, it seems more of an unimplemented path than a bug.