Different Textures per Thread Via Indirection? Can multiple textures be independently selected for e

Accelerated Computing CUDA CUDA Programming and Performance
1

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]