https://github.com/VigneshwarJ created https://github.com/llvm/llvm-project/pull/180314
…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. >From cad2ea652f4e88fb17f5718b390b19cf0a4278b7 Mon Sep 17 00:00:00 2001 From: vigneshwar jayakumar <[email protected]> Date: Fri, 6 Feb 2026 18:52:46 -0600 Subject: [PATCH] [HIP][Sema] Fix incorrect CK_NoOp for lvalue-to-rvalue conversion in 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. --- clang/lib/Sema/SemaExpr.cpp | 7 ++++--- .../SemaHIP/hip-builtin-lvalue-to-rvalue.hip | 19 +++++++++++++++++++ 2 files changed, 23 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaHIP/hip-builtin-lvalue-to-rvalue.hip 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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
