llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Steffen Larsen (steffenlarsen) <details> <summary>Changes</summary> 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. --- Full diff: https://github.com/llvm/llvm-project/pull/220576.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp (+8) - (added) clang/test/CIR/CodeGenHIP/builtin-bool-result.hip (+21) ``````````diff 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; +} `````````` </details> https://github.com/llvm/llvm-project/pull/220576 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
