CUDA使用纹理内存

纹理内存位于设备端，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) && (y < imageH))
	{
		uint i = __umul24(y, imageW) + x;
		d_output[i] = vox