[clang] [HIP][Sema] Fix incorrect CK_NoOp for lvalue-to-rvalue conversion in … (PR #180314)

[via cfe-commits]

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