是否有一个在CUDA内核中工作的memcpy()的等价物？

Question

我试图使用CUDA内核异步拆分并重新整形数组的结构.memcpy()在内核中不起作用,也不起作用cudaMemcpy()*; 我不知所措.

谁能告诉我从CUDA内核中复制内存的首选方法？

值得注意的是,它cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice)不能用于我想要做的事情,因为它只能从内核外部调用而不能异步执行.

Answer 1

是的,有一个相当于memcpy在cuda内核中工作.它被称为 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;
}

代码块at BB0_3是memcpy由编译器自动发出的字节大小的循环.从性能的角度来看,使用它可能不是一个好主意,但它完全受支持(并且已经在所有架构上使用了很长时间).

四年后编辑添加,因为设备端运行时API是作为CUDA 6发布周期的一部分发布的,所以也可以直接调用类似

cudaMemcpyAsync(void *to, void *from, size, cudaMemcpyDeviceToDevice)

在所有支持它的架构的设备代码中(Compute Capability 3.5和更新的硬件).

Answer 2

在我的测试中，最好的答案是编写自己的循环复制例程。就我而言：

__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来自主机的选项是有效的。

我需要在1600个复制调用的不同缓冲区中，将800个33000长度的连续向量划分为16500长度。用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不错。