Constant memory pointer to constant memory structure in constant memory with pointer to more constan

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi Everyone,

I have a strange problem and need some way around it:

[codebox]constant Lambert materials[1];

constant Sphere objects[1];

//-------------------------Set up materials-------------------------//

Lambert * h_materials = (Lambert*)malloc(sizeof(Lambert));

h_materials[0].set_cd(0.5, 0.5, 0);

h_materials[0].set_kd(0.3);

//--------------------------Set up objects--------------------------//

Sphere * h_objects= (Sphere*)malloc(sizeof(Sphere));

h_objects[0].set_center(0, 0, 0);

h_objects[0].set_radius(2.0f);

h_objects[0].setLambert(h_materials);

CUDA_SAFE_CALL(cudaMemcpyToSymbol(materials, h_materials, sizeof(Lambert)*h_numObjects[0]));

CUDA_SAFE_CALL(cudaMemcpyToSymbol(objects, h_objects, sizeof(Sphere)*h_numObjects[0]));

host device RGBColour lambertF(const ShadeRec& sr, const Vector3D& wo, const Vector3D& wi)

{

	return (sr.lambert->cd * sr.lambert->kd * 0.3183098861837906715f);

}

[/codebox]

The compiler says it can’t see what sr.lambert->cd points to assuming global memory, but it should assume constant memory. Is what I’m trying to do possible? If not, how can I solve this? Please help

2

A quick and dirty workaround - not tested yet:

Create a local variable that is a pointer to constant memory. Assign the pointer from the struct to the local variable. Use the local variable instead to dereference the pointer.

3

How do you create a pointer to constant memory? Any chance you could give an example with my code? External Media

4

Declare a constant memory variable (e.g. array of size 0), initialize a pointer to it. Come on, this can’t be that hard. ;)

This might only solve the compiler warning, but I also notice a second problem:

You are copying the structs to global memory using CudaMemCpyToSymbol(), however the pointer stored in the struct will continue pointing to HOST MEMORY, guaranteeing a crash of the kernel. You would also have to fix the pointer after copying the struct. Best by querying for the address of the struct on the device with the symbol name and writing that into the pointer in the other struct (CudaMemCpyToSymbol will accept an offset parameter). Depending on the number of structs, this may get a little slow.

Maybe instead of using pointers, store an integer array index only that identifies the struct being pointed to.

5

Thanks man, yes you are very clever. I did think to have an array index member field within the struct so I can access the other struct in the other array easier -I think I will have to do this.