undefined reference to Kernel error while using 3D texture for interpolation

Hi CUDA Gurus,

I encounter a weird undefined reference to my kernel function when i try to insert the 3D textures for interpolation in my code.

After looking at sample SDK’s (simpleTexture and simpleTexture3D) and CUDAIAP 3Dtexture example ,I have placed my texture related code in following way –


    // prepare texture

cudaArray *d_volumeArray = NULL ;
cudaChannelFormatDesc ca_descriptor = cudaCreateChannelDesc();

int const pwI = NX;
int const phI =Nmu ;
int const pdI = NSigma;

cudaExtent const ca_extent = {pwI,phI,pdI};

// CUDA_SAFE_CALL( cudaMemcpyToArray(d_volumeArray, 0, 0, V1, NXNmuNSigma, cudaMemcpyHostToDevice));

copy data to 3D array
cudaMemcpy3DParms copyParams = {0};
copyParams.extent = make_cudaExtent(NX,Nmu,NSigma);
copyParams.kind = cudaMemcpyHostToDevice;
copyParams.dstArray = d_volumeArray;
copyParams.srcPtr = make_cudaPitchedPtr((void*) V1,ca_extent.width*sizeof(float),ca_extent.width,ca_extent.h

// set texture parameters

// bind array to 3D texture
rray,ca_descriptor) );
"call to kernel and global function on device "


//defining texture reference
texture<float,3,cudaReadModeElementType> texinterpRef;

inside global func()

int tx = blockIdx.xblockDim.x + threadIdx.x;
int ty = blockIdx.y
blockDim.y + threadIdx.y;
int tz = threadIdx.z;

int idx = tzdc_NXdc_Nmu + ty*dc_NX + tx;

V[idx] = tex3D(texinterpRef,tx,ty,tz);


The precise error i get is

/tmp/tmpxft_00002015_00000000-14_factor_zero_cuda_single_tex.o: In function main': tmpxft_00002015_00000000-13_factor_zero_cuda_single_tex.i:(.text+0x179e): undefined reference to device_stub__Z24UpdateExpectedCTG_kernelPfS_S_S_S_S_S

collect2: ld returned 1 exit status

I m sure that while doing cudaMemcpy3D/cudaMalloc3DArray ,its going bonkers or may be just the way texture code is place in the two .cu files.

Any pointers/insight will be very helpful as i am still rudimentary with the concept.

thanks much in advance.

The texture declaration and the device code and host code that use it have to appear in the same compilation unit. You need to have the host and device side code in the same file as the texture declaration (that could be via an #include, I think) when at the phase when nvcc is working everything out.

I have already referenced the device.cu in host.cu pretty much like a standard include kernels

smthng like

— host.cu —

// include, kernels