[cfe-commits] r160050 - in /cfe/trunk: lib/Basic/Targets.cpp test/CodeGen/nvptx-inlineasm.c

[Justin Holewinski]

Author: jholewinski
Date: Wed Jul 11 10:34:55 2012
New Revision: 160050

URL: http://llvm.org/viewvc/llvm-project?rev=160050&view=rev
Log:
Fix handling of curly braces in NVPTX inline asm

Fixes bug 13322

Patch by Dmitry Mikushin

Added:
    cfe/trunk/test/CodeGen/nvptx-inlineasm.c
Modified:
    cfe/trunk/lib/Basic/Targets.cpp

Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=160050&r1=160049&r2=160050&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Wed Jul 11 10:34:55 2012
@@ -1049,6 +1049,7 @@
       AddrSpaceMap = &NVPTXAddrSpaceMap;
       // Define available target features
       // These must be defined in sorted order!
+      NoAsmVariants = true;
     }
     virtual void getTargetDefines(const LangOptions &Opts,
                                   MacroBuilder &Builder) const {

Added: cfe/trunk/test/CodeGen/nvptx-inlineasm.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/nvptx-inlineasm.c?rev=160050&view=auto
==============================================================================
--- cfe/trunk/test/CodeGen/nvptx-inlineasm.c (added)
+++ cfe/trunk/test/CodeGen/nvptx-inlineasm.c Wed Jul 11 10:34:55 2012
@@ -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;
+}