From 84f834b87f1f0b2eb56e94652eb3a4b440402a93 Mon Sep 17 00:00:00 2001
From: Ralf Jung <post@ralfj.de>
Date: Sat, 19 Jul 2025 23:23:40 +0200
Subject: [PATCH 1/8] atomicrmw on pointers: move integer-pointer cast hacks
 into backend

---
 src/builder.rs | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/src/builder.rs b/src/builder.rs
index a4ec4bf8dea..032f5d0a622 100644
--- a/src/builder.rs
+++ b/src/builder.rs
@@ -1656,6 +1656,7 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
         dst: RValue<'gcc>,
         src: RValue<'gcc>,
         order: AtomicOrdering,
+        ret_ptr: bool,
     ) -> RValue<'gcc> {
         let size = get_maybe_pointer_size(src);
         let name = match op {
@@ -1683,6 +1684,9 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
         let atomic_function = self.context.get_builtin_function(name);
         let order = self.context.new_rvalue_from_int(self.i32_type, order.to_gcc());
 
+        // FIXME: If `ret_ptr` is true and `src` is an integer, we should really tell GCC
+        // that this is a pointer operation that needs to preserve provenance -- but like LLVM,
+        // GCC does not currently seems to support that.
         let void_ptr_type = self.context.new_type::<*mut ()>();
         let volatile_void_ptr_type = void_ptr_type.make_volatile();
         let dst = self.context.new_cast(self.location, dst, volatile_void_ptr_type);
@@ -1690,7 +1694,8 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
         let new_src_type = atomic_function.get_param(1).to_rvalue().get_type();
         let src = self.context.new_bitcast(self.location, src, new_src_type);
         let res = self.context.new_call(self.location, atomic_function, &[dst, src, order]);
-        self.context.new_cast(self.location, res, src.get_type())
+        let res_type = if ret_ptr { void_ptr_type } else { src.get_type() };
+        self.context.new_cast(self.location, res, res_type)
     }
 
     fn atomic_fence(&mut self, order: AtomicOrdering, scope: SynchronizationScope) {

From c4b91bd9f87391c1dda3784205c0dd10b54ec2e8 Mon Sep 17 00:00:00 2001
From: Guillaume Gomez <guillaume1.gomez@gmail.com>
Date: Mon, 4 Aug 2025 10:49:43 +0200
Subject: [PATCH 2/8] Merge commit '482e8540a1b757ed7bccc2041c5400f051fdb01e'
 into subtree-update_cg_gcc_2025-08-04

---
 build_system/src/abi_test.rs                  |  2 +-
 build_system/src/fuzz.rs                      |  8 +-
 ...1-Add-stdarch-Cargo.toml-for-testing.patch | 23 +++---
 rust-toolchain                                |  2 +-
 src/allocator.rs                              | 14 ++--
 src/builder.rs                                | 26 ++++--
 src/int.rs                                    | 76 ++++++++++++++---
 src/intrinsic/archs.rs                        | 60 ++++++++++++++
 src/intrinsic/mod.rs                          |  9 ++-
 src/intrinsic/simd.rs                         | 81 +++++++++++++++++++
 src/lib.rs                                    |  5 +-
 tests/failing-lto-tests.txt                   |  2 +-
 tests/failing-run-make-tests.txt              |  1 -
 tests/failing-ui-tests.txt                    |  8 +-
 tools/cspell_dicts/rustc_codegen_gcc.txt      |  3 +
 15 files changed, 275 insertions(+), 45 deletions(-)

diff --git a/build_system/src/abi_test.rs b/build_system/src/abi_test.rs
index 3c1531be27a..a85886d87f3 100644
--- a/build_system/src/abi_test.rs
+++ b/build_system/src/abi_test.rs
@@ -31,7 +31,7 @@ pub fn run() -> Result<(), String> {
         Some("clones/abi-cafe".as_ref()),
         true,
     )
-    .map_err(|err| (format!("Git clone failed with message: {err:?}!")))?;
+    .map_err(|err| format!("Git clone failed with message: {err:?}!"))?;
     // Configure abi-cafe to use the exact same rustc version we use - this is crucial.
     // Otherwise, the concept of ABI compatibility becomes meanignless.
     std::fs::copy("rust-toolchain", "clones/abi-cafe/rust-toolchain")
diff --git a/build_system/src/fuzz.rs b/build_system/src/fuzz.rs
index 453211366b3..9714ce29af9 100644
--- a/build_system/src/fuzz.rs
+++ b/build_system/src/fuzz.rs
@@ -43,18 +43,18 @@ pub fn run() -> Result<(), String> {
             "--start" => {
                 start =
                     str::parse(&args.next().ok_or_else(|| "Fuzz start not provided!".to_string())?)
-                        .map_err(|err| (format!("Fuzz start not a number {err:?}!")))?;
+                        .map_err(|err| format!("Fuzz start not a number {err:?}!"))?;
             }
             "--count" => {
                 count =
                     str::parse(&args.next().ok_or_else(|| "Fuzz count not provided!".to_string())?)
-                        .map_err(|err| (format!("Fuzz count not a number {err:?}!")))?;
+                        .map_err(|err| format!("Fuzz count not a number {err:?}!"))?;
             }
             "-j" | "--jobs" => {
                 threads = str::parse(
                     &args.next().ok_or_else(|| "Fuzz thread count not provided!".to_string())?,
                 )
-                .map_err(|err| (format!("Fuzz thread count not a number {err:?}!")))?;
+                .map_err(|err| format!("Fuzz thread count not a number {err:?}!"))?;
             }
             _ => return Err(format!("Unknown option {arg}")),
         }
@@ -66,7 +66,7 @@ pub fn run() -> Result<(), String> {
         Some("clones/rustlantis".as_ref()),
         true,
     )
-    .map_err(|err| (format!("Git clone failed with message: {err:?}!")))?;
+    .map_err(|err| format!("Git clone failed with message: {err:?}!"))?;
 
     // Ensure that we are on the newest rustlantis commit.
     let cmd: &[&dyn AsRef<OsStr>] = &[&"git", &"pull", &"origin"];
diff --git a/patches/0001-Add-stdarch-Cargo.toml-for-testing.patch b/patches/0001-Add-stdarch-Cargo.toml-for-testing.patch
index 9cc377850b9..3a8c37a8b8d 100644
--- a/patches/0001-Add-stdarch-Cargo.toml-for-testing.patch
+++ b/patches/0001-Add-stdarch-Cargo.toml-for-testing.patch
@@ -1,29 +1,28 @@
-From b8f3eed3053c9333b5dfbeaeb2a6a65a4b3156df Mon Sep 17 00:00:00 2001
-From: Antoni Boucher <bouanto@zoho.com>
-Date: Tue, 29 Aug 2023 13:06:34 -0400
+From 190e26c9274b3c93a9ee3516b395590e6bd9213b Mon Sep 17 00:00:00 2001
+From: None <none@example.com>
+Date: Sun, 3 Aug 2025 19:54:56 -0400
 Subject: [PATCH] Patch 0001-Add-stdarch-Cargo.toml-for-testing.patch
 
 ---
- library/stdarch/Cargo.toml | 23 +++++++++++++++++++++++
- 1 file changed, 23 insertions(+)
+ library/stdarch/Cargo.toml | 20 ++++++++++++++++++++
+ 1 file changed, 20 insertions(+)
  create mode 100644 library/stdarch/Cargo.toml
 
 diff --git a/library/stdarch/Cargo.toml b/library/stdarch/Cargo.toml
 new file mode 100644
-index 0000000..4c63700
+index 0000000..bd6725c
 --- /dev/null
 +++ b/library/stdarch/Cargo.toml
-@@ -0,0 +1,21 @@
+@@ -0,0 +1,20 @@
 +[workspace]
 +resolver = "1"
 +members = [
-+  "crates/core_arch",
-+  "crates/std_detect",
-+  "crates/stdarch-gen-arm",
++  "crates/*",
 +  #"examples/"
 +]
 +exclude = [
-+  "crates/wasm-assert-instr-tests"
++  "crates/wasm-assert-instr-tests",
++  "rust_programs",
 +]
 +
 +[profile.release]
@@ -36,5 +35,5 @@ index 0000000..4c63700
 +opt-level = 3
 +incremental = true
 -- 
-2.42.0
+2.50.1
 
diff --git a/rust-toolchain b/rust-toolchain
index 2fe8ec4647f..058e734be5c 100644
--- a/rust-toolchain
+++ b/rust-toolchain
@@ -1,3 +1,3 @@
 [toolchain]
-channel = "nightly-2025-07-04"
+channel = "nightly-2025-08-03"
 components = ["rust-src", "rustc-dev", "llvm-tools-preview"]
diff --git a/src/allocator.rs b/src/allocator.rs
index 0d8dc93274f..2a95a7368aa 100644
--- a/src/allocator.rs
+++ b/src/allocator.rs
@@ -99,11 +99,15 @@ fn create_const_value_function(
     let func = context.new_function(None, FunctionType::Exported, output, &[], name, false);
 
     #[cfg(feature = "master")]
-    func.add_attribute(FnAttribute::Visibility(symbol_visibility_to_gcc(
-        tcx.sess.default_visibility(),
-    )));
-
-    func.add_attribute(FnAttribute::AlwaysInline);
+    {
+        func.add_attribute(FnAttribute::Visibility(symbol_visibility_to_gcc(
+            tcx.sess.default_visibility(),
+        )));
+
+        // FIXME(antoyo): cg_llvm sets AlwaysInline, but AlwaysInline is different in GCC and using
+        // it here will causes linking errors when using LTO.
+        func.add_attribute(FnAttribute::Inline);
+    }
 
     if tcx.sess.must_emit_unwind_tables() {
         // TODO(antoyo): emit unwind tables.
diff --git a/src/builder.rs b/src/builder.rs
index 4aee211e2ef..34ade3d025f 100644
--- a/src/builder.rs
+++ b/src/builder.rs
@@ -540,9 +540,15 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
 
     fn ret(&mut self, mut value: RValue<'gcc>) {
         let expected_return_type = self.current_func().get_return_type();
-        if !expected_return_type.is_compatible_with(value.get_type()) {
-            // NOTE: due to opaque pointers now being used, we need to cast here.
-            value = self.context.new_cast(self.location, value, expected_return_type);
+        let value_type = value.get_type();
+        if !expected_return_type.is_compatible_with(value_type) {
+            // NOTE: due to opaque pointers now being used, we need to (bit)cast here.
+            if self.is_native_int_type(value_type) && self.is_native_int_type(expected_return_type)
+            {
+                value = self.context.new_cast(self.location, value, expected_return_type);
+            } else {
+                value = self.context.new_bitcast(self.location, value, expected_return_type);
+            }
         }
         self.llbb().end_with_return(self.location, value);
     }
@@ -1279,11 +1285,19 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
 
     fn intcast(
         &mut self,
-        value: RValue<'gcc>,
+        mut value: RValue<'gcc>,
         dest_typ: Type<'gcc>,
-        _is_signed: bool,
+        is_signed: bool,
     ) -> RValue<'gcc> {
-        // NOTE: is_signed is for value, not dest_typ.
+        let value_type = value.get_type();
+        if is_signed && !value_type.is_signed(self.cx) {
+            let signed_type = value_type.to_signed(self.cx);
+            value = self.gcc_int_cast(value, signed_type);
+        } else if !is_signed && value_type.is_signed(self.cx) {
+            let unsigned_type = value_type.to_unsigned(self.cx);
+            value = self.gcc_int_cast(value, unsigned_type);
+        }
+
         self.gcc_int_cast(value, dest_typ)
     }
 
diff --git a/src/int.rs b/src/int.rs
index 6f21ce9352b..9fb7f6bad68 100644
--- a/src/int.rs
+++ b/src/int.rs
@@ -4,12 +4,15 @@
 
 // cSpell:words cmpti divti modti mulodi muloti udivti umodti
 
-use gccjit::{BinaryOp, ComparisonOp, FunctionType, Location, RValue, ToRValue, Type, UnaryOp};
+use gccjit::{
+    BinaryOp, CType, ComparisonOp, FunctionType, Location, RValue, ToRValue, Type, UnaryOp,
+};
 use rustc_abi::{CanonAbi, Endian, ExternAbi};
 use rustc_codegen_ssa::common::{IntPredicate, TypeKind};
 use rustc_codegen_ssa::traits::{BackendTypes, BaseTypeCodegenMethods, BuilderMethods, OverflowOp};
 use rustc_middle::ty::{self, Ty};
 use rustc_target::callconv::{ArgAbi, ArgAttributes, FnAbi, PassMode};
+use rustc_type_ir::{Interner, TyKind};
 
 use crate::builder::{Builder, ToGccComp};
 use crate::common::{SignType, TypeReflection};
@@ -167,9 +170,9 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
                 if a_type.is_vector() {
                     // Vector types need to be bitcast.
                     // TODO(antoyo): perhaps use __builtin_convertvector for vector casting.
-                    b = self.context.new_bitcast(self.location, b, a.get_type());
+                    b = self.context.new_bitcast(self.location, b, a_type);
                 } else {
-                    b = self.context.new_cast(self.location, b, a.get_type());
+                    b = self.context.new_cast(self.location, b, a_type);
                 }
             }
             self.context.new_binary_op(self.location, operation, a_type, a, b)
@@ -216,13 +219,22 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
         operation_name: &str,
         signed: bool,
         a: RValue<'gcc>,
-        b: RValue<'gcc>,
+        mut b: RValue<'gcc>,
     ) -> RValue<'gcc> {
         let a_type = a.get_type();
         let b_type = b.get_type();
         if (self.is_native_int_type_or_bool(a_type) && self.is_native_int_type_or_bool(b_type))
             || (a_type.is_vector() && b_type.is_vector())
         {
+            if !a_type.is_compatible_with(b_type) {
+                if a_type.is_vector() {
+                    // Vector types need to be bitcast.
+                    // TODO(antoyo): perhaps use __builtin_convertvector for vector casting.
+                    b = self.context.new_bitcast(self.location, b, a_type);
+                } else {
+                    b = self.context.new_cast(self.location, b, a_type);
+                }
+            }
             self.context.new_binary_op(self.location, operation, a_type, a, b)
         } else {
             debug_assert!(a_type.dyncast_array().is_some());
@@ -351,6 +363,11 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
             // TODO(antoyo): is it correct to use rhs type instead of the parameter typ?
             .new_local(self.location, rhs.get_type(), "binopResult")
             .get_address(self.location);
+        let new_type = type_kind_to_gcc_type(new_kind);
+        let new_type = self.context.new_c_type(new_type);
+        let lhs = self.context.new_cast(self.location, lhs, new_type);
+        let rhs = self.context.new_cast(self.location, rhs, new_type);
+        let res = self.context.new_cast(self.location, res, new_type.make_pointer());
         let overflow = self.overflow_call(intrinsic, &[lhs, rhs, res], None);
         (res.dereference(self.location).to_rvalue(), overflow)
     }
@@ -477,11 +494,27 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
             let lhs_low = self.context.new_cast(self.location, self.low(lhs), unsigned_type);
             let rhs_low = self.context.new_cast(self.location, self.low(rhs), unsigned_type);
 
+            let mut lhs_high = self.high(lhs);
+            let mut rhs_high = self.high(rhs);
+
+            match op {
+                IntPredicate::IntUGT
+                | IntPredicate::IntUGE
+                | IntPredicate::IntULT
+                | IntPredicate::IntULE => {
+                    lhs_high = self.context.new_cast(self.location, lhs_high, unsigned_type);
+                    rhs_high = self.context.new_cast(self.location, rhs_high, unsigned_type);
+                }
+                // TODO(antoyo): we probably need to handle signed comparison for unsigned
+                // integers.
+                _ => (),
+            }
+
             let condition = self.context.new_comparison(
                 self.location,
                 ComparisonOp::LessThan,
-                self.high(lhs),
-                self.high(rhs),
+                lhs_high,
+                rhs_high,
             );
             self.llbb().end_with_conditional(self.location, condition, block1, block2);
 
@@ -495,8 +528,8 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
             let condition = self.context.new_comparison(
                 self.location,
                 ComparisonOp::GreaterThan,
-                self.high(lhs),
-                self.high(rhs),
+                lhs_high,
+                rhs_high,
             );
             block2.end_with_conditional(self.location, condition, block3, block4);
 
@@ -620,7 +653,7 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
         }
     }
 
