ITADN

Inverted logic in warp_vote_any and warp_vote_all intrinsics

#356OpenSnehal-Reddy 创建于 2026-02-09
S
Snehal-Reddycommented
The warp_vote_any and warp_vote_all functions in cuda_std::warp return inverted boolean results. They currently return false when the vote succeeds and true when it fails. ### Details Looking at the implementation in crates/cuda_std/src/warp.rs: ``` pub unsafe fn warp_vote_all(mask: u32, predicate: bool) -> bool { let mut out: u32; unsafe { asm!( "{{", ".reg .pred %p<3>;", "setp.eq.u32 %p1, {}, 1;", "vote.sync.all.pred %p2, %p1, {};", "selp.u32 {}, 0, 1, %p2;", // <--- Bug here "}}", in(reg32) predicate as u32, in(reg32) mask, out(reg32) out ); } out != 0 } ``` 1. The `vote.sync.all.pred %p2, ..`. instruction sets the predicate register `%p2` to true if all threads agree. 2. The` selp.u32 {}, 0, 1, %p2` instruction selects the first operand (0) if `%p2` is true, and the second (1) if it is false. 3. The function returns out != 0. 4. Result: If the vote succeeds, it returns 0 != 0 which is false. The functions should return true if the vote condition is met. The `selp` instruction should be `selp.u32 {}, 1, 0, %p2`; so that success results in 1 (true). **This is also present similarly in `warp_vote_any`** For additional context, the logic in `crates/cuda_std/src/ptr.rs` (inside is_in_address_space) correctly uses the 1, 0 order for boolean conversion:
0 条评论