三维纹理怎么访问？

我现在有三张图像，分别读取后依次赋给float* h_data,然后由cudaMemcpy3D赋给一个texture
现在 将dim3 block(16,16,3); dims thread(32,32,1);
然后 在kernel函数中，如果想访问三张图像在x,y相同位置处的值，该怎么texture上提取呢？
先谢谢版主啦！

楼主您好，

首先感谢您的深夜来访。
但需要指出的是，您不能直接将数据复制到“Texture”，您需要将数据复制到一个CUDA Array，然后将此array和特定的texture reference绑定在一起，然后方可通过texture reference访问您的数据。

具体的说：
（1）您需要使用cudaMalloc3DArray()建立一个3D的CUDA Array或者一个多层的2D CUDA Array（这两者会影响数据的实际存放格式，在具体的应用中将展现不同的具体运行速度效果，请分别尝试哪个对您的算法来说更快。如果不想考虑，您可以直接只用一个3D的CUDA Array。
（2）传输前1/3的数据：您读取第一个图到float * h_data; 调用cudaMemcpy3D，并指定目标偏移为(0,0,0), 指定传输大小为(你的图的x上的元素数, y上元素数, 1)。这将传输数据到刚刚建立的CUDA Array的前1/3部分。
（3）继续中间1/3的数据：您继续读取第二个图，设定目标偏移为(0,0,1)，保持其他参数不变，调用cudaMemcpy3D()传输到cuda Array的中间1/3
（4）继续最后1/3的数据：您继续读取第三个图，设定目标偏移为(0,0,2），其他同上。这次操作完成后将完全填充好整个cuda Array。
（5）调用cudaBindTextureToArray(), 将您的纹理引用和存放数据的cuda array绑定在一起。
（6）在您的kernel中，
调用tex3D(您的纹理引用, X坐标，Y坐标, 0.0f);访问您的第1个图的(x,y)元素。
调用tex3D(您的纹理引用, X坐标，Y坐标, 1.0f);访问您的第2个图的(x,y)元素。
调用tex3D(您的纹理引用, X坐标，Y坐标, 2.0f);访问您的第3个图的(x,y)元素。

这是详细的全部流程，请直接参考到您的代码中，
感谢您的深夜来访。

横扫千军 发表于 2013-10-23 01:15
楼主您好，

首先感谢您的深夜来访。

应该是我感谢版主的深夜解答，不过我现在还有些问题：
1、你说的目标偏移保存是那个函数，我最近刚学，忘版主谅解，还有楼主看下我自己写得可以不？
char *hsrcData = (char*)malloc(width*height*3*sizeof(char));
将三幅图像内存连续赋给hsrcData，然后
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat);
        cudaExtent extent = make_cudaExtent(512,512,3);
        cudaArray *cuArray;
        cudaMalloc3DArray(&cuArray, &channelDesc, extent);
        //cudaMemcpy3D;
    checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, hData, size, cudaMemcpyHostToDevice));
    // Set texture parameters
    tex.addressMode[0] = cudaAddressModeWrap;
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.filterMode = cudaFilterModeLinear;
    tex.normalized = true;    // access with normalized texture coordinates
    // Bind the array to the texture
    checkCudaErrors(cudaBindTextureToArray(tex, cuArray, channelDesc)); 这样可以不？
2、我将线程这样划分可以不？

    dim3 dimBlock(16, 16, 3);
    dim3 dimGrid(32, 32, 1);
    transformKernel<<<dimGrid, dimBlock, 0>>>(dOutData, width, height);

xinliangyu 发表于 2013-10-23 10:26
应该是我感谢版主的深夜解答，不过我现在还有些问题：
1、你说的目标偏移保存是那个函数，我最近刚学，忘 ...

楼主我深刻无语您的回答，

请问您看了2#的回复了么？您真的看了么？？？
您都没看到提示您使用cudaMemcpy3D， 却在一个劲的问dstPos参数放在哪里。我深刻无语。

我和横扫千军版主将在7天内不再回答您的任何问题。请见谅。

尊重是相互的，几乎上千个字的回复被赤裸裸的无视，我也只能无视您。
请尊重版主们的劳动成果。

玫瑰幻想 发表于 2013-10-23 13:40
楼主我深刻无语您的回答，

请问您看了2#的回复了么？您真的看了么？？？

对不起版主辛苦的回答，是我自己没认真看，看来这几天没地方请教问题了！