Esiste un equivalente a memcpy() che funziona all'interno di un kernel CUDA?

Question

Sto cercando di separare e rimodellare la struttura di un array in modo asincrono utilizzando il kernel CUDA. memcpy() non funziona all'interno del kernel, e nemmeno cudaMemcpy() *; Sono in perdita.

qualcuno può dirmi il metodo preferito per la copia di memoria all'interno del kernel CUDA?

Vale la pena notare, cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) non funziona per quello che sto cercando di fare, perché può essere chiamato solo al di fuori del kernel e non esegue in modo asincrono.

Answer 1

sì, c'è un equivalente per memcpy che funziona all'interno kernel CUDA. Si chiama memcpy. Come esempio:

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

} 

che compila senza errori simili:

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

ed emette 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; 
} 

Il blocco di codice a BB0_3 è un byte dimensioni memcpy anello emessa automagicamente da il compilatore. Potrebbe non essere una buona idea dal punto di vista delle prestazioni usarlo, ma è pienamente supportato (ed è stato per molto tempo su tutte le architetture).

---

Modificato quattro anni dopo aggiungere che, poiché l'API runtime lato dispositivo è stato rilasciato come parte del ciclo di rilascio CUDA 6, è anche possibile chiamare direttamente qualcosa come

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

in codice del dispositivo per tutte le architetture che lo supportano (Compute Capability 3.5 e hardware più recente).

Answer 2

cudaMemcpy() esegue in modo asincrono ma hai ragione, non può essere eseguito da un kernel.

è la nuova forma della matrice determinato sulla base di un calcolo? Quindi, in genere si esegue lo stesso numero di thread quante sono le voci nell'array. Ogni thread eseguiva un calcolo per determinare l'origine e la destinazione di una singola voce nell'array e quindi copiarlo con un singolo compito. (dst[i] = src[j]). Se la nuova forma della matrice non si basa su calcoli, può essere più efficiente per eseguire una serie di cudaMemcpy() con cudaMemCpyDeviceToDevice dall'host.

Answer 3

Nel mio test la risposta migliore è scrivere una propria routine looping copia. Nel mio 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 opere in un kernel, ma può essere molto più lento. cudaMemcpyAsync dall'host è un'opzione valida.

avevo bisogno di partizionare 800 vettori contigui di ~ 33.000 lunghezza di 16.500 lunghezza del buffer diverso con 1.600 chiamate di copia.Timing con nvvp:

• memcpy nel kernel: 140 ms
• cudaMemcpy DtoD sull'host: 34 ms
• Loop copia in kernel: 8.6 ms

@talonmies riferisce che memcpy copie di byte da byte che è inefficiente con carichi e negozi. Il mio obiettivo è ancora il calcolo 3.0, quindi non posso testare cudaMemcpy sul dispositivo.