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 <[email protected]> Date: Tue, 12 May 2026 20:59:51 -0400 Subject: [PATCH 01/10] 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 <[email protected]> Date: Tue, 12 May 2026 22:10:55 -0400 Subject: [PATCH 02/10] 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 <[email protected]> Date: Wed, 13 May 2026 10:01:08 -0400 Subject: [PATCH 03/10] 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 <[email protected]> Date: Wed, 13 May 2026 10:11:43 -0400 Subject: [PATCH 04/10] 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 <[email protected]> Date: Wed, 13 May 2026 17:30:44 -0400 Subject: [PATCH 05/10] 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 <[email protected]> Date: Sat, 16 May 2026 15:27:06 -0400 Subject: [PATCH 06/10] 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 <[email protected]> Date: Sat, 16 May 2026 17:50:05 -0400 Subject: [PATCH 07/10] 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 <[email protected]> Date: Thu, 11 Jun 2026 17:54:30 -0400 Subject: [PATCH 08/10] 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 <[email protected]> Date: Fri, 10 Jul 2026 13:07:37 -0400 Subject: [PATCH 09/10] 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: >From 25e4ecd00eb6b422c5cd3901ee6cbec25b0e7e34 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 15 Jul 2026 10:09:01 -0400 Subject: [PATCH 10/10] Remove extra newline in test_sqrt_f16 function --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip index 0e7d2a8fa300a..6d4d14285cb57 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -87,4 +87,3 @@ __device__ void test_rsq_h(_Float16* out, _Float16 a) { __device__ void test_sqrt_f16(_Float16* out, _Float16 a) { *out = __builtin_amdgcn_sqrth(a); } - _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
