From: Justin Holewinski <justin.holewinski@gmail.com>
Date: Wed, 11 Jul 2012 15:34:55 +0000 (+0000)
Subject: Fix handling of curly braces in NVPTX inline asm
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=9903e94bee8fbbec6d44fe2de91ccee46f40187f;p=clang

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 5165193250..3df59c35d1 100644
--- 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
index 0000000000..860b50ff58
--- /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;
+}