Fix handling of curly braces in NVPTX inline asm

Fixes bug 13322

Patch by Dmitry Mikushin

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@160050 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index 51651932508e55fff0b68b76a23b5b3c8c28df30..3df59c35d12169edd25ca5b02596cea703673f3f 100644 (file)
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -1049,6 +1049,7 @@ namespace {
       AddrSpaceMap = &NVPTXAddrSpaceMap;
       // Define available target features
       // These must be defined in sorted order!
+      NoAsmVariants = true;
     }
     virtual void getTargetDefines(const LangOptions &Opts,
                                   MacroBuilder &Builder) const {

diff --git a/test/CodeGen/nvptx-inlineasm.c b/test/CodeGen/nvptx-inlineasm.c
new file mode 100644 (file)
index 0000000..860b50f
--- /dev/null
+++ b/test/CodeGen/nvptx-inlineasm.c
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s
+
+int bar(int a) {
+  int result;
+  // CHECK: call i32 asm sideeffect "{ {{.*}}
+  asm __volatile__ ("{ \n\t"
+                    ".reg .pred \t%%p1; \n\t"
+                    ".reg .pred \t%%p2; \n\t"
+                    "setp.ne.u32 \t%%p1, %1, 0; \n\t"
+                    "vote.any.pred \t%%p2, %%p1; \n\t"
+                    "selp.s32 \t%0, 1, 0, %%p2; \n\t"
+                    "}" : "=r"(result) : "r"(a));
+  return result;
+}