大家好，我在程序中需要对三维纹理进行采样，但是每次计算的时候，三维纹理的值需要更新。因此我使用了cudaMemcpy3D来将新的数值传输到三维纹理所绑定的cuda数组里面。但是我测试发现，该函数传输的速度非常的慢，大约只有同尺寸数据使用cudaMemcpy传输速度的十分之一左右（均是Device2Device）。不知道这是为什么？有什么办法可以解决呢？
或者是否有办法可以绕过这个函数，直接修改纹理的内容？似乎目前是不支持的。
楼主您好，我手头没有对cudaMemcpy3D的具体评测报告。
请问您是从cuda array复制到cuda array呢？还是从普通显存复制到cuda array呢？
后者需要进行数据重拍，慢一点正常的，但是否只有1/10我不清楚。
虽然从普通的线性显存复制到cuda array, 因为需要重拍元素而导致复制慢点，但是可以在使用的时候提供更好的性能。
请综合考虑此2因素，如果综合还是值得的，那么建议就这样使用。
如果综合考虑不值得，则您可以考虑如下建议：
（1）总是从cuda array复制到cuda array，或者从普通线性显存复制到普通线性显存。而不要从普通线性显存复制cuda array(或者反过来）.
（2）如果无法避免这个操作，建议考虑如下：
（2-1）将纹理绑定到线性显存，而不是cuda array, 这样您可以直接用普通cudaMemcpy更新纹理（下次启动有效）。甚至可以直接在kernel里更新纹理（下次kernel启动有效，通过绑定的普通指针，而不是纹理对象、引用）。
（2-2）如果还是想绑定到cuda array, 且可以的话，请尝试使用surface写入，它可以直接就地更新cuda array（依然是下次kernel启动有效）。
感谢您的来访，周末愉快。
感谢您的回答。我这个程序的简单一点的说，目的大概如下：
1.将三维数据A (实际是在global中是一维线性存放)，绑定到三维纹理，进行一些采样操作；
2.在显存中进行一些计算，并更新A的内容；
3.将新的A重新绑定到纹理，并转入1.
根据您的(2-1), 普通的线性显存也可以绑定到三维纹理是么？我的卡是K20, CUDA版本 5.5。如果是这样的话，那就太好了，我既能省下一个三维数组的空间，也能不需要大量的数据拷贝了。 我印象中线性内存只能绑定到一维纹理中（大概3.0以前是这样的），是否新版本的CUDA已经支持这种做法了？
楼主您好，我没有说过“可以将三维纹理和普通线性显存绑定”。
实际上上我的原文如此：
“请问您是从cuda array复制到cuda array呢？还是从普通显存复制到cuda array呢？”
以及“将纹理绑定到线性显存，而不是cuda array, 这样您可以直接用普通cudaMemcpy更新纹理（下次启动有效）。甚至可以直接在kernel里更新纹理（下次kernel启动有效，通过绑定的普通指针，而不是纹理对象、引用）。”
这两段。
只有1维和2维的能绑定到普通内存，所以我直接询问的是您是从“普通显存复制到cuda array"还是从"cuda array复制到cuda array".
并根据这个给出了尽量避免从线性内存复制到cuda array的建议。（以及其他建议）
但是我没有给出“普通的线性显存和3维纹理绑定”的建议。
我给出的将“纹理和线性内存绑定”是需要您手工计算坐标变换的。但绝不是直接可以这样。
以及，我在2#给出的建议已经基本考虑了所有情况了。
您可以考虑采用一些，或者当然也可以无视，这是您的自由选择。
我重申一下2#的话, 以及阐明点，避免你误会：
建议考虑使用array到array的复制。不建议使用普通线性内存到array的复制（如果你的综合代价不核算的话）。
建议考虑使用普通线性内存和纹理引用绑定。（需要你手工计算坐标）
建议使用surface就地修改cuda array. （如果你能写kernel来就地修改是可以的）
感谢来访，周末愉快。
对于三维纹理，主要是想用到三线性插值，比较方便也比较快，如果通过一维纹理或者直接数组操作的话，计算的开销会比较大。
根据您的建议“建议使用surface就地修改cuda array. （如果你能写kernel来就地修改是可以的）”，那是不是意味着：我可以避免掉将修改后的线性显存数据拷贝到3d cuda 数组的过程，而是直接更新与三维纹理绑定的3D cuda array？
如果你的目的是用texture sfu来进行插值，那么只能使用cuda array的3D纹理。
这个有单独的texture SFU来进行计算(不过请注意精度只有9位)。
想节省计算就只能这样，而采用1D/2D纹理就必须手工计算插值。这个需要你作出抉择。
此外，surface是写入的就是cuda array, 这也是目前我们直接修改cuda array的唯一办法。
这个是否可以“避免掉将修改后的线性显存数据拷贝到3d cuda 数组的过程”，要看你是否真的可以这样，因为这个修改只有下次kernel启动才有效（这次如果读取，会产生undefined的值的）。
而线性内存本次修改了可以直接读取，请根据你的需要来选择是这次修改，下次使用的surface方式。还是直接就能用的普通内存方式，但后者需要复制一次。这个也请你作出抉择。
感谢来访，夜安。
谢谢版主，我明天试一下 surface写入。
嗯嗯。欢迎尝试，别忘记注意事项哈。
夜安。亲。
版主您好，我尝试着使用surface来改写cuda数组的值，但是编译时报错，出错的相关代码和提示如下：
error : External calls are not supported (found non-inlined call to __surf3Dwriteu1)
float ProjSampel = tex2D(texProj2D, fProjAdd.x, fProjAdd.y);
surf3Dwrite(ProjSampel, SurRef3DImage, ImageX, ImageY, i, cudaBoundaryModeZero);
其中ImageX, ImageY, i 均为整形，此外我的Surface的声明如下：
surface<void, cudaSurfaceType3D> SurRef3DImage;
是创建的Surface类型不对么？
呵呵 找到原因了 将compute_和sm_的参数调高就行了。默认都是10
这两个参数是不是就根据当前显卡的最高计算能力设置就行了？
LZ您好：
建议按照您实际GPU的计算能力版本设置即可。
目前主流GPU中：
fermi核心的为2.x版的计算能力，具体为：GF100/110核心的为2.0的计算能力，其他为2.1的计算能力。在VS2010中这两种核心都设置为compute20，前者设置为sm20，后者设置为sm21；在VS2008中，前者直接设置sm20，后者直接设置sm21即可。
kepler核心的计算能力为3.0和3.5。具体为，telsa K20/K20X，Geforce TITAN，Geforce GTX780是3.5的，其他为3.0的。在VS2010中，前一类设置为compute35，sm35,后一类设置为compute30，sm30；在VS2008中前一类设置为sm35，后一类设置为sm30。
其他环境下需要命令行输入编译参数的，请参阅相关手册。
以及，如果您只是需要编译一下代码，而不需要在您的GPU上执行（比如编译出来给他人执行），那么您可以指定比您的GPU实际计算能力版本高的编译参数。
以及，计算能力版本越高，各方面的限制越少。
大致如此，祝您编码顺利~
恩 我改成sm_35和compute_35了，用的是K20C。此外还遇到个问题，surface和数组的绑定总是不成功：
[
ET0返回成功，ET2返回cudaErrorInvalidValue，不知道为何。
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
surface<void, cudaSurfaceType3D> SurRef3DImage;
另附上相关变量的声明
楼主您好，您需要将您的行：
ET0 = cudaMalloc3DArray(&cu_Image3D, &channelDesc, volumeSize);
改成：
ET0 = cudaMalloc3DArray(&cu_Image3D, &channelDesc, volumeSize, cudaArraySurfaceLoadStore);
方可使用surface。
默认的无load/store参数的array只能用于texture的。
谢谢版主，确实如此，这就是看文档部不仔细的结果啊…
另外还有一个问题，就是对于数据类型为32位浮点的时候，surface写入函数的参数的X，需要乘以数据字节数么？
我看手册上面没有强调，只是说坐标位置x,y,z。
surf3Dwrite(ProjSampel, SurRef3DImage, ImageX*sizeof(float), ImageY, i, cudaBoundaryModeTrap);
还是
surf3Dwrite(ProjSampel, SurRef3DImage, ImageX, ImageY, i, cudaBoundaryModeTrap);
楼主您客气了，服务您是我们的荣幸。
这个是需要的，您应该看到这里了。