CUDAカーネル内で動作するmemcpy（）と同等の機能はありますか？

Question

私は、CUDAカーネルを使用して、配列の構造を非同期的に分解し、再構成しようとしています。 memcpy()はカーネル内では機能しません。どちらもcudaMemcpy() *;私は迷っている。

誰でもCUDAカーネル内からメモリをコピーするのに適した方法を教えてもらえますか？

cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice)は、カーネルの外部からしか呼び出せず、非同期に実行されないため、私がやろうとしているところでは機能しません。

Answer 1

はい、キューダのカーネル内で動作するmemcpyに相当します。それは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リリースサイクルの一部としてリリースされて以来、直接、デバイス・コードで

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

のようなものを呼び出すことも可能であることを追加して4年後に編集された

---

それをサポートするすべてのアーキテクチャ（Compute Capability 3.5以降のハードウェア）で使用できます。

Answer 2

cudaMemcpy()実際には非同期に実行されますが、正しいですが、カーネル内では実行できません。

何らかの計算に基づいて配列の新しい形状が決定されますか？次に、配列内のエントリと同じ数のスレッドを実行します。各スレッドは計算を実行して、配列内の単一のエントリのソースとデスティネーションを決定し、そこに1回の割り当てでコピーします。 （dst[i] = src[j]）。配列の新しい形状が計算に基づいていない場合は、cudaMemcpy()をでホストから実行する方が効率的です。

Answer 3

私のテストでは、最良の答えはあなた自身のループコピールーチンを書くことです。私の場合：

__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は有効なオプションです。

長さが〜33,000、長さが16,500、長さが16,500の異なる連続したバッファを1,600のコピーコールで異なるバッファに分割する必要がありました。nvvpとタイミング：カーネルで

• のmemcpy：ホスト上の140ミリ秒
• cudaMemcpy DtoD：カーネルで34ミリ秒
• ループコピー：8.6ミリ秒

@talonmiesはmemcpyコピーをすることによってバイトと報告しバイトは、ロードとストアでは非効率的です。私はまだ3.0を対象にしているので、デバイス上でcudaMemcpyをテストすることはできません。