请考虑以下PTX代码:
//
// Generated by NVIDIA NVVM Compiler... sort of
//
// Compiler Build ID: CL-25769353
// Cuda compilation tools, release 10.1, V10.1.105
// Based on LLVM 3.4svn
//
.version 6.4
.target sm_30
.address_size 64
.func (.param .b32 func_retval0) foo(
.param .b32 foo_param_0,
.param .b32 foo_param_1,
.param .b32 foo_param_2
)
{
.reg .b16 %rs<3>;
.reg .b32 %r<3>;
ld.param.u16 %rs1, [foo_param_0];
ld.param.u16 %rs2, [foo_param_1];
ld.param.u32 %r2, [foo_param_2];
// inline asm
sad.s16 %r1, %rs1, %rs2, %r2;
// inline asm
st.param.b32 [func_retval0+0], %r1;
ret;
}
当尝试使用 ptxas (CUDA 10.1( 编译它时,我得到:
ptxas /tmp/a.ptx, line 27; error : Arguments mismatch for instruction 'sad'
ptxas fatal : Ptx assembly aborted due to errors
为什么?这种类型组合有什么问题?
PTX 参考资料说:
sad.type d, a, b, c; .type = { .u16, .u32, .u64, .s16, .s32, .s64 };
,似乎d
和c
总是u32
,type
适用于a
和b
。无论如何,这就是device_functions.h
中的__sad()
功能。
实际上,d
和a
的类型必须是"b
和c
类型的无符号版本"。所以,这应该有效:
.func (.param .b32 func_retval0) foo(
.param .b32 foo_param_0,
.param .b32 foo_param_1,
.param .b32 foo_param_2
)
{
.reg .b32 %r<5>;
ld.param.u32 %r4, [foo_param_2];
ld.param.s16 %r2, [foo_param_0];
ld.param.s16 %r3, [foo_param_1];
// inline asm
sad.s32 %r1, %r2, %r3, %r4;
// inline asm
st.param.b32 [func_retval0+0], %r1;
ret;
}