https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/220576
>From d54fdf0279d4a02b35f5efa8ddf1b9c635f23417 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 1 Sep 2026 09:00:07 -0500 Subject: [PATCH 1/3] [CIR] Convert i1 intrinsic results for bool-returning builtins A target builtin with no special case is emitted through the generic intrinsic path, which returns the intrinsic's own result type. For a builtin declared to return bool that type is i1, which CIR models as !cir.int<u, 1> rather than !cir.bool. cir.if accepts only !cir.bool, so CIR builtin handling needs to convert i1 results to !cir.bool when used in this context. This is verified using __builtin_amdgcn_is_shared. --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 8 +++++++ .../CIR/CodeGenHIP/builtin-bool-result.hip | 21 +++++++++++++++++++ 2 files changed, 29 insertions(+) create mode 100644 clang/test/CIR/CodeGenHIP/builtin-bool-result.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 2bab83ddc1f17..0dc240f7c2052 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -2723,6 +2723,14 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, if (isa<cir::VoidType>(correctedReturnType)) return RValue::get(nullptr); + // A bool-returning builtin may back an intrinsic that returns i1; CIR needs + // those as !cir.bool. + if (fd && fd->getReturnType()->isBooleanType() && + mlir::isa<cir::IntType>(intrinsicRes.getType())) + intrinsicRes = cir::CastOp::create( + builder, getLoc(e->getExprLoc()), convertType(fd->getReturnType()), + cir::CastKind::int_to_bool, intrinsicRes); + return RValue::get(intrinsicRes); } diff --git a/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip b/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip new file mode 100644 index 0000000000000..7a89e105b22dc --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip @@ -0,0 +1,21 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fcuda-is-device \ +// RUN: -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s \ +// RUN: -o - | FileCheck %s + +// Checks that builtins returning an i1 can be used as bool results implicitly. + +#define __device__ __attribute__((device)) + +// CHECK-LABEL: @_Z2shPv +// CHECK: call i1 @llvm.amdgcn.is.shared(ptr +__device__ bool sh(void *p) { return __builtin_amdgcn_is_shared(p); } + +// CHECK-LABEL: @_Z3usePv +// CHECK: call i1 @llvm.amdgcn.is.shared(ptr +__device__ int use(void *p) { + if (__builtin_amdgcn_is_shared(p)) + return 1; + return 0; +} >From 674adba2e99faaf98b5b147b08982e2b277b0bba Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Thu, 3 Sep 2026 00:09:22 -0500 Subject: [PATCH 2/3] Move return type conversion to return correction and add CIR test checks Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 37 ++++++++++++++----- .../CIR/CodeGenHIP/builtin-bool-result.hip | 19 +++++++--- 2 files changed, 40 insertions(+), 16 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 0dc240f7c2052..ace56363db56b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -975,6 +975,27 @@ static mlir::Type correctIntegerSignedness(mlir::Type iitType, QualType astType, return iitType; } +/// Helper function to correct the return type for intrinsic calls. This is +/// needed because the AST FunctionDecl may have a different return type than +/// the intrinsic's IIT descriptor. For example, builtins may need their +/// signedness corrected, or a builtin may return a bool while the intrinsic +/// returns an i1. +static mlir::Type correctReturnType(mlir::Type iitType, + const FunctionDecl *funcDecl, + mlir::MLIRContext *context) { + if (!funcDecl) + return iitType; + QualType astType = funcDecl->getReturnType(); + + // Relabel the return type to cir.bool if the builtin returns a bool and + // the intrinsic returns an i1. + auto intTy = mlir::dyn_cast<cir::IntType>(iitType); + if (intTy && intTy.getWidth() == 1 && astType->isBooleanType()) + return cir::BoolType::get(context); + + return correctIntegerSignedness(iitType, astType, context); +} + static mlir::Value getCorrectedPtr(mlir::Value argValue, mlir::Type expectedTy, CIRGenBuilderTy &builder) { auto ptrType = mlir::cast<cir::PointerType>(argValue.getType()); @@ -2705,14 +2726,10 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, args.push_back(argValue); } - // Correct return type signedness based on AST return type before creating - // the call, avoiding unnecessary casts in the IR. - mlir::Type correctedReturnType = intrinsicType.getReturnType(); - if (fd) { - correctedReturnType = - correctIntegerSignedness(intrinsicType.getReturnType(), - fd->getReturnType(), &getMLIRContext()); - } + // Correct the builtin type based on the AST function declaration's return + // type, if available. + mlir::Type correctedReturnType = + correctReturnType(intrinsicType.getReturnType(), fd, &getMLIRContext()); cir::LLVMIntrinsicCallOp intrinsicCall = cir::LLVMIntrinsicCallOp::create( builder, getLoc(e->getExprLoc()), builder.getStringAttr(name), @@ -2723,8 +2740,8 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, if (isa<cir::VoidType>(correctedReturnType)) return RValue::get(nullptr); - // A bool-returning builtin may back an intrinsic that returns i1; CIR needs - // those as !cir.bool. + // A bool-returning builtin may back an intrinsic that returns a wider + // int; CIR needs those as !cir.bool too. if (fd && fd->getReturnType()->isBooleanType() && mlir::isa<cir::IntType>(intrinsicRes.getType())) intrinsicRes = cir::CastOp::create( diff --git a/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip b/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip index 7a89e105b22dc..28efe28e4489f 100644 --- a/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip +++ b/clang/test/CIR/CodeGenHIP/builtin-bool-result.hip @@ -1,19 +1,26 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fcuda-is-device \ -// RUN: -emit-llvm %s -o - | FileCheck %s +// RUN: -emit-cir %s -o - | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fcuda-is-device \ +// RUN: -emit-llvm %s -o - | FileCheck --check-prefix=LLVM %s +// // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s \ -// RUN: -o - | FileCheck %s +// RUN: -o - | FileCheck --check-prefix=LLVM %s // Checks that builtins returning an i1 can be used as bool results implicitly. #define __device__ __attribute__((device)) -// CHECK-LABEL: @_Z2shPv -// CHECK: call i1 @llvm.amdgcn.is.shared(ptr +// CIR-LABEL: cir.func {{.*}} @_Z2shPv +// CIR: cir.call_llvm_intrinsic "amdgcn.is.shared" {{.*}} -> !cir.bool +// LLVM-LABEL: @_Z2shPv +// LLVM: call i1 @llvm.amdgcn.is.shared(ptr __device__ bool sh(void *p) { return __builtin_amdgcn_is_shared(p); } -// CHECK-LABEL: @_Z3usePv -// CHECK: call i1 @llvm.amdgcn.is.shared(ptr +// CIR-LABEL: cir.func {{.*}} @_Z3usePv +// CIR: cir.call_llvm_intrinsic "amdgcn.is.shared" {{.*}} -> !cir.bool +// LLVM-LABEL: @_Z3usePv +// LLVM: call i1 @llvm.amdgcn.is.shared(ptr __device__ int use(void *p) { if (__builtin_amdgcn_is_shared(p)) return 1; >From 6720e986c9d13a754bdf2227a79e814e7a862b91 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Thu, 3 Sep 2026 01:54:03 -0500 Subject: [PATCH 3/3] Remove the redundant conversion Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index ace56363db56b..5bd66c0be941e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -2740,14 +2740,6 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, if (isa<cir::VoidType>(correctedReturnType)) return RValue::get(nullptr); - // A bool-returning builtin may back an intrinsic that returns a wider - // int; CIR needs those as !cir.bool too. - if (fd && fd->getReturnType()->isBooleanType() && - mlir::isa<cir::IntType>(intrinsicRes.getType())) - intrinsicRes = cir::CastOp::create( - builder, getLoc(e->getExprLoc()), convertType(fd->getReturnType()), - cir::CastKind::int_to_bool, intrinsicRes); - return RValue::get(intrinsicRes); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
