Возможно ли, что memcpy(dst, src, 0) в CUDA пишет в *dst?

Question

У меня есть код ядра, который вызывает memcpy(my_dst, my_src, my_num_bytes) - а иногда у меня есть my_num_bytes равно 0. Как ни странно, некоторые спорадические эксперименты (с Titan X, CUDA 7.5, драйвер 358.16) предполагают, что данные записываются в пункт назначения, когда я выполняю такой вызов.

1. Сталкивались ли вы с таким поведением в CUDA?
2. Это указано где-нибудь? Запись руководства по программированию для memcpy() не говорит

Answer 1

Это может показаться ошибкой в ​​(текущей, то есть CUDA 7.5) реализации на стороне устройства. memcpy(),

Ядро вроде этого:

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

Мое чтение заключается в том, что этот код всегда будет копировать по крайней мере один байт, потому что значение аргумента длины не проверяется до окончания байтовой копии. Что-то вроде этого:

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:

вероятно, будет работать как ожидалось.