-    pub fn gcc_xor(&self, a: RValue<'gcc>, b: RValue<'gcc>) -> RValue<'gcc> {
+    pub fn gcc_xor(&self, a: RValue<'gcc>, mut b: RValue<'gcc>) -> RValue<'gcc> {
         let a_type = a.get_type();
         let b_type = b.get_type();
         if a_type.is_vector() && b_type.is_vector() {
@@ -628,6 +661,9 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
             a ^ b
         } else if self.is_native_int_type_or_bool(a_type) && self.is_native_int_type_or_bool(b_type)
         {
+            if !a_type.is_compatible_with(b_type) {
+                b = self.context.new_cast(self.location, b, a_type);
+            }
             a ^ b
         } else {
             self.concat_low_high_rvalues(
@@ -1042,3 +1078,25 @@ impl<'gcc, 'tcx> CodegenCx<'gcc, 'tcx> {
         self.context.new_array_constructor(None, typ, &values)
     }
 }
+
+fn type_kind_to_gcc_type<I: Interner>(kind: TyKind<I>) -> CType {
+    use rustc_middle::ty::IntTy::*;
+    use rustc_middle::ty::UintTy::*;
+    use rustc_middle::ty::{Int, Uint};
+
+    match kind {
+        Int(I8) => CType::Int8t,
+        Int(I16) => CType::Int16t,
+        Int(I32) => CType::Int32t,
+        Int(I64) => CType::Int64t,
+        Int(I128) => CType::Int128t,
+
+        Uint(U8) => CType::UInt8t,
+        Uint(U16) => CType::UInt16t,
+        Uint(U32) => CType::UInt32t,
+        Uint(U64) => CType::UInt64t,
+        Uint(U128) => CType::UInt128t,
+
+        _ => unimplemented!("Kind: {:?}", kind),
+    }
+}
diff --git a/src/intrinsic/archs.rs b/src/intrinsic/archs.rs
index 915ed875e32..d1b2a93243d 100644
--- a/src/intrinsic/archs.rs
+++ b/src/intrinsic/archs.rs
@@ -95,8 +95,11 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "cubema" => "__builtin_amdgcn_cubema",
                     "cubesc" => "__builtin_amdgcn_cubesc",
                     "cubetc" => "__builtin_amdgcn_cubetc",
+                    "cvt.f16.bf8" => "__builtin_amdgcn_cvt_f16_bf8",
+                    "cvt.f16.fp8" => "__builtin_amdgcn_cvt_f16_fp8",
                     "cvt.f32.bf8" => "__builtin_amdgcn_cvt_f32_bf8",
                     "cvt.f32.fp8" => "__builtin_amdgcn_cvt_f32_fp8",
+                    "cvt.f32.fp8.e5m3" => "__builtin_amdgcn_cvt_f32_fp8_e5m3",
                     "cvt.off.f32.i4" => "__builtin_amdgcn_cvt_off_f32_i4",
                     "cvt.pk.bf8.f32" => "__builtin_amdgcn_cvt_pk_bf8_f32",
                     "cvt.pk.f16.bf8" => "__builtin_amdgcn_cvt_pk_f16_bf8",
@@ -181,6 +184,12 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "dot4.f32.fp8.bf8" => "__builtin_amdgcn_dot4_f32_fp8_bf8",
                     "dot4.f32.fp8.fp8" => "__builtin_amdgcn_dot4_f32_fp8_fp8",
                     "ds.add.gs.reg.rtn" => "__builtin_amdgcn_ds_add_gs_reg_rtn",
+                    "ds.atomic.async.barrier.arrive.b64" => {
+                        "__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64"
+                    }
+                    "ds.atomic.barrier.arrive.rtn.b64" => {
+                        "__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64"
+                    }
                     "ds.bpermute" => "__builtin_amdgcn_ds_bpermute",
                     "ds.bpermute.fi.b32" => "__builtin_amdgcn_ds_bpermute_fi_b32",
                     "ds.gws.barrier" => "__builtin_amdgcn_ds_gws_barrier",
@@ -198,8 +207,32 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "fdot2.f16.f16" => "__builtin_amdgcn_fdot2_f16_f16",
                     "fdot2.f32.bf16" => "__builtin_amdgcn_fdot2_f32_bf16",
                     "fdot2c.f32.bf16" => "__builtin_amdgcn_fdot2c_f32_bf16",
+                    "flat.prefetch" => "__builtin_amdgcn_flat_prefetch",
                     "fmul.legacy" => "__builtin_amdgcn_fmul_legacy",
+                    "global.load.async.to.lds.b128" => {
+                        "__builtin_amdgcn_global_load_async_to_lds_b128"
+                    }
+                    "global.load.async.to.lds.b32" => {
+                        "__builtin_amdgcn_global_load_async_to_lds_b32"
+                    }
+                    "global.load.async.to.lds.b64" => {
+                        "__builtin_amdgcn_global_load_async_to_lds_b64"
+                    }
+                    "global.load.async.to.lds.b8" => "__builtin_amdgcn_global_load_async_to_lds_b8",
                     "global.load.lds" => "__builtin_amdgcn_global_load_lds",
+                    "global.prefetch" => "__builtin_amdgcn_global_prefetch",
+                    "global.store.async.from.lds.b128" => {
+                        "__builtin_amdgcn_global_store_async_from_lds_b128"
+                    }
+                    "global.store.async.from.lds.b32" => {
+                        "__builtin_amdgcn_global_store_async_from_lds_b32"
+                    }
+                    "global.store.async.from.lds.b64" => {
+                        "__builtin_amdgcn_global_store_async_from_lds_b64"
+                    }
+                    "global.store.async.from.lds.b8" => {
+                        "__builtin_amdgcn_global_store_async_from_lds_b8"
+                    }
                     "groupstaticsize" => "__builtin_amdgcn_groupstaticsize",
                     "iglp.opt" => "__builtin_amdgcn_iglp_opt",
                     "implicit.buffer.ptr" => "__builtin_amdgcn_implicit_buffer_ptr",
@@ -291,6 +324,7 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "s.incperflevel" => "__builtin_amdgcn_s_incperflevel",
                     "s.memrealtime" => "__builtin_amdgcn_s_memrealtime",
                     "s.memtime" => "__builtin_amdgcn_s_memtime",
+                    "s.monitor.sleep" => "__builtin_amdgcn_s_monitor_sleep",
                     "s.sendmsg" => "__builtin_amdgcn_s_sendmsg",
                     "s.sendmsghalt" => "__builtin_amdgcn_s_sendmsghalt",
                     "s.setprio" => "__builtin_amdgcn_s_setprio",
@@ -300,11 +334,15 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "s.sleep.var" => "__builtin_amdgcn_s_sleep_var",
                     "s.ttracedata" => "__builtin_amdgcn_s_ttracedata",
                     "s.ttracedata.imm" => "__builtin_amdgcn_s_ttracedata_imm",
+                    "s.wait.asynccnt" => "__builtin_amdgcn_s_wait_asynccnt",
                     "s.wait.event.export.ready" => "__builtin_amdgcn_s_wait_event_export_ready",
+                    "s.wait.tensorcnt" => "__builtin_amdgcn_s_wait_tensorcnt",
                     "s.waitcnt" => "__builtin_amdgcn_s_waitcnt",
                     "sad.hi.u8" => "__builtin_amdgcn_sad_hi_u8",
                     "sad.u16" => "__builtin_amdgcn_sad_u16",
                     "sad.u8" => "__builtin_amdgcn_sad_u8",
+                    "sat.pk4.i4.i8" => "__builtin_amdgcn_sat_pk4_i4_i8",
+                    "sat.pk4.u4.u8" => "__builtin_amdgcn_sat_pk4_u4_u8",
                     "sched.barrier" => "__builtin_amdgcn_sched_barrier",
                     "sched.group.barrier" => "__builtin_amdgcn_sched_group_barrier",
                     "sdot2" => "__builtin_amdgcn_sdot2",
@@ -346,8 +384,13 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
                     "smfmac.i32.16x16x64.i8" => "__builtin_amdgcn_smfmac_i32_16x16x64_i8",
                     "smfmac.i32.32x32x32.i8" => "__builtin_amdgcn_smfmac_i32_32x32x32_i8",
                     "smfmac.i32.32x32x64.i8" => "__builtin_amdgcn_smfmac_i32_32x32x64_i8",
+                    "struct.ptr.buffer.load.lds" => "__builtin_amdgcn_struct_ptr_buffer_load_lds",
                     "sudot4" => "__builtin_amdgcn_sudot4",
                     "sudot8" => "__builtin_amdgcn_sudot8",
+                    "tensor.load.to.lds" => "__builtin_amdgcn_tensor_load_to_lds",
+                    "tensor.load.to.lds.d2" => "__builtin_amdgcn_tensor_load_to_lds_d2",
+                    "tensor.store.from.lds" => "__builtin_amdgcn_tensor_store_from_lds",
+                    "tensor.store.from.lds.d2" => "__builtin_amdgcn_tensor_store_from_lds_d2",
                     "udot2" => "__builtin_amdgcn_udot2",
                     "udot4" => "__builtin_amdgcn_udot4",
                     "udot8" => "__builtin_amdgcn_udot8",
@@ -6326,6 +6369,23 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
             }
             s390(name, full_name)
         }
