Skip to content

Commit 3788dc7

Browse files
committed
fix(cuda_std): Fix inverted logic in warp_vote_any and warp_vote_all intrinsics
1 parent f76a232 commit 3788dc7

File tree

1 file changed

+2
-2
lines changed

1 file changed

+2
-2
lines changed

crates/cuda_std/src/warp.rs

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -351,7 +351,7 @@ pub unsafe fn warp_vote_all(mask: u32, predicate: bool) -> bool {
351351
".reg .pred %p<3>;",
352352
"setp.eq.u32 %p1, {}, 1;",
353353
"vote.sync.all.pred %p2, %p1, {};",
354-
"selp.u32 {}, 0, 1, %p2;",
354+
"selp.u32 {}, 1, 0, %p2;",
355355
"}}",
356356
in(reg32) predicate as u32,
357357
in(reg32) mask,
@@ -383,7 +383,7 @@ pub unsafe fn warp_vote_any(mask: u32, predicate: bool) -> bool {
383383
".reg .pred %p<3>;",
384384
"setp.eq.u32 %p1, {}, 1;",
385385
"vote.sync.any.pred %p2, %p1, {};",
386-
"selp.u32 {}, 0, 1, %p2;",
386+
"selp.u32 {}, 1, 0, %p2;",
387387
"}}",
388388
in(reg32) predicate as u32,
389389
in(reg32) mask,

0 commit comments

Comments
 (0)