[clang] [Clang] fix for heterogeneous chip's host side's MaybeODRUseExprs clear (PR #121085)

[via cfe-commits]

https://github.com/fuaq updated https://github.com/llvm/llvm-project/pull/121085

>From 77987b302bbd3a984504d161d8e8c12c4feebdba Mon Sep 17 00:00:00 2001
From: fuanqi <fuanqi at baidu.com>
Date: Tue, 24 Dec 2024 19:47:31 -0800
Subject: [PATCH] [Clang] fix for heterogeneous chip's host side's
 MaybeODRUseExprs clear

---
 clang/lib/Sema/SemaStmtAsm.cpp            |  6 ++++++
 clang/test/SemaCUDA/hostside_no_assert.cu | 13 +++++++++++++
 2 files changed, 19 insertions(+)
 create mode 100644 clang/test/SemaCUDA/hostside_no_assert.cu

diff --git a/clang/lib/Sema/SemaStmtAsm.cpp b/clang/lib/Sema/SemaStmtAsm.cpp
index cd8b98c7444eb..0d56583809edf 100644
--- a/clang/lib/Sema/SemaStmtAsm.cpp
+++ b/clang/lib/Sema/SemaStmtAsm.cpp
@@ -273,6 +273,12 @@ StmtResult Sema::ActOnGCCAsmStmt(SourceLocation AsmLoc, bool IsSimple,
                                  Expr *asmString, MultiExprArg clobbers,
                                  unsigned NumLabels,
                                  SourceLocation RParenLoc) {
+  struct _Cleaner {
+    Sema &S;
+    ~_Cleaner() {
+      S.DiscardCleanupsInEvaluationContext();
+    }
+  } _C{*this};
   unsigned NumClobbers = clobbers.size();
 
   SmallVector<TargetInfo::ConstraintInfo, 4> OutputConstraintInfos;
diff --git a/clang/test/SemaCUDA/hostside_no_assert.cu b/clang/test/SemaCUDA/hostside_no_assert.cu
new file mode 100644
index 0000000000000..9d8e0f376e9db
--- /dev/null
+++ b/clang/test/SemaCUDA/hostside_no_assert.cu
@@ -0,0 +1,13 @@
+// REQUIRES: nvptx-registered-target
+// REQUIRES: x86-registered-target
+
+// RUN: %clang_cc1 "-triple" "x86_64-unknown-linux-gnu" "-aux-triple" "nvptx64-nvidia-cuda" \
+// RUN:    "-target-cpu" "x86-64"  "-fsyntax-only" %s
+
+#include "Inputs/cuda.h"
+
+typedef __attribute__((ext_vector_type(16))) float float32x16_t;
+__device__ void test(float32x16_t& vodd) {
+  constexpr int pose = 16;
+  __asm__ __volatile__("vadd %0, %1, %2":"=&v"(vodd):"r"(pose),"v"(vodd));
+}