hi, i’d like to use texture for vectors which elements are complex struct types defined by me.
Is it possible?
Because i’ve seen from prog guide that texture defitnition supports only int, float and other cuda types defined in appendix B.
I tried to use CudaArrays, but i had problems for the same reason, and actually i can’t resolve the main problem. It seems everything is ok for normal types, but not for any special.
Probably i’m not catching all features, can anyone explain to me?
in particular, i’d like to tex a vector in which every element is a 3 float structure.
I’ still can’t compile. I tried to map it in float3. Is this the right way?
Read the programming guide line you mentioned more carefully:
3 element vectors are not mentioned in the list of supported types. As the guide says, see B.3.1 for a fully list of supported types (excepting those 3 element types, of course, and double precision numbers).
If you want to texture fetch some other object, you must align it on the same boundaries as one of the existing vector types, read that type with a vector and reconstruct your object with bit manipulation or other means. Search the forum for examples of this with doubles. An int2 (or int4) is read with a texture and two neighboring ints are combined to form a double.
Thank you for your answer, i think i had a little misunderstanding on it.
One last thing: since i’ve only to read these vectors, what about using the costant memory? it’s fully cached, isn’t it?
Constant memory is only useful when all threads in a warp are reading the same address from constant memory. If different threads in a warp read different values, then performance suffers and shared memory or textures are often faster, but you should benchmark to be certain.
Thanks for the complete answer. In my case texture would work better.
i’ve tried an approach with cudaArray, but i can’t find what i didn’t understand.
Hope somebody can explain me the mechanism.
My problem is to texture a vector of this type:
[codebox]typedef struct { float x, y, z; } VecR[/codebox]
This vector is called R_ in host memory, and his length is nSites.
So i defined a texture,
[codebox]texture<float3, 1, cudaReadModeElementType> texRef; [/codebox]
the cudaArray,
[codebox]cudaArray* cuArray;
cudaChannelFormatDesc channelDesc_V = cudaCreateChannelDesc();
cudaMallocArray (&cuArray, &channelDesc_V, nSites, 1);
cudaMemcpyToArray(cuArray, 0, 0, R_, sizeof(float3)*nSites, cudaMemcpyHostToDevice);
[/codebox]
and the texture bind:
[codebox]texRef.addressMode[0] = cudaAddressModeClamp;
texRef.filterMode = cudaFilterModePoint;
texRef.normalized = false;
cudaBindTextureToArray (texRef, cuArray, channelDesc_V );[/codebox]
Then in kernel i’ve tried the access:
[codebox]float3 Z;
Z=tex1D(texRef,i);
[/codebox]
where i is an integer defined previously with threadId.x and blockId.x.
When i try to compile it gives me this error:
error: no instance of overloaded function “tex1D” matches the argument list argument types are: (texture<float3, 1, cudaReadModeElementType>, int)
I didn’t understand if the texture definition is ok (from the prog guide, float3 is a supported type…); and if it is, what’s wrong with the tex1D call?
Thank for the patience
As was already pointed out to you, it seems there is no float3 texture support. Choose from float, float2, or float4.
ok you were right, i tried with float4 and it works. I misunderstood that line of programming guide.
Anyway, when i evaluated with gprof i had cudaErrors on functions cudaBindTexture & cudaUnbindTexture. I know that a cudaError on lauch is normal, what about these ones?
The doubt that somethin’s not working comes to me because i checked memory accesses on cudaprof, and i saw that 128 bits fetches are 0.
That’s strange because i thought there would be, since i defined a texture on a float4 cudaArray…