Is literal string supported? Using literal strings causes code gen error

Accelerated Computing CUDA CUDA Programming and Performance
1

I am hesistant to call this one a bug because the OpenCL spec is not very explicit when it comes to literal strings. But in any case, the following code doesn’t seem to compile.

int stringLen(const char *s)
{
int len = 0;
while(*s != 0)
{
s++;
len++;
}
return len;
}

__kernel void test_calc_kernel(__global void *heap, __global unsigned int *printbuf)
{
int len;
const char *s = “Hello World”;
len = stringLen(s);
printbuf[0] = 0;
}

clBuildProgram (CUDA 2.3) returns an error that says:

ptxas ptx input, line 53; fatal : Parsing error near ‘;’: syntax error
: Retrieving binary for ‘anonymous_jit_identity’, for gpu=‘sm_13’, usage mode=’

However, if I put the string on the __private char buffer, then it seems to compile and run. For example, if I change the second function to:

__kernel void test_calc_kernel(__global void *heap, __global unsigned int *printbuf)
{
int len;
const char s = {‘H’,‘e’,‘l’,‘l’,‘o’,’ ',‘W’,‘o’,‘r’,‘l’,‘d’};
len = stringLen(s);
printbuf[0] = 0;
}

For now, I made a workaround utility that converts all the literal strings in the code to char buffer. That’s not good since I suppose every invocation will involve a construction of the array.

Therefore, I’d like to know if the literal string problem is an expected behavior or something that will hopefully be fixed in a future release. If it is expected, i.e. literal string is for all intents and purposes not a supported feature in Nvidia’s implementation of OpenCL, I’d just resort to a different strategy, say pre-build a string table and put it on either __global or __constant. But if literal string is something that will be fixed, I’ll live with the workaround for now.

Thanks.

2

Hi,
3.0 gives much the same result.
I tried const char s = “Hello World”; within a function body, it breaks the compiler.
Then I tried the same declaration outside a function body, it compiled ok;
Then I tried char *s = “Hello World”; outside, this gives the same error you got.

Cautious conclusions:

  • String literals does not seem to be a problem when initializing global strings;
  • apparently you cannot declare a global variable inside a function body, even if you declare it with __global or static;
  • initializing pointers to a string literal does not appear to be supported;
  • initializing pointers to a string variable seems ok, also as a locally declared pointer.
  • generally, you have to be wary what address space a variable will be in, when using pointers.

I think your workaround could be simpler if you initialize your strings globally (within the kernel).

Hopes this helps somewhat.
Jan