diff options
| -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; +} |