¿Hay un equivalente a memcpy() que funcione dentro de un kernel CUDA?

Question

Estoy tratando de separarme y remodelar la estructura de una matriz de forma asincrónica utilizando el kernel CUDA. memcpy() no funciona dentro del kernel, y tampoco lo hace cudaMemcpy() *; Estoy perdido

¿Alguien puede decirme el método preferido para copiar memoria desde el kernel CUDA?

Vale la pena señalar, cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) no funcionará para lo que estoy tratando de hacer, ya que sólo se puede llamar desde fuera del núcleo y no se ejecuta de forma asíncrona.

Answer 1

Sí, hay un equivalente a memcpy que funciona dentro de los núcleos de cuda. Se llama memcpy. Como un ejemplo:

__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); 

} 

que compila sin error como este:

$ 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] 

y emite 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; 
} 

El bloque de código en BB0_3 es un byte de tamaño memcpy bucle emitida automágicamente por el compilador Puede que no sea una gran idea desde el punto de vista del rendimiento usarlo, pero es totalmente compatible (y lo ha sido durante mucho tiempo en todas las arquitecturas).

---

Editado cuatro años más tarde para agregar que desde la API de tiempo de ejecución lado del dispositivo fue lanzado como parte del ciclo de lanzamiento de CUDA 6, también es posible llamar directamente a algo así como

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

en código del dispositivo para todas las arquitecturas que lo soportan (Compute Capability 3.5 y hardware más nuevo).

Answer 2

cudaMemcpy() efectivamente se ejecuta de forma asíncrona, pero tiene razón, no se puede ejecutar desde un kernel.

¿La nueva forma de la matriz está determinada por algún cálculo? Entonces, normalmente ejecutará la misma cantidad de subprocesos que entradas en su matriz. Cada subproceso ejecutará un cálculo para determinar la fuente y el destino de una sola entrada en la matriz y luego copiarlo allí con una sola asignación. (dst[i] = src[j]). Si la nueva forma de la matriz no se basa en cálculos, podría ser más eficiente ejecutar una serie de cudaMemcpy() con cudaMemCpyDeviceToDevice desde el host.

Answer 3

En mis pruebas, la mejor respuesta es escribir su propia rutina de copia en bucle. En mi caso:

__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 obras en un núcleo, pero puede ser mucho más lento. cudaMemcpyAsync desde el host es una opción válida.

Necesitaba particionar 800 vectores contiguos de ~ 33,000 de longitud a 16,500 de longitud en un buffer diferente con 1,600 llamadas de copia.Timing con nvvp:

• memcpy en kernel: 140 ms
• cudaMemcpy DtoD en el host: 34 ms de copia
• bucle en kernel: 8.6 ms

@talonmies informa que memcpy copias byte por byte que es ineficiente con cargas y tiendas. Me estoy enfocando en calcular 3.0 aún así no puedo probar cudaMemcpy en el dispositivo.