Inverted logic in warp_vote_any and warp_vote_all intrinsics
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 条评论