纹理内存位于设备端,global memory也位于设备端,但是texture memory的访问速度较global memory要快。
因为纹理内存有cache, 只有当cache没有命中的时候才会去访问device memory,否则访问texture cache具有很小的延迟。
另外,texture cache的2D定位已经进行了优化,对于同一线程束的线程访问位置临近的texture memory时效率非常高。还有,texture memory的stream fetching 也进行了优化,所以即使cache没有命中,对texture memory的访问延迟也不会很高。
初始化纹理内存:
texture<uchar, 3, cudaReadModeNormalizedFloat> tex;
cudaArray *d_volumeArray = 0;
extern "C"
void initCudaTexture(const uchar *h_volume, cudaExtent volumeSize)
{
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
cutilSafeCall(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize));
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
copyParams.dstArray = d_volumeArray;
copyParams.extent = volumeSize;
copyParams.kind = cudaMemcpyHostToDevice;
cutilSafeCall(cudaMemcpy3D(©Params));
tex.normalized = true;
tex.filterMode = cudaFilterModeLinear;
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.addressMode[2] = cudaAddressModeWrap;
cutilSafeCall(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));
}
在上述代码中已经将h_volume中的数据拷贝到设备端的d_volumeArray中,然后又将其绑定到一个纹理内存。
下面在kernel函数中,对纹理进行访问,并将数据保存到PBO中,然后绘制。PBO的使用在前面已经写过了,可以参考之。
__global__ void kernel(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) && (