+        "spv" => {
+            #[allow(non_snake_case)]
+            fn spv(name: &str, full_name: &str) -> &'static str {
+                match name {
+                    // spv
+                    "num.subgroups" => "__builtin_spirv_num_subgroups",
+                    "subgroup.id" => "__builtin_spirv_subgroup_id",
+                    "subgroup.local.invocation.id" => {
+                        "__builtin_spirv_subgroup_local_invocation_id"
+                    }
+                    "subgroup.max.size" => "__builtin_spirv_subgroup_max_size",
+                    "subgroup.size" => "__builtin_spirv_subgroup_size",
+                    _ => unimplemented!("***** unsupported LLVM intrinsic {full_name}"),
+                }
+            }
+            spv(name, full_name)
+        }
         "ve" => {
             #[allow(non_snake_case)]
             fn ve(name: &str, full_name: &str) -> &'static str {
diff --git a/src/intrinsic/mod.rs b/src/intrinsic/mod.rs
index 0753ac1aeb8..eb0a5336a1f 100644
--- a/src/intrinsic/mod.rs
+++ b/src/intrinsic/mod.rs
@@ -925,10 +925,17 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
         // TODO(antoyo): use width?
         let arg_type = arg.get_type();
         let result_type = self.u32_type;
+        let arg = if arg_type.is_signed(self.cx) {
+            let new_type = arg_type.to_unsigned(self.cx);
+            self.gcc_int_cast(arg, new_type)
+        } else {
+            arg
+        };
+        let arg_type = arg.get_type();
         let count_leading_zeroes =
             // TODO(antoyo): write a new function Type::is_compatible_with(&Type) and use it here
             // instead of using is_uint().
-            if arg_type.is_uint(self.cx) {
+            if arg_type.is_uchar(self.cx) || arg_type.is_ushort(self.cx) || arg_type.is_uint(self.cx) {
                 "__builtin_clz"
             }
             else if arg_type.is_ulong(self.cx) {
diff --git a/src/intrinsic/simd.rs b/src/intrinsic/simd.rs
index 350915a277e..fdc15d580ef 100644
--- a/src/intrinsic/simd.rs
+++ b/src/intrinsic/simd.rs
@@ -206,6 +206,28 @@ pub fn generic_simd_intrinsic<'a, 'gcc, 'tcx>(
         );
     }
 
+    #[cfg(feature = "master")]
+    if name == sym::simd_funnel_shl {
+        return Ok(simd_funnel_shift(
+            bx,
+            args[0].immediate(),
+            args[1].immediate(),
+            args[2].immediate(),
+            true,
+        ));
+    }
+
+    #[cfg(feature = "master")]
+    if name == sym::simd_funnel_shr {
+        return Ok(simd_funnel_shift(
+            bx,
+            args[0].immediate(),
+            args[1].immediate(),
+            args[2].immediate(),
+            false,
+        ));
+    }
+
     if name == sym::simd_bswap {
         return Ok(simd_bswap(bx, args[0].immediate()));
     }
@@ -1434,3 +1456,62 @@ pub fn generic_simd_intrinsic<'a, 'gcc, 'tcx>(
 
     unimplemented!("simd {}", name);
 }
