Hi ,
I have been trying to Bind Texture to a 3d Array ,making use of cudaMalloc3DArray (for allocation) and cudaMemcpy3D.
It works fine when used with single precision(floats),but for doubles,it throws an invalid texture reference.
Can some one please help and point out what I am doing wrong when using this for doubles
Following is how the code looks like ( Only relevant texture bit )
///////////////////////////////////////////////////////////////////
// Test texture with doubles
////////////////////////////////////////////////////////////////////
// Declare Texture
texture<int2,3,cudaReadModeElementType> TexRef;
//////////////////////////////////////////////////
// Bind Texture to 3D Array
/////////////////////////////////////////////////
// Set Texture Mutable Reference Parameters
cudaChannelFormatDesc ca_descriptor;
cudaExtent ca_extent;
cudaArray *cuArray=0;
ca_descriptor = cudaCreateChannelDesc();
ca_extent.width = dimx;
ca_extent.height = dimy;
ca_extent.depth = dimz;
// Allocate 3D Cuda Array
cudaMalloc3DArray( &cuArray, &ca_descriptor, ca_extent);
// create a cudaMemcpy3DParms structures
// Structure that tells cuda how to copy data using the cudaMemcpy3D function.
cudaMemcpy3DParms cpy_params = {0};
cpy_params.extent = ca_extent;
cpy_params.kind = cudaMemcpyHostToDevice;
cpy_params.dstArray = cuArray;
cpy_params.srcPtr = make_cudaPitchedPtr((void*)h_data,dimx*sizeof(double),dimx,dimy );
// Copy 3D Data from Host to Device
cudaMemcpy3D( &cpy_params );
// Run Time Texture Reference Attributes
TexRef.normalized = false;
// Filter Mode For Fetch , presuming interpolation doesnt work with doubles
TexRef.filterMode = cudaFilterModePoint;
TexRef.addressMode[0] = cudaAddressModeClamp;
TexRef.addressMode[1] = cudaAddressModeClamp;
TexRef.addressMode[2] = cudaAddressModeClamp;
// Bind Texture to the Cuda Array
cudaBindTextureToArray(TexRef,cuArray,ca_descriptor );
printf(“CUDA Error Post Binding Texture: %s \n”, cudaGetErrorString(cudaGetLastError()));
// Fails here to Bind Texture throwing an invalid texture reference
// //////////////////////////////////////
// for the record ,Fetching Texture code
// /////////////////////////////////////
inline device double fetch_textures_double(texture<int2, 3> t, int x,int y,int z)
{
int2 v = tex3D(t,x,y,z);
return __hiloint2double(v.y,v.x);
}