From e1fee48cd0d1e515f247fe3bceceb0f854623f73 Mon Sep 17 00:00:00 2001 From: Justin Holewinski Date: Wed, 20 Apr 2011 15:37:17 +0000 Subject: [PATCH] PTX: Add intrinsics to list of built-in intrinsics, which allows them to be used by Clang. To help Clang integration, the PTX target has been split into two targets: ptx32 and ptx64, depending on the desired pointer size. - Add GCCBuiltin class to all intrinsics - Split PTX target into ptx32 and ptx64 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@129851 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/llvm/ADT/Triple.h | 3 +- include/llvm/IntrinsicsPTX.td | 90 ++++++++++++++------- lib/Support/Triple.cpp | 30 ++++--- lib/Target/PTX/PTX.h | 3 +- lib/Target/PTX/PTX.td | 3 - lib/Target/PTX/PTXAsmPrinter.cpp | 3 +- lib/Target/PTX/PTXInstrInfo.td | 4 +- lib/Target/PTX/PTXSubtarget.cpp | 5 +- lib/Target/PTX/PTXSubtarget.h | 6 +- lib/Target/PTX/PTXTargetMachine.cpp | 34 +++++--- lib/Target/PTX/PTXTargetMachine.h | 18 ++++- lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp | 8 +- test/CodeGen/PTX/add.ll | 2 +- test/CodeGen/PTX/bra.ll | 2 +- test/CodeGen/PTX/exit.ll | 2 +- test/CodeGen/PTX/fdiv-sm10.ll | 2 +- test/CodeGen/PTX/fdiv-sm13.ll | 2 +- test/CodeGen/PTX/intrinsic.ll | 2 +- test/CodeGen/PTX/ld.ll | 2 +- test/CodeGen/PTX/llvm-intrinsic.ll | 2 +- test/CodeGen/PTX/mad.ll | 2 +- test/CodeGen/PTX/mov.ll | 2 +- test/CodeGen/PTX/mul.ll | 2 +- test/CodeGen/PTX/options.ll | 12 +-- test/CodeGen/PTX/parameter-order.ll | 2 +- test/CodeGen/PTX/ret.ll | 2 +- test/CodeGen/PTX/setp.ll | 2 +- test/CodeGen/PTX/shl.ll | 2 +- test/CodeGen/PTX/shr.ll | 2 +- test/CodeGen/PTX/st.ll | 2 +- test/CodeGen/PTX/sub.ll | 2 +- 31 files changed, 165 insertions(+), 90 deletions(-) diff --git a/include/llvm/ADT/Triple.h b/include/llvm/ADT/Triple.h index 4cfad31df53..2659bce6175 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 01241fe4d48..28379c918de 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 { // 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; + def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin; + def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin; + def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin; } - class PTXReadSpecialRegisterIntrinsic_r32 - : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; + class PTXReadSpecialRegisterIntrinsic_r32 + : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, + GCCBuiltin; - class PTXReadSpecialRegisterIntrinsic_r64 - : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>; + class PTXReadSpecialRegisterIntrinsic_r64 + : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>, + GCCBuiltin; } -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 ad93121d02d..dbdb303a4fd 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 ""; @@ -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 49045cdc01e..ec2be9291a0 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 dbc6f579a29..ae8326e3199 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 27c96053bfe..3363c7300f8 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 X(ThePTXTarget); + RegisterAsmPrinter X(ThePTX32Target); + RegisterAsmPrinter Y(ThePTX64Target); } diff --git a/lib/Target/PTX/PTXInstrInfo.td b/lib/Target/PTX/PTXInstrInfo.td index 972002cf6ae..c124c03896b 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 527622d0c78..a224f2b8be1 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 57cd43da476..47d98424065 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 4701a941d18..78a7b0daf15 100644 --- a/lib/Target/PTX/PTXTargetMachine.cpp +++ b/lib/Target/PTX/PTXTargetMachine.cpp @@ -30,9 +30,15 @@ namespace llvm { } extern "C" void LLVMInitializePTXTarget() { - RegisterTargetMachine X(ThePTXTarget); - RegisterAsmInfo Y(ThePTXTarget); - TargetRegistry::RegisterAsmStreamer(ThePTXTarget, createPTXAsmStreamer); + + RegisterTargetMachine X(ThePTX32Target); + RegisterTargetMachine Y(ThePTX64Target); + + RegisterAsmInfo Z(ThePTX32Target); + RegisterAsmInfo 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 a5dba537d1d..149be8e3b7e 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 a577d7755af..9df6c7567bd 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 X(ThePTXTarget, "ptx", "PTX"); + RegisterTarget X32(ThePTX32Target, "ptx32", + "PTX (32-bit) [Experimental]"); + RegisterTarget X64(ThePTX64Target, "ptx64", + "PTX (64-bit) [Experimental]"); } diff --git a/test/CodeGen/PTX/add.ll b/test/CodeGen/PTX/add.ll index 598591c0fcb..235b00e8782 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 0506a990668..49383eb3cf9 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 4071babb80c..7816c801728 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 42f615d0c8d..121360ce9be 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 eb20f787639..0ec7bae8030 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 7405dd6f5e5..cea41827ca4 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 1119aa46944..58e16a20a45 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 3ce4c29f9f5..1e265f5b7b3 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 786345b2913..0c25f2c0030 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 00dcf19f1da..120572a0e86 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 fd0788fce66..5ce042675dc 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 6576a6d8bbb..ac33fef0d6e 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 dbbbb67a140..8131f13a6e8 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 d5037f25fd3..ba0523f6424 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 5348482e093..5836122049e 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 b564b43ab93..6e72c922132 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 3f8ade862b7..8693e0ecf49 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 4e9b08a33a2..dee5c61abe6 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 4810e4fc055..7dd2c6f6ac7 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;