Có tương đương với memcpy() hoạt động bên trong hạt nhân CUDA không?

Question

Tôi đang cố gắng tách rời và định hình lại cấu trúc của một mảng không đồng bộ bằng hạt nhân CUDA. memcpy() không hoạt động bên trong hạt nhân và không hoạt động cudaMemcpy() *; Tôi đang thua lỗ.

Bất cứ ai có thể cho tôi biết phương pháp ưu tiên để sao chép bộ nhớ từ bên trong hạt nhân CUDA?

Điều đáng chú ý là cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice) sẽ KHÔNG hoạt động cho những gì tôi đang cố gắng làm, vì nó chỉ có thể được gọi từ bên ngoài hạt nhân và không thực thi không đồng bộ.

Answer 1

Có, tương đương với memcpy hoạt động bên trong hạt nhân cuda. Nó được gọi là memcpy. Như một ví dụ:

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

} 

mà biên dịch mà không có lỗi như thế này:

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

và phát ra 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; 
} 

Khối mã tại BB0_3 là một byte có kích thước memcpy vòng lặp phát ra automagically bởi trình biên dịch. Nó có thể không phải là một ý tưởng tuyệt vời từ một quan điểm hiệu suất để sử dụng nó, nhưng nó được hỗ trợ đầy đủ (và đã được một thời gian dài trên tất cả các kiến ​​trúc).

---

Sửa bốn năm sau để thêm rằng kể từ khi API runtime bên thiết bị đã được phát hành như là một phần của chu kỳ phát hành CUDA 6, nó cũng có thể trực tiếp gọi một cái gì đó giống như

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

trong mã thiết bị cho tất cả các kiến ​​trúc hỗ trợ nó (Khả năng tính toán 3.5 và phần cứng mới hơn).

Answer 2

cudaMemcpy() thực sự chạy không đồng bộ nhưng bạn nói đúng, nó không thể được thực hiện từ bên trong hạt nhân.

Hình dạng mới của mảng được xác định dựa trên một số phép tính? Sau đó, bạn thường sẽ chạy cùng một số luồng vì có các mục trong mảng của bạn. Mỗi luồng sẽ chạy một phép tính để xác định nguồn và đích của một mục duy nhất trong mảng và sau đó sao chép nó vào đó với một phép gán đơn. (dst[i] = src[j]). Nếu hình dạng mới của mảng không dựa trên các phép tính, thì có thể hiệu quả hơn khi chạy một loạt các cudaMemcpy() với cudaMemCpyDeviceToDevice từ máy chủ.

Answer 3

Trong thử nghiệm của tôi, câu trả lời hay nhất là viết thói quen sao chép lặp của riêng bạn. Trong trường hợp của tôi:

__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 công trình trong một hạt nhân nhưng nó có thể chậm hơn nhiều. cudaMemcpyAsync từ máy chủ lưu trữ là một tùy chọn hợp lệ.

Tôi cần phân vùng 800 vectơ liền kề có chiều dài ~ 33.000 đến 16.500 độ dài trong bộ đệm khác nhau với 1.600 cuộc gọi sao chép.Thời gian với NVVP:

• memcpy trong kernel: 140 ms
• cudaMemcpy DtoD trên máy chủ: 34 ms
• loop bản sao trong kernel: 8.6 ms

@talonmies báo cáo rằng memcpy bản byte bởi byte không hiệu quả với tải và lưu trữ. Tôi đang nhắm mục tiêu tính 3.0 vẫn còn vì vậy tôi không thể kiểm tra cudaMemcpy trên thiết bị.