Есть ли эквивалент memcpy(), который работает внутри ядра CUDA?

Ответ 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;
}

Блок кода в BB0_3 представляет собой цикл memcpy байтового размера, автоматически генерируемый компилятором. Возможно, с точки зрения производительности не очень хорошая идея использовать его, но он полностью поддерживается (и долгое время использовался на всех архитектурах).

Отредактировано четыре года спустя, чтобы добавить, что, поскольку API времени выполнения на стороне устройства был выпущен как часть цикла выпуска CUDA 6, также можно напрямую вызывать что-то вроде

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

в коде устройства для всех архитектур, которые его поддерживают (Compute Capability 3.5 и более новое оборудование, использующее отдельную компиляцию и связывание устройств).

Ответ 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 с хоста является допустимым вариантом.

Мне нужно было разделить 800 смежных векторов длиной ~ 33 000 и длиной 16 500 в другом буфере с 1600 вызовами копирования. Время с 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 хорошо.

Ответ 3

cudaMemcpy() действительно выполняется асинхронно, но вы правы, он не может быть выполнен изнутри ядра.

Определена ли новая форма массива на основе какого-то вычисления? Затем вы обычно запускаете то же количество потоков, что и записи в вашем массиве. Каждый поток запускает вычисление, чтобы определить источник и назначение одной записи в массиве, а затем скопировать ее там с одним присваиванием. (dst[i] = src[j]). Если новая форма массива не основана на вычислениях, может быть более эффективным запустить серию cudaMemcpy() с cudaMemCpyDeviceToDevice с хоста.