[clang] [CIR][AMDGPU] Add support for AMDGCN s_prefetch builtins (PR #223228)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Sun Sep 13 03:55:52 PDT 2026
https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/223228
Adds codegen for the following AMDGCN s_prefetch builtins:
- __builtin_amdgcn_s_prefetch_data
- __builtin_amdgcn_s_prefetch_inst
These are lowered to the corresponding `llvm.amdgcn.s.prefetch` intrinsics.
This also removes `builtins-amdgcn-s-prefetch-inst-nyi.hip`, which pinned the NYI diagnostic that no longer fires.
Assisted by: Claude Opus 5
>From 9d62bd4488daee4f89d2e1b6063429aeb1ce3277 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sun, 13 Sep 2026 12:54:28 +0200
Subject: [PATCH] [CIR][AMDGPU] Add support for AMDGCN s_prefetch builtins
Adds codegen for the following AMDGCN s_prefetch builtins:
__builtin_amdgcn_s_prefetch_data
__builtin_amdgcn_s_prefetch_inst
These are lowered to the corresponding llvm.amdgcn.s.prefetch intrinsics.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 22 ++++----
.../CodeGenHIP/builtins-amdgcn-prefetch.hip | 51 +++++++++++++++++++
.../builtins-amdgcn-s-prefetch-inst-nyi.hip | 10 ----
3 files changed, 61 insertions(+), 22 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip
delete mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 1c23ea142bf8d..1a26d8b1f6bee 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1050,18 +1050,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
}
- case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
- case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "amdgcn.s.prefetch.data",
+ cir::VoidType::get(builder.getContext()))
+ .getValue();
+ case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst:
+ return emitBuiltinWithOneOverloadedType<2>(
+ expr, "amdgcn.s.prefetch.inst",
+ cir::VoidType::get(builder.getContext()))
+ .getValue();
case Builtin::BIlogbf:
case Builtin::BI__builtin_logbf:
return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip
new file mode 100644
index 0000000000000..1b58d69838ef4
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip
@@ -0,0 +1,51 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+#define __device__ __attribute__((device))
+
+typedef __attribute__((address_space(1))) float *global_ptr_t;
+typedef __attribute__((address_space(4))) const char *constant_ptr_t;
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z20test_s_prefetch_dataPiPU3AS1fPU3AS4Kcj
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void>, !u32i) -> !void
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void, target_address_space(1)>, !u32i) -> !void
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u32i) -> !void
+// LLVM: define{{.*}} void @_Z20test_s_prefetch_dataPiPU3AS1fPU3AS4Kcj
+// LLVM: call void @llvm.amdgcn.s.prefetch.data.p0(ptr %{{.*}}, i32 0)
+// LLVM: call void @llvm.amdgcn.s.prefetch.data.p1(ptr addrspace(1) %{{.*}}, i32 %{{.*}})
+// LLVM: call void @llvm.amdgcn.s.prefetch.data.p4(ptr addrspace(4) %{{.*}}, i32 31)
+__device__ void test_s_prefetch_data(int *fp, global_ptr_t gp,
+ constant_ptr_t cp, unsigned int len) {
+ __builtin_amdgcn_s_prefetch_data(fp, 0);
+ __builtin_amdgcn_s_prefetch_data(gp, len);
+ __builtin_amdgcn_s_prefetch_data(cp, 31);
+}
+
+// CIR-LABEL: @_Z20test_s_prefetch_instPiPU3AS1fPU3AS4Kcj
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void>, !u32i) -> !void
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void, target_address_space(1)>, !u32i) -> !void
+// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u32i) -> !void
+// LLVM: define{{.*}} void @_Z20test_s_prefetch_instPiPU3AS1fPU3AS4Kcj
+// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %{{.*}}, i32 0)
+// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %{{.*}}, i32 %{{.*}})
+// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %{{.*}}, i32 31)
+__device__ void test_s_prefetch_inst(int *fp, global_ptr_t gp,
+ constant_ptr_t cp, unsigned int len) {
+ __builtin_amdgcn_s_prefetch_inst(fp, 0);
+ __builtin_amdgcn_s_prefetch_inst(gp, len);
+ __builtin_amdgcn_s_prefetch_inst(cp, 31);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip
deleted file mode 100644
index 74d058a725857..0000000000000
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip
+++ /dev/null
@@ -1,10 +0,0 @@
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -x hip -std=c++11 -fclangir \
-// RUN: -fcuda-is-device -emit-cir %s -verify -o %t.cir
-
-#define __device__ __attribute__((device))
-
-// expected-error at +2 {{ClangIR code gen Not Yet Implemented: unimplemented AMDGPU builtin call: __builtin_amdgcn_s_prefetch_inst}}
-__device__ void test_s_prefetch_inst(const void *p, unsigned int len) {
- __builtin_amdgcn_s_prefetch_inst(p, len);
-}
More information about the cfe-commits
mailing list