how to use tex3D with multi-channel texture data

Accelerated Computing CUDA CUDA Programming and Performance
1

I need to use multi-channel texture. In the documentation (tex3D), text3D supports the operation:
float4 tex3D(sampler3D samp, float3 s)
It returns a four channel value. However, I am not clear how to define the sampler3D correctly. Here are some fragments of what I did:

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include “cutil_math.h”

texture <float4, cudaTextureType3D, cudaReadModeElementType> samp;

global void readDataKernel(float *saver)
{
int xIdx = blockIdx.x * blockDim.x + threadIdx.x;
int yIdx = blockIdx.y * blockDim.y + threadIdx.y;
…
float3 pos = make_float3(posx, posy, pposz);
float4 val;
val = text3D(samp, pos); //FIXME
…
}

extern “C” void readData(cudaArray *d_arrayBuf, float *saver)
{
samp.addressMode[0] = cudaAddressModeBorder;
samp.addressMode[1] = cudaAddressModeBorder;
samp.filterMode = cudaFilterModeLinear;
samp.normalized = false;
cudaChannelFormateDesc channelDesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
//cudaChannelFormateDesc channelDesc = cudaCreateChannelDesc();
cudaMalloc3DArray(&d_arrayBuf, &channelDesc, make_cudaExtent(xNum, yNum, zNum));
…
cudaBindTextureToArray(samp, d_arrayBuf, channelDesc);
unsigned int blockSzX = 8;
unsigned int blockSzY = 8;
dim3 dGrid = {(xNum+blockSzX-1)/blockSzX, (yNum+blockSzY-1)/blockSzY, 1};
dim3 dBlock={blockSzX, blockSzY, 1};
readDataKernel<<<dGrid, dBlock>>>(saver);
cudaUnbindTexture(samp);
}

When I compile the above code, I got an error at the line of val = text3D(samp, pos):
“no instance of overloaded function “text3D” matches the argument list”.
How can I fix this problem? Anyone can help? Thanks!

2

That isn’t CUDA documentation, it is Cg documentation. Cg is obsolete.

CUDA documentation is here:

https://docs.nvidia.com/cuda/index.html

3

Many thanks to Robert!
Now I have figured it out.
For 4-channel texture, the following works,
texture <float4, cudaTextureType3D, cudaReadModeElementType> samp4Channel;
float4 val = text3D(samp4Channel, pos.x, pos.y, pos.z);

For 2-Channel texture,
texture <float2, cudaTextureType3D, cudaReadModeElementType> samp2Channel;
float2 val = text3D(samp2Channel, pos.x, pos.y, pos.z);