+
+#[cfg(feature = "master")]
+fn simd_funnel_shift<'a, 'gcc, 'tcx>(
+    bx: &mut Builder<'a, 'gcc, 'tcx>,
+    a: RValue<'gcc>,
+    b: RValue<'gcc>,
+    shift: RValue<'gcc>,
+    shift_left: bool,
+) -> RValue<'gcc> {
+    use crate::common::SignType;
+
+    let a_type = a.get_type();
+    let vector_type = a_type.unqualified().dyncast_vector().expect("vector type");
+    let num_units = vector_type.get_num_units();
+    let elem_type = vector_type.get_element_type();
+
+    let (new_int_type, int_shift_val, int_mask) = if elem_type.is_compatible_with(bx.u8_type)
+        || elem_type.is_compatible_with(bx.i8_type)
+    {
+        (bx.u16_type, 8, u8::MAX as u64)
+    } else if elem_type.is_compatible_with(bx.u16_type) || elem_type.is_compatible_with(bx.i16_type)
+    {
+        (bx.u32_type, 16, u16::MAX as u64)
+    } else if elem_type.is_compatible_with(bx.u32_type) || elem_type.is_compatible_with(bx.i32_type)
+    {
+        (bx.u64_type, 32, u32::MAX as u64)
+    } else if elem_type.is_compatible_with(bx.u64_type) || elem_type.is_compatible_with(bx.i64_type)
+    {
+        (bx.u128_type, 64, u64::MAX)
+    } else {
+        unimplemented!("funnel shift on {:?}", elem_type);
+    };
+
+    let int_mask = bx.context.new_rvalue_from_long(new_int_type, int_mask as i64);
+    let int_shift_val = bx.context.new_rvalue_from_int(new_int_type, int_shift_val);
+    let mut elements = vec![];
+    let unsigned_type = elem_type.to_unsigned(bx);
+    for i in 0..num_units {
+        let index = bx.context.new_rvalue_from_int(bx.int_type, i as i32);
+        let a_val = bx.context.new_vector_access(None, a, index).to_rvalue();
+        let a_val = bx.context.new_bitcast(None, a_val, unsigned_type);
+        // TODO: we probably need to use gcc_int_cast instead.
+        let a_val = bx.gcc_int_cast(a_val, new_int_type);
+        let b_val = bx.context.new_vector_access(None, b, index).to_rvalue();
+        let b_val = bx.context.new_bitcast(None, b_val, unsigned_type);
+        let b_val = bx.gcc_int_cast(b_val, new_int_type);
+        let shift_val = bx.context.new_vector_access(None, shift, index).to_rvalue();
+        let shift_val = bx.gcc_int_cast(shift_val, new_int_type);
+        let mut val = a_val << int_shift_val | b_val;
+        if shift_left {
+            val = (val << shift_val) >> int_shift_val;
+        } else {
+            val = (val >> shift_val) & int_mask;
+        }
+        let val = bx.gcc_int_cast(val, elem_type);
+        elements.push(val);
+    }
+    bx.context.new_rvalue_from_vector(None, a_type, &elements)
+}
diff --git a/src/lib.rs b/src/lib.rs
index 613315f77a6..b11f11d38e3 100644
--- a/src/lib.rs
+++ b/src/lib.rs
@@ -50,6 +50,7 @@ extern crate rustc_session;
 extern crate rustc_span;
 extern crate rustc_symbol_mangling;
 extern crate rustc_target;
