如何解决如何通过在一个方法中调用一个函数两次来编译 cuda 代码?
我正在尝试编译这段代码:
struct foo {
unsigned long long x0;
};
//__device__ __noinline__ foo bar(foo a,foo b){ // << try this
__device__ foo bar(foo a,foo b){
foo r;
asm(
".reg .u64 my_cool_var;\n\t"
"add.cc.u64 %0,%1,%2;\n\t"
: "=l"(r.x0)
: "l"(a.x0)
"l"(b.x0)
);
return r;
}
__device__ foo func_with2call(foo x,foo y){
foo res = bar(x,y);
foo iy = bar(x,res);
return iy;
}
__global__ void cuda_test1(foo *a,foo *b,foo *r) {
*r = func_with2call(*a,*b);
}
编译器输出:
ptxas /tmp/tmpxft_000010f5_00000000-6_main.ptx,line 38; error : Duplicate deFinition of variable 'my_cool_var'
ptxas fatal : Ptx assembly aborted due to errors
我知道,这是代码内联的问题。例如,如果我用 bar
属性编译 __noinline__
函数,则没有错误。有没有办法保持内联(除了用不同的内部变量名复制 bar
函数代码),但仍然调用 bar
函数两次?
解决方法
此限制在 Inline PTX Assembly Guide 中讨论。您可以通过强制每个定义进入其自己的范围来解决它,例如:
__device__ foo bar(foo a,foo b){
foo r;
asm(
"{.reg .u64 my_cool_var;\n\t"
"add.cc.u64 %0,%1,%2;\n\t"
"}"
: "=l"(r.x0)
: "l"(a.x0)
"l"(b.x0)
);
return r;
}
这将安全地内联而不会发生冲突。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。