[clang] [CIR][AMDGPU] Add support for AMDGCN sqrt builtins (PR #197342)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 15 05:51:52 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197342
>From 703be1efd1cee0c7fc32cbfd25bf721c051cb331 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Tue, 12 May 2026 20:59:51 -0400
Subject: [PATCH 1/9] add amdgcn sqrt builtin
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 8 ++++----
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 ++++++++++++++++
2 files changed, 20 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..eabf56ce09afd 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -260,10 +260,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_sqrtf:
case AMDGPU::BI__builtin_amdgcn_sqrth:
case AMDGPU::BI__builtin_amdgcn_sqrt_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.sqrt", src.getType(),
+ mlir::ValueRange{src});
}
case AMDGPU::BI__builtin_amdgcn_rsq:
case AMDGPU::BI__builtin_amdgcn_rsqf:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index ca708bca8587b..67b1ba91c438f 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -111,3 +111,19 @@ __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_sqrt_f32Pff
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z13test_sqrt_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
+__device__ void test_sqrt_f32(float* out, float a) {
+ *out = __builtin_amdgcn_sqrtf(a);
+}
+
+// CIR-LABEL: @_Z13test_sqrt_f64Pdd
+// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @_Z13test_sqrt_f64Pdd
+// LLVM: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
+__device__ void test_sqrt_f64(double* out, double a) {
+ *out = __builtin_amdgcn_sqrt(a);
+}
>From e43f5faeefa5455da608630944893d8c6b281553 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Tue, 12 May 2026 22:10:55 -0400
Subject: [PATCH 2/9] fix intinsic call instruction
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 67b1ba91c438f..ff024cf5eead1 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: @_Z13test_sqrt_f32Pff
-// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.float) -> !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.float) -> !cir.float
// LLVM: define{{.*}} void @_Z13test_sqrt_f32Pff
// LLVM: call{{.*}} float @llvm.amdgcn.sqrt.f32(float %{{.*}})
__device__ void test_sqrt_f32(float* out, float a) {
@@ -121,7 +121,7 @@ __device__ void test_sqrt_f32(float* out, float a) {
}
// CIR-LABEL: @_Z13test_sqrt_f64Pdd
-// CIR: cir.llvm.intrinsic "amdgcn.sqrt" {{.*}} : (!cir.double) -> !cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.double) -> !cir.double
// LLVM: define{{.*}} void @_Z13test_sqrt_f64Pdd
// LLVM: call{{.*}} double @llvm.amdgcn.sqrt.f64(double %{{.*}})
__device__ void test_sqrt_f64(double* out, double a) {
>From 1577d20d69bfad0e11a490ad321783a8df16b9e8 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 10:01:08 -0400
Subject: [PATCH 3/9] add test for bf16 type
---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 +++++++++++++++++++
1 file changed, 26 insertions(+)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.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..44bdb2c666204
--- /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_sqrt_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z14test_sqrt_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+__device__ void test_sqrt_bf16(__bf16* out, __bf16 a) {
+ *out = __builtin_amdgcn_sqrt_bf16(a);
+}
>From a39ec718288f1a92220db04dc183f56515d0db23 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 10:11:43 -0400
Subject: [PATCH 4/9] delete bf16 test
---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 -------------------
1 file changed, 26 deletions(-)
delete mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
deleted file mode 100644
index 44bdb2c666204..0000000000000
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ /dev/null
@@ -1,26 +0,0 @@
-#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_sqrt_bf16PDF16bDF16b
-// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.bf16) -> !cir.bf16
-// LLVM: define{{.*}} void @_Z14test_sqrt_bf16PDF16bDF16b
-// LLVM: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
-__device__ void test_sqrt_bf16(__bf16* out, __bf16 a) {
- *out = __builtin_amdgcn_sqrt_bf16(a);
-}
>From e1475d7cbb3cfdf500dc25001be05ed8d589fb42 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 17:30:44 -0400
Subject: [PATCH 5/9] add tests for f16 and bf16 types
---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 ++++++++
.../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 +++++++++++++++++++
2 files changed, 91 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..2e132ce291d32
--- /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_sqrt_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z14test_sqrt_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sqrt.bf16(bfloat %{{.*}})
+__device__ void test_sqrt_bf16(__bf16* out, __bf16 a) {
+ *out = __builtin_amdgcn_sqrt_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..325ae2dd97237
--- /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= --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: @_Z13test_sqrt_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+__device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
+ *out = __builtin_amdgcn_sqrth(a);
+}
>From ff3f8adb5944df0d92518f1cdfc18780dab8c06b Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 16 May 2026 15:27:06 -0400
Subject: [PATCH 6/9] add missing test prefix
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
index 325ae2dd97237..ac4da55c4105d 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
@@ -38,7 +38,7 @@
// 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= --input-file=%t.ll %s
+// 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
>From 3250fbd7171ce28044b8d9abe16cdbf6facd1184 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 16 May 2026 17:50:05 -0400
Subject: [PATCH 7/9] remove regex matching in CHECK line
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
index ac4da55c4105d..3c966d16272eb 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
@@ -59,7 +59,7 @@
// CIR-LABEL: @_Z13test_sqrt_f16PDF16_DF16_
// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
// LLVM: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
-// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}sqrt.f16(half %{{.*}})
+// LLVM: call{{.*}} half @llvm.amdgcn.sqrt.f16(half %{{.*}})
__device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
*out = __builtin_amdgcn_sqrth(a);
}
>From e5d0d603d7fa13cc0fcb4a0927320276834d7b7d Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 11 Jun 2026 17:54:30 -0400
Subject: [PATCH 8/9] switch out header for macro (__device specifier__)
---
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 4 +-
.../CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip | 8 +++
.../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 -------------------
3 files changed, 10 insertions(+), 67 deletions(-)
delete 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
index 2e132ce291d32..59b86fbf05f77 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
//===----------------------------------------------------------------------===//
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index c21359ff7a88b..96b8d26e6cf80 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -63,3 +63,11 @@
__device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
*out = __builtin_amdgcn_div_fixuph(a, b, c);
}
+
+// CIR-LABEL: @_Z13test_sqrt_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.sqrt.f16(half %{{.*}})
+__device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
+ *out = __builtin_amdgcn_sqrth(a);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
deleted file mode 100644
index 3c966d16272eb..0000000000000
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip
+++ /dev/null
@@ -1,65 +0,0 @@
-#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: @_Z13test_sqrt_f16PDF16_DF16_
-// CIR: cir.call_llvm_intrinsic "amdgcn.sqrt" {{.*}} : (!cir.f16) -> !cir.f16
-// LLVM: define{{.*}} void @_Z13test_sqrt_f16PDF16_DF16_
-// LLVM: call{{.*}} half @llvm.amdgcn.sqrt.f16(half %{{.*}})
-__device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
- *out = __builtin_amdgcn_sqrth(a);
-}
>From f17647e859311f8d621bbc68c45afc113e8c7f6a Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 13:07:37 -0400
Subject: [PATCH 9/9] switch manual codegen to 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 eabf56ce09afd..c3ca1d760be60 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -260,10 +260,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_sqrtf:
case AMDGPU::BI__builtin_amdgcn_sqrth:
case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
- mlir::Value src = emitScalarExpr(expr->getArg(0));
- return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
- "amdgcn.sqrt", src.getType(),
- mlir::ValueRange{src});
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt").getValue();
}
case AMDGPU::BI__builtin_amdgcn_rsq:
case AMDGPU::BI__builtin_amdgcn_rsqf:
More information about the cfe-commits
mailing list