+extern crate rustc_type_ir;
 
 // This prevents duplicating functions and statics that are already part of the host rustc process.
 #[allow(unused_extern_crates)]
@@ -362,9 +363,9 @@ impl WriteBackendMethods for GccCodegenBackend {
         _exported_symbols_for_lto: &[String],
         each_linked_rlib_for_lto: &[PathBuf],
         modules: Vec<FatLtoInput<Self>>,
-        diff_fncs: Vec<AutoDiffItem>,
+        diff_functions: Vec<AutoDiffItem>,
     ) -> Result<ModuleCodegen<Self::Module>, FatalError> {
-        if !diff_fncs.is_empty() {
+        if !diff_functions.is_empty() {
             unimplemented!();
         }
 
diff --git a/tests/failing-lto-tests.txt b/tests/failing-lto-tests.txt
index b9126fb73a7..bf0633f7320 100644
--- a/tests/failing-lto-tests.txt
+++ b/tests/failing-lto-tests.txt
@@ -28,6 +28,6 @@ tests/ui/macros/macro-comma-behavior-rpass.rs
 tests/ui/macros/rfc-2011-nicer-assert-messages/assert-with-custom-errors-does-not-create-unnecessary-code.rs
 tests/ui/macros/rfc-2011-nicer-assert-messages/feature-gate-generic_assert.rs
 tests/ui/macros/stringify.rs
-tests/ui/reexport-test-harness-main.rs
 tests/ui/rfcs/rfc-1937-termination-trait/termination-trait-in-test.rs
 tests/ui/binding/fn-arg-incomplete-pattern-drop-order.rs
+tests/ui/lto/debuginfo-lto-alloc.rs
diff --git a/tests/failing-run-make-tests.txt b/tests/failing-run-make-tests.txt
index 842533cd3c6..29032b321fa 100644
--- a/tests/failing-run-make-tests.txt
+++ b/tests/failing-run-make-tests.txt
@@ -6,7 +6,6 @@ tests/run-make/doctests-keep-binaries/
 tests/run-make/doctests-runtool/
 tests/run-make/emit-shared-files/
 tests/run-make/exit-code/
-tests/run-make/issue-22131/
 tests/run-make/issue-64153/
 tests/run-make/llvm-ident/
 tests/run-make/native-link-modifier-bundle/
diff --git a/tests/failing-ui-tests.txt b/tests/failing-ui-tests.txt
index 6979c04d534..41fb4729c07 100644
--- a/tests/failing-ui-tests.txt
+++ b/tests/failing-ui-tests.txt
@@ -10,11 +10,10 @@ tests/ui/iterators/iter-sum-overflow-overflow-checks.rs
 tests/ui/mir/mir_drop_order.rs
 tests/ui/mir/mir_let_chains_drop_order.rs
 tests/ui/mir/mir_match_guard_let_chains_drop_order.rs
-tests/ui/oom_unwind.rs
+tests/ui/panics/oom-panic-unwind.rs
 tests/ui/panic-runtime/abort-link-to-unwinding-crates.rs
 tests/ui/panic-runtime/abort.rs
 tests/ui/panic-runtime/link-to-abort.rs
-tests/ui/unwind-no-uwtable.rs
 tests/ui/parser/unclosed-delimiter-in-dep.rs
 tests/ui/consts/missing_span_in_backtrace.rs
 tests/ui/drop/dynamic-drop.rs
@@ -82,3 +81,8 @@ tests/ui/coroutine/panic-drops.rs
 tests/ui/coroutine/panic-safe.rs
 tests/ui/process/nofile-limit.rs
 tests/ui/simd/intrinsic/generic-arithmetic-pass.rs
+tests/ui/linking/no-gc-encapsulation-symbols.rs
+tests/ui/panics/unwind-force-no-unwind-tables.rs
+tests/ui/attributes/fn-align-dyn.rs
+tests/ui/linkage-attr/raw-dylib/elf/glibc-x86_64.rs
+tests/ui/explicit-tail-calls/recursion-etc.rs
diff --git a/tools/cspell_dicts/rustc_codegen_gcc.txt b/tools/cspell_dicts/rustc_codegen_gcc.txt
index 31023e50ffa..4fb018b3ecd 100644
--- a/tools/cspell_dicts/rustc_codegen_gcc.txt
+++ b/tools/cspell_dicts/rustc_codegen_gcc.txt
@@ -8,6 +8,7 @@ clzll
 cmse
 codegened
 csky
+ctfe
 ctlz
 ctpop
 cttz
@@ -25,6 +26,7 @@ fwrapv
 gimple
 hrtb
 immediates
+interner
 liblto
 llbb
 llcx
@@ -47,6 +49,7 @@ mavx
 mcmodel
 minimumf
 minnumf
+miri
 monomorphization
 monomorphizations
 monomorphized

From 2136c500cdbdd08fe7c25927b5d3ce790d75f68b Mon Sep 17 00:00:00 2001
From: Sebastian Poeplau <poeplau@adacore.com>
Date: Fri, 1 Aug 2025 12:39:50 +0200
Subject: [PATCH 3/8] Preserve the .debug_gdb_scripts section

Make sure that compiler and linker don't optimize the section's contents
away by adding the global holding the data to "llvm.used". The volatile
load in the main shim is retained because "llvm.used", which translates
to SHF_GNU_RETAIN on ELF targets, requires a reasonably recent linker;
emitting the volatile load ensures compatibility with older linkers, at
least when libstd is used.

Pretty printers in dylib dependencies are now emitted by the main crate
instead of the dylib; apart from matching how rlibs are handled, this
approach has the advantage that `omit_gdb_pretty_printer_section` keeps
working with dylib dependencies.
---
 src/debuginfo.rs | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/src/debuginfo.rs b/src/debuginfo.rs
index 4c8585192a1..4c0b6439523 100644
--- a/src/debuginfo.rs
+++ b/src/debuginfo.rs
@@ -254,7 +254,8 @@ impl<'gcc, 'tcx> DebugInfoCodegenMethods<'tcx> for CodegenCx<'gcc, 'tcx> {
         // TODO(antoyo): implement.
     }
 
-    fn debuginfo_finalize(&self) {
+    fn debuginfo_finalize(&mut self) {
+        // TODO: emit section `.debug_gdb_scripts`.
         self.context.set_debug_info(true)
     }
 

From 82aac02c5267a427fdc5e8e1abbba98fc11c832c Mon Sep 17 00:00:00 2001
From: bjorn3 <17426603+bjorn3@users.noreply.github.com>
Date: Wed, 6 Aug 2025 18:01:07 +0000
Subject: [PATCH 4/8] Revert "Preserve the .debug_gdb_scripts section"

This reverts commit 868bdde25b030e0b71a29a5dbc04a891036e702e.
---
 src/debuginfo.rs | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/src/debuginfo.rs b/src/debuginfo.rs
index 4c0b6439523..4c8585192a1 100644
--- a/src/debuginfo.rs
+++ b/src/debuginfo.rs
@@ -254,8 +254,7 @@ impl<'gcc, 'tcx> DebugInfoCodegenMethods<'tcx> for CodegenCx<'gcc, 'tcx> {
         // TODO(antoyo): implement.
     }
 
-    fn debuginfo_finalize(&mut self) {
-        // TODO: emit section `.debug_gdb_scripts`.
+    fn debuginfo_finalize(&self) {
         self.context.set_debug_info(true)
     }
 

From 098d06b0300e7ec3d4174904c2d6a9d4444a30a2 Mon Sep 17 00:00:00 2001
From: bjorn3 <17426603+bjorn3@users.noreply.github.com>
Date: Thu, 7 Aug 2025 13:13:04 +0000
Subject: [PATCH 5/8] Prevent name collisions with internal implementation
 details

The implementation of the linkage attribute inside extern blocks defines
symbols starting with _rust_extern_with_linkage_. If someone tries to
also define this symbol you will get a symbol conflict or even an ICE.
By adding an unpredictable component to the symbol name, this becomes
less of an issue.
---
 src/consts.rs | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/src/consts.rs b/src/consts.rs
index c04c75e1b11..873f1f1951c 100644
--- a/src/consts.rs
+++ b/src/consts.rs
@@ -6,6 +6,7 @@ use rustc_codegen_ssa::traits::{
     BaseTypeCodegenMethods, ConstCodegenMethods, StaticCodegenMethods,
 };
 use rustc_hir::def::DefKind;
+use rustc_hir::def_id::LOCAL_CRATE;
 use rustc_middle::middle::codegen_fn_attrs::{CodegenFnAttrFlags, CodegenFnAttrs};
 use rustc_middle::mir::interpret::{
     self, ConstAllocation, ErrorHandled, Scalar as InterpScalar, read_target_uint,
@@ -384,8 +385,8 @@ fn check_and_apply_linkage<'gcc, 'tcx>(
         // linkage and there are no definitions), then
         // `extern_with_linkage_foo` will instead be initialized to
         // zero.
-        let mut real_name = "_rust_extern_with_linkage_".to_string();
-        real_name.push_str(sym);
+        let real_name =
+            format!("_rust_extern_with_linkage_{:016x}_{sym}", cx.tcx.stable_crate_id(LOCAL_CRATE));
         let global2 = cx.define_global(&real_name, gcc_type, is_tls, attrs.link_section);
         // TODO(antoyo): set linkage.
         let value = cx.const_ptrcast(global1.get_address(None), gcc_type);

From 26183d574197b2605a559097e604f4906c4892fc Mon Sep 17 00:00:00 2001
From: Antoni Boucher <bouanto@zoho.com>
Date: Thu, 14 Aug 2025 09:05:09 -0400
Subject: [PATCH 6/8] Update to nightly-2025-08-14

---
 rust-toolchain | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/rust-toolchain b/rust-toolchain
index 058e734be5c..059d469e8d4 100644
--- a/rust-toolchain
+++ b/rust-toolchain
@@ -1,3 +1,3 @@
 [toolchain]
-channel = "nightly-2025-08-03"
+channel = "nightly-2025-08-14"
 components = ["rust-src", "rustc-dev", "llvm-tools-preview"]

From 3f476d309853e98fee3663089fe7842ad5d3733a Mon Sep 17 00:00:00 2001
From: Antoni Boucher <bouanto@zoho.com>
Date: Thu, 14 Aug 2025 09:51:01 -0400
Subject: [PATCH 7/8] Rename/remove some UI and run-make tests

---
 tests/failing-run-make-tests.txt | 3 ---
 tests/failing-ui-tests.txt       | 4 +++-
 2 files changed, 3 insertions(+), 4 deletions(-)

diff --git a/tests/failing-run-make-tests.txt b/tests/failing-run-make-tests.txt
index 29032b321fa..c5e22970c66 100644
--- a/tests/failing-run-make-tests.txt
+++ b/tests/failing-run-make-tests.txt
@@ -6,7 +6,6 @@ tests/run-make/doctests-keep-binaries/
 tests/run-make/doctests-runtool/
 tests/run-make/emit-shared-files/
 tests/run-make/exit-code/
-tests/run-make/issue-64153/
 tests/run-make/llvm-ident/
 tests/run-make/native-link-modifier-bundle/
 tests/run-make/remap-path-prefix-dwarf/
@@ -34,8 +33,6 @@ tests/run-make/c-link-to-rust-staticlib/
 tests/run-make/foreign-double-unwind/
 tests/run-make/foreign-exceptions/
 tests/run-make/glibc-staticlib-args/
-tests/run-make/issue-36710/
-tests/run-make/issue-68794-textrel-on-minimal-lib/
 tests/run-make/lto-smoke-c/
 tests/run-make/return-non-c-like-enum/
 
diff --git a/tests/failing-ui-tests.txt b/tests/failing-ui-tests.txt
index 41fb4729c07..e2615bce190 100644
--- a/tests/failing-ui-tests.txt
+++ b/tests/failing-ui-tests.txt
@@ -20,7 +20,7 @@ tests/ui/drop/dynamic-drop.rs
 tests/ui/rfcs/rfc-2091-track-caller/std-panic-locations.rs
 tests/ui/simd/issue-17170.rs
 tests/ui/simd/issue-39720.rs
-tests/ui/issues/issue-14875.rs
+tests/ui/drop/panic-during-drop-14875.rs
 tests/ui/issues/issue-29948.rs
 tests/ui/process/println-with-broken-pipe.rs
 tests/ui/lto/thin-lto-inlines2.rs
@@ -86,3 +86,5 @@ tests/ui/panics/unwind-force-no-unwind-tables.rs
 tests/ui/attributes/fn-align-dyn.rs
 tests/ui/linkage-attr/raw-dylib/elf/glibc-x86_64.rs
 tests/ui/explicit-tail-calls/recursion-etc.rs
+tests/ui/explicit-tail-calls/indexer.rs
+tests/ui/explicit-tail-calls/drop-order.rs

From eff47b545a4065f87cc14943bc638096cf929325 Mon Sep 17 00:00:00 2001
From: Antoni Boucher <bouanto@zoho.com>
Date: Thu, 14 Aug 2025 11:46:34 -0400
Subject: [PATCH 8/8] Don't set feature compiler-builtins-no-f16-f128

---
 .github/workflows/m68k.yml | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/.github/workflows/m68k.yml b/.github/workflows/m68k.yml
index 176ee8628b0..e49c62d6c93 100644
--- a/.github/workflows/m68k.yml
+++ b/.github/workflows/m68k.yml
@@ -82,14 +82,14 @@ jobs:
     - name: Build sample project with target defined as JSON spec
       run: |
         ./y.sh prepare --only-libcore --cross
-        ./y.sh build --sysroot --features compiler-builtins-no-f16-f128 --target-triple m68k-unknown-linux-gnu --target ${{ github.workspace }}/target_specs/m68k-unknown-linux-gnu.json
+        ./y.sh build --sysroot --target-triple m68k-unknown-linux-gnu --target ${{ github.workspace }}/target_specs/m68k-unknown-linux-gnu.json
         CG_RUSTFLAGS="-Clinker=m68k-unknown-linux-gnu-gcc" ./y.sh cargo build --manifest-path=./tests/hello-world/Cargo.toml --target ${{ github.workspace }}/target_specs/m68k-unknown-linux-gnu.json
         ./y.sh clean all
 
     - name: Build
       run: |
         ./y.sh prepare --only-libcore --cross
-        ./y.sh build --sysroot --features compiler-builtins-no-f16-f128 --target-triple m68k-unknown-linux-gnu
+        ./y.sh build --sysroot --target-triple m68k-unknown-linux-gnu
         ./y.sh test --mini-tests --target-triple m68k-unknown-linux-gnu
         CG_GCC_TEST_TARGET=m68k-unknown-linux-gnu ./y.sh test --cargo-tests --target-triple m68k-unknown-linux-gnu
         ./y.sh clean all
@@ -102,7 +102,7 @@ jobs:
 
     - name: Run tests
       run: |
-        ./y.sh test --target-triple m68k-unknown-linux-gnu --release --clean --build-sysroot --sysroot-features compiler-builtins-no-f16-f128 ${{ matrix.commands }}
+        ./y.sh test --target-triple m68k-unknown-linux-gnu --release --clean --build-sysroot ${{ matrix.commands }}
 
     - name: Run Hello World!
       run: |