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