https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198087
>From 26de77918be5edef3efc5180064e45b298024b3f Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 16 May 2026 09:23:03 -0400 Subject: [PATCH 1/5] add ldexp builtin --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 22 ++++-- .../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 67 +++++++++++++++++++ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 +++++ 3 files changed, 100 insertions(+), 5 deletions(-) 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..d7fb6251a696e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -320,12 +320,24 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_ldexp: - case AMDGPU::BI__builtin_amdgcn_ldexpf: + case AMDGPU::BI__builtin_amdgcn_ldexpf: { + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp", + src0.getType(), + mlir::ValueRange{src0, src1}); + } case AMDGPU::BI__builtin_amdgcn_ldexph: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + // The raw instruction has a different behavior for out of bounds exponent + // values (implicit truncation instead of saturate to short_min/short_max). + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + mlir::Value truncSrc1 = cir::CastOp::create( + builder, getLoc(expr->getExprLoc()), builder.getSInt16Ty(), + cir::CastKind::integral, src1); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp", + src0.getType(), + mlir::ValueRange{src0, truncSrc1}); } case AMDGPU::BI__builtin_amdgcn_frexp_mant: case AMDGPU::BI__builtin_amdgcn_frexp_mantf: 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..dcde886c205dc --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip @@ -0,0 +1,67 @@ +#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: @_Z14test_ldexp_f16PDF16_DF16_i +// CIR: [[TRUNC:%.*]] = cir.cast integral {{.*}} : !s32i -> !s16i +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}}, [[TRUNC]] : (!cir.f16, !s16i) -> !cir.f16 +// LLVM: define{{.*}} void @_Z14test_ldexp_f16PDF16_DF16_i +// LLVM: [[TRUNC:%.*]] = trunc i32 {{.*}} to i16 +// LLVM: call{{.*}} half @llvm.ldexp.f16.i16(half %{{.*}}, i16 [[TRUNC]]) +__device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) { + *out = __builtin_amdgcn_ldexph(a, b); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index ca708bca8587b..07ed93dd4ded3 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: @_Z14test_ldexp_f32Pffi +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.float, !s32i) -> !cir.float +// LLVM: define{{.*}} void @_Z14test_ldexp_f32Pffi +// LLVM: call{{.*}} float @llvm.ldexp.f32.i32(float %{{.*}}, i32 %{{.*}}) +__device__ void test_ldexp_f32(float* out, float a, int b) { + *out = __builtin_amdgcn_ldexpf(a, b); +} + +// CIR-LABEL: @_Z14test_ldexp_f64Pddi +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.double, !s32i) -> !cir.double +// LLVM: define{{.*}} void @_Z14test_ldexp_f64Pddi +// LLVM: call{{.*}} double @llvm.ldexp.f64.i32(double %{{.*}}, i32 %{{.*}}) +__device__ void test_ldexp_f64(double* out, double a, int b) { + *out = __builtin_amdgcn_ldexp(a, b); +} >From 2f4b8abe936581d661d8765d0676b4737501aa6a Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 12:35:07 -0400 Subject: [PATCH 2/5] switch out manual codegen for function --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 7 +- .../CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip | 10 +++ .../CIR/CodeGenHIP/builtins-amdgcn-vi.hip | 67 ------------------- 3 files changed, 12 insertions(+), 72 deletions(-) delete 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 d7fb6251a696e..c5426a18d74dd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -321,11 +321,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_ldexp: case AMDGPU::BI__builtin_amdgcn_ldexpf: { - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - mlir::Value src1 = emitScalarExpr(expr->getArg(1)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp", - src0.getType(), - mlir::ValueRange{src0, src1}); + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") + .getValue(); } case AMDGPU::BI__builtin_amdgcn_ldexph: { // The raw instruction has a different behavior for out of bounds exponent diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip index c21359ff7a88b..bc672d07bcbbc 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -63,3 +63,13 @@ __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) { *out = __builtin_amdgcn_div_fixuph(a, b, c); } + +// CIR-LABEL: @_Z14test_ldexp_f16PDF16_DF16_i +// CIR: [[TRUNC:%.*]] = cir.cast integral {{.*}} : !s32i -> !s16i +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}}, [[TRUNC]] : (!cir.f16, !s16i) -> !cir.f16 +// LLVM: define{{.*}} void @_Z14test_ldexp_f16PDF16_DF16_i +// LLVM: [[TRUNC:%.*]] = trunc i32 {{.*}} to i16 +// LLVM: call{{.*}} half @llvm.ldexp.f16.i16(half %{{.*}}, i16 [[TRUNC]]) +__device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) { + *out = __builtin_amdgcn_ldexph(a, b); +} 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 dcde886c205dc..0000000000000 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi.hip +++ /dev/null @@ -1,67 +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: @_Z14test_ldexp_f16PDF16_DF16_i -// CIR: [[TRUNC:%.*]] = cir.cast integral {{.*}} : !s32i -> !s16i -// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}}, [[TRUNC]] : (!cir.f16, !s16i) -> !cir.f16 -// LLVM: define{{.*}} void @_Z14test_ldexp_f16PDF16_DF16_i -// LLVM: [[TRUNC:%.*]] = trunc i32 {{.*}} to i16 -// LLVM: call{{.*}} half @llvm.ldexp.f16.i16(half %{{.*}}, i16 [[TRUNC]]) -__device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) { - *out = __builtin_amdgcn_ldexph(a, b); -} >From 9a7d6611f30298820d795a7c034cb82cb9e4ab4a Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 12:56:33 -0400 Subject: [PATCH 3/5] correct builtin name --- 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 c5426a18d74dd..c41a3d6d42ea7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -321,7 +321,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_ldexp: case AMDGPU::BI__builtin_amdgcn_ldexpf: { - return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ldexp") .getValue(); } case AMDGPU::BI__builtin_amdgcn_ldexph: { >From f1554732cd2727ccbae54b7af368b024e6e4438d Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 17:08:11 -0400 Subject: [PATCH 4/5] fix format --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index c41a3d6d42ea7..4be3f96af95b7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -321,8 +321,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_ldexp: case AMDGPU::BI__builtin_amdgcn_ldexpf: { - return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ldexp") - .getValue(); + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ldexp").getValue(); } case AMDGPU::BI__builtin_amdgcn_ldexph: { // The raw instruction has a different behavior for out of bounds exponent >From 3ab9a27b8f0cca16a438aebda87c58b3140856f4 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 20:44:11 -0400 Subject: [PATCH 5/5] switch from function to manual codegen --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 4be3f96af95b7..d7fb6251a696e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -321,7 +321,11 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } case AMDGPU::BI__builtin_amdgcn_ldexp: case AMDGPU::BI__builtin_amdgcn_ldexpf: { - return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ldexp").getValue(); + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp", + src0.getType(), + mlir::ValueRange{src0, src1}); } case AMDGPU::BI__builtin_amdgcn_ldexph: { // The raw instruction has a different behavior for out of bounds exponent _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
