I’m trying to figure out if there is any way of organizing several texture references into an array, so I could at least enumerate them and fetch from a dynamically selected texture.
I’ve tried a couple of obvious ways (texture<>*, std::vector), but none of them compiles. Using one 1D texture instead of several 3D’s has one serious disadvantage - I want to use built-in trilinear interpolation in order to accelerate my program.
I’m using CUDA 4.0 RC2.
Thanks in advance
rainman,
Hey, I’ve been wanting to do the same thing. The texture interpolation is way faster than using regular global memory for me, because of the extra reads required to do the interpolation manually. For CUDA 4.0 RC they have the layered textures, but looks like these are only 1D and 2D, so no luck for those of us who would like 3d arrays of textures. I remember a long time ago someone described using preprocessor macros functions to manage an array of texture references. This looked dangerous, but the guys said it worked. What I’ve been doing since getting a Fermi was to use function pointers to setup the textures I’m going to need, then call the appropriate function. I would really like to have layered 3d textures though…
pojken
Thanks for the tip about function pointers :)
Can you post a brief example for me?
Sure. So my situation is I have electric potentials defined for particular sets of electrode configurations. Depending on what simulation I’m doing, the textures that need to be fetched are different. Hopefully your situation maps onto this one. I do something like:
/* this constant gets set in the host code, for the particular configuration needed, with:
int all3dfunc = ALL3DFUNC_TRANSFEROPTICS;
cudaMemcpyToSymbol(“d_all3dFuncType”,&all3dfunc,sizeof(int),0,cudaMemcpyHostToDevice);
*/
constant int d_all3dFuncType;
// typedef for the function we will be calling
typedef float(all3dGetPotentials_t)(float x,float y,float z,float volts,int& status);
// and here is the function pointer
device all3dGetPotentials_t Get3dPotentials;
// now the different configurations I have are defined in a set of functions. for example:
device inline float
PotentialsTransferOptics(float x,float y,float z,float* volts,int& status)
{
float4 paPotentials;
float total = 0;
paPotentials = tex3D(tex3d0,x,y,z);
total+= paPotentials.x * volts[0];
total+= paPotentials.y * volts[1];
total+= paPotentials.z * volts[2];
total+= paPotentials.w * volts[3];
paPotentials = tex3D(tex3d1,x,y,z);
total+= paPotentials.x * volts[4];
total+= paPotentials.y * volts[5];
total+= paPotentials.z * volts[6];
total+= paPotentials.w * volts[7];
// check for electrode hit
if(paPotentials.x > 1.0f)
status = HITELECTRODE;
return total;
}
// this enum would be somewhere visible to host and this file
enum all3dFunction_t
{
ALL3DFUNC_DUALCELLNORESEJ,
ALL3DFUNC_DUALCELLTK,
ALL3DFUNC_GATEASSEMBLY,
ALL3DFUNC_TRANSFEROPTICS,
ALL3DFUNC_OPTICSANDTRAP,
ALL3DFUNC_NUMBEROFFUNCTIONS
};
// array of all possible functions, each with unique set of textures they have to call
device all3dGetPotentials_t all3dInterpolate_table[ALL3DFUNC_NUMBEROFFUNCTIONS] = {Potential3dDualCellNoResej,
Potential3dDualCellTurnerKruger,Potential3dGateAssembly,PotentialsTransferOptics,Potential3dTransferOpticsAndTrap};
// finally, there is some device function that needs to use the function pointer
device float3
ComputeAccel_All3d(float* volts, float4 pos_pix,int& status,float invmass_amu,float charge)
{
// set the function pointer for getting the potentials
Get3dPotentials = all3dInterpolate_table[d_all3dFuncType];
// convert to grid units, position units of texture file
// positions are assumed to be valid
float3 pos_gu;
pos_gu.x = abs(pos_pix.x*d_elecParams[XDIM].guPerPixx);
pos_gu.y = abs(pos_pix.y*d_elecParams[YDIM].guPerPixy);
pos_gu.z = abs(pos_pix.z*d_elecParams[ZDIM].guPerPixz);
float xm = pos_gu.x - 0.5f;
float xp = pos_gu.x + 0.5f;
float ym = pos_gu.y - 0.5f;
float yp = pos_gu.y + 0.5f;
float zm = pos_gu.z - 0.5f;
float zp = pos_gu.z + 0.5f;
// use function pointer
float V2 = (*Get3dPotentials)(xp,pos_gu.y,pos_gu.z,volts,status);
etc…