[Clang] fix for heterogeneous chip's host side's MaybeODRUseExprs clear

Author: fuaq

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;

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));
+}