5
编译器
新人贴
Segmentation fault (core dumped)
2025-12-18 14:44:13
142
老师好,在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文件的报错: ![1766039929217.png](/api-forum/file/Redirect?path=forum/mgsubgg6anum/1766039929561/1766039929217.png '1766039929217.png') 通过搜索__shfl.*_sync和shfl.sync,发现还有很多sync的接口,这种dcu上是不是不兼容?需要都改成不带_sync的形式吗? ![1766040065452.png](/api-forum/file/Redirect?path=forum/mgsubgg6anum/1766040065704/1766040065452.png '1766040065452.png') 请问老师修改路线是怎么样的,谢谢。
全部评论(8)
只看楼主
王
王君浩
楼主
1楼
2025/12/18
S
suily
回复 @王君浩
麻烦您把错误汇总那一块的日志放出来,比如第一个error出现的地方、以及结尾显示
2025/12/18
王
王君浩
回复 @suily
老师您好,这是第一次出现报错的信息,带有具体的错误信息:ptx error, (10, 3): parsed invalid instruction: wave mode:Wave64 size of operand b 4 must be 8 bytes
2025/12/18
S
suily
回复 @王君浩
您使用fastpt了吗?可以使用fastpt+修改报错处的洗牌函数(去掉_sync后缀)
2025/12/19
全部7条回复
没有更多内容了
--
楼主
Level.3
0 帖子
楼主热帖
暂无数据