PDA

View Full Version : how to declare a constant array of float2 vectors?



ksi
04-14-2011, 02:38 PM
In my kernel code I am trying to declare a constant array of float2 vectors like this:


__constant float2 grads[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};


But this gives a rather obscure error when attempting to build the kernel:

ptxas application ptx input, line 11; fatal : Parsing error near ',': syntax error
ptxas fatal : Ptx assembly aborted due to errors
error : Ptx compilation failed: gpu='sm_20', device code='cuModuleLoadDataEx_4'
: Considering profile 'compute_20' for gpu='sm_20' in 'cuModuleLoadDataEx_4'
: Retrieving binary for 'cuModuleLoadDataEx_4', for gpu='sm_20', usage mode=' '
: Considering profile 'compute_20' for gpu='sm_20' in 'cuModuleLoadDataEx_4'
: Control flags for 'cuModuleLoadDataEx_4' disable search path
: Ptx binary found for 'cuModuleLoadDataEx_4', architecture='compute_20'
: Ptx compilation for 'cuModuleLoadDataEx_4', for gpu='sm_20', ocg options=' '

Any suggestions for the correct way to do this?

I'm using the Linux Nvidia OpenCL implementation.
This code did previously work ok on Mac, so I was thinking it was correct syntax...

Thanks!

david.garcia
04-14-2011, 03:40 PM
Did you declare grads at program scope or at function scope? That is, does it look like "foo" or like "bar" in the code below:



// Program scope. This is valid and should compile.
__constant float2 foo[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};

__kernel void fubar()
{
// Function scope. This is not valid and should not compile.
__constant float2 foo[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};
}


Section 6.5.3. of the spec says:

Variables allocated in the __constant address space can only be defined as program scope variables and are required to be initialized.

ksi
04-14-2011, 05:23 PM
Yes it is at program scope. (not function scope)

ksi
04-17-2011, 11:53 AM
...so it should be correct syntax and should work, but doesn't. Any ideas? Bug?

david.garcia
04-17-2011, 03:05 PM
Yeah, as I mentioned earlier, if it is in program scope it should work. There's a compiler bug.

AngryLittleBird
04-18-2011, 03:10 AM
Thanks - had the same problem now it's resolved!

ksi
04-18-2011, 02:11 PM
Care to share with us how it got resolved?

ksi
04-27-2011, 11:45 AM
just an update on this, in case anybody else is running into the same issue:
Nvidia says it should be fixed in their driver r280 release.

grabner
11-20-2012, 11:53 AM
just an update on this, in case anybody else is running into the same issue:
Nvidia says it should be fixed in their driver r280 release.
I did run into the same issue (with driver 306.94) and did some further investigations, see https://devtalk.nvidia.com/default/topic/524095.

Kind regards,
Markus