https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198025
>From aef923942bc764a15bda9b8da7c0287332cedcc6 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 15 May 2026 17:13:59 -0400 Subject: [PATCH 1/3] add amdgcn exp2 buitins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 8 +++--- .../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 +++++++++++++++++++ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 ++++++ 3 files changed, 38 insertions(+), 4 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 0a7ba0c194400..e7aee0dc9bbcc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -308,10 +308,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_exp2f: case AMDGPU::BI__builtin_amdgcn_exp2_bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + mlir::Value src = emitScalarExpr(expr->getArg(0)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.exp2", src.getType(), + mlir::ValueRange{src}); } case AMDGPU::BI__builtin_amdgcn_log_clampf: { cgm.errorNYI(expr->getSourceRange(), diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip new file mode 100644 index 0000000000000..efac022872f16 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -0,0 +1,26 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +//===----------------------------------------------------------------------===// +// Test AMDGPU builtins +//===----------------------------------------------------------------------===// + +// CIR-LABEL: @_Z14test_exp2_bf16PDF16bDF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.exp2" {{.*}} : (!cir.bf16) -> !cir.bf16 +// LLVM: define{{.*}} void @_Z14test_exp2_bf16PDF16bDF16b +// LLVM: call{{.*}} bfloat @llvm.amdgcn.exp2.bf16(bfloat %{{.*}}) +__device__ void test_exp2_bf16(__bf16* out, __bf16 a) { + *out = __builtin_amdgcn_exp2_bf16(a); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index ca708bca8587b..0468842fc57b9 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -111,3 +111,11 @@ __device__ void test_readfirstlane(int* out, int a) { __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) { *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr(); } + +// CIR-LABEL: @_Z13test_exp2_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.exp2" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z13test_exp2_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.exp2.f32(float %{{.*}}) +__device__ void test_exp2_f32(float* out, float a) { + *out = __builtin_amdgcn_exp2f(a); +} >From b3e8b598c088714c0fa60bddfeb86557ad3dff27 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:51:25 -0400 Subject: [PATCH 2/3] switch manual codegen for function --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index efac022872f16..d0d380bb580b8 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -1,5 +1,3 @@ -#include "../CodeGenCUDA/Inputs/cuda.h" - // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ // RUN: -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -o %t.cir @@ -13,6 +11,8 @@ // RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +#define __device__ __attribute__((device)) + //===----------------------------------------------------------------------===// // Test AMDGPU builtins //===----------------------------------------------------------------------===// >From 77440541cc9ab751908e04f70544148e626edd91 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:51:41 -0400 Subject: [PATCH 3/3] switch manual codegen for function --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index e7aee0dc9bbcc..61c376757b08e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -308,10 +308,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_exp2f: case AMDGPU::BI__builtin_amdgcn_exp2_bf16: { - mlir::Value src = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.exp2", src.getType(), - mlir::ValueRange{src}); + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.exp2").getValue(); } case AMDGPU::BI__builtin_amdgcn_log_clampf: { cgm.errorNYI(expr->getSourceRange(), _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
