Hi,
Can I use two levels of indirection for the texture reference in tex1D, thereby causing each thread to have a different texture reference?
I have multiple textures declared and wish to independently select them for each thread. To do this I first made an array of texture references. This works when I use a constant integer in the code as the index to the array, but it selects the same texture for all threads. When I replace the constant integer index with an array which resolves to a constant integer, but a different constant integer for each thread, I get the following error message:
ptxas /tmp/tmpxft_00007160_00000000-2_test2.ptx, line 158; error : Unknown symbol ‘m’
ptxas fatal : Ptx assembly aborted due to errors
make: *** [obj/release/test.cu.o] error 255
Is this possible or must I use one large texture containing all data sets with offsets to select a different data set for each thread?
Thanks
[codebox]// KERNEL
texture<float, 1, cudaReadModeElementType> tex_data_a;
texture<float, 1, cudaReadModeElementType> tex_data_b;
global void
testKernel( float g_idata, float g_odata, int g_sel_data )
{
extern shared float sdata;
texture<float, 1, cudaReadModeElementType> data_ptrs[ 2 ] = { tex_data_a, tex_data_b };
const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
float *Acc;
int tmpc = g_sel_data[ tid ];
Acc = &sdata[ 0 ];
*( Acc + tid ) = g_idata[ 0 * numB + tid ];
__syncthreads();
for ( int i = 0; i < numA; i++ ) {
g_odata[ tid * numA + i ] = // following are the alternatives I have tried
tex1D ( tex_data_a, *( Acc + tid ) ); // Works
tex1D ( data_ptrs[ 1 ], *( Acc + tid ) ); // Works
tex1D ( ( data_ptrs[ ( g_sel_data[ tid ] ) ] ), *( Acc + tid ) ); // Gives above error with symbol 'f'
tex1D ( data_ptrs[ tmpc ], *( Acc + tid ) ); // Gives above error with symbol 'm'
}
__syncthreads();
g_idata[ 0 * numB + tid ] = *( Acc + tid );
}[/codebox]