CUDA 的 memcpy(dst, src, 0) 是否有可能写入 *dst?
Is it possible that CUDA's memcpy(dst, src, 0) writes to *dst?
我有一些调用 memcpy(my_dst, my_src, my_num_bytes)
的内核代码 - 有时我 my_num_bytes
等于 0。奇怪的是,一些零星的实验(使用 Titan X,CUDA 7.5,驱动程序 358.16)建议数据当我使用这样的调用时确实会写入目的地。
- 您在 CUDA 中遇到过这种行为吗?
- 这是在任何地方指定的吗?
memcpy()
doesn't say. 的编程指南条目
这似乎是设备端实施(当前,即 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:
可能会按预期工作。
我有一些调用 memcpy(my_dst, my_src, my_num_bytes)
的内核代码 - 有时我 my_num_bytes
等于 0。奇怪的是,一些零星的实验(使用 Titan X,CUDA 7.5,驱动程序 358.16)建议数据当我使用这样的调用时确实会写入目的地。
- 您在 CUDA 中遇到过这种行为吗?
- 这是在任何地方指定的吗?
memcpy()
doesn't say. 的编程指南条目
这似乎是设备端实施(当前,即 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:
可能会按预期工作。