llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Vigneshwar Jayakumar (VigneshwarJ) <details> <summary>Changes</summary> …builtin args The HIP implicit address space cast for builtin pointer arguments used CK_NoOp to convert lvalue args to rvalues. This caused an assertion failure in LifetimeSafety analysis: Assertion `Dst->getLength() == Src->getLength()` failed in FactsGenerator::flow() in some cases. Use DefaultLvalueConversion which correctly emits CK_LValueToRValue. --- Full diff: https://github.com/llvm/llvm-project/pull/180314.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaExpr.cpp (+4-3) - (added) clang/test/SemaHIP/hip-builtin-lvalue-to-rvalue.hip (+19) ``````````diff diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 77d3c4ef63f1f..f88563d834c26 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6823,9 +6823,10 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc, // First, ensure that the Arg is an RValue. if (ArgExprs[Idx]->isGLValue()) { - ArgExprs[Idx] = ImplicitCastExpr::Create( - Context, ArgExprs[Idx]->getType(), CK_NoOp, ArgExprs[Idx], - nullptr, VK_PRValue, FPOptionsOverride()); + ExprResult Res = DefaultLvalueConversion(ArgExprs[Idx]); + if (Res.isInvalid()) + return ExprError(); + ArgExprs[Idx] = Res.get(); } // Construct a new arg type with address space of Param diff --git a/clang/test/SemaHIP/hip-builtin-lvalue-to-rvalue.hip b/clang/test/SemaHIP/hip-builtin-lvalue-to-rvalue.hip new file mode 100644 index 0000000000000..b3238d7b29d3e --- /dev/null +++ b/clang/test/SemaHIP/hip-builtin-lvalue-to-rvalue.hip @@ -0,0 +1,19 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx1200 -fsyntax-only -fcuda-is-device -ast-dump %s 2>&1 | FileCheck %s + +// expected-no-diagnostics + +// Test that HIP builtins correctly perform lvalue-to-rvalue conversion on pointer arg. + +#define __device__ __attribute__((device)) + +typedef __attribute__((__vector_size__(8 * sizeof(__fp16)))) __fp16 fp16x8_t; + +// CHECK: ImplicitCastExpr {{.*}} <AddressSpaceConversion> +// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue> +// CHECK-NEXT: DeclRefExpr {{.*}} lvalue Var {{.*}} 'glb_ptr' +__device__ fp16x8_t test_global_load_transpose(const _Float16 *in_ptr) { + fp16x8_t *glb_ptr = reinterpret_cast<fp16x8_t *>( + reinterpret_cast<__UINTPTR_TYPE__>(in_ptr)); + return __builtin_amdgcn_global_load_tr_b128_v8f16(glb_ptr); +} `````````` </details> https://github.com/llvm/llvm-project/pull/180314 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
