[clang] [CIR][AMDGPU] Add support for AMDGCN rsq builtins (PR #197349)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 10 09:54:51 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197349
>From 329d0fcee811127eb3e4a5c76467f70770231e32 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Tue, 12 May 2026 21:52:28 -0400
Subject: [PATCH 1/5] add amdgcn rsq builtin
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 15 ++++----
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 34 +++++++++++++++++++
2 files changed, 41 insertions(+), 8 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..1cf6672f6e86e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -269,17 +269,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_rsqf:
case AMDGPU::BI__builtin_amdgcn_rsqh:
case AMDGPU::BI__builtin_amdgcn_rsq_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.rsq",
+ src.getType(), mlir::ValueRange{src});
}
case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
- 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.rsq.clamp", src.getType(),
+ mlir::ValueRange{src});
}
case AMDGPU::BI__builtin_amdgcn_sinf:
case AMDGPU::BI__builtin_amdgcn_sinh:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index ca708bca8587b..2f2ae287e19a8 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -111,3 +111,37 @@ __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: @_Z12test_rsq_f32Pff
+// CIR: cir.llvm.intrinsic "amdgcn.rsq" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z12test_rsq_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.rsq.f32(float %{{.*}})
+__device__ void test_rsq_f32(float* out, float a)
+{
+ *out = __builtin_amdgcn_rsqf(a);
+}
+
+// CIR-LABEL: @_Z12test_rsq_f64Pdd
+// CIR: cir.llvm.intrinsic "amdgcn.rsq" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @_Z12test_rsq_f64Pdd
+// LLVM: call{{.*}} double @llvm.amdgcn.rsq.f64(double %{{.*}})
+__device__ void test_rsq_f64(double* out, double a) {
+ *out = __builtin_amdgcn_rsq(a);
+}
+
+// CIR-LABEL: @_Z18test_rsq_clamp_f32Pff
+// CIR: cir.llvm.intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.rsq.clamp.f32(float %{{.*}})
+__device__ void test_rsq_clamp_f32(float* out, float a)
+{
+ *out = __builtin_amdgcn_rsq_clampf(a);
+}
+
+// CIR-LABEL: @_Z18test_rsq_clamp_f64Pdd
+// CIR: cir.llvm.intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f64Pdd
+// LLVM: call{{.*}} double @llvm.amdgcn.rsq.clamp.f64(double %{{.*}})
+__device__ void test_rsq_clamp_f64(double* out, double a) {
+ *out = __builtin_amdgcn_rsq_clamp(a);
+}
>From 18fe54fa6a6dee0219becb3a343fc254837c94e1 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Tue, 12 May 2026 22:16:43 -0400
Subject: [PATCH 2/5] fix intinsic call instruction
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 2f2ae287e19a8..c2de3108719f8 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -113,7 +113,7 @@ __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out)
}
// CIR-LABEL: @_Z12test_rsq_f32Pff
-// CIR: cir.llvm.intrinsic "amdgcn.rsq" {{.*}} : (!cir.float) -> !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.float) -> !cir.float
// LLVM: define{{.*}} void @_Z12test_rsq_f32Pff
// LLVM: call{{.*}} float @llvm.amdgcn.rsq.f32(float %{{.*}})
__device__ void test_rsq_f32(float* out, float a)
@@ -122,7 +122,7 @@ __device__ void test_rsq_f32(float* out, float a)
}
// CIR-LABEL: @_Z12test_rsq_f64Pdd
-// CIR: cir.llvm.intrinsic "amdgcn.rsq" {{.*}} : (!cir.double) -> !cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.double) -> !cir.double
// LLVM: define{{.*}} void @_Z12test_rsq_f64Pdd
// LLVM: call{{.*}} double @llvm.amdgcn.rsq.f64(double %{{.*}})
__device__ void test_rsq_f64(double* out, double a) {
@@ -130,7 +130,7 @@ __device__ void test_rsq_f64(double* out, double a) {
}
// CIR-LABEL: @_Z18test_rsq_clamp_f32Pff
-// CIR: cir.llvm.intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.float) -> !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.float) -> !cir.float
// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f32Pff
// LLVM: call{{.*}} float @llvm.amdgcn.rsq.clamp.f32(float %{{.*}})
__device__ void test_rsq_clamp_f32(float* out, float a)
@@ -139,7 +139,7 @@ __device__ void test_rsq_clamp_f32(float* out, float a)
}
// CIR-LABEL: @_Z18test_rsq_clamp_f64Pdd
-// CIR: cir.llvm.intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.double) -> !cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.double) -> !cir.double
// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f64Pdd
// LLVM: call{{.*}} double @llvm.amdgcn.rsq.clamp.f64(double %{{.*}})
__device__ void test_rsq_clamp_f64(double* out, double a) {
>From b0ead0d9ec1e7f335667d340064bdbbe6fb92050 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 17:55:48 -0400
Subject: [PATCH 3/5] add tests for f16 and bf16 types
---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 28 ++++++++
.../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 +++++++++++++++++++
2 files changed, 93 insertions(+)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
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..2f3152e2de503
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -0,0 +1,28 @@
+#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: @_Z13test_rsq_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z13test_rsq_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.rsq.bf16(bfloat %{{.*}})
+// OGCG: define{{.*}} void @_Z13test_rsq_bf16PDF16bDF16b
+// OGCG: call{{.*}} bfloat @llvm.amdgcn.rsq.bf16(bfloat %{{.*}})
+__device__ void test_rsq_bf16(__bf16* out, __bf16 a) {
+ *out = __builtin_amdgcn_rsq_bf16(a);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
new file mode 100644
index 0000000000000..aee11f1ee73fc
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
@@ -0,0 +1,65 @@
+#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 tonga -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx900 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1012 -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 tonga -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 -fclangir \
+// RUN: -target-cpu gfx900 -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 -fclangir \
+// RUN: -target-cpu gfx1010 -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 -fclangir \
+// RUN: -target-cpu gfx1012 -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 tonga -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 gfx900 -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 gfx1010 -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 gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z10test_rsq_hPDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z10test_rsq_hPDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.rsq.f16(half %{{.*}})
+__device__ void test_rsq_h(_Float16* out, _Float16 a) {
+ *out = __builtin_amdgcn_rsqh(a);
+}
>From 2861d25e99fc3bcdb98637283a202e7ca012a59b Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 15:19:59 -0400
Subject: [PATCH 4/5] remove OGCG checks
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index 2f3152e2de503..90508b51b464c 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -21,8 +21,6 @@
// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.bf16) -> !cir.bf16
// LLVM: define{{.*}} void @_Z13test_rsq_bf16PDF16bDF16b
// LLVM: call{{.*}} bfloat @llvm.amdgcn.rsq.bf16(bfloat %{{.*}})
-// OGCG: define{{.*}} void @_Z13test_rsq_bf16PDF16bDF16b
-// OGCG: call{{.*}} bfloat @llvm.amdgcn.rsq.bf16(bfloat %{{.*}})
__device__ void test_rsq_bf16(__bf16* out, __bf16 a) {
*out = __builtin_amdgcn_rsq_bf16(a);
}
>From 6e55f8db4efc88ca6d0d7b4330c05df5109be6e1 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 12:54:29 -0400
Subject: [PATCH 5/5] switch out manual codgen out for function
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 10 +++-------
1 file changed, 3 insertions(+), 7 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 1cf6672f6e86e..4b2724b9ec42b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -269,16 +269,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_rsqf:
case AMDGPU::BI__builtin_amdgcn_rsqh:
case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
- mlir::Value src = emitScalarExpr(expr->getArg(0));
- return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "amdgcn.rsq",
- src.getType(), mlir::ValueRange{src});
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue();
}
case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
- mlir::Value src = emitScalarExpr(expr->getArg(0));
- return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
- "amdgcn.rsq.clamp", src.getType(),
- mlir::ValueRange{src});
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp")
+ .getValue();
}
case AMDGPU::BI__builtin_amdgcn_sinf:
case AMDGPU::BI__builtin_amdgcn_sinh:
More information about the cfe-commits
mailing list