From c88e2f49854724fe373865221485f5e071eed2bc Mon Sep 17 00:00:00 2001
From: Manuel Drehwald <git@manuel.drehwald.info>
Date: Wed, 3 Dec 2025 11:32:54 -0800
Subject: [PATCH] start adding amdgpu intrinsics

---
 compiler/rustc_codegen_llvm/src/intrinsic.rs       | 3 +++
 compiler/rustc_hir_analysis/src/check/intrinsic.rs | 7 +++++++
 compiler/rustc_span/src/symbol.rs                  | 3 +++
 3 files changed, 13 insertions(+)

diff --git a/compiler/rustc_codegen_llvm/src/intrinsic.rs b/compiler/rustc_codegen_llvm/src/intrinsic.rs
index 33541f7b695f8..6d89d0977755c 100644
--- a/compiler/rustc_codegen_llvm/src/intrinsic.rs
+++ b/compiler/rustc_codegen_llvm/src/intrinsic.rs
@@ -41,6 +41,9 @@ fn call_simple_intrinsic<'ll, 'tcx>(
     args: &[OperandRef<'tcx, &'ll Value>],
 ) -> Option<&'ll Value> {
     let (base_name, type_params): (&'static str, &[&'ll Type]) = match name {
+        sym::amd_workitem => ("llvm.amdgcn.workitem.id.x", &[]),
+        sym::amd_workgroup_id => ("llvm.amdgcn.workgroup.id.x", &[]),
+        sym::amd_workgroup_size => ("llvm.amdgcn.workgroup.size.x", &[]),
         sym::sqrtf16 => ("llvm.sqrt", &[bx.type_f16()]),
         sym::sqrtf32 => ("llvm.sqrt", &[bx.type_f32()]),
         sym::sqrtf64 => ("llvm.sqrt", &[bx.type_f64()]),
diff --git a/compiler/rustc_hir_analysis/src/check/intrinsic.rs b/compiler/rustc_hir_analysis/src/check/intrinsic.rs
index 676c9a980afff..e110f34290e9a 100644
--- a/compiler/rustc_hir_analysis/src/check/intrinsic.rs
+++ b/compiler/rustc_hir_analysis/src/check/intrinsic.rs
@@ -200,6 +200,9 @@ fn intrinsic_operation_unsafety(tcx: TyCtxt<'_>, intrinsic_id: LocalDefId) -> hi
         | sym::sinf64
         | sym::sinf128
         | sym::size_of
+        | sym::amd_workitem
+        | sym::amd_workgroup_id
+        | sym::amd_workgroup_size
         | sym::sqrtf16
         | sym::sqrtf32
         | sym::sqrtf64
@@ -353,6 +356,10 @@ pub(crate) fn check_intrinsic_type(
             tcx.types.unit,
         ),
 
+        sym::amd_workitem => (0, 0, vec![], tcx.types.i32),
+        sym::amd_workgroup_id => (0, 0, vec![], tcx.types.i32),
+        sym::amd_workgroup_size => (0, 0, vec![], tcx.types.i32),
+
         sym::sqrtf16 => (0, 0, vec![tcx.types.f16], tcx.types.f16),
         sym::sqrtf32 => (0, 0, vec![tcx.types.f32], tcx.types.f32),
         sym::sqrtf64 => (0, 0, vec![tcx.types.f64], tcx.types.f64),
diff --git a/compiler/rustc_span/src/symbol.rs b/compiler/rustc_span/src/symbol.rs
index f2b13dad1fd90..928c8b54a7931 100644
--- a/compiler/rustc_span/src/symbol.rs
+++ b/compiler/rustc_span/src/symbol.rs
@@ -2147,6 +2147,9 @@ symbols! {
         speed,
         spirv,
         spotlight,
+        amd_workitem,
+        amd_workgroup_id,
+        amd_workgroup_size,
         sqrtf16,
         sqrtf32,
         sqrtf64,