[clang] 2c9afb4 - [CIR][AMDGPU] Add support for AMDGCN bitfield extract builtins (#222186)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-09-09T14:56:29-04:00
New Revision: 2c9afb46c32ccf19fe048688e273cad4c7b3dfc6

URL: https://github.com/llvm/llvm-project/commit/2c9afb46c32ccf19fe048688e273cad4c7b3dfc6
DIFF: https://github.com/llvm/llvm-project/commit/2c9afb46c32ccf19fe048688e273cad4c7b3dfc6.diff

LOG: [CIR][AMDGPU] Add support for AMDGCN bitfield extract builtins (#222186)

Adds codegen for the following AMDGCN bitfield extract builtins:

* __builtin_amdgcn_ubfe (unsigned)
* __builtin_amdgcn_sbfe (signed)

These are lowered to the corresponding `llvm.amdgcn.ubfe` and
`llvm.amdgcn.sbfe` intrinsics.

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 94d5a835d60fa..c31ac8fa946c9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -346,18 +346,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
                      getContext().BuiltinInfo.getName(builtinId));
     return mlir::Value{};
   }
-  case AMDGPU::BI__builtin_amdgcn_ubfe: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
-  }
-  case AMDGPU::BI__builtin_amdgcn_sbfe: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
-  }
+  case AMDGPU::BI__builtin_amdgcn_ubfe:
+    return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.ubfe").getValue();
+  case AMDGPU::BI__builtin_amdgcn_sbfe:
+    return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.sbfe").getValue();
   case AMDGPU::BI__builtin_amdgcn_ballot_w32:
   case AMDGPU::BI__builtin_amdgcn_ballot_w64:
     return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot",

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 9bbfb6747cbf8..9b687e98d1195 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -331,3 +331,21 @@ __device__ void test_class_f32(bool* out, float a, int b) {
 __device__ void test_class_f64(bool* out, double a, int b) {
   *out = __builtin_amdgcn_class(a, b);
 }
+
+// CIR-LABEL: @_Z9test_ubfePjjjj
+// CIR: cir.call_llvm_intrinsic "amdgcn.ubfe" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i
+// LLVM: define{{.*}} void @_Z9test_ubfePjjjj
+// LLVM: call{{.*}} i32 @llvm.amdgcn.ubfe.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}})
+__device__ void test_ubfe(unsigned int* out, unsigned int src0, unsigned int src1,
+                          unsigned int src2) {
+  *out = __builtin_amdgcn_ubfe(src0, src1, src2);
+}
+
+// CIR-LABEL: @_Z9test_sbfePjjjj
+// CIR: cir.call_llvm_intrinsic "amdgcn.sbfe" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i
+// LLVM: define{{.*}} void @_Z9test_sbfePjjjj
+// LLVM: call{{.*}} i32 @llvm.amdgcn.sbfe.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}})
+__device__ void test_sbfe(unsigned int* out, unsigned int src0, unsigned int src1,
+                          unsigned int src2) {
+  *out = __builtin_amdgcn_sbfe(src0, src1, src2);
+}