Texture & dynamic parallelism

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,
is it possible to read a texture within a child kernel? tex3D(…) always returns 0. But if I perform the same read of the texture within the parent kernel… it works.

texture<unsigned char, 3, cudaReadModeNormalizedFloat> tex;

__global__ void childKernel(float* arr)
{
  *arr = tex3D(tex, 0, 0, 0);
}

__global__ void parentKernel(float* arr)
{
  childKernel<<<1, 1>>>(arr);
  cudaDeviceSynchronize();
}

*arr contains 0.

texture<unsigned char, 3, cudaReadModeNormalizedFloat> tex;

__global__ void parentKernel(float* arr)
{
  *arr = tex3D(tex, 0, 0, 0);
}

*arr contains the right value.

thanks

2

I believe you need to use texture objects (and not legacy textures) in dynamic parallelism.

3

I tried to use texture object but the result was the same. I do another try anyway.

but… does this sentence mean that texture created in the host can’t be used in child kernel?
“…textures may be created from the host and used in device code as for any kernel, but may only be used by a top-level kernel (i.e. the one which is launched from the host).”

4

it works! I used texture object passed as argument to parent kernel and then to child. thanks a lot!