https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197852
>From dcaf439dbe833dd33ed2b9ae052dedf8130a1c2c Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 14 May 2026 20:42:44 -0400 Subject: [PATCH 01/10] add amdgcn tanh builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 7 +- .../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 ++++++++ .../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 +++++++++++++++++++ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 14 +++- 4 files changed, 105 insertions(+), 7 deletions(-) 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/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 0a7ba0c194400..3a5f66be24920 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -386,10 +386,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_tanhf: case AMDGPU::BI__builtin_amdgcn_tanhh: case AMDGPU::BI__builtin_amdgcn_tanh_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.tanh", + src.getType(), mlir::ValueRange{src}); } case AMDGPU::BI__builtin_amdgcn_uicmp: case AMDGPU::BI__builtin_amdgcn_uicmpl: 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..aca9cd45ee234 --- /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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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_tanh_bf16PDF16bDF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.bf16) -> !cir.bf16 +// LLVM: define{{.*}} void @_Z14test_tanh_bf16PDF16bDF16b +// LLVM: call{{.*}} bfloat @llvm.amdgcn.tanh.bf16(bfloat %{{.*}}) +__device__ void test_tanh_bf16(__bf16* out, __bf16 a) { + *out = __builtin_amdgcn_tanh_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..8cbd3c48a58d5 --- /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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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_tanh_f16PDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_ +// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}tanh.f16(half %{{.*}}) +__device__ void test_tanh_f16(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_tanhh(a); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index ca708bca8587b..29d41f5560686 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -1,14 +1,14 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ -// RUN: -target-cpu tahiti -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: -target-cpu tahiti -target-feature +tanh-insts -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 tahiti -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: -target-cpu tahiti --target-feature +tanh-insts fcuda-is-device -emit-llvm %s -o %t-cir.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ -// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: -target-cpu tahiti -target-feature +tanh-insts -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s #define __device__ __attribute__((device)) @@ -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: @_Z14test_tanhf_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.tanh.f32(float %{{.*}}) +__device__ void test_tanhf_f32(float* out, float a) { + *out = __builtin_amdgcn_tanhf(a); +} >From de0508a7cb25ba7ea6c955eb8d559dcce7e1d32a Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 14 May 2026 20:50:20 -0400 Subject: [PATCH 02/10] fix style --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 3a5f66be24920..6ca59612fede4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -387,8 +387,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_tanhh: case AMDGPU::BI__builtin_amdgcn_tanh_bf16: { mlir::Value src = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "amdgcn.tanh", - src.getType(), mlir::ValueRange{src}); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.tanh", src.getType(), + mlir::ValueRange{src}); } case AMDGPU::BI__builtin_amdgcn_uicmp: case AMDGPU::BI__builtin_amdgcn_uicmpl: >From a5ef1cb0bdec0f3f9e92adcadee716c10cb8f44f Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 15 May 2026 14:47:16 -0400 Subject: [PATCH 03/10] isolated tanh test --- .../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 26 -------- .../CIR/CodeGenHIP/builtins-amdgcn-tanh.hip | 38 +++++++++++ .../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 65 ------------------- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 6 +- 4 files changed, 41 insertions(+), 94 deletions(-) delete mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip 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 deleted file mode 100644 index aca9cd45ee234..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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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_tanh_bf16PDF16bDF16b -// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.bf16) -> !cir.bf16 -// LLVM: define{{.*}} void @_Z14test_tanh_bf16PDF16bDF16b -// LLVM: call{{.*}} bfloat @llvm.amdgcn.tanh.bf16(bfloat %{{.*}}) -__device__ void test_tanh_bf16(__bf16* out, __bf16 a) { - *out = __builtin_amdgcn_tanh_bf16(a); -} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip new file mode 100644 index 0000000000000..5024c0d7a9d81 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip @@ -0,0 +1,38 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir -target-cpu gfx1250 \ +// RUN: -target-feature +tanh-insts -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 -target-cpu gfx1250 \ +// RUN: -target-feature +tanh-insts -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 -target-cpu gfx1250 \ +// RUN: -target-feature +tanh-insts -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// CIR-LABEL: @_Z14test_tanhf_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.tanh.f32(float %{{.*}}) +__device__ void test_tanhf_f32(float* out, float a) { + *out = __builtin_amdgcn_tanhf(a); +} + +// CIR-LABEL: @_Z13test_tanh_f16PDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_ +// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}tanh.f16(half %{{.*}}) +__device__ void test_tanh_f16(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_tanhh(a); +} + +// CIR-LABEL: @_Z14test_tanh_bf16PDF16bDF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.bf16) -> !cir.bf16 +// LLVM: define{{.*}} void @_Z14test_tanh_bf16PDF16bDF16b +// LLVM: call{{.*}} bfloat @llvm.amdgcn.tanh.bf16(bfloat %{{.*}}) +__device__ void test_tanh_bf16(__bf16* out, __bf16 a) { + *out = __builtin_amdgcn_tanh_bf16(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 8cbd3c48a58d5..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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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 -target-feature +tanh-insts -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_tanh_f16PDF16_DF16_ -// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16 -// LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_ -// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}tanh.f16(half %{{.*}}) -__device__ void test_tanh_f16(_Float16* out, _Float16 a) { - *out = __builtin_amdgcn_tanhh(a); -} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 29d41f5560686..904f571ddcb41 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -1,14 +1,14 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ -// RUN: -target-cpu tahiti -target-feature +tanh-insts -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: -target-cpu tahiti -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 tahiti --target-feature +tanh-insts fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t-cir.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ -// RUN: -target-cpu tahiti -target-feature +tanh-insts -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s #define __device__ __attribute__((device)) >From 43e130a4e5397d5dfe67974d2ef944355f47d609 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 15 May 2026 16:34:52 -0400 Subject: [PATCH 04/10] put tests in gfx1250 test --- ...gcn-tanh.hip => builtins-amdgcn-gfx1250.hip} | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) rename clang/test/CIR/CodeGenHIP/{builtins-amdgcn-tanh.hip => builtins-amdgcn-gfx1250.hip} (73%) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip similarity index 73% rename from clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip rename to clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index 5024c0d7a9d81..cc60b266fe78a 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-tanh.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -1,18 +1,23 @@ #include "../CodeGenCUDA/Inputs/cuda.h" // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir -target-cpu gfx1250 \ -// RUN: -target-feature +tanh-insts -fcuda-is-device -emit-cir %s -o %t.cir +// 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 -target-cpu gfx1250 \ -// RUN: -target-feature +tanh-insts -fcuda-is-device -emit-llvm %s -o %t.ll +// 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 -target-cpu gfx1250 \ -// RUN: -target-feature +tanh-insts -fcuda-is-device -emit-llvm %s -o %t.ll +// 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_tanhf_f32Pff // CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) -> !cir.float // LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff >From 8eae0aa47f4f8c5cfdf8967fc2eea90ed0d14df0 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 15 May 2026 16:36:27 -0400 Subject: [PATCH 05/10] remove unneccessary space in test --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index cc60b266fe78a..77a97eaf0433b 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -13,7 +13,6 @@ // 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 //===----------------------------------------------------------------------===// >From 3efdc4cad6d7ff450055f7f29cd6bfe822c29cb2 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:13:06 -0400 Subject: [PATCH 06/10] 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 6ca59612fede4..e9a7d1d44335e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -386,10 +386,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_tanhf: case AMDGPU::BI__builtin_amdgcn_tanhh: case AMDGPU::BI__builtin_amdgcn_tanh_bf16: { - mlir::Value src = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.tanh", src.getType(), - mlir::ValueRange{src}); + return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.tanh").getValue(); } case AMDGPU::BI__builtin_amdgcn_uicmp: case AMDGPU::BI__builtin_amdgcn_uicmpl: >From dca72e2f66528f5d18085f8d22643845e295423c Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:13:54 -0400 Subject: [PATCH 07/10] switch header for macro in test --- 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 77a97eaf0433b..a981a250757c3 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 4856582044c63fc2cc0f87ea7747bfeeeadf7baf Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:17:09 -0400 Subject: [PATCH 08/10] remove f32 test from wrong file --- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 -------- 1 file changed, 8 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 904f571ddcb41..ca708bca8587b 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -111,11 +111,3 @@ __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: @_Z14test_tanhf_f32Pff -// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) -> !cir.float -// LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff -// LLVM: call{{.*}} float @llvm.amdgcn.tanh.f32(float %{{.*}}) -__device__ void test_tanhf_f32(float* out, float a) { - *out = __builtin_amdgcn_tanhf(a); -} >From 70d910fe87d24be50870f9ea0f7c13b93dfec158 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 16:54:23 -0400 Subject: [PATCH 09/10] correct number of args --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index e9a7d1d44335e..8ad8cbcef051e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -386,7 +386,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_tanhf: case AMDGPU::BI__builtin_amdgcn_tanhh: case AMDGPU::BI__builtin_amdgcn_tanh_bf16: { - return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.tanh").getValue(); + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue(); } case AMDGPU::BI__builtin_amdgcn_uicmp: case AMDGPU::BI__builtin_amdgcn_uicmpl: >From 896128b1c697dfc68f7fbeab3c824990c91d2838 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 18 Jul 2026 13:13:47 -0400 Subject: [PATCH 10/10] Fix LLVM check for tanh intrinsic in builtins test --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index a981a250757c3..8a6a0f392e581 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -28,7 +28,7 @@ __device__ void test_tanhf_f32(float* out, float a) { // CIR-LABEL: @_Z13test_tanh_f16PDF16_DF16_ // CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16 // LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_ -// LLVM: call{{.*}} half @llvm.{{((amdgcn.){0,1})}}tanh.f16(half %{{.*}}) +// LLVM: call{{.*}} half @llvm.amdgcn.tanh.f16(half %{{.*}}) __device__ void test_tanh_f16(_Float16* out, _Float16 a) { *out = __builtin_amdgcn_tanhh(a); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
