老师好,在ubuntu22.04-dtk25.04.2,cuda版本是12.6上编译GPU程序出现:
编译命令:RUSTFLAGS="-C target-cpu=native" cargo build --release --features cuda
> cargo:warning=dcc: error: unable to execute command: Segmentation fault (core dumped)
cargo:warning=ptx error, (10, 3): parsed invalid instruction: wave mode:Wave64 size of operand b 4 must be 8 bytes
cargo:warning= 10 | shfl.sync.up.b32 v|did, arg0, arg1, 0, 0xffffffff;
cargo:warning= | ^, instruction: $( .reg.b32 $v; .reg.pred $did; shfl.sync.up.b32 $v$|$did, $0, $1, 0, 0xffffffff; @$did add.u32 $0, $0, $v;$)
根据出现的PTX指令检索到三处出现shfl.sync.up.b32的代码,注释并改成下面的__shfl_up方式:
```
//asm("{ .reg.pred %%did;"
// " shfl.sync.up.b32 %0|%%did, %1, %2, 0, 0xffffffff;"//wang 2025 0xffffffffffffffff
// " @!%%did mov.b32 %0, 0;"
// "}" : "=r"(idx) : "r"(len), "r"(degree));
idx = __shfl_up(len, degree, WARP_SZ);//
if (laneid < degree) idx = 0;//wang 2025
```
```language
53 __device__ __forceinline__
54 uint32_t sum_up(uint32_t sum, const uint32_t limit = WARP_SZ)
55 {
56 #pragma unroll
57 for (uint32_t off = 1; off < limit; off <<= 1){
58 //asm("{ .reg.b32 %v; .reg.pred %did;"
59 //" shfl.sync.up.b32 %v|%did, %0, %1, 0, 0xffffffff;"
60 // " shfl.sync.up.b32 %v|%did, %0, %1, 0, 0xffffffffffffffff;"//wang 2025
61 // " @%did add.u32 %0, %0, %v;"
62 // "}" : "+r"(sum) : "r"(off));
63 uint32_t v = __shfl_up(sum, off, limit);//
64 if (__lane_id() >= off) sum += v;// wang 2025
65 }
66
67 return sum;
68 }
```
```
174 //asm("{ .reg.b32 %%v; .reg.pred %%did;");
175 //asm("shfl.sync.up.b32 %v|%did, %0, 1, 0, 0xffffffff;" :: "r"(sum));
176 //asm("shfl.sync.up.b32 %%v|%%did, %0, 1, 0, 0xffffffffffffffff;" :: "r"(sum));//wang 2025
177 //asm("@%%did mov.b32 %0, %%v;" : "+r"(h.x));
178 //asm("}");
179 uint32_t v = __shfl_up(sum, 1, WARP_SZ);//
180 if (__lane_id() >= 1) h.x = v;// wang 2025
181
182 //h.y = __shfl_down_sync(0xffffffff, sum, gridDim.x-1) - h.x;
183 h.y = __shfl_down(sum, gridDim.x-1, WARP_SZ) - h.x;// wang 2025
184
185 if (blockIdx.x == 0 && sub_laneid == 0 && warp_off < 1< cargo:warning=11 dcc 0x0000561284522b73 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) + 867
cargo:warning=12 dcc 0x000056128213b214 cc1_main(llvm::ArrayRef, char const*, void*) + 3572
cargo:warning=13 dcc 0x0000561282137aff
cargo:warning=14 dcc 0x0000561282137f91 BuildInClang::buildInClang(int, char**, llvm::ToolContext const&) + 785
cargo:warning=15 dcc 0x00005612820ee42b dcc_main(int, char**, llvm::ToolContext const&) + 507
cargo:warning=16 dcc 0x000056128202eda5 main + 53
cargo:warning=17 libc.so.6 0x00007fe9600f3d90
cargo:warning=18 libc.so.6 0x00007fe9600f3e40 __libc_start_main + 128
cargo:warning=19 dcc 0x00005612820ee16e _start + 46
cargo:warning=dcc: error: unable to execute command: Segmentation fault (core dumped)
还出现dtk文件的报错:

通过搜索__shfl.*_sync和shfl.sync,发现还有很多sync的接口,这种dcu上是不是不兼容?需要都改成不带_sync的形式吗?

请问老师修改路线是怎么样的,谢谢。