CUDAのmemcpy（dst、src、0）が* dstに書き込む可能性はありますか？

Question

私はmemcpy(my_dst, my_src, my_num_bytes)を呼び出すいくつかのカーネルコードを持っていますが、時には0に等しいmy_num_bytesがあります。奇妙なことに、（Titan X、CUDA 7.5、ドライバ358.16の）いくつかの散発的な実験では、 。

1. この現象はCUDAで発生しましたか？
2. これはどこに指定されていますか？ memcpy()doesn't sayのプログラミングガイドエントリ。

Answer 1

これはデバイス側memcpy()の（現在の、つまりCUDA 7.5）実装のバグであるようです。このような

カーネル：

__global__ void kernel(char* source, char* dst, int len, int sz) 
{ 

    int i = threadIdx.x * len; 

    memcpy(source+i, dst+i, sz); 
} 

はこのようにPTXを放出するツールチェーンをリード：

 // .globl  _Z6kernelPcS_ii 
.visible .entry _Z6kernelPcS_ii(
     .param .u64 _Z6kernelPcS_ii_param_0, 
     .param .u64 _Z6kernelPcS_ii_param_1, 
     .param .u32 _Z6kernelPcS_ii_param_2, 
     .param .u32 _Z6kernelPcS_ii_param_3 
) 
{ 
     .reg .pred  %p<2>; 
     .reg .b16  %rs<2>; 
     .reg .b32  %r<4>; 
     .reg .b64  %rd<15>; 


     ld.param.u64 %rd7, [_Z6kernelPcS_ii_param_0]; 
     ld.param.u64 %rd8, [_Z6kernelPcS_ii_param_1]; 
     ld.param.u32 %r1, [_Z6kernelPcS_ii_param_2]; 
     cvta.to.global.u64  %rd9, %rd8; 
     cvta.to.global.u64  %rd10, %rd7; 
     mov.u32   %r2, %tid.x; 
     mul.lo.s32  %r3, %r2, %r1; 
     cvt.s64.s32  %rd11, %r3; 
     add.s64   %rd1, %rd10, %rd11; 
     add.s64   %rd2, %rd9, %rd11; 
     mov.u64   %rd14, 0; 
     ld.param.s32 %rd3, [_Z6kernelPcS_ii_param_3]; 

BB6_1: 
     add.s64   %rd12, %rd2, %rd14; 
     ld.global.u8 %rs1, [%rd12]; 
     add.s64   %rd13, %rd1, %rd14; 
     st.global.u8 [%rd13], %rs1; 
     add.s64   %rd14, %rd14, 1; 
     setp.lt.u64  %p1, %rd14, %rd3; 
     @%p1 bra  BB6_1; 

     ret; 
} 

私の読書は、このコードは常にの値があるため、少なくとも1つのバイトをコピーすることです長さ引数は、バイトコピーの後までテストされません。このようなもの：

BB6_1: 
     setp.ge.u64  %p1, %rd14, %rd3; 
     @%p1 bra  Done; 
     add.s64   %rd12, %rd2, %rd14; 
     ld.global.u8 %rs1, [%rd12]; 
     add.s64   %rd13, %rd1, %rd14; 
     st.global.u8 [%rd13], %rs1; 
     add.s64   %rd14, %rd14, 1; 
     bra    BB6_1; 
Done: 

おそらく、期待どおりに動作します。