[llvm-commits] [llvm] r129851 - in /llvm/trunk: include/llvm/ include/llvm/ADT/ lib/Support/ lib/Target/PTX/ lib/Target/PTX/TargetInfo/ test/CodeGen/PTX/
Justin Holewinski
justin.holewinski at gmail.com
Wed Apr 20 08:37:17 PDT 2011
Author: jholewinski
Date: Wed Apr 20 10:37:17 2011
New Revision: 129851
URL: http://llvm.org/viewvc/llvm-project?rev=129851&view=rev
Log:
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
Modified:
llvm/trunk/include/llvm/ADT/Triple.h
llvm/trunk/include/llvm/IntrinsicsPTX.td
llvm/trunk/lib/Support/Triple.cpp
llvm/trunk/lib/Target/PTX/PTX.h
llvm/trunk/lib/Target/PTX/PTX.td
llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp
llvm/trunk/lib/Target/PTX/PTXInstrInfo.td
llvm/trunk/lib/Target/PTX/PTXSubtarget.cpp
llvm/trunk/lib/Target/PTX/PTXSubtarget.h
llvm/trunk/lib/Target/PTX/PTXTargetMachine.cpp
llvm/trunk/lib/Target/PTX/PTXTargetMachine.h
llvm/trunk/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp
llvm/trunk/test/CodeGen/PTX/add.ll
llvm/trunk/test/CodeGen/PTX/bra.ll
llvm/trunk/test/CodeGen/PTX/exit.ll
llvm/trunk/test/CodeGen/PTX/fdiv-sm10.ll
llvm/trunk/test/CodeGen/PTX/fdiv-sm13.ll
llvm/trunk/test/CodeGen/PTX/intrinsic.ll
llvm/trunk/test/CodeGen/PTX/ld.ll
llvm/trunk/test/CodeGen/PTX/llvm-intrinsic.ll
llvm/trunk/test/CodeGen/PTX/mad.ll
llvm/trunk/test/CodeGen/PTX/mov.ll
llvm/trunk/test/CodeGen/PTX/mul.ll
llvm/trunk/test/CodeGen/PTX/options.ll
llvm/trunk/test/CodeGen/PTX/parameter-order.ll
llvm/trunk/test/CodeGen/PTX/ret.ll
llvm/trunk/test/CodeGen/PTX/setp.ll
llvm/trunk/test/CodeGen/PTX/shl.ll
llvm/trunk/test/CodeGen/PTX/shr.ll
llvm/trunk/test/CodeGen/PTX/st.ll
llvm/trunk/test/CodeGen/PTX/sub.ll
Modified: llvm/trunk/include/llvm/ADT/Triple.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/ADT/Triple.h?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/include/llvm/ADT/Triple.h (original)
+++ llvm/trunk/include/llvm/ADT/Triple.h Wed Apr 20 10:37:17 2011
@@ -64,7 +64,8 @@
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
};
Modified: llvm/trunk/include/llvm/IntrinsicsPTX.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IntrinsicsPTX.td?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IntrinsicsPTX.td (original)
+++ llvm/trunk/include/llvm/IntrinsicsPTX.td Wed Apr 20 10:37:17 2011
@@ -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_r64
- : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>;
+ class PTXReadSpecialRegisterIntrinsic_r32<string name>
+ : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
+ GCCBuiltin<name>;
+
+ 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;
-
-def int_ptx_read_laneid : PTXReadSpecialRegisterIntrinsic_r32;
-def int_ptx_read_warpid : PTXReadSpecialRegisterIntrinsic_r32;
-def int_ptx_read_nwarpid : PTXReadSpecialRegisterIntrinsic_r32;
-
-defm int_ptx_read_ctaid : PTXReadSpecialRegisterIntrinsic_v4i32;
-defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32;
-
-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_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_clock : PTXReadSpecialRegisterIntrinsic_r32;
-def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64;
-
-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;
+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
+ <"__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
+ <"__builtin_ptx_read_ctaid">;
+defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32
+ <"__builtin_ptx_read_nctaid">;
+
+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
+ <"__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
+ <"__builtin_ptx_read_clock">;
+def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64
+ <"__builtin_ptx_read_clock64">;
+
+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">;
Modified: llvm/trunk/lib/Support/Triple.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Support/Triple.cpp?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Support/Triple.cpp (original)
+++ llvm/trunk/lib/Support/Triple.cpp Wed Apr 20 10:37:17 2011
@@ -41,7 +41,8 @@
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 @@
case xcore: return "xcore";
- case ptx: return "ptx";
+ case ptx32: return "ptx";
+ case ptx64: return "ptx";
}
}
@@ -165,8 +167,10 @@
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 @@
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 @@
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 @@
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;
}
Modified: llvm/trunk/lib/Target/PTX/PTX.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTX.h?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTX.h (original)
+++ llvm/trunk/lib/Target/PTX/PTX.h Wed Apr 20 10:37:17 2011
@@ -42,7 +42,8 @@
FunctionPass *createPTXMFInfoExtract(PTXTargetMachine &TM,
CodeGenOpt::Level OptLevel);
- extern Target ThePTXTarget;
+ extern Target ThePTX32Target;
+ extern Target ThePTX64Target;
} // namespace llvm;
// Defines symbolic names for PTX registers.
Modified: llvm/trunk/lib/Target/PTX/PTX.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTX.td?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTX.td (original)
+++ llvm/trunk/lib/Target/PTX/PTX.td Wed Apr 20 10:37:17 2011
@@ -24,9 +24,6 @@
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",
Modified: llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp Wed Apr 20 10:37:17 2011
@@ -447,5 +447,6 @@
// Force static initialization.
extern "C" void LLVMInitializePTXAsmPrinter() {
- RegisterAsmPrinter<PTXAsmPrinter> X(ThePTXTarget);
+ RegisterAsmPrinter<PTXAsmPrinter> X(ThePTX32Target);
+ RegisterAsmPrinter<PTXAsmPrinter> Y(ThePTX64Target);
}
Modified: llvm/trunk/lib/Target/PTX/PTXInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.td?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXInstrInfo.td (original)
+++ llvm/trunk/lib/Target/PTX/PTXInstrInfo.td Wed Apr 20 10:37:17 2011
@@ -22,8 +22,8 @@
//===----------------------------------------------------------------------===//
// 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()">;
Modified: llvm/trunk/lib/Target/PTX/PTXSubtarget.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXSubtarget.cpp?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXSubtarget.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXSubtarget.cpp Wed Apr 20 10:37:17 2011
@@ -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);
}
Modified: llvm/trunk/lib/Target/PTX/PTXSubtarget.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXSubtarget.h?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXSubtarget.h (original)
+++ llvm/trunk/lib/Target/PTX/PTXSubtarget.h Wed Apr 20 10:37:17 2011
@@ -50,10 +50,10 @@
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 @@
bool supportsDouble() const { return SupportsDouble; }
- bool use64BitAddresses() const { return Use64BitAddresses; }
+ bool is64Bit() const { return Is64Bit; }
bool supportsSM13() const { return PTXShaderModel >= PTX_SM_1_3; }
Modified: llvm/trunk/lib/Target/PTX/PTXTargetMachine.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXTargetMachine.cpp?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXTargetMachine.cpp (original)
+++ llvm/trunk/lib/Target/PTX/PTXTargetMachine.cpp Wed Apr 20 10:37:17 2011
@@ -30,9 +30,15 @@
}
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 @@
// 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));
Modified: llvm/trunk/lib/Target/PTX/PTXTargetMachine.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXTargetMachine.h?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/PTXTargetMachine.h (original)
+++ llvm/trunk/lib/Target/PTX/PTXTargetMachine.h Wed Apr 20 10:37:17 2011
@@ -33,7 +33,7 @@
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 @@
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
Modified: llvm/trunk/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp (original)
+++ llvm/trunk/lib/Target/PTX/TargetInfo/PTXTargetInfo.cpp Wed Apr 20 10:37:17 2011
@@ -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]");
}
Modified: llvm/trunk/test/CodeGen/PTX/add.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/add.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/add.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/add.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/bra.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/bra.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/bra.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/bra.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/exit.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/exit.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/exit.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/exit.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/fdiv-sm10.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/fdiv-sm10.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/fdiv-sm10.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/fdiv-sm10.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/fdiv-sm13.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/fdiv-sm13.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/fdiv-sm13.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/fdiv-sm13.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/intrinsic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/intrinsic.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/intrinsic.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/intrinsic.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/ld.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/ld.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/ld.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/ld.ll Wed Apr 20 10:37:17 2011
@@ -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]
Modified: llvm/trunk/test/CodeGen/PTX/llvm-intrinsic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/llvm-intrinsic.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/llvm-intrinsic.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/llvm-intrinsic.ll Wed Apr 20 10:37:17 2011
@@ -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:
Modified: llvm/trunk/test/CodeGen/PTX/mad.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mad.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/mad.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/mad.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/mov.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mov.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/mov.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/mov.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/mul.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mul.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/mul.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/mul.ll Wed Apr 20 10:37:17 2011
@@ -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
Modified: llvm/trunk/test/CodeGen/PTX/options.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/options.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/options.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/options.ll Wed Apr 20 10:37:17 2011
@@ -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
Modified: llvm/trunk/test/CodeGen/PTX/parameter-order.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/parameter-order.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/parameter-order.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/parameter-order.ll Wed Apr 20 10:37:17 2011
@@ -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) {
Modified: llvm/trunk/test/CodeGen/PTX/ret.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/ret.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/ret.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/ret.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/setp.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/setp.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/setp.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/setp.ll Wed Apr 20 10:37:17 2011
@@ -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;
Modified: llvm/trunk/test/CodeGen/PTX/shl.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/shl.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/shl.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/shl.ll Wed Apr 20 10:37:17 2011
@@ -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
Modified: llvm/trunk/test/CodeGen/PTX/shr.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/shr.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/shr.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/shr.ll Wed Apr 20 10:37:17 2011
@@ -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
Modified: llvm/trunk/test/CodeGen/PTX/st.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/st.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/st.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/st.ll Wed Apr 20 10:37:17 2011
@@ -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]
Modified: llvm/trunk/test/CodeGen/PTX/sub.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/sub.ll?rev=129851&r1=129850&r2=129851&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/PTX/sub.ll (original)
+++ llvm/trunk/test/CodeGen/PTX/sub.ll Wed Apr 20 10:37:17 2011
@@ -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;
More information about the llvm-commits
mailing list