diff options
31 files changed, 165 insertions, 90 deletions
diff --git a/include/llvm/ADT/Triple.h b/include/llvm/ADT/Triple.h index 4cfad31..2659bce 100644 --- a/include/llvm/ADT/Triple.h +++ b/include/llvm/ADT/Triple.h @@ -64,7 +64,8 @@ public: x86_64, // X86-64: amd64, x86_64 xcore, // XCore: xcore mblaze, // MBlaze: mblaze - ptx, // PTX: ptx + ptx32, // PTX: ptx (32-bit) + ptx64, // PTX: ptx (64-bit) InvalidArch }; diff --git a/include/llvm/IntrinsicsPTX.td b/include/llvm/IntrinsicsPTX.td index 01241fe..28379c9 100644 --- a/include/llvm/IntrinsicsPTX.td +++ b/include/llvm/IntrinsicsPTX.td @@ -12,53 +12,81 @@ //===----------------------------------------------------------------------===// let TargetPrefix = "ptx" in { - multiclass PTXReadSpecialRegisterIntrinsic_v4i32 { + multiclass PTXReadSpecialRegisterIntrinsic_v4i32<string prefix> { // FIXME: Do we need the 128-bit integer type version? // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem]>; // FIXME: Enable this once v4i32 support is enabled in back-end. // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem]>; - def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; - def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; - def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; - def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; + def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<!strconcat(prefix, "_x")>; + def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<!strconcat(prefix, "_y")>; + def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<!strconcat(prefix, "_z")>; + def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<!strconcat(prefix, "_w")>; } - class PTXReadSpecialRegisterIntrinsic_r32 - : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; + class PTXReadSpecialRegisterIntrinsic_r32<string name> + : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin<name>; - class PTXReadSpecialRegisterIntrinsic_r64 - : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>; + class PTXReadSpecialRegisterIntrinsic_r64<string name> + : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>, + GCCBuiltin<name>; } -defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic_v4i32; -defm int_ptx_read_ntid : PTXReadSpecialRegisterIntrinsic_v4i32; +defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic_v4i32 + <"__builtin_ptx_read_tid">; +defm int_ptx_read_ntid : PTXReadSpecialRegisterIntrinsic_v4i32 + <"__builtin_ptx_read_ntid">; -def int_ptx_read_laneid : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_warpid : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_nwarpid : PTXReadSpecialRegisterIntrinsic_r32; +def int_ptx_read_laneid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_laneid">; +def int_ptx_read_warpid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_warpid">; +def int_ptx_read_nwarpid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_nwarpid">; -defm int_ptx_read_ctaid : PTXReadSpecialRegisterIntrinsic_v4i32; -defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32; +defm int_ptx_read_ctaid : PTXReadSpecialRegisterIntrinsic_v4i32 + <"__builtin_ptx_read_ctaid">; +defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32 + <"__builtin_ptx_read_nctaid">; -def int_ptx_read_smid : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_nsmid : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_gridid : PTXReadSpecialRegisterIntrinsic_r32; +def int_ptx_read_smid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_smid">; +def int_ptx_read_nsmid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_nsmid">; +def int_ptx_read_gridid : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_gridid">; -def int_ptx_read_lanemask_eq : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_lanemask_le : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_lanemask_lt : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_lanemask_ge : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_lanemask_gt : PTXReadSpecialRegisterIntrinsic_r32; +def int_ptx_read_lanemask_eq : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_lanemask_eq">; +def int_ptx_read_lanemask_le : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_lanemask_le">; +def int_ptx_read_lanemask_lt : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_lanemask_lt">; +def int_ptx_read_lanemask_ge : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_lanemask_ge">; +def int_ptx_read_lanemask_gt : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_lanemask_gt">; -def int_ptx_read_clock : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64; +def int_ptx_read_clock : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_clock">; +def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64 + <"__builtin_ptx_read_clock64">; -def int_ptx_read_pm0 : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32; -def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32; +def int_ptx_read_pm0 : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_pm0">; +def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_pm1">; +def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_pm2">; +def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32 + <"__builtin_ptx_read_pm3">; let TargetPrefix = "ptx" in - def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>; + def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>, + GCCBuiltin<"__builtin_ptx_bar_sync">; diff --git a/lib/Support/Triple.cpp b/lib/Support/Triple.cpp index ad93121..dbdb303 100644 --- a/lib/Support/Triple.cpp +++ b/lib/Support/Triple.cpp @@ -41,7 +41,8 @@ const char *Triple::getArchTypeName(ArchType Kind) { case x86_64: return "x86_64"; case xcore: return "xcore"; case mblaze: return "mblaze"; - case ptx: return "ptx"; + case ptx32: return "ptx32"; + case ptx64: return "ptx64"; } return "<invalid>"; @@ -74,7 +75,8 @@ const char *Triple::getArchTypePrefix(ArchType Kind) { case xcore: return "xcore"; - case ptx: return "ptx"; + case ptx32: return "ptx"; + case ptx64: return "ptx"; } } @@ -165,8 +167,10 @@ Triple::ArchType Triple::getArchTypeForLLVMName(StringRef Name) { return x86_64; if (Name == "xcore") return xcore; - if (Name == "ptx") - return ptx; + if (Name == "ptx32") + return ptx32; + if (Name == "ptx64") + return ptx64; return UnknownArch; } @@ -205,8 +209,10 @@ Triple::ArchType Triple::getArchTypeForDarwinArchName(StringRef Str) { Str == "armv6" || Str == "armv7") return Triple::arm; - if (Str == "ptx") - return Triple::ptx; + if (Str == "ptx32") + return Triple::ptx32; + if (Str == "ptx64") + return Triple::ptx64; return Triple::UnknownArch; } @@ -238,8 +244,10 @@ const char *Triple::getArchNameForAssembler() { return "armv6"; if (Str == "armv7" || Str == "thumbv7") return "armv7"; - if (Str == "ptx") - return "ptx"; + if (Str == "ptx32") + return "ptx32"; + if (Str == "ptx64") + return "ptx64"; return NULL; } @@ -288,8 +296,10 @@ Triple::ArchType Triple::ParseArch(StringRef ArchName) { return tce; else if (ArchName == "xcore") return xcore; - else if (ArchName == "ptx") - return ptx; + else if (ArchName == "ptx32") + return ptx32; + else if (ArchName == "ptx64") + return ptx64; else return UnknownArch; } diff --git a/lib/Target/PTX/PTX.h b/lib/Target/PTX/PTX.h index 49045cd..ec2be92 100644 --- a/lib/Target/PTX/PTX.h +++ b/lib/Target/PTX/PTX.h @@ -42,7 +42,8 @@ namespace llvm { FunctionPass *createPTXMFInfoExtract(PTXTargetMachine &TM, CodeGenOpt::Level OptLevel); - extern Target ThePTXTarget; + extern Target ThePTX32Target; + extern Target ThePTX64Target; } // namespace llvm; // Defines symbolic names for PTX registers. diff --git a/lib/Target/PTX/PTX.td b/lib/Target/PTX/PTX.td index dbc6f57..ae8326e 100644 --- a/lib/Target/PTX/PTX.td +++ b/lib/Target/PTX/PTX.td @@ -24,9 +24,6 @@ include "llvm/Target/Target.td" def FeatureDouble : SubtargetFeature<"double", "SupportsDouble", "true", "Do not demote .f64 to .f32">; -def Feature64Bit : SubtargetFeature<"64bit", "Use64BitAddresses", "true", - "Use 64-bit integer types for addresses.">; - //===- PTX Version --------------------------------------------------------===// def FeaturePTX20 : SubtargetFeature<"ptx20", "PTXVersion", "PTX_VERSION_2_0", diff --git a/lib/Target/PTX/PTXAsmPrinter.cpp b/lib/Target/PTX/PTXAsmPrinter.cpp index 27c9605..3363c73 100644 --- a/lib/Target/PTX/PTXAsmPrinter.cpp +++ b/lib/Target/PTX/PTXAsmPrinter.cpp @@ -447,5 +447,6 @@ printPredicateOperand(const MachineInstr *MI, raw_ostream &O) { // Force static initialization. extern "C" void LLVMInitializePTXAsmPrinter() { - RegisterAsmPrinter<PTXAsmPrinter> X(ThePTXTarget); + RegisterAsmPrinter<PTXAsmPrinter> X(ThePTX32Target); + RegisterAsmPrinter<PTXAsmPrinter> Y(ThePTX64Target); } diff --git a/lib/Target/PTX/PTXInstrInfo.td b/lib/Target/PTX/PTXInstrInfo.td index 972002c..c124c03 100644 --- a/lib/Target/PTX/PTXInstrInfo.td +++ b/lib/Target/PTX/PTXInstrInfo.td @@ -22,8 +22,8 @@ include "PTXInstrFormats.td" //===----------------------------------------------------------------------===// // Addressing -def Use32BitAddresses : Predicate<"!getSubtarget().use64BitAddresses()">; -def Use64BitAddresses : Predicate<"getSubtarget().use64BitAddresses()">; +def Use32BitAddresses : Predicate<"!getSubtarget().is64Bit()">; +def Use64BitAddresses : Predicate<"getSubtarget().is64Bit()">; // Shader Model Support def SupportsSM13 : Predicate<"getSubtarget().supportsSM13()">; diff --git a/lib/Target/PTX/PTXSubtarget.cpp b/lib/Target/PTX/PTXSubtarget.cpp index 527622d..a224f2b 100644 --- a/lib/Target/PTX/PTXSubtarget.cpp +++ b/lib/Target/PTX/PTXSubtarget.cpp @@ -16,11 +16,12 @@ using namespace llvm; -PTXSubtarget::PTXSubtarget(const std::string &TT, const std::string &FS) +PTXSubtarget::PTXSubtarget(const std::string &TT, const std::string &FS, + bool is64Bit) : PTXShaderModel(PTX_SM_1_0), PTXVersion(PTX_VERSION_2_0), SupportsDouble(false), - Use64BitAddresses(false) { + Is64Bit(is64Bit) { std::string TARGET = "generic"; ParseSubtargetFeatures(FS, TARGET); } diff --git a/lib/Target/PTX/PTXSubtarget.h b/lib/Target/PTX/PTXSubtarget.h index 57cd43d..47d9842 100644 --- a/lib/Target/PTX/PTXSubtarget.h +++ b/lib/Target/PTX/PTXSubtarget.h @@ -50,10 +50,10 @@ namespace llvm { bool SupportsDouble; // Use .u64 instead of .u32 for addresses. - bool Use64BitAddresses; + bool Is64Bit; public: - PTXSubtarget(const std::string &TT, const std::string &FS); + PTXSubtarget(const std::string &TT, const std::string &FS, bool is64Bit); std::string getTargetString() const; @@ -61,7 +61,7 @@ namespace llvm { bool supportsDouble() const { return SupportsDouble; } - bool use64BitAddresses() const { return Use64BitAddresses; } + bool is64Bit() const { return Is64Bit; } bool supportsSM13() const { return PTXShaderModel >= PTX_SM_1_3; } diff --git a/lib/Target/PTX/PTXTargetMachine.cpp b/lib/Target/PTX/PTXTargetMachine.cpp index 4701a94..78a7b0d 100644 --- a/lib/Target/PTX/PTXTargetMachine.cpp +++ b/lib/Target/PTX/PTXTargetMachine.cpp @@ -30,9 +30,15 @@ namespace llvm { } extern "C" void LLVMInitializePTXTarget() { - RegisterTargetMachine<PTXTargetMachine> X(ThePTXTarget); - RegisterAsmInfo<PTXMCAsmInfo> Y(ThePTXTarget); - TargetRegistry::RegisterAsmStreamer(ThePTXTarget, createPTXAsmStreamer); + + RegisterTargetMachine<PTX32TargetMachine> X(ThePTX32Target); + RegisterTargetMachine<PTX64TargetMachine> Y(ThePTX64Target); + + RegisterAsmInfo<PTXMCAsmInfo> Z(ThePTX32Target); + RegisterAsmInfo<PTXMCAsmInfo> W(ThePTX64Target); + + TargetRegistry::RegisterAsmStreamer(ThePTX32Target, createPTXAsmStreamer); + TargetRegistry::RegisterAsmStreamer(ThePTX64Target, createPTXAsmStreamer); } namespace { @@ -45,18 +51,28 @@ namespace { // DataLayout and FrameLowering are filled with dummy data PTXTargetMachine::PTXTargetMachine(const Target &T, const std::string &TT, - const std::string &FS) + const std::string &FS, + bool is64Bit) : LLVMTargetMachine(T, TT), - // FIXME: This feels like a dirty hack, but Subtarget does not appear to be - // initialized at this point, and we need to finish initialization of - // DataLayout. - DataLayout((FS.find("64bit") != FS.npos) ? DataLayout64 : DataLayout32), - Subtarget(TT, FS), + DataLayout(is64Bit ? DataLayout64 : DataLayout32), + Subtarget(TT, FS, is64Bit), FrameLowering(Subtarget), InstrInfo(*this), TLInfo(*this) { } +PTX32TargetMachine::PTX32TargetMachine(const Target &T, + const std::string& TT, + const std::string& FS) + : PTXTargetMachine(T, TT, FS, false) { +} + +PTX64TargetMachine::PTX64TargetMachine(const Target &T, + const std::string& TT, + const std::string& FS) + : PTXTargetMachine(T, TT, FS, true) { +} + bool PTXTargetMachine::addInstSelector(PassManagerBase &PM, CodeGenOpt::Level OptLevel) { PM.add(createPTXISelDag(*this, OptLevel)); diff --git a/lib/Target/PTX/PTXTargetMachine.h b/lib/Target/PTX/PTXTargetMachine.h index a5dba53..149be8e 100644 --- a/lib/Target/PTX/PTXTargetMachine.h +++ b/lib/Target/PTX/PTXTargetMachine.h @@ -33,7 +33,7 @@ class PTXTargetMachine : public LLVMTargetMachine { public: PTXTargetMachine(const Target &T, const std::string &TT, - const std::string &FS); + const std::string &FS, bool is64Bit); virtual const TargetData *getTargetData() const { return &DataLayout; } @@ -55,6 +55,22 @@ class PTXTargetMachine : public LLVMTargetMachine { virtual bool addPostRegAlloc(PassManagerBase &PM, CodeGenOpt::Level OptLevel); }; // class PTXTargetMachine + + +class PTX32TargetMachine : public PTXTargetMachine { +public: + + PTX32TargetMachine(const Target &T, const std::string &TT, + const std::string& FS); +}; // class PTX32TargetMachine + +class PTX64TargetMachine : public PTXTargetMachine { +public: + + PTX64TargetMachine(const Target &T, const std::string &TT, + const std::string& FS); +}; // class PTX32TargetMachine + } // namespace llvm #endif // PTX_TARGET_MACHINE_H diff --git a/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp b/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp index a577d77..9df6c75 100644 --- a/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp +++ b/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp @@ -13,9 +13,13 @@ using namespace llvm; -Target llvm::ThePTXTarget; +Target llvm::ThePTX32Target; +Target llvm::ThePTX64Target; extern "C" void LLVMInitializePTXTargetInfo() { // see llvm/ADT/Triple.h - RegisterTarget<Triple::ptx> X(ThePTXTarget, "ptx", "PTX"); + RegisterTarget<Triple::ptx32> X32(ThePTX32Target, "ptx32", + "PTX (32-bit) [Experimental]"); + RegisterTarget<Triple::ptx64> X64(ThePTX64Target, "ptx64", + "PTX (64-bit) [Experimental]"); } diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll index 598591c..235b00e 100644 --- a/test/CodeGen/PTX/add.ll +++ b/test/CodeGen/PTX/add.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i16 @t1_u16(i16 %x, i16 %y) { ; CHECK: add.u16 rh0, rh1, rh2; diff --git a/test/CodeGen/PTX/bra.ll b/test/CodeGen/PTX/bra.ll index 0506a99..49383eb 100644 --- a/test/CodeGen/PTX/bra.ll +++ b/test/CodeGen/PTX/bra.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device void @test_bra_direct() { ; CHECK: bra $L__BB0_1; diff --git a/test/CodeGen/PTX/exit.ll b/test/CodeGen/PTX/exit.ll index 4071bab..7816c80 100644 --- a/test/CodeGen/PTX/exit.ll +++ b/test/CodeGen/PTX/exit.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_kernel void @t1() { ; CHECK: exit; diff --git a/test/CodeGen/PTX/fdiv-sm10.ll b/test/CodeGen/PTX/fdiv-sm10.ll index 42f615d..121360c 100644 --- a/test/CodeGen/PTX/fdiv-sm10.ll +++ b/test/CodeGen/PTX/fdiv-sm10.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx -mattr=+sm10 | FileCheck %s +; RUN: llc < %s -march=ptx32 -mattr=+sm10 | FileCheck %s define ptx_device float @t1_f32(float %x, float %y) { ; CHECK: div.approx.f32 f0, f1, f2; diff --git a/test/CodeGen/PTX/fdiv-sm13.ll b/test/CodeGen/PTX/fdiv-sm13.ll index eb20f78..0ec7bae 100644 --- a/test/CodeGen/PTX/fdiv-sm13.ll +++ b/test/CodeGen/PTX/fdiv-sm13.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx -mattr=+sm13 | FileCheck %s +; RUN: llc < %s -march=ptx32 -mattr=+sm13 | FileCheck %s define ptx_device float @t1_f32(float %x, float %y) { ; CHECK: div.approx.f32 f0, f1, f2; diff --git a/test/CodeGen/PTX/intrinsic.ll b/test/CodeGen/PTX/intrinsic.ll index 7405dd6..cea4182 100644 --- a/test/CodeGen/PTX/intrinsic.ll +++ b/test/CodeGen/PTX/intrinsic.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx -mattr=+ptx20,+sm20 | FileCheck %s +; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | FileCheck %s define ptx_device i32 @test_tid_x() { ; CHECK: mov.u32 r0, %tid.x; diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll index 1119aa4..58e16a2 100644 --- a/test/CodeGen/PTX/ld.ll +++ b/test/CodeGen/PTX/ld.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s ;CHECK: .extern .global .b8 array_i16[20]; @array_i16 = external global [10 x i16] diff --git a/test/CodeGen/PTX/llvm-intrinsic.ll b/test/CodeGen/PTX/llvm-intrinsic.ll index 3ce4c29..1e265f5 100644 --- a/test/CodeGen/PTX/llvm-intrinsic.ll +++ b/test/CodeGen/PTX/llvm-intrinsic.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx -mattr=+ptx20,+sm20 | FileCheck %s +; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | FileCheck %s define ptx_device float @test_sqrt_f32(float %x) { entry: diff --git a/test/CodeGen/PTX/mad.ll b/test/CodeGen/PTX/mad.ll index 786345b..0c25f2c 100644 --- a/test/CodeGen/PTX/mad.ll +++ b/test/CodeGen/PTX/mad.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx -mattr=+sm13 | FileCheck %s +; RUN: llc < %s -march=ptx32 -mattr=+sm13 | FileCheck %s define ptx_device float @t1_f32(float %x, float %y, float %z) { ; CHECK: mad.rn.f32 f0, f1, f2, f3; diff --git a/test/CodeGen/PTX/mov.ll b/test/CodeGen/PTX/mov.ll index 00dcf19..120572a 100644 --- a/test/CodeGen/PTX/mov.ll +++ b/test/CodeGen/PTX/mov.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i16 @t1_u16() { ; CHECK: mov.u16 rh0, 0; diff --git a/test/CodeGen/PTX/mul.ll b/test/CodeGen/PTX/mul.ll index fd0788f..5ce0426 100644 --- a/test/CodeGen/PTX/mul.ll +++ b/test/CodeGen/PTX/mul.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s ;define ptx_device i32 @t1(i32 %x, i32 %y) { ; %z = mul i32 %x, %y diff --git a/test/CodeGen/PTX/options.ll b/test/CodeGen/PTX/options.ll index 6576a6d..ac33fef 100644 --- a/test/CodeGen/PTX/options.ll +++ b/test/CodeGen/PTX/options.ll @@ -1,9 +1,9 @@ -; RUN: llc < %s -march=ptx -mattr=ptx20 | grep ".version 2.0" -; RUN: llc < %s -march=ptx -mattr=ptx21 | grep ".version 2.1" -; RUN: llc < %s -march=ptx -mattr=ptx22 | grep ".version 2.2" -; RUN: llc < %s -march=ptx -mattr=sm10 | grep ".target sm_10" -; RUN: llc < %s -march=ptx -mattr=sm13 | grep ".target sm_13" -; RUN: llc < %s -march=ptx -mattr=sm20 | grep ".target sm_20" +; RUN: llc < %s -march=ptx32 -mattr=ptx20 | grep ".version 2.0" +; RUN: llc < %s -march=ptx32 -mattr=ptx21 | grep ".version 2.1" +; RUN: llc < %s -march=ptx32 -mattr=ptx22 | grep ".version 2.2" +; RUN: llc < %s -march=ptx32 -mattr=sm10 | grep ".target sm_10" +; RUN: llc < %s -march=ptx32 -mattr=sm13 | grep ".target sm_13" +; RUN: llc < %s -march=ptx32 -mattr=sm20 | grep ".target sm_20" define ptx_device void @t1() { ret void diff --git a/test/CodeGen/PTX/parameter-order.ll b/test/CodeGen/PTX/parameter-order.ll index dbbbb67..8131f13 100644 --- a/test/CodeGen/PTX/parameter-order.ll +++ b/test/CodeGen/PTX/parameter-order.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s ; CHECK: .func (.reg .u32 r0) test_parameter_order (.reg .u32 r1, .reg .u32 r2) define ptx_device i32 @test_parameter_order(i32 %x, i32 %y) { diff --git a/test/CodeGen/PTX/ret.ll b/test/CodeGen/PTX/ret.ll index d5037f2..ba0523f 100644 --- a/test/CodeGen/PTX/ret.ll +++ b/test/CodeGen/PTX/ret.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device void @t1() { ; CHECK: ret; diff --git a/test/CodeGen/PTX/setp.ll b/test/CodeGen/PTX/setp.ll index 5348482..5836122 100644 --- a/test/CodeGen/PTX/setp.ll +++ b/test/CodeGen/PTX/setp.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i32 @test_setp_eq_u32_rr(i32 %x, i32 %y) { ; CHECK: setp.eq.u32 p0, r1, r2; diff --git a/test/CodeGen/PTX/shl.ll b/test/CodeGen/PTX/shl.ll index b564b43..6e72c92 100644 --- a/test/CodeGen/PTX/shl.ll +++ b/test/CodeGen/PTX/shl.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i32 @t1(i32 %x, i32 %y) { ; CHECK: shl.b32 r0, r1, r2 diff --git a/test/CodeGen/PTX/shr.ll b/test/CodeGen/PTX/shr.ll index 3f8ade8..8693e0e 100644 --- a/test/CodeGen/PTX/shr.ll +++ b/test/CodeGen/PTX/shr.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i32 @t1(i32 %x, i32 %y) { ; CHECK: shr.u32 r0, r1, r2 diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll index 4e9b08a..dee5c61 100644 --- a/test/CodeGen/PTX/st.ll +++ b/test/CodeGen/PTX/st.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s ;CHECK: .extern .global .b8 array_i16[20]; @array_i16 = external global [10 x i16] diff --git a/test/CodeGen/PTX/sub.ll b/test/CodeGen/PTX/sub.ll index 4810e4f..7dd2c6f 100644 --- a/test/CodeGen/PTX/sub.ll +++ b/test/CodeGen/PTX/sub.ll @@ -1,4 +1,4 @@ -; RUN: llc < %s -march=ptx | FileCheck %s +; RUN: llc < %s -march=ptx32 | FileCheck %s define ptx_device i16 @t1_u16(i16 %x, i16 %y) { ; CHECK: sub.u16 rh0, rh1, rh2; |