3d texture fetching

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello everyone,

I have a problem of 3d texture fetching. I get the code of simpleTexture3d in SDK and have some modification of it. First, the voxel type is unsigned int, not unsigned char.
And I use tex3d with cudaFilterModePoint, but I got the unexpexted output. And I do not know why, I have searched this forum but maybe I do not get the right one.
Attached below are the code of this,

Thanks for any advice,
Hai

#ifndef SIMPLETEXTURE3D_KERNEL_CU
#define SIMPLETEXTURE3D_KERNEL_CU

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include

#include <cutil_inline.h>
#include <cutil_math.h>

typedef unsigned int uint;

texture<uint, 3, cudaReadModeElementType> tex; // 3D texture
cudaArray *d_volumeArray = 0;

global void
d_render(uint *d_output, uint imageW, uint imageH, float w)
{
uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;

float u = x / (float) imageW ;
float v = y / (float) imageH ;
float voxel = tex3D(tex, u, v, w);

if ((x < imageW) && (y < imageH)) {
    // write output color
    uint i = __umul24(y, imageW) + x;
    d_output[i] = (uint)(voxel);
}

}

extern “C”
void setTextureFilterMode(bool bLinearFilter)
{
tex.filterMode = bLinearFilter ? cudaFilterModeLinear : cudaFilterModePoint;
}

extern “C”
void initCuda(const uint *h_volume, cudaExtent volumeSize)
{
// create 3D array
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc();
cutilSafeCall( cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize) );

// copy data to 3D array
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume, volumeSize.width * sizeof(uint), volumeSize.width, volumeSize.height);
copyParams.dstArray = d_volumeArray;
copyParams.extent   = volumeSize;
copyParams.kind     = cudaMemcpyHostToDevice;
cutilSafeCall( cudaMemcpy3D(&copyParams) );

// set texture parameters
tex.normalized = true;                      // access with normalized texture coordinates
tex.filterMode = cudaFilterModePoint;      // Point texel
tex.addressMode[0] = cudaAddressModeWrap;   // wrap texture coordinates
tex.addressMode[1] = cudaAddressModeWrap;
tex.addressMode[2] = cudaAddressModeWrap;

// bind array to 3D texture
cutilSafeCall(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));

}

extern “C”
void render_kernel(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH, float w)
{
d_render<<<gridSize, blockSize>>>(d_output, imageW, imageH, w);
}

#endif // #ifndef SIMPLETEXTURE3D_KERNEL_CU
simpleTexture3D.cpp (5.49 KB)