CUDAカーネルを使用して、非同期で配列の構造を分解して再形成しようとしています。memcpy()カーネル内では機能せず、cudaMemcpy()*;も機能しません。私は途方に暮れています。
誰かがCUDAカーネル内からメモリをコピーするための好ましい方法を教えてもらえますか？
cudaMemcpy(void *to, void *from, size, cudaMemcpyDeviceToDevice)カーネルの外部からのみ呼び出すことができ、非同期で実行されないため、私がやろうとしていることには機能しないことに注意してください。
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パフォーマンスの観点からはそれを使用するのは良い考えではないかもしれませんが、完全にサポートされています（そして、すべてのアーキテクチャで長い間使用されてきました）。
4年後に編集され、デバイス側のランタイムAPIがCUDA 6リリースサイクルの一部としてリリースされたため、次のようなものを直接呼び出すことも可能です。
cudaMemcpyAsync(void *to, void *from, size, cudaMemcpyDeviceToDevice)
それをサポートするすべてのアーキテクチャのデバイスコード（個別のコンパイルとデバイスリンクを使用したCompute Capability 3.5以降のハードウェア）。
私のテストでは、最良の答えは、独自のループコピールーチンを作成することです。私の場合：
__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ホストからは有効なオプションです。
1,600のコピー呼び出しを使用して、長さが約33,000から16,500の長さの連続する800個のベクトルを異なるバッファーに分割する必要がありました。nvvpを使用したタイミング：
@talonmiesは、memcpyバイトごとにコピーすることを報告しますが、これはロードとストアでは非効率的です。まだコンピューティング3.0をターゲットにしているため、デバイスでcudaMemcpyをテストできません。
編集：新しいデバイスでテスト済み。デバイスのランタイムcudaMemcpyAsync(out, in, bytes, cudaMemcpyDeviceToDevice, 0)は、良いコピーループに匹敵し、悪いコピーループよりも優れています。デバイスランタイムAPIを使用すると、コンパイルの変更が必要になる場合があることに注意してください（sm> = 3.5、個別のコンパイル）。コンパイルについては、プログラミングガイドとnvccドキュメントを参照してください。
デバイスmemcpyが不良です。ホストcudaMemcpyAsyncは大丈夫です。デバイスはcudaMemcpyAsync良好です。
cudaMemcpy()確かに非同期で実行されますが、その通り、カーネル内から実行することはできません。
アレイの新しい形状は、何らかの計算に基づいて決定されていますか？次に、通常、配列にエントリがあるのと同じ数のスレッドを実行します。各スレッドは計算を実行して、配列内の単一のエントリのソースと宛先を決定し、単一の割り当てでそこにコピーします。（dst[i] = src[j]）。配列の新しい形状が計算に基づいていない場合は、ホストから一連のcudaMemcpy()withを実行する方が効率的である可能性があります。cudaMemCpyDeviceToDevice