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@+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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
