Hi all,
Is there any way to pass a texture type, or textureReference into a kernel function? I have one kernel function that takes two buffers (same size) as input. Each time, it updates the first buffer in-place, and every second time I call the function it swaps the buffer pointers.
__global__ MyKernel(float* bufA, float* bufB)
{
get linear index -> idx
bufA[idx] = somefunctionof(bufA, bufB);
}
If I was just using global memory, this code on the host would work:
float *bufA, *bufB;
for( int iter=0; iter<10; iter++ )
{
MyKernel<<<grid,threads>>>(bufA, bufB);
swap(bufA, bufB);
}
Instead, I’m using texture memory, and therefore have declared two global texture<> variables:
texture<float, 1, cudaReadModeElementType> aTex;
texture<float, 1, cudaReadModeElementType> bTex;
and bound them to the respective buffers in the host code above, before the for loop.
For reads from bufA or bufB, MyKernel now uses tex1Dfetch on aTex and bTex (though for writes, it still does bufA[ind] = blah)… therefore swapping the pointers in the host code has no effect (since aTex and bTex always refer to bufA and bufB).
Solution 1.
Code two MyKernels, one that uses texA/texB, the other identical, but swaps texA for texB. ← ugly from a code maintenance standpoint
Solution 2.
One MyKernel, with an if statement ← no thanks
Solution 3.
Re-bind the textures during every iteration of the for loop ← performance impact unknown
Solution 4.
(Hopeful) Pass the texture<…> variables in as parameters to MyKernel (doesn’t appear to work) … or, obtain a runtime textureReference* to the textures, and pass those in.
It doesn’t appear that tex1Dfetch supports Solution 4 though…
Has anyone dealt with this before? Any solutions?
Thanks :)
Dan