[pull] main from Rust-GPU:main

crates/cuda_std/src/atomic/intrinsics.rs

@@ -52,11 +52,11 @@ pub unsafe fn fence_acqrel_system() {
 
 #[allow(unused_macros)]
 macro_rules! load_scope {
-    (volatile, $scope:ident) => {
+    (volatile, $scope_asm:ident) => {
         ""
     };
-    ($ordering:ident, $scope:ident) => {
-        concat!(".", stringify!($scope))
+    ($ordering:ident, $scope_asm:ident) => {
+        concat!(".", stringify!($scope_asm))
     };
 }
 
@@ -70,7 +70,7 @@ macro_rules! load {
                 pub unsafe fn [<atomic_load_ $ordering _ $width _ $scope>](ptr: *const [<u $width>]) -> [<u $width>] {
                     let mut out;
                     asm!(
-                        concat!("ld.", stringify!($ordering), load_scope!($ordering, $scope), ".", stringify!([<u $width>]), " {}, [{}];"),
+                        concat!("ld.", stringify!($ordering), load_scope!($ordering, $scope_asm), ".", stringify!([<u $width>]), " {}, [{}];"),
                         out([<reg $width>]) out,
                         in(reg64) ptr
                     );
@@ -116,7 +116,7 @@ macro_rules! store {
                 #[doc = concat!("Performs a ", stringify!($ordering), " atomic store at the ", stringify!($scope), " level with a width of ", stringify!($width), " bits")]
                 pub unsafe fn [<atomic_store_ $ordering _ $width _ $scope>](ptr: *mut [<u $width>], val: [<u $width>]) {
                     asm!(
-                        concat!("st.", stringify!($ordering), load_scope!($ordering, $scope), ".", stringify!([<u $width>]), " [{}], {};"),
+                        concat!("st.", stringify!($ordering), load_scope!($ordering, $scope_asm), ".", stringify!([<u $width>]), " [{}], {};"),
                         in(reg64) ptr,
                         in([<reg $width>]) val,
                     );

crates/cuda_std/src/warp.rs

@@ -351,7 +351,7 @@ pub unsafe fn warp_vote_all(mask: u32, predicate: bool) -> bool {
             ".reg .pred %p<3>;",
             "setp.eq.u32 %p1, {}, 1;",
             "vote.sync.all.pred %p2, %p1, {};",
-            "selp.u32 {}, 0, 1, %p2;",
+            "selp.u32 {}, 1, 0, %p2;",
             "}}",
             in(reg32) predicate as u32,
             in(reg32) mask,
@@ -383,7 +383,7 @@ pub unsafe fn warp_vote_any(mask: u32, predicate: bool) -> bool {
             ".reg .pred %p<3>;",
             "setp.eq.u32 %p1, {}, 1;",
             "vote.sync.any.pred %p2, %p1, {};",
-            "selp.u32 {}, 0, 1, %p2;",
+            "selp.u32 {}, 1, 0, %p2;",
             "}}",
             in(reg32) predicate as u32,
             in(reg32) mask,
@@ -735,10 +735,10 @@ impl WarpShuffleValue for bf16 {
 #[repr(u32)]
 #[derive(Clone, Copy)]
 pub enum WarpShuffleMode {
-    Up,
-    Down,
-    Idx,
-    Xor,
+    Idx = 0,
+    Up = 1,
+    Down = 2,
+    Xor = 3,
 }
 
 // C-compatible struct to match LLVM IR's {i32, i8} return type

crates/rustc_codegen_nvvm/src/abi.rs

@@ -59,8 +59,32 @@ pub(crate) fn readjust_fn_abi<'tcx>(
 
         // pass all adts directly as values, ptx wants them to be passed all by value, but rustc's
         // ptx-kernel abi seems to be wrong, and it's unstable.
-        if arg.layout.ty.is_adt() && !matches!(arg.mode, PassMode::Direct { .. }) {
-            arg.mode = PassMode::Direct(ArgAttributes::new());
+        if arg.layout.ty.is_adt() {
+            let align = arg.layout.align.abi.bytes();
+            if align >= 16 && align.is_power_of_two() {
+                let unit = Reg {
+                    kind: RegKind::Integer,
+                    size: Size::from_bytes(16),
+                };
+                let cast = CastTarget {
+                    prefix: Default::default(),
+                    rest: rustc_target::callconv::Uniform {
+                        unit,
+                        total: arg.layout.size,
+                        is_consecutive: false,
+                    },
+                    rest_offset: Some(Size::ZERO),
+                    attrs: ArgAttributes::new(),
+                };
+                arg.mode = PassMode::Cast {
+                    cast: Box::new(cast),
+                    pad_i32: false,
+                };
+            } else if !matches!(arg.mode, PassMode::Direct { .. }) {
+                let mut attrs = ArgAttributes::new();
+                attrs.pointee_align = Some(arg.layout.align.abi);
+                arg.mode = PassMode::Direct(attrs);
+            }
         }
         arg
     };