[clang] [CIR][AMDGPU] Add support for AMDGCN exp2 builtins (PR #198025)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 16 08:44:25 PDT 2026
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 <ayokunle321 at gmail.com>
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 <ayokunle321 at gmail.com>
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 <ayokunle321 at gmail.com>
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(),
More information about the cfe-commits
mailing list