Passing a texture object to kernel Problems from passing texture template to kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

To read a cuda array from within the kernel, I am binding it to a texture on the host side.

So something like this:

texture<int, 1, cudaReadModeElementType> texRefA;

cudaBindTextureToArray(texRefA, cuArray);

But the problem is, apparently one can not pass a texture template to the kernel! So I can not do this:

my_kernel <<< n_blocks, block_size >>> (cuArray, texRefA);

__global__ void my_kernel(cudaArray *a, texture<int, 1, cudaReadModeElementType> texRefA)

This problem is also somewhat documented here.

This implies that I need to declare the texture as a global variable, in my main file, and keep the kernel in that same file so it has access to the texture.

Has this problem been fixed?

Is there any way I can pass a texture object? How else can the kernel access a texture object without having to be defined in the same file as where the texture object is defined?

Thanks for your help!

2

As far as i know, textures are supposed to be declared global. Cant say ive actually tried to do otherwise yet but i could see why you would need this.
As far as being in the same file, all you need to do is #include another file with the texture in it if thats what youre looking for.
I usually put struct definitions, textures, and device constant in a cuda_shared.cu file that all my kernel files #include.

Hope this helps.

3

Hmm…I see. Thank you!

4

One (very newbie) question: Do you have to build a custom build step for the cuda_shared.cu file as well? I assume so, but in that case would it be the same nvcc command that you use for the main.cu?

For example, my main.cu is

“C:\CUDA\bin\nvcc.exe” -ccbin “$(VCInstallDir)bin” -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/O2,/Zi,/MT -I"C:\CUDA\include" -I./ -I…/…/common/inc -o $(ConfigurationName)\main.obj main.cu

For cuda_shared I would think it would be the same as above except direct it to cuda_shared.obj

“C:\CUDA\bin\nvcc.exe” -ccbin “$(VCInstallDir)bin” -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/O2,/Zi,/MT -I"C:\CUDA\include" -I./ -I…/…/common/inc -o $(ConfigurationName)\cuda_shared.obj cuda_shared.cu

But then I get this unusual error that everything I have defined in the shared file has already been defined, although they have not.

So I must be doing something wrong with the nvcc command line.

Any ideas?

Thanks for your help!

5

If in visual studio, i only compile the files with the kernel calls in them (not the actual global function definitions but the host code calling them) having #included those files.
Everything but the host code is set to “excluded from build”, including my cuda_shared file.

6

So, I’ve done exactly that. I’ve added

extern const texture<int, 1, cudaReadModeElementType> texRefA;

to a common.cu and excluded it from build, and I’m including common.cu in my main.cu and kernel.cu.

Except now I get the error

“error: attempt to take address of texture variable `texRefA’”

When I try to use texRefA in my kernel.

Do you know why that could be?

I don’t get that error when I take the const away, but then I get the error that I am redefining texRefA

7

I define it as neither extern nor const but i have no clue if thats your problem or not.