请问一个有关cudaMemcpy3D的问题

大家好,我在程序中需要对三维纹理进行采样,但是每次计算的时候,三维纹理的值需要更新。因此我使用了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);

楼主您客气了,服务您是我们的荣幸。

这个是需要的,您应该看到这里了。