PDA

View Full Version : nvidia 197.45 crash : global initialized with literal



matrem
06-08-2010, 07:45 AM
With last nVidia release, initializing a global vector type variable with a literal crash the compiler.

this example crash in clBuildProgram :


float4 f = (float4)(0.f, 0.f, 0.f, 0.f);

__kernel void test()
{
}

Ati stream on its side absolutely want f to be declared in constant memory space... I don't see this requirement in the specification.

david.garcia
06-08-2010, 08:48 AM
Ati stream on its side absolutely want f to be declared in constant memory space... I don't see this requirement in the specification.

AMD is doing the right thing here. See section 6.5 of the spec: "All program scope variables must be declared in the __constant address space."

matrem
06-08-2010, 09:14 AM
A global variable must be a program scoped variable?
Can't we have global kernel scoped variables (initialized at kernel execution)?

In fact I read this sentence but it was not really clear for me...

matrem
06-08-2010, 09:17 AM
By the way using __constant don't change the crash with 197.45.

With 257.15 beta drivers, without constant we have the same error than ati, and with constant there is an opencl error (it's a bit better than a crash) with a message speaking about "Ptx compilation failed".

david.garcia
06-08-2010, 03:36 PM
A global variable must be a program scoped variable?

All variables declared outside of a kernel (i.e. at program scope) must be constant. The term "global variable" is ambiguous and is better avoided.

You may get better feedback if you report the problems you are seeing to the company that is shipping the particular SDK you are using.

matrem
06-09-2010, 01:29 AM
Thanks.
I'll report to nVidia forum, but I always think that it should be good to have a section on khronos forum to report implementation bugs. :)