Do I really have to get rid of all my constructors and destructors if i want to use my structs in constant memory?
struct Test
{
float a;
void init()
{
a=1.0f;
}
};
__constant__ Test _test[5]; <-- GOOD
struct Test
{
float a;
Test
: a(1.0f)
{}
};
__constant__ Test _test[5]; <-- BAD
error: can't generate code for non empty constructors or destructors on device
Why is it not allowed to use constructors for the constant memory? As far as I know do constructors exist in plain-vanilla C.
My guess is that there is no hook to run the constructor code on the device for an object with file scope (like a constant object). I see no reason why NVIDIA could not implement this in the future, but it does take some extra work. For example, with the runtime API, the compiler would need to generate code to register with the CUDA library all the device-side constructors that needed to be called. Then, when the device context is establish by one of the CUDA API calls, at that time all the constructor code could be called. I’m sure something similar could be worked out for the driver API as well.
Incidentally, calling constructors for objects declared with file scope on the CPU side is tricky business, especially when you want to make it work for shared libraries. All the details have been long since worked out, but people used to see problems on various platforms. Ex:
My guess is that there is no hook to run the constructor code on the device for an object with file scope (like a constant object). I see no reason why NVIDIA could not implement this in the future, but it does take some extra work. For example, with the runtime API, the compiler would need to generate code to register with the CUDA library all the device-side constructors that needed to be called. Then, when the device context is establish by one of the CUDA API calls, at that time all the constructor code could be called. I’m sure something similar could be worked out for the driver API as well.
Incidentally, calling constructors for objects declared with file scope on the CPU side is tricky business, especially when you want to make it work for shared libraries. All the details have been long since worked out, but people used to see problems on various platforms. Ex:
thank you very much for your post.
So if I want to use my structs/classes in that way in the constant memory, i definetely have to get rid of the constructors - mhmhmh…that’s bad - but a solution ;)
thanks!
thank you very much for your post.
So if I want to use my structs/classes in that way in the constant memory, i definetely have to get rid of the constructors - mhmhmh…that’s bad - but a solution ;)
thanks!
Unless someone from NVIDIA jumps in this thread and tells us that this is fixed in CUDA 3.2, yes. Even then, you probably won’t want to wait, so separating your initialization function and calling it yourself is your best solution at the moment.
Unless someone from NVIDIA jumps in this thread and tells us that this is fixed in CUDA 3.2, yes. Even then, you probably won’t want to wait, so separating your initialization function and calling it yourself is your best solution at the moment.
If it saves you some work, you could instantiate your structs on the host (if they are host and device compatible) and then use cudaMemcpyToSymbol() to get them to constant memory.
If it saves you some work, you could instantiate your structs on the host (if they are host and device compatible) and then use cudaMemcpyToSymbol() to get them to constant memory.
but thanks anyway - i will get rid of my utility classes (Point3D, Vector3D, Normal, …) It really makes no sense anymore to use them in cuda, because i can use float3, float4 instead. The other objects will get an ini() function instead of a constructor and then i should be able to store them in the constant memory → more total speed ;)
but thanks anyway - i will get rid of my utility classes (Point3D, Vector3D, Normal, …) It really makes no sense anymore to use them in cuda, because i can use float3, float4 instead. The other objects will get an ini() function instead of a constructor and then i should be able to store them in the constant memory → more total speed ;)