diff options
| author | Justin Holewinski <justin.holewinski@gmail.com> | 2012-07-11 15:34:55 +0000 |
|---|---|---|
| committer | Justin Holewinski <justin.holewinski@gmail.com> | 2012-07-11 15:34:55 +0000 |
| commit | 9903e94bee8fbbec6d44fe2de91ccee46f40187f (patch) | |
| tree | 28728f8d9e760280a8ebbcfe2df2adcd01ef357e | |
| parent | c852e9ff06f7089405265c53af33cf06fdc36e52 (diff) | |
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
| -rw-r--r-- | lib/Basic/Targets.cpp | 1 | ||||
| -rw-r--r-- | test/CodeGen/nvptx-inlineasm.c | 15 |
2 files changed, 16 insertions, 0 deletions
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; +} |