[clang] [CIR][AMDGPU] Add support for AMDGCN s_prefetch builtins (PR #223228)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clangir

Author: Ayokunle Amodu (ayokunle321)

<details>
<summary>Changes</summary>

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

---
Full diff: https://github.com/llvm/llvm-project/pull/223228.diff


3 Files Affected:

- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+10-12) 
- (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip (+51) 
- (removed) clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip (-10) 


``````````diff
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);
-}

``````````

</details>


https://github.com/llvm/llvm-project/pull/223228