我正在尝试使用CUDA内核以异步方式拆分和重塑数组的结构。 memcpy()在内核中无法工作，cudaMemcpy()也是如此；我很困惑。
有人能告诉我从CUDA内核内部复制内存的首选方法吗？
值得注意的是，cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice)对于我所要做的事情不起作用，因为它只能从内核外部调用并且不会异步执行。
CUDA内核中是否有类似于memcpy()的等效函数？
22
- Zak
3
3个回答
40
是的，在CUDA内核中有一个等同于memcpy的函数，它也叫做memcpy。以下是一个示例：
__global__ void kernel(int **in, int **out, int len, int N)
{
int idx = threadIdx.x + blockIdx.x*blockDim.x;
for(; idx<N; idx+=gridDim.x*blockDim.x)
memcpy(out[idx], in[idx], sizeof(int)*len);
}
能够像这样无错误地编译:
$ nvcc -Xptxas="-v" -arch=sm_20 -c memcpy.cu
ptxas info : Compiling entry function '_Z6kernelPPiS0_ii' for 'sm_20'
ptxas info : Function properties for _Z6kernelPPiS0_ii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 11 registers, 48 bytes cmem[0]
并发出 PTX：
.version 3.0
.target sm_20
.address_size 32
.file 1 "/tmp/tmpxft_00000407_00000000-9_memcpy.cpp3.i"
.file 2 "memcpy.cu"
.file 3 "/usr/local/cuda/nvvm/ci_include.h"
.entry _Z6kernelPPiS0_ii(
.param .u32 _Z6kernelPPiS0_ii_param_0,
.param .u32 _Z6kernelPPiS0_ii_param_1,
.param .u32 _Z6kernelPPiS0_ii_param_2,
.param .u32 _Z6kernelPPiS0_ii_param_3
)
{
.reg .pred %p<4>;
.reg .s32 %r<32>;
.reg .s16 %rc<2>;
ld.param.u32 %r15, [_Z6kernelPPiS0_ii_param_0];
ld.param.u32 %r16, [_Z6kernelPPiS0_ii_param_1];
ld.param.u32 %r2, [_Z6kernelPPiS0_ii_param_3];
cvta.to.global.u32 %r3, %r15;
cvta.to.global.u32 %r4, %r16;
.loc 2 4 1
mov.u32 %r5, %ntid.x;
mov.u32 %r17, %ctaid.x;
mov.u32 %r18, %tid.x;
mad.lo.s32 %r30, %r5, %r17, %r18;
.loc 2 6 1
setp.ge.s32 %p1, %r30, %r2;
@%p1 bra BB0_5;
ld.param.u32 %r26, [_Z6kernelPPiS0_ii_param_2];
shl.b32 %r7, %r26, 2;
.loc 2 6 54
mov.u32 %r19, %nctaid.x;
.loc 2 4 1
mov.u32 %r29, %ntid.x;
.loc 2 6 54
mul.lo.s32 %r8, %r29, %r19;
BB0_2:
.loc 2 7 1
shl.b32 %r21, %r30, 2;
add.s32 %r22, %r4, %r21;
ld.global.u32 %r11, [%r22];
add.s32 %r23, %r3, %r21;
ld.global.u32 %r10, [%r23];
mov.u32 %r31, 0;
BB0_3:
add.s32 %r24, %r10, %r31;
ld.u8 %rc1, [%r24];
add.s32 %r25, %r11, %r31;
st.u8 [%r25], %rc1;
add.s32 %r31, %r31, 1;
setp.lt.u32 %p2, %r31, %r7;
@%p2 bra BB0_3;
.loc 2 6 54
add.s32 %r30, %r8, %r30;
ld.param.u32 %r27, [_Z6kernelPPiS0_ii_param_3];
.loc 2 6 1
setp.lt.s32 %p3, %r30, %r27;
@%p3 bra BB0_2;
BB0_5:
.loc 2 9 2
ret;
}
BB0_3处的代码块是编译器自动生成的一个字节大小的memcpy循环。从性能角度考虑，使用它可能不是一个好主意，但它得到了充分支持（在所有架构上长期存在）。
四年后编辑以补充信息：自CUDA 6发布周期以来，由于设备端运行时API的发布，也可以直接调用类似下面的内容：
cudaMemcpyAsync(void *to, void *from, size, cudaMemcpyDeviceToDevice)
对于所有支持的架构，包括使用单独编译和设备链接的计算能力3.5及更高版本硬件，在设备代码中都可以使用。
- talonmies
2
2从性能角度来看，使用它可能不是一个好主意。你的意思是最好使用for循环来复制数组的每个位置吗？如果不是，你能告诉我在哪些可能的数组长度下，使用memcpy更有效吗？ - kon psych
太好了--我曾经担心会有函数调用开销，但nVidia的人比那聪明多了。 - wcochran
10
在我的测试中，最好的解决方案是编写自己的循环复制例程。 在我的情况下：
__device__
void devCpyCplx(const thrust::complex<float> *in, thrust::complex<float> *out, int len){
// Casting for improved loads and stores
for (int i=0; i<len/2; ++i) {
((float4*) out)[i] = ((float4*) out)[i];
}
if (len%2) {
((float2*) out)[len-1] = ((float2*) in)[len-1];
}
}
memcpy 可以在内核中使用，但速度可能会慢很多。从主机使用 cudaMemcpyAsync 是一个有效的选择。
我需要将800个连续的长度约为33,000的向量分割成不同缓冲区的长度为16,500，并进行1,600次复制调用。使用nvvp计时：
- 在内核中使用 memcpy：140毫秒
- 在主机上使用 cudaMemcpy DtoD：34毫秒
- 在内核中使用循环复制：8.6毫秒
@talonmies 称，memcpy 逐字节复制效率低下。我仍然针对 compute 3.0，因此无法在设备上测试cudaMemcpy。
编辑： 在新设备上测试。设备运行时 cudaMemcpyAsync(out, in, bytes, cudaMemcpyDeviceToDevice, 0) 与好的复制循环相当，比糟糕的复制循环更好。注意使用设备运行时 API 可能需要编译更改（sm>=3.5，单独编译）。请参阅编程指南和nvcc文档进行编译。
设备 memcpy 不好。主机 cudaMemcpyAsync 可以。设备 cudaMemcpyAsync 很好。
- plswork04
1
cudaMemcpy()确实是异步运行的，但你说得对，它不能在内核中执行。
新数组的形状是基于某种计算确定的吗？如果是这样，您通常会运行与数组中条目数量相同的线程数。每个线程将运行一个计算以确定数组中单个条目的源和目标，然后使用一个单一的赋值将其复制到那里（dst[i] = src[j]）。 如果新数组的形状不是基于计算的，则从主机运行一系列使用cudaMemCpyDeviceToDevice的cudaMemcpy()可能更有效率。
- Roger Dahl
网页内容由stack overflow 提供, 点击上面的可以查看英文原文，
原文链接
原文链接
cudaMemcpy进行设备到设备的拷贝。 - talonmies