llvm.org GIT mirror llvm / 42a0b48
Remove the PTX back-end and all of its artifacts (triple, etc.) This back-end was deprecated in favor of the NVPTX back-end. NV_CONTRIB git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@157417 91177308-0d34-0410-b5e6-96231b3b80d8 Justin Holewinski 8 years ago
95 changed file(s) with 97 addition(s) and 9086 deletion(s). Raw diff Collapse all Expand all
8282 MSP430
8383 NVPTX
8484 PowerPC
85 PTX
8685 Sparc
8786 X86
8887 XCore
368368 msp430-*) llvm_cv_target_arch="MSP430" ;;
369369 hexagon-*) llvm_cv_target_arch="Hexagon" ;;
370370 mblaze-*) llvm_cv_target_arch="MBlaze" ;;
371 ptx-*) llvm_cv_target_arch="PTX" ;;
372371 nvptx-*) llvm_cv_target_arch="NVPTX" ;;
373372 *) llvm_cv_target_arch="Unknown" ;;
374373 esac])
402401 msp430-*) host_arch="MSP430" ;;
403402 hexagon-*) host_arch="Hexagon" ;;
404403 mblaze-*) host_arch="MBlaze" ;;
405 ptx-*) host_arch="PTX" ;;
406404 *) host_arch="Unknown" ;;
407405 esac
408406
541539 MSP430) AC_SUBST(TARGET_HAS_JIT,0) ;;
542540 Hexagon) AC_SUBST(TARGET_HAS_JIT,0) ;;
543541 MBlaze) AC_SUBST(TARGET_HAS_JIT,0) ;;
544 PTX) AC_SUBST(TARGET_HAS_JIT,0) ;;
545542 NVPTX) AC_SUBST(TARGET_HAS_JIT,0) ;;
546543 *) AC_SUBST(TARGET_HAS_JIT,0) ;;
547544 esac
654651 AC_ARG_ENABLE([targets],AS_HELP_STRING([--enable-targets],
655652 [Build specific host targets: all or target1,target2,... Valid targets are:
656653 host, x86, x86_64, sparc, powerpc, arm, mips, spu, hexagon,
657 xcore, msp430, ptx, nvptx, and cpp (default=all)]),,
654 xcore, msp430, nvptx, and cpp (default=all)]),,
658655 enableval=all)
659656 if test "$enableval" = host-only ; then
660657 enableval=host
661658 fi
662659 case "$enableval" in
663 all) TARGETS_TO_BUILD="X86 Sparc PowerPC ARM Mips CellSPU XCore MSP430 CppBackend MBlaze PTX NVPTX Hexagon" ;;
660 all) TARGETS_TO_BUILD="X86 Sparc PowerPC ARM Mips CellSPU XCore MSP430 CppBackend MBlaze NVPTX Hexagon" ;;
664661 *)for a_target in `echo $enableval|sed -e 's/,/ /g' ` ; do
665662 case "$a_target" in
666663 x86) TARGETS_TO_BUILD="X86 $TARGETS_TO_BUILD" ;;
676673 cpp) TARGETS_TO_BUILD="CppBackend $TARGETS_TO_BUILD" ;;
677674 hexagon) TARGETS_TO_BUILD="Hexagon $TARGETS_TO_BUILD" ;;
678675 mblaze) TARGETS_TO_BUILD="MBlaze $TARGETS_TO_BUILD" ;;
679 ptx) TARGETS_TO_BUILD="PTX $TARGETS_TO_BUILD" ;;
680676 nvptx) TARGETS_TO_BUILD="NVPTX $TARGETS_TO_BUILD" ;;
681677 host) case "$llvm_cv_target_arch" in
682678 x86) TARGETS_TO_BUILD="X86 $TARGETS_TO_BUILD" ;;
690686 XCore) TARGETS_TO_BUILD="XCore $TARGETS_TO_BUILD" ;;
691687 MSP430) TARGETS_TO_BUILD="MSP430 $TARGETS_TO_BUILD" ;;
692688 Hexagon) TARGETS_TO_BUILD="Hexagon $TARGETS_TO_BUILD" ;;
693 PTX) TARGETS_TO_BUILD="PTX $TARGETS_TO_BUILD" ;;
694689 NVPTX) TARGETS_TO_BUILD="NVPTX $TARGETS_TO_BUILD" ;;
695690 *) AC_MSG_ERROR([Can not set target to build]) ;;
696691 esac ;;
14181418 --enable-targets Build specific host targets: all or
14191419 target1,target2,... Valid targets are: host, x86,
14201420 x86_64, sparc, powerpc, arm, mips, spu, hexagon,
1421 xcore, msp430, ptx, nvptx, and cpp (default=all)
1421 xcore, msp430, nvptx, and cpp (default=all)
14221422 --enable-bindings Build specific language bindings:
14231423 all,auto,none,{binding-name} (default=auto)
14241424 --enable-libffi Check for the presence of libffi (default is NO)
39003900 msp430-*) llvm_cv_target_arch="MSP430" ;;
39013901 hexagon-*) llvm_cv_target_arch="Hexagon" ;;
39023902 mblaze-*) llvm_cv_target_arch="MBlaze" ;;
3903 ptx-*) llvm_cv_target_arch="PTX" ;;
39043903 nvptx-*) llvm_cv_target_arch="NVPTX" ;;
39053904 *) llvm_cv_target_arch="Unknown" ;;
39063905 esac
39343933 msp430-*) host_arch="MSP430" ;;
39353934 hexagon-*) host_arch="Hexagon" ;;
39363935 mblaze-*) host_arch="MBlaze" ;;
3937 ptx-*) host_arch="PTX" ;;
39383936 *) host_arch="Unknown" ;;
39393937 esac
39403938
51475145 ;;
51485146 MBlaze) TARGET_HAS_JIT=0
51495147 ;;
5150 PTX) TARGET_HAS_JIT=0
5151 ;;
51525148 NVPTX) TARGET_HAS_JIT=0
51535149 ;;
51545150 *) TARGET_HAS_JIT=0
53355331 enableval=host
53365332 fi
53375333 case "$enableval" in
5338 all) TARGETS_TO_BUILD="X86 Sparc PowerPC ARM Mips CellSPU XCore MSP430 CppBackend MBlaze PTX NVPTX Hexagon" ;;
5334 all) TARGETS_TO_BUILD="X86 Sparc PowerPC ARM Mips CellSPU XCore MSP430 CppBackend MBlaze NVPTX Hexagon" ;;
53395335 *)for a_target in `echo $enableval|sed -e 's/,/ /g' ` ; do
53405336 case "$a_target" in
53415337 x86) TARGETS_TO_BUILD="X86 $TARGETS_TO_BUILD" ;;
53515347 cpp) TARGETS_TO_BUILD="CppBackend $TARGETS_TO_BUILD" ;;
53525348 hexagon) TARGETS_TO_BUILD="Hexagon $TARGETS_TO_BUILD" ;;
53535349 mblaze) TARGETS_TO_BUILD="MBlaze $TARGETS_TO_BUILD" ;;
5354 ptx) TARGETS_TO_BUILD="PTX $TARGETS_TO_BUILD" ;;
53555350 nvptx) TARGETS_TO_BUILD="NVPTX $TARGETS_TO_BUILD" ;;
53565351 host) case "$llvm_cv_target_arch" in
53575352 x86) TARGETS_TO_BUILD="X86 $TARGETS_TO_BUILD" ;;
53655360 XCore) TARGETS_TO_BUILD="XCore $TARGETS_TO_BUILD" ;;
53665361 MSP430) TARGETS_TO_BUILD="MSP430 $TARGETS_TO_BUILD" ;;
53675362 Hexagon) TARGETS_TO_BUILD="Hexagon $TARGETS_TO_BUILD" ;;
5368 PTX) TARGETS_TO_BUILD="PTX $TARGETS_TO_BUILD" ;;
53695363 NVPTX) TARGETS_TO_BUILD="NVPTX $TARGETS_TO_BUILD" ;;
53705364 *) { { echo "$as_me:$LINENO: error: Can not set target to build" >&5
53715365 echo "$as_me: error: Can not set target to build" >&2;}
1034810342 lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
1034910343 lt_status=$lt_dlunknown
1035010344 cat > conftest.$ac_ext <
10351 #line 10352 "configure"
10345 #line 10346 "configure"
1035210346 #include "confdefs.h"
1035310347
1035410348 #if HAVE_DLFCN_H
6161 x86_64, // X86-64: amd64, x86_64
6262 xcore, // XCore: xcore
6363 mblaze, // MBlaze: mblaze
64 ptx32, // PTX: ptx (32-bit)
65 ptx64, // PTX: ptx (64-bit)
6664 nvptx, // NVPTX: 32-bit
6765 nvptx64, // NVPTX: 64-bit
6866 le32, // le32: generic little-endian 32-bit CPU (PNaCl / Emscripten)
444444 include "llvm/IntrinsicsARM.td"
445445 include "llvm/IntrinsicsCellSPU.td"
446446 include "llvm/IntrinsicsXCore.td"
447 include "llvm/IntrinsicsPTX.td"
448447 include "llvm/IntrinsicsHexagon.td"
449448 include "llvm/IntrinsicsNVVM.td"
869869 Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.error">;
870870 def int_nvvm_compiler_warn :
871871 Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.warn">;
872
873
874 // Old PTX back-end intrinsics retained here for backwards-compatibility
875
876 multiclass PTXReadSpecialRegisterIntrinsic_v4i32 {
877 // FIXME: Do we need the 128-bit integer type version?
878 // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem]>;
879
880 // FIXME: Enable this once v4i32 support is enabled in back-end.
881 // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem]>;
882
883 def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
884 GCCBuiltin;
885 def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
886 GCCBuiltin;
887 def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
888 GCCBuiltin;
889 def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
890 GCCBuiltin;
891 }
892
893 class PTXReadSpecialRegisterIntrinsic_r32
894 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
895 GCCBuiltin;
896
897 class PTXReadSpecialRegisterIntrinsic_r64
898 : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>,
899 GCCBuiltin;
900
901 defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic_v4i32
902 <"__builtin_ptx_read_tid">;
903 defm int_ptx_read_ntid : PTXReadSpecialRegisterIntrinsic_v4i32
904 <"__builtin_ptx_read_ntid">;
905
906 def int_ptx_read_laneid : PTXReadSpecialRegisterIntrinsic_r32
907 <"__builtin_ptx_read_laneid">;
908 def int_ptx_read_warpid : PTXReadSpecialRegisterIntrinsic_r32
909 <"__builtin_ptx_read_warpid">;
910 def int_ptx_read_nwarpid : PTXReadSpecialRegisterIntrinsic_r32
911 <"__builtin_ptx_read_nwarpid">;
912
913 defm int_ptx_read_ctaid : PTXReadSpecialRegisterIntrinsic_v4i32
914 <"__builtin_ptx_read_ctaid">;
915 defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32
916 <"__builtin_ptx_read_nctaid">;
917
918 def int_ptx_read_smid : PTXReadSpecialRegisterIntrinsic_r32
919 <"__builtin_ptx_read_smid">;
920 def int_ptx_read_nsmid : PTXReadSpecialRegisterIntrinsic_r32
921 <"__builtin_ptx_read_nsmid">;
922 def int_ptx_read_gridid : PTXReadSpecialRegisterIntrinsic_r32
923 <"__builtin_ptx_read_gridid">;
924
925 def int_ptx_read_lanemask_eq : PTXReadSpecialRegisterIntrinsic_r32
926 <"__builtin_ptx_read_lanemask_eq">;
927 def int_ptx_read_lanemask_le : PTXReadSpecialRegisterIntrinsic_r32
928 <"__builtin_ptx_read_lanemask_le">;
929 def int_ptx_read_lanemask_lt : PTXReadSpecialRegisterIntrinsic_r32
930 <"__builtin_ptx_read_lanemask_lt">;
931 def int_ptx_read_lanemask_ge : PTXReadSpecialRegisterIntrinsic_r32
932 <"__builtin_ptx_read_lanemask_ge">;
933 def int_ptx_read_lanemask_gt : PTXReadSpecialRegisterIntrinsic_r32
934 <"__builtin_ptx_read_lanemask_gt">;
935
936 def int_ptx_read_clock : PTXReadSpecialRegisterIntrinsic_r32
937 <"__builtin_ptx_read_clock">;
938 def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64
939 <"__builtin_ptx_read_clock64">;
940
941 def int_ptx_read_pm0 : PTXReadSpecialRegisterIntrinsic_r32
942 <"__builtin_ptx_read_pm0">;
943 def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32
944 <"__builtin_ptx_read_pm1">;
945 def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32
946 <"__builtin_ptx_read_pm2">;
947 def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32
948 <"__builtin_ptx_read_pm3">;
949
950 def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>,
951 GCCBuiltin<"__builtin_ptx_bar_sync">;
+0
-92
include/llvm/IntrinsicsPTX.td less more
None //===- IntrinsicsPTX.td - Defines PTX intrinsics -----------*- tablegen -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines all of the PTX-specific intrinsics.
10 //
11 //===----------------------------------------------------------------------===//
12
13 let TargetPrefix = "ptx" in {
14 multiclass PTXReadSpecialRegisterIntrinsic_v4i32 {
15 // FIXME: Do we need the 128-bit integer type version?
16 // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem]>;
17
18 // FIXME: Enable this once v4i32 support is enabled in back-end.
19 // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem]>;
20
21 def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
22 GCCBuiltin;
23 def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
24 GCCBuiltin;
25 def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
26 GCCBuiltin;
27 def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
28 GCCBuiltin;
29 }
30
31 class PTXReadSpecialRegisterIntrinsic_r32
32 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
33 GCCBuiltin;
34
35 class PTXReadSpecialRegisterIntrinsic_r64
36 : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>,
37 GCCBuiltin;
38 }
39
40 defm int_ptx_read_tid : PTXReadSpecialRegisterIntrinsic_v4i32
41 <"__builtin_ptx_read_tid">;
42 defm int_ptx_read_ntid : PTXReadSpecialRegisterIntrinsic_v4i32
43 <"__builtin_ptx_read_ntid">;
44
45 def int_ptx_read_laneid : PTXReadSpecialRegisterIntrinsic_r32
46 <"__builtin_ptx_read_laneid">;
47 def int_ptx_read_warpid : PTXReadSpecialRegisterIntrinsic_r32
48 <"__builtin_ptx_read_warpid">;
49 def int_ptx_read_nwarpid : PTXReadSpecialRegisterIntrinsic_r32
50 <"__builtin_ptx_read_nwarpid">;
51
52 defm int_ptx_read_ctaid : PTXReadSpecialRegisterIntrinsic_v4i32
53 <"__builtin_ptx_read_ctaid">;
54 defm int_ptx_read_nctaid : PTXReadSpecialRegisterIntrinsic_v4i32
55 <"__builtin_ptx_read_nctaid">;
56
57 def int_ptx_read_smid : PTXReadSpecialRegisterIntrinsic_r32
58 <"__builtin_ptx_read_smid">;
59 def int_ptx_read_nsmid : PTXReadSpecialRegisterIntrinsic_r32
60 <"__builtin_ptx_read_nsmid">;
61 def int_ptx_read_gridid : PTXReadSpecialRegisterIntrinsic_r32
62 <"__builtin_ptx_read_gridid">;
63
64 def int_ptx_read_lanemask_eq : PTXReadSpecialRegisterIntrinsic_r32
65 <"__builtin_ptx_read_lanemask_eq">;
66 def int_ptx_read_lanemask_le : PTXReadSpecialRegisterIntrinsic_r32
67 <"__builtin_ptx_read_lanemask_le">;
68 def int_ptx_read_lanemask_lt : PTXReadSpecialRegisterIntrinsic_r32
69 <"__builtin_ptx_read_lanemask_lt">;
70 def int_ptx_read_lanemask_ge : PTXReadSpecialRegisterIntrinsic_r32
71 <"__builtin_ptx_read_lanemask_ge">;
72 def int_ptx_read_lanemask_gt : PTXReadSpecialRegisterIntrinsic_r32
73 <"__builtin_ptx_read_lanemask_gt">;
74
75 def int_ptx_read_clock : PTXReadSpecialRegisterIntrinsic_r32
76 <"__builtin_ptx_read_clock">;
77 def int_ptx_read_clock64 : PTXReadSpecialRegisterIntrinsic_r64
78 <"__builtin_ptx_read_clock64">;
79
80 def int_ptx_read_pm0 : PTXReadSpecialRegisterIntrinsic_r32
81 <"__builtin_ptx_read_pm0">;
82 def int_ptx_read_pm1 : PTXReadSpecialRegisterIntrinsic_r32
83 <"__builtin_ptx_read_pm1">;
84 def int_ptx_read_pm2 : PTXReadSpecialRegisterIntrinsic_r32
85 <"__builtin_ptx_read_pm2">;
86 def int_ptx_read_pm3 : PTXReadSpecialRegisterIntrinsic_r32
87 <"__builtin_ptx_read_pm3">;
88
89 let TargetPrefix = "ptx" in
90 def int_ptx_bar_sync : Intrinsic<[], [llvm_i32_ty], []>,
91 GCCBuiltin<"__builtin_ptx_bar_sync">;
3737 case x86_64: return "x86_64";
3838 case xcore: return "xcore";
3939 case mblaze: return "mblaze";
40 case ptx32: return "ptx32";
41 case ptx64: return "ptx64";
4240 case nvptx: return "nvptx";
4341 case nvptx64: return "nvptx64";
4442 case le32: return "le32";
7573
7674 case xcore: return "xcore";
7775
78 case ptx32: return "ptx";
79 case ptx64: return "ptx";
8076 case nvptx: return "nvptx";
8177 case nvptx64: return "nvptx";
8278 case le32: return "le32";
163159 .Case("x86", x86)
164160 .Case("x86-64", x86_64)
165161 .Case("xcore", xcore)
166 .Case("ptx32", ptx32)
167 .Case("ptx64", ptx64)
168162 .Case("nvptx", nvptx)
169163 .Case("nvptx64", nvptx64)
170164 .Case("le32", le32)
197191 .Cases("arm", "armv4t", "armv5", "armv6", Triple::arm)
198192 .Cases("armv7", "armv7f", "armv7k", "armv7s", "xscale", Triple::arm)
199193 .Case("r600", Triple::r600)
200 .Case("ptx32", Triple::ptx32)
201 .Case("ptx64", Triple::ptx64)
202194 .Case("nvptx", Triple::nvptx)
203195 .Case("nvptx64", Triple::nvptx64)
204196 .Case("amdil", Triple::amdil)
222214 .Cases("armv6", "thumbv6", "armv6")
223215 .Cases("armv7", "thumbv7", "armv7")
224216 .Case("r600", "r600")
225 .Case("ptx32", "ptx32")
226 .Case("ptx64", "ptx64")
227217 .Case("nvptx", "nvptx")
228218 .Case("nvptx64", "nvptx64")
229219 .Case("le32", "le32")
258248 .Case("sparcv9", Triple::sparcv9)
259249 .Case("tce", Triple::tce)
260250 .Case("xcore", Triple::xcore)
261 .Case("ptx32", Triple::ptx32)
262 .Case("ptx64", Triple::ptx64)
263251 .Case("nvptx", Triple::nvptx)
264252 .Case("nvptx64", Triple::nvptx64)
265253 .Case("le32", Triple::le32)
688676 case llvm::Triple::mipsel:
689677 case llvm::Triple::nvptx:
690678 case llvm::Triple::ppc:
691 case llvm::Triple::ptx32:
692679 case llvm::Triple::r600:
693680 case llvm::Triple::sparc:
694681 case llvm::Triple::tce:
701688 case llvm::Triple::mips64el:
702689 case llvm::Triple::nvptx64:
703690 case llvm::Triple::ppc64:
704 case llvm::Triple::ptx64:
705691 case llvm::Triple::sparcv9:
706692 case llvm::Triple::x86_64:
707693 return 64;
739725 case Triple::mipsel:
740726 case Triple::nvptx:
741727 case Triple::ppc:
742 case Triple::ptx32:
743728 case Triple::r600:
744729 case Triple::sparc:
745730 case Triple::tce:
753738 case Triple::mips64el: T.setArch(Triple::mipsel); break;
754739 case Triple::nvptx64: T.setArch(Triple::nvptx); break;
755740 case Triple::ppc64: T.setArch(Triple::ppc); break;
756 case Triple::ptx64: T.setArch(Triple::ptx32); break;
757741 case Triple::sparcv9: T.setArch(Triple::sparc); break;
758742 case Triple::x86_64: T.setArch(Triple::x86); break;
759743 }
782766 case Triple::mips64el:
783767 case Triple::nvptx64:
784768 case Triple::ppc64:
785 case Triple::ptx64:
786769 case Triple::sparcv9:
787770 case Triple::x86_64:
788771 // Already 64-bit.
792775 case Triple::mipsel: T.setArch(Triple::mips64el); break;
793776 case Triple::nvptx: T.setArch(Triple::nvptx64); break;
794777 case Triple::ppc: T.setArch(Triple::ppc64); break;
795 case Triple::ptx32: T.setArch(Triple::ptx64); break;
796778 case Triple::sparc: T.setArch(Triple::sparcv9); break;
797779 case Triple::x86: T.setArch(Triple::x86_64); break;
798780 }
1515 ;===------------------------------------------------------------------------===;
1616
1717 [common]
18 subdirectories = ARM CellSPU CppBackend Hexagon MBlaze MSP430 NVPTX Mips PTX PowerPC Sparc X86 XCore
18 subdirectories = ARM CellSPU CppBackend Hexagon MBlaze MSP430 NVPTX Mips PowerPC Sparc X86 XCore
1919
2020 ; This is a special group whose required libraries are extended (by llvm-build)
2121 ; with the best execution engine (the native JIT, if available, or the
+0
-32
lib/Target/PTX/CMakeLists.txt less more
None set(LLVM_TARGET_DEFINITIONS PTX.td)
1
2 tablegen(LLVM PTXGenAsmWriter.inc -gen-asm-writer)
3 tablegen(LLVM PTXGenDAGISel.inc -gen-dag-isel)
4 tablegen(LLVM PTXGenInstrInfo.inc -gen-instr-info)
5 tablegen(LLVM PTXGenRegisterInfo.inc -gen-register-info)
6 tablegen(LLVM PTXGenSubtargetInfo.inc -gen-subtarget)
7 add_public_tablegen_target(PTXCommonTableGen)
8
9 add_llvm_target(PTXCodeGen
10 PTXAsmPrinter.cpp
11 PTXISelDAGToDAG.cpp
12 PTXISelLowering.cpp
13 PTXInstrInfo.cpp
14 PTXFPRoundingModePass.cpp
15 PTXFrameLowering.cpp
16 PTXMCAsmStreamer.cpp
17 PTXMCInstLower.cpp
18 PTXMFInfoExtract.cpp
19 PTXMachineFunctionInfo.cpp
20 PTXParamManager.cpp
21 PTXRegAlloc.cpp
22 PTXRegisterInfo.cpp
23 PTXSelectionDAGInfo.cpp
24 PTXSubtarget.cpp
25 PTXTargetMachine.cpp
26 )
27
28 add_subdirectory(TargetInfo)
29 add_subdirectory(InstPrinter)
30 add_subdirectory(MCTargetDesc)
31
+0
-8
lib/Target/PTX/InstPrinter/CMakeLists.txt less more
None include_directories( ${CMAKE_CURRENT_BINARY_DIR}/.. ${CMAKE_CURRENT_SOURCE_DIR}/.. )
1
2 add_llvm_library(LLVMPTXAsmPrinter
3 PTXInstPrinter.cpp
4 )
5
6 add_dependencies(LLVMPTXAsmPrinter PTXCommonTableGen)
7
+0
-23
lib/Target/PTX/InstPrinter/LLVMBuild.txt less more
None ;===- ./lib/Target/PTX/InstPrinter/LLVMBuild.txt ---------------*- Conf -*--===;
1 ;
2 ; The LLVM Compiler Infrastructure
3 ;
4 ; This file is distributed under the University of Illinois Open Source
5 ; License. See LICENSE.TXT for details.
6 ;
7 ;===------------------------------------------------------------------------===;
8 ;
9 ; This is an LLVMBuild description file for the components in this subdirectory.
10 ;
11 ; For more information on the LLVMBuild system, please see:
12 ;
13 ; http://llvm.org/docs/LLVMBuild.html
14 ;
15 ;===------------------------------------------------------------------------===;
16
17 [component_0]
18 type = Library
19 name = PTXAsmPrinter
20 parent = PTX
21 required_libraries = MC Support
22 add_to_library_groups = PTX
+0
-16
lib/Target/PTX/InstPrinter/Makefile less more
None ##===- lib/Target/PTX/AsmPrinter/Makefile ------------------*- Makefile -*-===##
1 #
2 # The LLVM Compiler Infrastructure
3 #
4 # This file is distributed under the University of Illinois Open Source
5 # License. See LICENSE.TXT for details.
6 #
7 ##===----------------------------------------------------------------------===##
8 LEVEL = ../../../..
9 LIBRARYNAME = LLVMPTXAsmPrinter
10
11 # Hack: we need to include 'main' ptx target directory to grab private headers
12 CPP.Flags += -I$(PROJ_OBJ_DIR)/.. -I$(PROJ_SRC_DIR)/..
13
14 include $(LEVEL)/Makefile.common
15
+0
-249
lib/Target/PTX/InstPrinter/PTXInstPrinter.cpp less more
None //===-- PTXInstPrinter.cpp - Convert PTX MCInst to assembly syntax --------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This class prints a PTX MCInst to a .ptx file.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #define DEBUG_TYPE "asm-printer"
14 #include "PTXInstPrinter.h"
15 #include "MCTargetDesc/PTXBaseInfo.h"
16 #include "llvm/MC/MCAsmInfo.h"
17 #include "llvm/MC/MCExpr.h"
18 #include "llvm/MC/MCInst.h"
19 #include "llvm/MC/MCSymbol.h"
20 #include "llvm/MC/MCInstrInfo.h"
21 #include "llvm/ADT/APFloat.h"
22 #include "llvm/ADT/StringExtras.h"
23 #include "llvm/Support/ErrorHandling.h"
24 #include "llvm/Support/raw_ostream.h"
25 using namespace llvm;
26
27 #include "PTXGenAsmWriter.inc"
28
29 PTXInstPrinter::PTXInstPrinter(const MCAsmInfo &MAI,
30 const MCInstrInfo &MII,
31 const MCRegisterInfo &MRI,
32 const MCSubtargetInfo &STI) :
33 MCInstPrinter(MAI, MII, MRI) {
34 // Initialize the set of available features.
35 setAvailableFeatures(STI.getFeatureBits());
36 }
37
38 void PTXInstPrinter::printRegName(raw_ostream &OS, unsigned RegNo) const {
39 // Decode the register number into type and offset
40 unsigned RegSpace = RegNo & 0x7;
41 unsigned RegType = (RegNo >> 3) & 0x7;
42 unsigned RegOffset = RegNo >> 6;
43
44 // Print the register
45 OS << "%";
46
47 switch (RegSpace) {
48 default:
49 llvm_unreachable("Unknown register space!");
50 case PTXRegisterSpace::Reg:
51 switch (RegType) {
52 default:
53 llvm_unreachable("Unknown register type!");
54 case PTXRegisterType::Pred:
55 OS << "p";
56 break;
57 case PTXRegisterType::B16:
58 OS << "rh";
59 break;
60 case PTXRegisterType::B32:
61 OS << "r";
62 break;
63 case PTXRegisterType::B64:
64 OS << "rd";
65 break;
66 case PTXRegisterType::F32:
67 OS << "f";
68 break;
69 case PTXRegisterType::F64:
70 OS << "fd";
71 break;
72 }
73 break;
74 case PTXRegisterSpace::Return:
75 OS << "ret";
76 break;
77 case PTXRegisterSpace::Argument:
78 OS << "arg";
79 break;
80 }
81
82 OS << RegOffset;
83 }
84
85 void PTXInstPrinter::printInst(const MCInst *MI, raw_ostream &O,
86 StringRef Annot) {
87 printPredicate(MI, O);
88 switch (MI->getOpcode()) {
89 default:
90 printInstruction(MI, O);
91 break;
92 case PTX::CALL:
93 printCall(MI, O);
94 }
95 O << ";";
96 printAnnotation(O, Annot);
97 }
98
99 void PTXInstPrinter::printPredicate(const MCInst *MI, raw_ostream &O) {
100 // The last two operands are the predicate operands
101 int RegIndex;
102 int OpIndex;
103
104 if (MI->getOpcode() == PTX::CALL) {
105 RegIndex = 0;
106 OpIndex = 1;
107 } else {
108 RegIndex = MI->getNumOperands()-2;
109 OpIndex = MI->getNumOperands()-1;
110 }
111
112 int PredOp = MI->getOperand(OpIndex).getImm();
113 if (PredOp == PTXPredicate::None)
114 return;
115
116 if (PredOp == PTXPredicate::Negate)
117 O << '!';
118 else
119 O << '@';
120
121 printOperand(MI, RegIndex, O);
122 }
123
124 void PTXInstPrinter::printCall(const MCInst *MI, raw_ostream &O) {
125 O << "\tcall.uni\t";
126 // The first two operands are the predicate slot
127 unsigned Index = 2;
128 unsigned NumRets = MI->getOperand(Index++).getImm();
129
130 if (NumRets > 0) {
131 O << "(";
132 printOperand(MI, Index++, O);
133 for (unsigned i = 1; i < NumRets; ++i) {
134 O << ", ";
135 printOperand(MI, Index++, O);
136 }
137 O << "), ";
138 }
139
140 const MCExpr* Expr = MI->getOperand(Index++).getExpr();
141 unsigned NumArgs = MI->getOperand(Index++).getImm();
142
143 // if the function call is to printf or puts, change to vprintf
144 if (const MCSymbolRefExpr *SymRefExpr = dyn_cast(Expr)) {
145 const MCSymbol &Sym = SymRefExpr->getSymbol();
146 if (Sym.getName() == "printf" || Sym.getName() == "puts") {
147 O << "vprintf";
148 } else {
149 O << Sym.getName();
150 }
151 } else {
152 O << *Expr;
153 }
154
155 O << ", (";
156
157 if (NumArgs > 0) {
158 printOperand(MI, Index++, O);
159 for (unsigned i = 1; i < NumArgs; ++i) {
160 O << ", ";
161 printOperand(MI, Index++, O);
162 }
163 }
164 O << ")";
165 }
166
167 void PTXInstPrinter::printOperand(const MCInst *MI, unsigned OpNo,
168 raw_ostream &O) {
169 const MCOperand &Op = MI->getOperand(OpNo);
170 if (Op.isImm()) {
171 O << Op.getImm();
172 } else if (Op.isFPImm()) {
173 double Imm = Op.getFPImm();
174 APFloat FPImm(Imm);
175 APInt FPIntImm = FPImm.bitcastToAPInt();
176 O << "0D";
177 // PTX requires us to output the full 64 bits, even if the number is zero
178 if (FPIntImm.getZExtValue() > 0) {
179 O << FPIntImm.toString(16, false);
180 } else {
181 O << "0000000000000000";
182 }
183 } else if (Op.isReg()) {
184 printRegName(O, Op.getReg());
185 } else {
186 assert(Op.isExpr() && "unknown operand kind in printOperand");
187 const MCExpr *Expr = Op.getExpr();
188 if (const MCSymbolRefExpr *SymRefExpr = dyn_cast(Expr)) {
189 const MCSymbol &Sym = SymRefExpr->getSymbol();
190 O << Sym.getName();
191 } else {
192 O << *Op.getExpr();
193 }
194 }
195 }
196
197 void PTXInstPrinter::printMemOperand(const MCInst *MI, unsigned OpNo,
198 raw_ostream &O) {
199 // By definition, operand OpNo+1 is an i32imm
200 const MCOperand &Op2 = MI->getOperand(OpNo+1);
201 printOperand(MI, OpNo, O);
202 if (Op2.getImm() == 0)
203 return; // don't print "+0"
204 O << "+" << Op2.getImm();
205 }
206
207 void PTXInstPrinter::printRoundingMode(const MCInst *MI, unsigned OpNo,
208 raw_ostream &O) {
209 const MCOperand &Op = MI->getOperand(OpNo);
210 assert (Op.isImm() && "Rounding modes must be immediate values");
211 switch (Op.getImm()) {
212 default:
213 llvm_unreachable("Unknown rounding mode!");
214 case PTXRoundingMode::RndDefault:
215 llvm_unreachable("FP rounding-mode pass did not handle instruction!");
216 case PTXRoundingMode::RndNone:
217 // Do not print anything.
218 break;
219 case PTXRoundingMode::RndNearestEven:
220 O << ".rn";
221 break;
222 case PTXRoundingMode::RndTowardsZero:
223 O << ".rz";
224 break;
225 case PTXRoundingMode::RndNegInf:
226 O << ".rm";
227 break;
228 case PTXRoundingMode::RndPosInf:
229 O << ".rp";
230 break;
231 case PTXRoundingMode::RndApprox:
232 O << ".approx";
233 break;
234 case PTXRoundingMode::RndNearestEvenInt:
235 O << ".rni";
236 break;
237 case PTXRoundingMode::RndTowardsZeroInt:
238 O << ".rzi";
239 break;
240 case PTXRoundingMode::RndNegInfInt:
241 O << ".rmi";
242 break;
243 case PTXRoundingMode::RndPosInfInt:
244 O << ".rpi";
245 break;
246 }
247 }
248
+0
-45
lib/Target/PTX/InstPrinter/PTXInstPrinter.h less more
None //===- PTXInstPrinter.h - Convert PTX MCInst to assembly syntax -*- C++ -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This class prints n PTX MCInst to a .ptx file.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTXINSTPRINTER_H
14 #define PTXINSTPRINTER_H
15
16 #include "llvm/MC/MCInstPrinter.h"
17 #include "llvm/MC/MCSubtargetInfo.h"
18
19 namespace llvm {
20
21 class MCOperand;
22
23 class PTXInstPrinter : public MCInstPrinter {
24 public:
25 PTXInstPrinter(const MCAsmInfo &MAI, const MCInstrInfo &MII,
26 const MCRegisterInfo &MRI, const MCSubtargetInfo &STI);
27
28 virtual void printInst(const MCInst *MI, raw_ostream &O, StringRef Annot);
29 virtual void printRegName(raw_ostream &OS, unsigned RegNo) const;
30
31 // Autogenerated by tblgen.
32 void printInstruction(const MCInst *MI, raw_ostream &O);
33 static const char *getRegisterName(unsigned RegNo);
34
35 void printPredicate(const MCInst *MI, raw_ostream &O);
36 void printCall(const MCInst *MI, raw_ostream &O);
37 void printOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
38 void printMemOperand(const MCInst *MI, unsigned OpNo, raw_ostream &O);
39 void printRoundingMode(const MCInst *MI, unsigned OpNo, raw_ostream &O);
40 };
41 }
42
43 #endif
44
+0
-32
lib/Target/PTX/LLVMBuild.txt less more
None ;===- ./lib/Target/PTX/LLVMBuild.txt ---------------------------*- Conf -*--===;
1 ;
2 ; The LLVM Compiler Infrastructure
3 ;
4 ; This file is distributed under the University of Illinois Open Source
5 ; License. See LICENSE.TXT for details.
6 ;
7 ;===------------------------------------------------------------------------===;
8 ;
9 ; This is an LLVMBuild description file for the components in this subdirectory.
10 ;
11 ; For more information on the LLVMBuild system, please see:
12 ;
13 ; http://llvm.org/docs/LLVMBuild.html
14 ;
15 ;===------------------------------------------------------------------------===;
16
17 [common]
18 subdirectories = InstPrinter MCTargetDesc TargetInfo
19
20 [component_0]
21 type = TargetGroup
22 name = PTX
23 parent = Target
24 has_asmprinter = 1
25
26 [component_1]
27 type = Library
28 name = PTXCodeGen
29 parent = PTX
30 required_libraries = Analysis AsmPrinter CodeGen Core MC PTXDesc PTXInfo SelectionDAG Support Target TransformUtils
31 add_to_library_groups = PTX
+0
-6
lib/Target/PTX/MCTargetDesc/CMakeLists.txt less more
None add_llvm_library(LLVMPTXDesc
1 PTXMCTargetDesc.cpp
2 PTXMCAsmInfo.cpp
3 )
4
5 add_dependencies(LLVMPTXDesc PTXCommonTableGen)
+0
-23
lib/Target/PTX/MCTargetDesc/LLVMBuild.txt less more
None ;===- ./lib/Target/PTX/MCTargetDesc/LLVMBuild.txt --------------*- Conf -*--===;
1 ;
2 ; The LLVM Compiler Infrastructure
3 ;
4 ; This file is distributed under the University of Illinois Open Source
5 ; License. See LICENSE.TXT for details.
6 ;
7 ;===------------------------------------------------------------------------===;
8 ;
9 ; This is an LLVMBuild description file for the components in this subdirectory.
10 ;
11 ; For more information on the LLVMBuild system, please see:
12 ;
13 ; http://llvm.org/docs/LLVMBuild.html
14 ;
15 ;===------------------------------------------------------------------------===;
16
17 [component_0]
18 type = Library
19 name = PTXDesc
20 parent = PTX
21 required_libraries = MC PTXAsmPrinter PTXInfo Support
22 add_to_library_groups = PTX
+0
-16
lib/Target/PTX/MCTargetDesc/Makefile less more
None ##===- lib/Target/PTX/TargetDesc/Makefile ------------------*- Makefile -*-===##
1 #
2 # The LLVM Compiler Infrastructure
3 #
4 # This file is distributed under the University of Illinois Open Source
5 # License. See LICENSE.TXT for details.
6 #
7 ##===----------------------------------------------------------------------===##
8
9 LEVEL = ../../../..
10 LIBRARYNAME = LLVMPTXDesc
11
12 # Hack: we need to include 'main' target directory to grab private headers
13 CPP.Flags += -I$(PROJ_OBJ_DIR)/.. -I$(PROJ_SRC_DIR)/..
14
15 include $(LEVEL)/Makefile.common
+0
-134
lib/Target/PTX/MCTargetDesc/PTXBaseInfo.h less more
None //===-- PTXBaseInfo.h - Top level definitions for PTX -------- --*- C++ -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains small standalone helper functions and enum definitions for
10 // the PTX target useful for the compiler back-end and the MC libraries.
11 // As such, it deliberately does not include references to LLVM core
12 // code gen types, passes, etc..
13 //
14 //===----------------------------------------------------------------------===//
15
16 #ifndef PTXBASEINFO_H
17 #define PTXBASEINFO_H
18
19 #include "PTXMCTargetDesc.h"
20 #include "llvm/Support/ErrorHandling.h"
21 #include "llvm/Support/raw_ostream.h"
22
23 namespace llvm {
24 namespace PTXStateSpace {
25 enum {
26 Global = 0, // default to global state space
27 Constant = 1,
28 Local = 2,
29 Parameter = 3,
30 Shared = 4
31 };
32 } // namespace PTXStateSpace
33
34 namespace PTXPredicate {
35 enum {
36 Normal = 0,
37 Negate = 1,
38 None = 2
39 };
40 } // namespace PTXPredicate
41
42 /// Namespace to hold all target-specific flags.
43 namespace PTXRoundingMode {
44 // Instruction Flags
45 enum {
46 // Rounding Mode Flags
47 RndMask = 15,
48 RndDefault = 0, // ---
49 RndNone = 1, //
50 RndNearestEven = 2, // .rn
51 RndTowardsZero = 3, // .rz
52 RndNegInf = 4, // .rm
53 RndPosInf = 5, // .rp
54 RndApprox = 6, // .approx
55 RndNearestEvenInt = 7, // .rni
56 RndTowardsZeroInt = 8, // .rzi
57 RndNegInfInt = 9, // .rmi
58 RndPosInfInt = 10 // .rpi
59 };
60 } // namespace PTXII
61
62 namespace PTXRegisterType {
63 // Register type encoded in MCOperands
64 enum {
65 Pred = 0,
66 B16,
67 B32,
68 B64,
69 F32,
70 F64
71 };
72 } // namespace PTXRegisterType
73
74 namespace PTXRegisterSpace {
75 // Register space encoded in MCOperands
76 enum {
77 Reg = 0,
78 Local,
79 Param,
80 Argument,
81 Return
82 };
83 }
84
85 inline static void decodeRegisterName(raw_ostream &OS,
86 unsigned EncodedReg) {
87 OS << "%";
88
89 unsigned RegSpace = EncodedReg & 0x7;
90 unsigned RegType = (EncodedReg >> 3) & 0x7;
91 unsigned RegOffset = EncodedReg >> 6;
92
93 switch (RegSpace) {
94 default:
95 llvm_unreachable("Unknown register space!");
96 case PTXRegisterSpace::Reg:
97 switch (RegType) {
98 default:
99 llvm_unreachable("Unknown register type!");
100 case PTXRegisterType::Pred:
101 OS << "p";
102 break;
103 case PTXRegisterType::B16:
104 OS << "rh";
105 break;
106 case PTXRegisterType::B32:
107 OS << "r";
108 break;
109 case PTXRegisterType::B64:
110 OS << "rd";
111 break;
112 case PTXRegisterType::F32:
113 OS << "f";
114 break;
115 case PTXRegisterType::F64:
116 OS << "fd";
117 break;
118 }
119 break;
120 case PTXRegisterSpace::Return:
121 OS << "ret";
122 break;
123 case PTXRegisterSpace::Argument:
124 OS << "arg";
125 break;
126 }
127
128 OS << RegOffset;
129 }
130 } // namespace llvm
131
132 #endif
133
+0
-37
lib/Target/PTX/MCTargetDesc/PTXMCAsmInfo.cpp less more
None //===-- PTXMCAsmInfo.cpp - PTX asm properties -----------------------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains the declarations of the PTXMCAsmInfo properties.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #include "PTXMCAsmInfo.h"
14 #include "llvm/ADT/Triple.h"
15
16 using namespace llvm;
17
18 void PTXMCAsmInfo::anchor() { }
19
20 PTXMCAsmInfo::PTXMCAsmInfo(const Target &T, const StringRef &TT) {
21 Triple TheTriple(TT);
22 if (TheTriple.getArch() == Triple::ptx64)
23 PointerSize = 8;
24
25 CommentString = "//";
26
27 PrivateGlobalPrefix = "$L__";
28
29 AllowPeriodsInName = false;
30
31 HasSetDirective = false;
32
33 HasDotTypeDotSizeDirective = false;
34
35 HasSingleParameterDotFile = false;
36 }
+0
-30
lib/Target/PTX/MCTargetDesc/PTXMCAsmInfo.h less more
None //===-- PTXMCAsmInfo.h - PTX asm properties --------------------*- C++ -*--===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains the declaration of the PTXMCAsmInfo class.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTX_MCASM_INFO_H
14 #define PTX_MCASM_INFO_H
15
16 #include "llvm/MC/MCAsmInfo.h"
17
18 namespace llvm {
19 class Target;
20 class StringRef;
21
22 class PTXMCAsmInfo : public MCAsmInfo {
23 virtual void anchor();
24 public:
25 explicit PTXMCAsmInfo(const Target &T, const StringRef &TT);
26 };
27 } // namespace llvm
28
29 #endif // PTX_MCASM_INFO_H
+0
-98
lib/Target/PTX/MCTargetDesc/PTXMCTargetDesc.cpp less more
None //===-- PTXMCTargetDesc.cpp - PTX Target Descriptions ---------------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file provides PTX specific target descriptions.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #include "PTXMCTargetDesc.h"
14 #include "PTXMCAsmInfo.h"
15 #include "InstPrinter/PTXInstPrinter.h"
16 #include "llvm/MC/MCCodeGenInfo.h"
17 #include "llvm/MC/MCInstrInfo.h"
18 #include "llvm/MC/MCRegisterInfo.h"
19 #include "llvm/MC/MCSubtargetInfo.h"
20 #include "llvm/Support/TargetRegistry.h"
21
22 #define GET_INSTRINFO_MC_DESC
23 #include "PTXGenInstrInfo.inc"
24
25 #define GET_SUBTARGETINFO_MC_DESC
26 #include "PTXGenSubtargetInfo.inc"
27
28 #define GET_REGINFO_MC_DESC
29 #include "PTXGenRegisterInfo.inc"
30
31 using namespace llvm;
32
33 static MCInstrInfo *createPTXMCInstrInfo() {
34 MCInstrInfo *X = new MCInstrInfo();
35 InitPTXMCInstrInfo(X);
36 return X;
37 }
38
39 static MCRegisterInfo *createPTXMCRegisterInfo(StringRef TT) {
40 MCRegisterInfo *X = new MCRegisterInfo();
41 // PTX does not have a return address register.
42 InitPTXMCRegisterInfo(X, 0);
43 return X;
44 }
45
46 static MCSubtargetInfo *createPTXMCSubtargetInfo(StringRef TT, StringRef CPU,
47 StringRef FS) {
48 MCSubtargetInfo *X = new MCSubtargetInfo();
49 InitPTXMCSubtargetInfo(X, TT, CPU, FS);
50 return X;
51 }
52
53 static MCCodeGenInfo *createPTXMCCodeGenInfo(StringRef TT, Reloc::Model RM,
54 CodeModel::Model CM,
55 CodeGenOpt::Level OL) {
56 MCCodeGenInfo *X = new MCCodeGenInfo();
57 X->InitMCCodeGenInfo(RM, CM, OL);
58 return X;
59 }
60
61 static MCInstPrinter *createPTXMCInstPrinter(const Target &T,
62 unsigned SyntaxVariant,
63 const MCAsmInfo &MAI,
64 const MCInstrInfo &MII,
65 const MCRegisterInfo &MRI,
66 const MCSubtargetInfo &STI) {
67 assert(SyntaxVariant == 0 && "We only have one syntax variant");
68 return new PTXInstPrinter(MAI, MII, MRI, STI);
69 }
70
71 extern "C" void LLVMInitializePTXTargetMC() {
72 // Register the MC asm info.
73 RegisterMCAsmInfo X(ThePTX32Target);
74 RegisterMCAsmInfo Y(ThePTX64Target);
75
76 // Register the MC codegen info.
77 TargetRegistry::RegisterMCCodeGenInfo(ThePTX32Target, createPTXMCCodeGenInfo);
78 TargetRegistry::RegisterMCCodeGenInfo(ThePTX64Target, createPTXMCCodeGenInfo);
79
80 // Register the MC instruction info.
81 TargetRegistry::RegisterMCInstrInfo(ThePTX32Target, createPTXMCInstrInfo);
82 TargetRegistry::RegisterMCInstrInfo(ThePTX64Target, createPTXMCInstrInfo);
83
84 // Register the MC register info.
85 TargetRegistry::RegisterMCRegInfo(ThePTX32Target, createPTXMCRegisterInfo);
86 TargetRegistry::RegisterMCRegInfo(ThePTX64Target, createPTXMCRegisterInfo);
87
88 // Register the MC subtarget info.
89 TargetRegistry::RegisterMCSubtargetInfo(ThePTX32Target,
90 createPTXMCSubtargetInfo);
91 TargetRegistry::RegisterMCSubtargetInfo(ThePTX64Target,
92 createPTXMCSubtargetInfo);
93
94 // Register the MCInstPrinter.
95 TargetRegistry::RegisterMCInstPrinter(ThePTX32Target, createPTXMCInstPrinter);
96 TargetRegistry::RegisterMCInstPrinter(ThePTX64Target, createPTXMCInstPrinter);
97 }
+0
-36
lib/Target/PTX/MCTargetDesc/PTXMCTargetDesc.h less more
None //===-- PTXMCTargetDesc.h - PTX Target Descriptions ------------*- C++ -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file provides PTX specific target descriptions.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTXMCTARGETDESC_H
14 #define PTXMCTARGETDESC_H
15
16 namespace llvm {
17 class Target;
18
19 extern Target ThePTX32Target;
20 extern Target ThePTX64Target;
21
22 } // End llvm namespace
23
24 // Defines symbolic names for PTX registers.
25 #define GET_REGINFO_ENUM
26 #include "PTXGenRegisterInfo.inc"
27
28 // Defines symbolic names for the PTX instructions.
29 #define GET_INSTRINFO_ENUM
30 #include "PTXGenInstrInfo.inc"
31
32 #define GET_SUBTARGETINFO_ENUM
33 #include "PTXGenSubtargetInfo.inc"
34
35 #endif
+0
-23
lib/Target/PTX/Makefile less more
None ##===- lib/Target/PTX/Makefile -----------------------------*- Makefile -*-===##
1 #
2 # The LLVM Compiler Infrastructure
3 #
4 # This file is distributed under the University of Illinois Open Source
5 # License. See LICENSE.TXT for details.
6 #
7 ##===----------------------------------------------------------------------===##
8
9 LEVEL = ../../..
10 LIBRARYNAME = LLVMPTXCodeGen
11 TARGET = PTX
12
13 # Make sure that tblgen is run, first thing.
14 BUILT_SOURCES = PTXGenAsmWriter.inc \
15 PTXGenDAGISel.inc \
16 PTXGenInstrInfo.inc \
17 PTXGenRegisterInfo.inc \
18 PTXGenSubtargetInfo.inc
19
20 DIRS = InstPrinter TargetInfo MCTargetDesc
21
22 include $(LEVEL)/Makefile.common
+0
-43
lib/Target/PTX/PTX.h less more
None //
1 // The LLVM Compiler Infrastructure
2 //
3 // This file is distributed under the University of Illinois Open Source
4 // License. See LICENSE.TXT for details.
5 //
6 //===----------------------------------------------------------------------===//
7 //
8 // This file contains the entry points for global functions defined in the LLVM
9 // PTX back-end.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTX_H
14 #define PTX_H
15
16 #include "MCTargetDesc/PTXBaseInfo.h"
17 #include "llvm/Target/TargetMachine.h"
18
19 namespace llvm {
20 class MachineInstr;
21 class MCInst;
22 class PTXAsmPrinter;
23 class PTXTargetMachine;
24 class FunctionPass;
25
26 FunctionPass *createPTXISelDag(PTXTargetMachine &TM,
27 CodeGenOpt::Level OptLevel);
28
29 FunctionPass *createPTXMFInfoExtract(PTXTargetMachine &TM,
30 CodeGenOpt::Level OptLevel);
31
32 FunctionPass *createPTXFPRoundingModePass(PTXTargetMachine &TM,
33 CodeGenOpt::Level OptLevel);
34
35 FunctionPass *createPTXRegisterAllocator();
36
37 void LowerPTXMachineInstrToMCInst(const MachineInstr *MI, MCInst &OutMI,
38 PTXAsmPrinter &AP);
39
40 } // namespace llvm;
41
42 #endif // PTX_H
+0
-141
lib/Target/PTX/PTX.td less more
None //===-- PTX.td - Describe the PTX Target Machine -----------*- tablegen -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 // This is the top level entry point for the PTX target.
9 //===----------------------------------------------------------------------===//
10
11 //===----------------------------------------------------------------------===//
12 // Target-independent interfaces
13 //===----------------------------------------------------------------------===//
14
15 include "llvm/Target/Target.td"
16
17 //===----------------------------------------------------------------------===//
18 // Subtarget Features
19 //===----------------------------------------------------------------------===//
20
21 //===- Architectural Features ---------------------------------------------===//
22
23 def FeatureDouble : SubtargetFeature<"double", "SupportsDouble", "true",
24 "Do not demote .f64 to .f32">;
25
26 def FeatureNoFMA : SubtargetFeature<"no-fma","SupportsFMA", "false",
27 "Disable Fused-Multiply Add">;
28
29 //===- PTX Version --------------------------------------------------------===//
30
31 def FeaturePTX20 : SubtargetFeature<"ptx20", "PTXVersion", "PTX_VERSION_2_0",
32 "Use PTX Language Version 2.0">;
33
34 def FeaturePTX21 : SubtargetFeature<"ptx21", "PTXVersion", "PTX_VERSION_2_1",
35 "Use PTX Language Version 2.1">;
36
37 def FeaturePTX22 : SubtargetFeature<"ptx22", "PTXVersion", "PTX_VERSION_2_2",
38 "Use PTX Language Version 2.2">;
39
40 def FeaturePTX23 : SubtargetFeature<"ptx23", "PTXVersion", "PTX_VERSION_2_3",
41 "Use PTX Language Version 2.3">;
42
43 //===- PTX Target ---------------------------------------------------------===//
44
45 def FeatureSM10 : SubtargetFeature<"sm10", "PTXTarget", "PTX_SM_1_0",
46 "Use Shader Model 1.0">;
47 def FeatureSM11 : SubtargetFeature<"sm11", "PTXTarget", "PTX_SM_1_1",
48 "Use Shader Model 1.1">;
49 def FeatureSM12 : SubtargetFeature<"sm12", "PTXTarget", "PTX_SM_1_2",
50 "Use Shader Model 1.2">;
51 def FeatureSM13 : SubtargetFeature<"sm13", "PTXTarget", "PTX_SM_1_3",
52 "Use Shader Model 1.3">;
53 def FeatureSM20 : SubtargetFeature<"sm20", "PTXTarget", "PTX_SM_2_0",
54 "Use Shader Model 2.0", [FeatureDouble]>;
55 def FeatureSM21 : SubtargetFeature<"sm21", "PTXTarget", "PTX_SM_2_1",
56 "Use Shader Model 2.1", [FeatureDouble]>;
57 def FeatureSM22 : SubtargetFeature<"sm22", "PTXTarget", "PTX_SM_2_2",
58 "Use Shader Model 2.2", [FeatureDouble]>;
59 def FeatureSM23 : SubtargetFeature<"sm23", "PTXTarget", "PTX_SM_2_3",
60 "Use Shader Model 2.3", [FeatureDouble]>;
61
62 def FeatureCOMPUTE10 : SubtargetFeature<"compute10", "PTXTarget",
63 "PTX_COMPUTE_1_0",
64 "Use Compute Compatibility 1.0">;
65 def FeatureCOMPUTE11 : SubtargetFeature<"compute11", "PTXTarget",
66 "PTX_COMPUTE_1_1",
67 "Use Compute Compatibility 1.1">;
68 def FeatureCOMPUTE12 : SubtargetFeature<"compute12", "PTXTarget",
69 "PTX_COMPUTE_1_2",
70 "Use Compute Compatibility 1.2">;
71 def FeatureCOMPUTE13 : SubtargetFeature<"compute13", "PTXTarget",
72 "PTX_COMPUTE_1_3",
73 "Use Compute Compatibility 1.3">;
74 def FeatureCOMPUTE20 : SubtargetFeature<"compute20", "PTXTarget",
75 "PTX_COMPUTE_2_0",
76 "Use Compute Compatibility 2.0",
77 [FeatureDouble]>;
78
79 //===----------------------------------------------------------------------===//
80 // PTX supported processors
81 //===----------------------------------------------------------------------===//
82
83 class Proc Features>
84 : Processor;
85
86 def : Proc<"generic", []>;
87
88 // Processor definitions for compute/shader models
89 def : Proc<"compute_10", [FeatureCOMPUTE10]>;
90 def : Proc<"compute_11", [FeatureCOMPUTE11]>;
91 def : Proc<"compute_12", [FeatureCOMPUTE12]>;
92 def : Proc<"compute_13", [FeatureCOMPUTE13]>;
93 def : Proc<"compute_20", [FeatureCOMPUTE20]>;
94 def : Proc<"sm_10", [FeatureSM10]>;
95 def : Proc<"sm_11", [FeatureSM11]>;
96 def : Proc<"sm_12", [FeatureSM12]>;
97 def : Proc<"sm_13", [FeatureSM13]>;
98 def : Proc<"sm_20", [FeatureSM20]>;
99 def : Proc<"sm_21", [FeatureSM21]>;
100 def : Proc<"sm_22", [FeatureSM22]>;
101 def : Proc<"sm_23", [FeatureSM23]>;
102
103 // Processor definitions for common GPU architectures
104 def : Proc<"g80", [FeatureSM10]>;
105 def : Proc<"gt200", [FeatureSM13]>;
106 def : Proc<"gf100", [FeatureSM20, FeatureDouble]>;
107 def : Proc<"fermi", [FeatureSM20, FeatureDouble]>;
108
109 //===----------------------------------------------------------------------===//
110 // Register File Description
111 //===----------------------------------------------------------------------===//
112
113 include "PTXRegisterInfo.td"
114
115 //===----------------------------------------------------------------------===//
116 // Instruction Descriptions
117 //===----------------------------------------------------------------------===//
118
119 include "PTXInstrInfo.td"
120
121 def PTXInstrInfo : InstrInfo;
122
123 //===----------------------------------------------------------------------===//
124 // Assembly printer
125 //===----------------------------------------------------------------------===//
126 // PTX uses the MC printer for asm output, so make sure the TableGen
127 // AsmWriter bits get associated with the correct class.
128 def PTXAsmWriter : AsmWriter {
129 string AsmWriterClassName = "InstPrinter";
130 bit isMCAsmWriter = 1;
131 }
132
133 //===----------------------------------------------------------------------===//
134 // Target Declaration
135 //===----------------------------------------------------------------------===//
136
137 def PTX : Target {
138 let InstructionSet = PTXInstrInfo;
139 let AssemblyWriters = [PTXAsmWriter];
140 }
+0
-561
lib/Target/PTX/PTXAsmPrinter.cpp less more
None //===-- PTXAsmPrinter.cpp - PTX LLVM assembly writer ----------------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains a printer that converts from our internal representation
10 // of machine-dependent LLVM code to PTX assembly language.
11 //
12 //===----------------------------------------------------------------------===//
13
14 #define DEBUG_TYPE "ptx-asm-printer"
15
16 #include "PTXAsmPrinter.h"
17 #include "PTX.h"
18 #include "PTXMachineFunctionInfo.h"
19 #include "PTXParamManager.h"
20 #include "PTXRegisterInfo.h"
21 #include "PTXTargetMachine.h"
22 #include "llvm/Argument.h"
23 #include "llvm/DerivedTypes.h"
24 #include "llvm/Function.h"
25 #include "llvm/Module.h"
26 #include "llvm/ADT/SmallString.h"
27 #include "llvm/ADT/Twine.h"
28 #include "llvm/Analysis/DebugInfo.h"
29 #include "llvm/CodeGen/AsmPrinter.h"
30 #include "llvm/CodeGen/MachineFrameInfo.h"
31 #include "llvm/CodeGen/MachineInstr.h"
32 #include "llvm/CodeGen/MachineRegisterInfo.h"
33 #include "llvm/MC/MCContext.h"
34 #include "llvm/MC/MCExpr.h"
35 #include "llvm/MC/MCInst.h"
36 #include "llvm/MC/MCStreamer.h"
37 #include "llvm/MC/MCSymbol.h"
38 #include "llvm/Target/Mangler.h"
39 #include "llvm/Target/TargetLoweringObjectFile.h"
40 #include "llvm/Support/CommandLine.h"
41 #include "llvm/Support/Debug.h"
42 #include "llvm/Support/ErrorHandling.h"
43 #include "llvm/Support/MathExtras.h"
44 #include "llvm/Support/Path.h"
45 #include "llvm/Support/TargetRegistry.h"
46 #include "llvm/Support/raw_ostream.h"
47
48 using namespace llvm;
49
50 static const char PARAM_PREFIX[] = "__param_";
51 static const char RETURN_PREFIX[] = "__ret_";
52
53 static const char *getRegisterTypeName(unsigned RegType) {
54 switch (RegType) {
55 default:
56 llvm_unreachable("Unknown register type");
57 case PTXRegisterType::Pred:
58 return ".pred";
59 case PTXRegisterType::B16:
60 return ".b16";
61 case PTXRegisterType::B32:
62 return ".b32";
63 case PTXRegisterType::B64:
64 return ".b64";
65 case PTXRegisterType::F32:
66 return ".f32";
67 case PTXRegisterType::F64:
68 return ".f64";
69 }
70 }
71
72 static const char *getStateSpaceName(unsigned addressSpace) {
73 switch (addressSpace) {
74 default: llvm_unreachable("Unknown state space");
75 case PTXStateSpace::Global: return "global";
76 case PTXStateSpace::Constant: return "const";
77 case PTXStateSpace::Local: return "local";
78 case PTXStateSpace::Parameter: return "param";
79 case PTXStateSpace::Shared: return "shared";
80 }
81 }
82
83 static const char *getTypeName(Type* type) {
84 while (true) {
85 switch (type->getTypeID()) {
86 default: llvm_unreachable("Unknown type");
87 case Type::FloatTyID: return ".f32";
88 case Type::DoubleTyID: return ".f64";
89 case Type::IntegerTyID:
90 switch (type->getPrimitiveSizeInBits()) {
91 default: llvm_unreachable("Unknown integer bit-width");
92 case 16: return ".u16";
93 case 32: return ".u32";
94 case 64: return ".u64";
95 }
96 case Type::ArrayTyID:
97 case Type::PointerTyID:
98 type = dyn_cast(type)->getElementType();
99 break;
100 }
101 }
102 return NULL;
103 }
104
105 bool PTXAsmPrinter::doFinalization(Module &M) {
106 // XXX Temproarily remove global variables so that doFinalization() will not
107 // emit them again (global variables are emitted at beginning).
108
109 Module::GlobalListType &global_list = M.getGlobalList();
110 int i, n = global_list.size();
111 GlobalVariable **gv_array = new GlobalVariable* [n];
112
113 // first, back-up GlobalVariable in gv_array
114 i = 0;
115 for (Module::global_iterator I = global_list.begin(), E = global_list.end();
116 I != E; ++I)
117 gv_array[i++] = &*I;
118
119 // second, empty global_list
120 while (!global_list.empty())
121 global_list.remove(global_list.begin());
122
123 // call doFinalization
124 bool ret = AsmPrinter::doFinalization(M);
125
126 // now we restore global variables
127 for (i = 0; i < n; i ++)
128 global_list.insert(global_list.end(), gv_array[i]);
129
130 delete[] gv_array;
131 return ret;
132 }
133
134 void PTXAsmPrinter::EmitStartOfAsmFile(Module &M)
135 {
136 const PTXSubtarget& ST = TM.getSubtarget();
137
138 // Emit the PTX .version and .target attributes
139 OutStreamer.EmitRawText(Twine("\t.version ") + ST.getPTXVersionString());
140 OutStreamer.EmitRawText(Twine("\t.target ") + ST.getTargetString() +
141 (ST.supportsDouble() ? ""
142 : ", map_f64_to_f32"));
143 // .address_size directive is optional, but it must immediately follow
144 // the .target directive if present within a module
145 if (ST.supportsPTX23()) {
146 const char *addrSize = ST.is64Bit() ? "64" : "32";
147 OutStreamer.EmitRawText(Twine("\t.address_size ") + addrSize);
148 }
149
150 OutStreamer.AddBlankLine();
151
152 // Define any .file directives
153 DebugInfoFinder DbgFinder;
154 DbgFinder.processModule(M);
155
156 for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(),
157 E = DbgFinder.compile_unit_end(); I != E; ++I) {
158 DICompileUnit DIUnit(*I);
159 StringRef FN = DIUnit.getFilename();
160 StringRef Dir = DIUnit.getDirectory();
161 GetOrCreateSourceID(FN, Dir);
162 }
163
164 OutStreamer.AddBlankLine();
165
166 // declare external functions
167 for (Module::const_iterator i = M.begin(), e = M.end();
168 i != e; ++i)
169 EmitFunctionDeclaration(i);
170
171 // declare global variables
172 for (Module::const_global_iterator i = M.global_begin(), e = M.global_end();
173 i != e; ++i)
174 EmitVariableDeclaration(i);
175 }
176
177 void PTXAsmPrinter::EmitFunctionBodyStart() {
178 OutStreamer.EmitRawText(Twine("{"));
179
180 const PTXMachineFunctionInfo *MFI = MF->getInfo();
181 const PTXParamManager &PM = MFI->getParamManager();
182
183 // Print register definitions
184 SmallString<128> regDefs;
185 raw_svector_ostream os(regDefs);
186 unsigned numRegs;
187
188 // pred
189 numRegs = MFI->countRegisters(PTXRegisterType::Pred, PTXRegisterSpace::Reg);
190 if(numRegs > 0)
191 os << "\t.reg .pred %p<" << numRegs << ">;\n";
192
193 // i16
194 numRegs = MFI->countRegisters(PTXRegisterType::B16, PTXRegisterSpace::Reg);
195 if(numRegs > 0)
196 os << "\t.reg .b16 %rh<" << numRegs << ">;\n";
197
198 // i32
199 numRegs = MFI->countRegisters(PTXRegisterType::B32, PTXRegisterSpace::Reg);
200 if(numRegs > 0)
201 os << "\t.reg .b32 %r<" << numRegs << ">;\n";
202
203 // i64
204 numRegs = MFI->countRegisters(PTXRegisterType::B64, PTXRegisterSpace::Reg);
205 if(numRegs > 0)
206 os << "\t.reg .b64 %rd<" << numRegs << ">;\n";
207
208 // f32
209 numRegs = MFI->countRegisters(PTXRegisterType::F32, PTXRegisterSpace::Reg);
210 if(numRegs > 0)
211 os << "\t.reg .f32 %f<" << numRegs << ">;\n";
212
213 // f64
214 numRegs = MFI->countRegisters(PTXRegisterType::F64, PTXRegisterSpace::Reg);
215 if(numRegs > 0)
216 os << "\t.reg .f64 %fd<" << numRegs << ">;\n";
217
218 // Local params
219 for (PTXParamManager::param_iterator i = PM.local_begin(), e = PM.local_end();
220 i != e; ++i)
221 os << "\t.param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i)
222 << ";\n";
223
224 OutStreamer.EmitRawText(os.str());
225
226
227 const MachineFrameInfo* FrameInfo = MF->getFrameInfo();
228 DEBUG(dbgs() << "Have " << FrameInfo->getNumObjects()
229 << " frame object(s)\n");
230 for (unsigned i = 0, e = FrameInfo->getNumObjects(); i != e; ++i) {
231 DEBUG(dbgs() << "Size of object: " << FrameInfo->getObjectSize(i) << "\n");
232 if (FrameInfo->getObjectSize(i) > 0) {
233 OutStreamer.EmitRawText("\t.local .align " +
234 Twine(FrameInfo->getObjectAlignment(i)) +
235 " .b8 __local" +
236 Twine(i) +
237 "[" +
238 Twine(FrameInfo->getObjectSize(i)) +
239 "];");
240 }
241 }
242
243 //unsigned Index = 1;
244 // Print parameter passing params
245 //for (PTXMachineFunctionInfo::param_iterator
246 // i = MFI->paramBegin(), e = MFI->paramEnd(); i != e; ++i) {
247 // std::string def = "\t.param .b";
248 // def += utostr(*i);
249 // def += " __ret_";
250 // def += utostr(Index);
251 // Index++;
252 // def += ";";
253 // OutStreamer.EmitRawText(Twine(def));
254 //}
255 }
256
257 void PTXAsmPrinter::EmitFunctionBodyEnd() {
258 OutStreamer.EmitRawText(Twine("}"));
259 }
260
261 void PTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {
262 MCInst TmpInst;
263 LowerPTXMachineInstrToMCInst(MI, TmpInst, *this);
264 OutStreamer.EmitInstruction(TmpInst);
265 }
266
267 void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) {
268 // Check to see if this is a special global used by LLVM, if so, emit it.
269 if (EmitSpecialLLVMGlobal(gv))
270 return;
271
272 MCSymbol *gvsym = Mang->getSymbol(gv);
273
274 assert(gvsym->isUndefined() && "Cannot define a symbol twice!");
275
276 SmallString<128> decl;
277 raw_svector_ostream os(decl);
278
279 // check if it is defined in some other translation unit
280 if (gv->isDeclaration())
281 os << ".extern ";
282
283 // state space: e.g., .global
284 os << '.' << getStateSpaceName(gv->getType()->getAddressSpace()) << ' ';
285
286 // alignment (optional)
287 unsigned alignment = gv->getAlignment();
288 if (alignment != 0)
289 os << ".align " << gv->getAlignment() << ' ';
290
291
292 if (PointerType::classof(gv->getType())) {
293 PointerType* pointerTy = dyn_cast(gv->getType());
294 Type* elementTy = pointerTy->getElementType();
295
296 if (elementTy->isArrayTy()) {
297 assert(elementTy->isArrayTy() && "Only pointers to arrays are supported");
298
299 ArrayType* arrayTy = dyn_cast(elementTy);
300 elementTy = arrayTy->getElementType();
301
302 unsigned numElements = arrayTy->getNumElements();
303
304 while (elementTy->isArrayTy()) {
305 arrayTy = dyn_cast(elementTy);
306 elementTy = arrayTy->getElementType();
307
308 numElements *= arrayTy->getNumElements();
309 }
310
311 // FIXME: isPrimitiveType() == false for i16?
312 assert(elementTy->isSingleValueType() &&
313 "Non-primitive types are not handled");
314
315 // Find the size of the element in bits
316 unsigned elementSize = elementTy->getPrimitiveSizeInBits();
317
318 os << ".b" << elementSize << ' ' << gvsym->getName()
319 << '[' << numElements << ']';
320 } else {
321 os << ".b8" << gvsym->getName() << "[]";
322 }
323
324 // handle string constants (assume ConstantArray means string)
325 if (gv->hasInitializer()) {
326 const Constant *C = gv->getInitializer();
327 if (const ConstantArray *CA = dyn_cast(C)) {
328 os << " = {";
329
330 for (unsigned i = 0, e = C->getNumOperands(); i != e; ++i) {
331 if (i > 0)
332 os << ',';
333
334 os << "0x";
335 os.write_hex(cast(CA->getOperand(i))->getZExtValue());
336 }
337
338 os << '}';
339 }
340 }
341 } else {
342 // Note: this is currently the fall-through case and most likely generates
343 // incorrect code.
344 os << getTypeName(gv->getType()) << ' ' << gvsym->getName();
345
346 if (isa(gv->getType()) || isa(gv->getType()))
347 os << "[]";
348 }
349
350 os << ';';
351
352 OutStreamer.EmitRawText(os.str());
353 OutStreamer.AddBlankLine();
354 }
355
356 void PTXAsmPrinter::EmitFunctionEntryLabel() {
357 // The function label could have already been emitted if two symbols end up
358 // conflicting due to asm renaming. Detect this and emit an error.
359 if (!CurrentFnSym->isUndefined())
360 report_fatal_error("'" + Twine(CurrentFnSym->getName()) +
361 "' label emitted multiple times to assembly file");
362
363 const PTXMachineFunctionInfo *MFI = MF->getInfo();
364 const PTXParamManager &PM = MFI->getParamManager();
365 const bool isKernel = MFI->isKernel();
366 const PTXSubtarget& ST = TM.getSubtarget();
367
368 SmallString<128> decl;
369 raw_svector_ostream os(decl);
370 os << (isKernel ? ".entry" : ".func");
371
372 if (!isKernel) {
373 os << " (";
374 if (ST.useParamSpaceForDeviceArgs()) {
375 for (PTXParamManager::param_iterator i = PM.ret_begin(), e = PM.ret_end(),
376 b = i; i != e; ++i) {
377 if (i != b)
378 os << ", ";
379
380 os << ".param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i);
381 }
382 } else {
383 for (PTXMachineFunctionInfo::reg_iterator
384 i = MFI->retreg_begin(), e = MFI->retreg_end(), b = i;
385 i != e; ++i) {
386 if (i != b)
387 os << ", ";
388
389 os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
390 << MFI->getRegisterName(*i);
391 }
392 }
393 os << ')';
394 }
395
396 // Print function name
397 os << ' ' << CurrentFnSym->getName() << " (";
398
399 const Function *F = MF->getFunction();
400
401 // Print parameters
402 if (isKernel || ST.useParamSpaceForDeviceArgs()) {
403 /*for (PTXParamManager::param_iterator i = PM.arg_begin(), e = PM.arg_end(),
404 b = i; i != e; ++i) {
405 if (i != b)
406 os << ", ";
407
408 os << ".param .b" << PM.getParamSize(*i) << ' ' << PM.getParamName(*i);
409 }*/
410 int Counter = 1;
411 for (Function::const_arg_iterator i = F->arg_begin(), e = F->arg_end(),
412 b = i; i != e; ++i) {
413 if (i != b)
414 os << ", ";
415 const Type *ArgType = (*i).getType();
416 os << ".param .b";
417 if (ArgType->isPointerTy()) {
418 if (ST.is64Bit())
419 os << "64";
420 else
421 os << "32";
422 } else {
423 os << ArgType->getPrimitiveSizeInBits();
424 }
425 if (ArgType->isPointerTy() && ST.emitPtrAttribute()) {
426 const PointerType *PtrType = dyn_cast(ArgType);
427 os << " .ptr";
428 switch (PtrType->getAddressSpace()) {
429 default:
430 llvm_unreachable("Unknown address space in argument");
431 case PTXStateSpace::Global:
432 os << " .global";
433 break;
434 case PTXStateSpace::Shared:
435 os << " .shared";
436 break;
437 }
438 }
439 os << " __param_" << Counter++;
440 }
441 } else {
442 for (PTXMachineFunctionInfo::reg_iterator
443 i = MFI->argreg_begin(), e = MFI->argreg_end(), b = i;
444 i != e; ++i) {
445 if (i != b)
446 os << ", ";
447
448 os << ".reg " << getRegisterTypeName(MFI->getRegisterType(*i)) << ' '
449 << MFI->getRegisterName(*i);
450 }
451 }
452 os << ')';
453
454 OutStreamer.EmitRawText(os.str());
455 }
456
457 void PTXAsmPrinter::EmitFunctionDeclaration(const Function* func)
458 {
459 const PTXSubtarget& ST = TM.getSubtarget();
460
461 std::string decl = "";
462
463 // hard-coded emission of extern vprintf function
464
465 if (func->getName() == "printf" || func->getName() == "puts") {
466 decl += ".extern .func (.param .b32 __param_1) vprintf (.param .b";
467 if (ST.is64Bit())
468 decl += "64";
469 else
470 decl += "32";
471 decl += " __param_2, .param .b";
472 if (ST.is64Bit())
473 decl += "64";
474 else
475 decl += "32";
476 decl += " __param_3)\n";
477 }
478
479 OutStreamer.EmitRawText(Twine(decl));
480 }
481
482 unsigned PTXAsmPrinter::GetOrCreateSourceID(StringRef FileName,
483 StringRef DirName) {
484 // If FE did not provide a file name, then assume stdin.
485 if (FileName.empty())
486 return GetOrCreateSourceID("", StringRef());
487
488 // MCStream expects full path name as filename.
489 if (!DirName.empty() && !sys::path::is_absolute(FileName)) {
490 SmallString<128> FullPathName = DirName;
491 sys::path::append(FullPathName, FileName);
492 // Here FullPathName will be copied into StringMap by GetOrCreateSourceID.
493 return GetOrCreateSourceID(StringRef(FullPathName), StringRef());
494 }
495
496 StringMapEntry &Entry = SourceIdMap.GetOrCreateValue(FileName);
497 if (Entry.getValue())
498 return Entry.getValue();
499
500 unsigned SrcId = SourceIdMap.size();
501 Entry.setValue(SrcId);
502
503 // Print out a .file directive to specify files for .loc directives.
504 OutStreamer.EmitDwarfFileDirective(SrcId, "", Entry.getKey());
505
506 return SrcId;
507 }
508
509 MCOperand PTXAsmPrinter::GetSymbolRef(const MachineOperand &MO,
510 const MCSymbol *Symbol) {
511 const MCExpr *Expr;
512 Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, OutContext);
513 return MCOperand::CreateExpr(Expr);
514 }
515
516 MCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
517 MCOperand MCOp;
518 const PTXMachineFunctionInfo *MFI = MF->getInfo();
519 unsigned EncodedReg;
520 switch (MO.getType()) {
521 default:
522 llvm_unreachable("Unknown operand type");
523 case MachineOperand::MO_Register:
524 if (MO.getReg() > 0) {
525 // Encode the register
526 EncodedReg = MFI->getEncodedRegister(MO.getReg());
527 } else {
528 EncodedReg = 0;
529 }
530 MCOp = MCOperand::CreateReg(EncodedReg);
531 break;
532 case MachineOperand::MO_Immediate:
533 MCOp = MCOperand::CreateImm(MO.getImm());
534 break;
535 case MachineOperand::MO_MachineBasicBlock:
536 MCOp = MCOperand::CreateExpr(MCSymbolRefExpr::Create(
537 MO.getMBB()->getSymbol(), OutContext));
538 break;
539 case MachineOperand::MO_GlobalAddress:
540 MCOp = GetSymbolRef(MO, Mang->getSymbol(MO.getGlobal()));
541 break;
542 case MachineOperand::MO_ExternalSymbol:
543 MCOp = GetSymbolRef(MO, GetExternalSymbolSymbol(MO.getSymbolName()));
544 break;
545 case MachineOperand::MO_FPImmediate:
546 APFloat Val = MO.getFPImm()->getValueAPF();
547 bool ignored;
548 Val.convert(APFloat::IEEEdouble, APFloat::rmTowardZero, &ignored);
549 MCOp = MCOperand::CreateFPImm(Val.convertToDouble());
550 break;
551 }
552
553 return MCOp;
554 }
555
556 // Force static initialization.
557 extern "C" void LLVMInitializePTXAsmPrinter() {
558 RegisterAsmPrinter X(ThePTX32Target);
559 RegisterAsmPrinter Y(ThePTX64Target);
560 }
+0
-57
lib/Target/PTX/PTXAsmPrinter.h less more
None //===-- PTXAsmPrinter.h - Print machine code to a PTX file ------*- C++ -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // PTX Assembly printer class.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTXASMPRINTER_H
14 #define PTXASMPRINTER_H
15
16 #include "PTX.h"
17 #include "PTXTargetMachine.h"
18 #include "llvm/ADT/StringMap.h"
19 #include "llvm/CodeGen/AsmPrinter.h"
20 #include "llvm/Support/Compiler.h"
21
22 namespace llvm {
23
24 class MCOperand;
25
26 class LLVM_LIBRARY_VISIBILITY PTXAsmPrinter : public AsmPrinter {
27 public:
28 explicit PTXAsmPrinter(TargetMachine &TM, MCStreamer &Streamer)
29 : AsmPrinter(TM, Streamer) {}
30
31 const char *getPassName() const { return "PTX Assembly Printer"; }
32
33 bool doFinalization(Module &M);
34
35 virtual void EmitStartOfAsmFile(Module &M);
36 virtual void EmitFunctionBodyStart();
37 virtual void EmitFunctionBodyEnd();
38 virtual void EmitFunctionEntryLabel();
39 virtual void EmitInstruction(const MachineInstr *MI);
40
41 unsigned GetOrCreateSourceID(StringRef FileName,
42 StringRef DirName);
43
44 MCOperand GetSymbolRef(const MachineOperand &MO, const MCSymbol *Symbol);
45 MCOperand lowerOperand(const MachineOperand &MO);
46
47 private:
48 void EmitVariableDeclaration(const GlobalVariable *gv);
49 void EmitFunctionDeclaration(const Function* func);
50
51 StringMap SourceIdMap;
52 }; // class PTXAsmPrinter
53 } // namespace llvm
54
55 #endif
56
+0
-181
lib/Target/PTX/PTXFPRoundingModePass.cpp less more
None //===-- PTXFPRoundingModePass.cpp - Assign rounding modes pass ------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines a machine function pass that sets appropriate FP rounding
10 // modes for all relevant instructions.
11 //
12 //===----------------------------------------------------------------------===//
13
14 #define DEBUG_TYPE "ptx-fp-rounding-mode"
15
16 #include "PTX.h"
17 #include "PTXTargetMachine.h"
18 #include "llvm/ADT/DenseMap.h"
19 #include "llvm/CodeGen/MachineFunctionPass.h"
20 #include "llvm/CodeGen/MachineRegisterInfo.h"
21 #include "llvm/Support/Debug.h"
22 #include "llvm/Support/ErrorHandling.h"
23 #include "llvm/Support/raw_ostream.h"
24
25 using namespace llvm;
26
27 // NOTE: PTXFPRoundingModePass should be executed just before emission.
28
29 namespace {
30 /// PTXFPRoundingModePass - Pass to assign appropriate FP rounding modes to
31 /// all FP instructions. Essentially, this pass just looks for all FP
32 /// instructions that have a rounding mode set to RndDefault, and sets an
33 /// appropriate rounding mode based on the target device.
34 ///
35 class PTXFPRoundingModePass : public MachineFunctionPass {
36 private:
37 static char ID;
38
39 typedef std::pair RndModeDesc;
40
41 PTXTargetMachine& TargetMachine;
42 DenseMap Instrs;
43
44 public:
45 PTXFPRoundingModePass(PTXTargetMachine &TM, CodeGenOpt::Level OptLevel)
46 : MachineFunctionPass(ID),
47 TargetMachine(TM) {
48 initializeMap();
49 }
50
51 virtual bool runOnMachineFunction(MachineFunction &MF);
52
53 virtual const char *getPassName() const {
54 return "PTX FP Rounding Mode Pass";
55 }
56
57 private:
58
59 void initializeMap();
60 void processInstruction(MachineInstr &MI);
61 }; // class PTXFPRoundingModePass
62 } // end anonymous namespace
63
64 using namespace llvm;
65
66 char PTXFPRoundingModePass::ID = 0;
67
68 bool PTXFPRoundingModePass::runOnMachineFunction(MachineFunction &MF) {
69 // Look at each basic block
70 for (MachineFunction::iterator bbi = MF.begin(), bbe = MF.end(); bbi != bbe;
71 ++bbi) {
72 MachineBasicBlock &MBB = *bbi;
73 // Look at each instruction
74 for (MachineBasicBlock::iterator ii = MBB.begin(), ie = MBB.end();
75 ii != ie; ++ii) {
76 MachineInstr &MI = *ii;
77 processInstruction(MI);
78 }
79 }
80 return false;
81 }
82
83 void PTXFPRoundingModePass::initializeMap() {
84 using namespace PTXRoundingMode;
85 const PTXSubtarget& ST = TargetMachine.getSubtarget();
86
87 // Build a map of default rounding mode for all instructions that need a
88 // rounding mode.
89 Instrs[PTX::FADDrr32] = std::make_pair(1U, (unsigned)RndNearestEven);
90 Instrs[PTX::FADDri32] = std::make_pair(1U, (unsigned)RndNearestEven);
91 Instrs[PTX::FADDrr64] = std::make_pair(1U, (unsigned)RndNearestEven);
92 Instrs[PTX::FADDri64] = std::make_pair(1U, (unsigned)RndNearestEven);
93 Instrs[PTX::FSUBrr32] = std::make_pair(1U, (unsigned)RndNearestEven);
94 Instrs[PTX::FSUBri32] = std::make_pair(1U, (unsigned)RndNearestEven);
95 Instrs[PTX::FSUBrr64] = std::make_pair(1U, (unsigned)RndNearestEven);
96 Instrs[PTX::FSUBri64] = std::make_pair(1U, (unsigned)RndNearestEven);
97 Instrs[PTX::FMULrr32] = std::make_pair(1U, (unsigned)RndNearestEven);
98 Instrs[PTX::FMULri32] = std::make_pair(1U, (unsigned)RndNearestEven);
99 Instrs[PTX::FMULrr64] = std::make_pair(1U, (unsigned)RndNearestEven);
100 Instrs[PTX::FMULri64] = std::make_pair(1U, (unsigned)RndNearestEven);
101
102 Instrs[PTX::FNEGrr32] = std::make_pair(1U, (unsigned)RndNone);
103 Instrs[PTX::FNEGri32] = std::make_pair(1U, (unsigned)RndNone);
104 Instrs[PTX::FNEGrr64] = std::make_pair(1U, (unsigned)RndNone);
105 Instrs[PTX::FNEGri64] = std::make_pair(1U, (unsigned)RndNone);
106
107 unsigned FDivRndMode = ST.fdivNeedsRoundingMode() ? RndNearestEven : RndNone;
108 Instrs[PTX::FDIVrr32] = std::make_pair(1U, FDivRndMode);
109 Instrs[PTX::FDIVri32] = std::make_pair(1U, FDivRndMode);
110 Instrs[PTX::FDIVrr64] = std::make_pair(1U, FDivRndMode);
111 Instrs[PTX::FDIVri64] = std::make_pair(1U, FDivRndMode);
112
113 unsigned FMADRndMode = ST.fmadNeedsRoundingMode() ? RndNearestEven : RndNone;
114 Instrs[PTX::FMADrrr32] = std::make_pair(1U, FMADRndMode);
115 Instrs[PTX::FMADrri32] = std::make_pair(1U, FMADRndMode);
116 Instrs[PTX::FMADrii32] = std::make_pair(1U, FMADRndMode);
117 Instrs[PTX::FMADrrr64] = std::make_pair(1U, FMADRndMode);
118 Instrs[PTX::FMADrri64] = std::make_pair(1U, FMADRndMode);
119 Instrs[PTX::FMADrii64] = std::make_pair(1U, FMADRndMode);
120
121 Instrs[PTX::FSQRTrr32] = std::make_pair(1U, (unsigned)RndNearestEven);
122 Instrs[PTX::FSQRTri32] = std::make_pair(1U, (unsigned)RndNearestEven);
123 Instrs[PTX::FSQRTrr64] = std::make_pair(1U, (unsigned)RndNearestEven);
124 Instrs[PTX::FSQRTri64] = std::make_pair(1U, (unsigned)RndNearestEven);
125
126 Instrs[PTX::FSINrr32] = std::make_pair(1U, (unsigned)RndApprox);
127 Instrs[PTX::FSINri32] = std::make_pair(1U, (unsigned)RndApprox);
128 Instrs[PTX::FSINrr64] = std::make_pair(1U, (unsigned)RndApprox);
129 Instrs[PTX::FSINri64] = std::make_pair(1U, (unsigned)RndApprox);
130 Instrs[PTX::FCOSrr32] = std::make_pair(1U, (unsigned)RndApprox);
131 Instrs[PTX::FCOSri32] = std::make_pair(1U, (unsigned)RndApprox);
132 Instrs[PTX::FCOSrr64] = std::make_pair(1U, (unsigned)RndApprox);
133 Instrs[PTX::FCOSri64] = std::make_pair(1U, (unsigned)RndApprox);
134
135 Instrs[PTX::CVTu16f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
136 Instrs[PTX::CVTs16f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
137 Instrs[PTX::CVTu16f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
138 Instrs[PTX::CVTs16f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
139 Instrs[PTX::CVTu32f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
140 Instrs[PTX::CVTs32f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
141 Instrs[PTX::CVTu32f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
142 Instrs[PTX::CVTs32f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
143 Instrs[PTX::CVTu64f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
144 Instrs[PTX::CVTs64f32] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
145 Instrs[PTX::CVTu64f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
146 Instrs[PTX::CVTs64f64] = std::make_pair(1U, (unsigned)RndTowardsZeroInt);
147
148 Instrs[PTX::CVTf32u16] = std::make_pair(1U, (unsigned)RndNearestEven);
149 Instrs[PTX::CVTf32s16] = std::make_pair(1U, (unsigned)RndNearestEven);
150 Instrs[PTX::CVTf32u32] = std::make_pair(1U, (unsigned)RndNearestEven);
151 Instrs[PTX::CVTf32s32] = std::make_pair(1U, (unsigned)RndNearestEven);
152 Instrs[PTX::CVTf32u64] = std::make_pair(1U, (unsigned)RndNearestEven);
153 Instrs[PTX::CVTf32s64] = std::make_pair(1U, (unsigned)RndNearestEven);
154 Instrs[PTX::CVTf32f64] = std::make_pair(1U, (unsigned)RndNearestEven);
155 Instrs[PTX::CVTf64u16] = std::make_pair(1U, (unsigned)RndNearestEven);
156 Instrs[PTX::CVTf64s16] = std::make_pair(1U, (unsigned)RndNearestEven);
157 Instrs[PTX::CVTf64u32] = std::make_pair(1U, (unsigned)RndNearestEven);
158 Instrs[PTX::CVTf64s32] = std::make_pair(1U, (unsigned)RndNearestEven);
159 Instrs[PTX::CVTf64u64] = std::make_pair(1U, (unsigned)RndNearestEven);
160 Instrs[PTX::CVTf64s64] = std::make_pair(1U, (unsigned)RndNearestEven);
161 }
162
163 void PTXFPRoundingModePass::processInstruction(MachineInstr &MI) {
164 // Is this an instruction that needs a rounding mode?
165 if (Instrs.count(MI.getOpcode())) {
166 const RndModeDesc &Desc = Instrs[MI.getOpcode()];
167 // Get the rounding mode operand
168 MachineOperand &Op = MI.getOperand(Desc.first);
169 // Update the rounding mode if needed
170 if (Op.getImm() == PTXRoundingMode::RndDefault) {
171 Op.setImm(Desc.second);
172 }
173 }
174 }
175
176 FunctionPass *llvm::createPTXFPRoundingModePass(PTXTargetMachine &TM,
177 CodeGenOpt::Level OptLevel) {
178 return new PTXFPRoundingModePass(TM, OptLevel);
179 }
180
+0
-24
lib/Target/PTX/PTXFrameLowering.cpp less more
None //===-- PTXFrameLowering.cpp - PTX Frame Information ----------------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains the PTX implementation of TargetFrameLowering class.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #include "PTXFrameLowering.h"
14 #include "llvm/CodeGen/MachineFunction.h"
15
16 using namespace llvm;
17
18 void PTXFrameLowering::emitPrologue(MachineFunction &MF) const {
19 }
20
21 void PTXFrameLowering::emitEpilogue(MachineFunction &MF,
22 MachineBasicBlock &MBB) const {
23 }
+0
-44
lib/Target/PTX/PTXFrameLowering.h less more
None //===-- PTXFrameLowering.h - Define frame lowering for PTX -----*- C++ -*--===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 //
10 //
11 //===----------------------------------------------------------------------===//
12
13 #ifndef PTX_FRAMEINFO_H
14 #define PTX_FRAMEINFO_H
15
16 #include "PTX.h"
17 #include "PTXSubtarget.h"
18 #include "llvm/Target/TargetFrameLowering.h"
19
20 namespace llvm {
21 class PTXSubtarget;
22
23 class PTXFrameLowering : public TargetFrameLowering {
24 protected:
25 const PTXSubtarget &STI;
26
27 public:
28 explicit PTXFrameLowering(const PTXSubtarget &sti)
29 : TargetFrameLowering(TargetFrameLowering::StackGrowsDown, 2, -2),
30 STI(sti) {
31 }
32
33 /// emitProlog/emitEpilog - These methods insert prolog and epilog code into
34 /// the function.
35 void emitPrologue(MachineFunction &MF) const;
36 void emitEpilogue(MachineFunction &MF, MachineBasicBlock &MBB) const;
37
38 bool hasFP(const MachineFunction &MF) const { return false; }
39 };
40
41 } // End llvm namespace
42
43 #endif
+0
-356
lib/Target/PTX/PTXISelDAGToDAG.cpp less more
None //===-- PTXISelDAGToDAG.cpp - A dag to dag inst selector for PTX ----------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines an instruction selector for the PTX target.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #include "PTX.h"
14 #include "PTXMachineFunctionInfo.h"
15 #include "PTXTargetMachine.h"
16 #include "llvm/ADT/StringExtras.h"
17 #include "llvm/CodeGen/SelectionDAGISel.h"
18 #include "llvm/DerivedTypes.h"
19 #include "llvm/Support/Debug.h"
20 #include "llvm/Support/raw_ostream.h"
21
22 using namespace llvm;
23
24 namespace {
25 // PTXDAGToDAGISel - PTX specific code to select PTX machine
26 // instructions for SelectionDAG operations.
27 class PTXDAGToDAGISel : public SelectionDAGISel {
28 public:
29 PTXDAGToDAGISel(PTXTargetMachine &TM, CodeGenOpt::Level OptLevel);
30
31 virtual const char *getPassName() const {
32 return "PTX DAG->DAG Pattern Instruction Selection";
33 }
34
35 SDNode *Select(SDNode *Node);
36
37 // Complex Pattern Selectors.
38 bool SelectADDRrr(SDValue &Addr, SDValue &R1, SDValue &R2);
39 bool SelectADDRri(SDValue &Addr, SDValue &Base, SDValue &Offset);
40 bool SelectADDRii(SDValue &Addr, SDValue &Base, SDValue &Offset);
41 bool SelectADDRlocal(SDValue &Addr, SDValue &Base, SDValue &Offset);
42
43 // Include the pieces auto'gened from the target description
44 #include "PTXGenDAGISel.inc"
45
46 private:
47 // We need this only because we can't match intruction BRAdp
48 // pattern (PTXbrcond bb:$d, ...) in PTXInstrInfo.td
49 SDNode *SelectBRCOND(SDNode *Node);
50
51 SDNode *SelectREADPARAM(SDNode *Node);
52 SDNode *SelectWRITEPARAM(SDNode *Node);
53 SDNode *SelectFrameIndex(SDNode *Node);
54
55 bool isImm(const SDValue &operand);
56 bool SelectImm(const SDValue &operand, SDValue &imm);
57
58 const PTXSubtarget& getSubtarget() const;
59 }; // class PTXDAGToDAGISel
60 } // namespace
61
62 // createPTXISelDag - This pass converts a legalized DAG into a
63 // PTX-specific DAG, ready for instruction scheduling
64 FunctionPass *llvm::createPTXISelDag(PTXTargetMachine &TM,
65 CodeGenOpt::Level OptLevel) {
66 return new PTXDAGToDAGISel(TM, OptLevel);
67 }
68
69 PTXDAGToDAGISel::PTXDAGToDAGISel(PTXTargetMachine &TM,
70 CodeGenOpt::Level OptLevel)
71 : SelectionDAGISel(TM, OptLevel) {}
72
73 SDNode *PTXDAGToDAGISel::Select(SDNode *Node) {
74 switch (Node->getOpcode()) {
75 case ISD::BRCOND:
76 return SelectBRCOND(Node);
77 case PTXISD::READ_PARAM:
78 return SelectREADPARAM(Node);
79 case PTXISD::WRITE_PARAM:
80 return SelectWRITEPARAM(Node);
81 case ISD::FrameIndex:
82 return SelectFrameIndex(Node);
83 default:
84 return SelectCode(Node);
85 }
86 }
87
88 SDNode *PTXDAGToDAGISel::SelectBRCOND(SDNode *Node) {
89 assert(Node->getNumOperands() >= 3);
90
91 SDValue Chain = Node->getOperand(0);
92 SDValue Pred = Node->getOperand(1);
93 SDValue Target = Node->getOperand(2); // branch target
94 SDValue PredOp = CurDAG->getTargetConstant(PTXPredicate::Normal, MVT::i32);
95 DebugLoc dl = Node->getDebugLoc();
96
97 assert(Target.getOpcode() == ISD::BasicBlock);
98 assert(Pred.getValueType() == MVT::i1);
99
100 // Emit BRAdp
101 SDValue Ops[] = { Target, Pred, PredOp, Chain };
102 return CurDAG->getMachineNode(PTX::BRAdp, dl, MVT::Other, Ops, 4);
103 }
104
105 SDNode *PTXDAGToDAGISel::SelectREADPARAM(SDNode *Node) {
106 SDValue Chain = Node->getOperand(0);
107 SDValue Index = Node->getOperand(1);
108
109 int OpCode;
110
111 // Get the type of parameter we are reading
112 EVT VT = Node->getValueType(0);
113 assert(VT.isSimple() && "READ_PARAM only implemented for MVT types");
114
115 MVT Type = VT.getSimpleVT();
116
117 if (Type == MVT::i1)
118 OpCode = PTX::READPARAMPRED;
119 else if (Type == MVT::i16)
120 OpCode = PTX::READPARAMI16;
121 else if (Type == MVT::i32)
122 OpCode = PTX::READPARAMI32;
123 else if (Type == MVT::i64)
124 OpCode = PTX::READPARAMI64;
125 else if (Type == MVT::f32)
126 OpCode = PTX::READPARAMF32;
127 else {
128 assert(Type == MVT::f64 && "Unexpected type!");
129 OpCode = PTX::READPARAMF64;
130 }
131
132 SDValue Pred = CurDAG->getRegister(PTX::NoRegister, MVT::i1);
133 SDValue PredOp = CurDAG->getTargetConstant(PTXPredicate::None, MVT::i32);
134 DebugLoc dl = Node->getDebugLoc();
135
136 SDValue Ops[] = { Index, Pred, PredOp, Chain };
137 return CurDAG->getMachineNode(OpCode, dl, VT, Ops, 4);
138 }
139
140 SDNode *PTXDAGToDAGISel::SelectWRITEPARAM(SDNode *Node) {
141
142 SDValue Chain = Node->getOperand(0);
143 SDValue Value = Node->getOperand(1);
144
145 int OpCode;
146
147 //Node->dumpr(CurDAG);
148
149 // Get the type of parameter we are writing
150 EVT VT = Value->getValueType(0);
151 assert(VT.isSimple() && "WRITE_PARAM only implemented for MVT types");
152
153 MVT Type = VT.getSimpleVT();
154
155 if (Type == MVT::i1)
156 OpCode = PTX::WRITEPARAMPRED;
157 else if (Type == MVT::i16)
158 OpCode = PTX::WRITEPARAMI16;
159 else if (Type == MVT::i32)
160 OpCode = PTX::WRITEPARAMI32;
161 else if (Type == MVT::i64)
162 OpCode = PTX::WRITEPARAMI64;
163 else if (Type == MVT::f32)
164 OpCode = PTX::WRITEPARAMF32;
165 else if (Type == MVT::f64)
166 OpCode = PTX::WRITEPARAMF64;
167 else
168 llvm_unreachable("Invalid type in SelectWRITEPARAM");
169
170 SDValue Pred = CurDAG->getRegister(PTX::NoRegister, MVT::i1);
171 SDValue PredOp = CurDAG->getTargetConstant(PTXPredicate::None, MVT::i32);
172 DebugLoc dl = Node->getDebugLoc();
173
174 SDValue Ops[] = { Value, Pred, PredOp, Chain };
175 SDNode* Ret = CurDAG->getMachineNode(OpCode, dl, MVT::Other, Ops, 4);
176
177 //dbgs() << "SelectWRITEPARAM produced:\n\t";
178 //Ret->dumpr(CurDAG);
179
180 return Ret;
181 }
182
183 SDNode *PTXDAGToDAGISel::SelectFrameIndex(SDNode *Node) {
184 int FI = cast(Node)->getIndex();
185 //dbgs() << "Selecting FrameIndex at index " << FI << "\n";
186 //SDValue TFI = CurDAG->getTargetFrameIndex(FI, Node->getValueType(0));
187
188 PTXMachineFunctionInfo *MFI = MF->getInfo();
189
190 SDValue FrameSymbol = CurDAG->getTargetExternalSymbol(MFI->getFrameSymbol(FI),
191 Node->getValueType(0));
192
193 return FrameSymbol.getNode();
194 }
195
196 // Match memory operand of the form [reg+reg]
197 bool PTXDAGToDAGISel::SelectADDRrr(SDValue &Addr, SDValue &R1, SDValue &R2) {
198 if (Addr.getOpcode() != ISD::ADD || Addr.getNumOperands() < 2 ||
199 isImm(Addr.getOperand(0)) || isImm(Addr.getOperand(1)))
200 return false;
201
202 assert(Addr.getValueType().isSimple() && "Type must be simple");
203
204 R1 = Addr;
205 R2 = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
206
207 return true;
208 }
209
210 // Match memory operand of the form [reg], [imm+reg], and [reg+imm]
211 bool PTXDAGToDAGISel::SelectADDRri(SDValue &Addr, SDValue &Base,
212 SDValue &Offset) {
213 // FrameIndex addresses are handled separately
214 //errs() << "SelectADDRri: ";
215 //Addr.getNode()->dumpr();
216 if (isa(Addr)) {
217 //errs() << "Failure\n";
218 return false;
219 }
220
221 if (CurDAG->isBaseWithConstantOffset(Addr)) {
222 Base = Addr.getOperand(0);
223 if (isa(Base)) {
224 //errs() << "Failure\n";
225 return false;
226 }
227 ConstantSDNode *CN = dyn_cast(Addr.getOperand(1));
228 Offset = CurDAG->getTargetConstant(CN->getZExtValue(), MVT::i32);
229 //errs() << "Success\n";
230 return true;
231 }
232
233 /*if (Addr.getNumOperands() == 1) {
234 Base = Addr;
235 Offset = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
236 errs() << "Success\n";
237 return true;
238 }*/
239
240 //errs() << "SelectADDRri fails on: ";
241 //Addr.getNode()->dumpr();
242
243 if (isImm(Addr)) {
244 //errs() << "Failure\n";
245 return false;
246 }
247
248 Base = Addr;
249 Offset = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
250
251 //errs() << "Success\n";
252 return true;
253
254 /*if (Addr.getOpcode() != ISD::ADD) {
255 // let SelectADDRii handle the [imm] case
256 if (isImm(Addr))
257 return false;
258 // it is [reg]
259
260 assert(Addr.getValueType().isSimple() && "Type must be simple");
261 Base = Addr;
262 Offset = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
263
264 return true;
265 }
266
267 if (Addr.getNumOperands() < 2)
268 return false;
269
270 // let SelectADDRii handle the [imm+imm] case
271 if (isImm(Addr.getOperand(0)) && isImm(Addr.getOperand(1)))
272 return false;
273
274 // try [reg+imm] and [imm+reg]
275 for (int i = 0; i < 2; i ++)
276 if (SelectImm(Addr.getOperand(1-i), Offset)) {
277 Base = Addr.getOperand(i);
278 return true;
279 }
280
281 // neither [reg+imm] nor [imm+reg]
282 return false;*/
283 }
284
285 // Match memory operand of the form [imm+imm] and [imm]
286 bool PTXDAGToDAGISel::SelectADDRii(SDValue &Addr, SDValue &Base,
287 SDValue &Offset) {
288 // is [imm+imm]?
289 if (Addr.getOpcode() == ISD::ADD) {
290 return SelectImm(Addr.getOperand(0), Base) &&
291 SelectImm(Addr.getOperand(1), Offset);
292 }
293
294 // is [imm]?
295 if (SelectImm(Addr, Base)) {
296 assert(Addr.getValueType().isSimple() && "Type must be simple");
297
298 Offset = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
299
300 return true;
301 }
302
303 return false;
304 }
305
306 // Match memory operand of the form [reg], [imm+reg], and [reg+imm]
307 bool PTXDAGToDAGISel::SelectADDRlocal(SDValue &Addr, SDValue &Base,
308 SDValue &Offset) {
309 //errs() << "SelectADDRlocal: ";
310 //Addr.getNode()->dumpr();
311 if (isa(Addr)) {
312 Base = Addr;
313 Offset = CurDAG->getTargetConstant(0, Addr.getValueType().getSimpleVT());
314 //errs() << "Success\n";
315 return true;
316 }
317
318 if (CurDAG->isBaseWithConstantOffset(Addr)) {
319 Base = Addr.getOperand(0);
320 if (!isa(Base)) {
321 //errs() << "Failure\n";
322 return false;
323 }
324 ConstantSDNode *CN = dyn_cast(Addr.getOperand(1));
325 Offset = CurDAG->getTargetConstant(CN->getZExtValue(), MVT::i32);
326 //errs() << "Offset: ";
327 //Offset.getNode()->dumpr();
328 //errs() << "Success\n";
329 return true;
330 }
331
332 //errs() << "Failure\n";
333 return false;
334 }
335
336 bool PTXDAGToDAGISel::isImm(const SDValue &operand) {
337 return ConstantSDNode::classof(operand.getNode());
338 }
339
340 bool PTXDAGToDAGISel::SelectImm(const SDValue &operand, SDValue &imm) {
341 SDNode *node = operand.getNode();
342 if (!ConstantSDNode::classof(node))
343 return false;
344
345 ConstantSDNode *CN = cast(node);
346 imm = CurDAG->getTargetConstant(*CN->getConstantIntValue(),
347 operand.getValueType());
348 return true;
349 }
350
351 const PTXSubtarget& PTXDAGToDAGISel::getSubtarget() const
352 {
353 return TM.getSubtarget();
354 }
355
+0
-516
lib/Target/PTX/PTXISelLowering.cpp less more
None //===-- PTXISelLowering.cpp - PTX DAG Lowering Implementation -------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file implements the PTXTargetLowering class.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #include "PTXISelLowering.h"
14 #include "PTX.h"
15 #include "PTXMachineFunctionInfo.h"
16 #include "PTXRegisterInfo.h"
17 #include "PTXSubtarget.h"
18 #include "llvm/Function.h"
19 #include "llvm/Support/ErrorHandling.h"
20 #include "llvm/CodeGen/CallingConvLower.h"
21 #include "llvm/CodeGen/MachineFunction.h"
22 #include "llvm/CodeGen/MachineFrameInfo.h"
23 #include "llvm/CodeGen/MachineRegisterInfo.h"
24 #include "llvm/CodeGen/SelectionDAG.h"
25 #include "llvm/CodeGen/TargetLoweringObjectFileImpl.h"
26 #include "llvm/Support/Debug.h"
27 #include "llvm/Support/raw_ostream.h"
28
29 using namespace llvm;
30
31 //===----------------------------------------------------------------------===//
32 // TargetLowering Implementation
33 //===----------------------------------------------------------------------===//
34
35 PTXTargetLowering::PTXTargetLowering(TargetMachine &TM)
36 : TargetLowering(TM, new TargetLoweringObjectFileELF()) {
37 // Set up the register classes.
38 addRegisterClass(MVT::i1, &PTX::RegPredRegClass);
39 addRegisterClass(MVT::i16, &PTX::RegI16RegClass);
40 addRegisterClass(MVT::i32, &PTX::RegI32RegClass);
41 addRegisterClass(MVT::i64, &PTX::RegI64RegClass);
42 addRegisterClass(MVT::f32, &PTX::RegF32RegClass);
43 addRegisterClass(MVT::f64, &PTX::RegF64RegClass);
44
45 setBooleanContents(ZeroOrOneBooleanContent);
46 setBooleanVectorContents(ZeroOrOneBooleanContent); // FIXME: Is this correct?
47 setMinFunctionAlignment(2);
48
49 // Let LLVM use loads/stores for all mem* operations
50 maxStoresPerMemcpy = 4096;
51 maxStoresPerMemmove = 4096;
52 maxStoresPerMemset = 4096;
53
54 ////////////////////////////////////
55 /////////// Expansion //////////////
56 ////////////////////////////////////
57
58 // (any/zero/sign) extload => load + (any/zero/sign) extend
59
60 setLoadExtAction(ISD::EXTLOAD, MVT::i16, Expand);
61 setLoadExtAction(ISD::ZEXTLOAD, MVT::i16, Expand);
62 setLoadExtAction(ISD::SEXTLOAD, MVT::i16, Expand);
63
64 // f32 extload => load + fextend
65
66 setLoadExtAction(ISD::EXTLOAD, MVT::f32, Expand);
67
68 // f64 truncstore => trunc + store
69
70 setTruncStoreAction(MVT::f64, MVT::f32, Expand);
71
72 // sign_extend_inreg => sign_extend
73
74 setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i1, Expand);
75
76 // br_cc => brcond
77
78 setOperationAction(ISD::BR_CC, MVT::Other, Expand);
79
80 // select_cc => setcc
81
82 setOperationAction(ISD::SELECT_CC, MVT::Other, Expand);
83 setOperationAction(ISD::SELECT_CC, MVT::f32, Expand);
84 setOperationAction(ISD::SELECT_CC, MVT::f64, Expand);
85
86 ////////////////////////////////////
87 //////////// Legal /////////////////
88 ////////////////////////////////////
89
90 setOperationAction(ISD::ConstantFP, MVT::f32, Legal);
91 setOperationAction(ISD::ConstantFP, MVT::f64, Legal);
92
93 ////////////////////////////////////
94 //////////// Custom ////////////////
95 ////////////////////////////////////
96
97 // customise setcc to use bitwise logic if possible
98
99 //setOperationAction(ISD::SETCC, MVT::i1, Custom);
100 setOperationAction(ISD::SETCC, MVT::i1, Legal);
101
102 // customize translation of memory addresses
103
104 setOperationAction(ISD::GlobalAddress, MVT::i32, Custom);
105 setOperationAction(ISD::GlobalAddress, MVT::i64, Custom);
106
107 // Compute derived properties from the register classes
108 computeRegisterProperties();
109 }
110
111 EVT PTXTargetLowering::getSetCCResultType(EVT VT) const {
112 return MVT::i1;
113 }
114
115 SDValue PTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
116 switch (Op.getOpcode()) {
117 default:
118 llvm_unreachable("Unimplemented operand");
119 case ISD::SETCC:
120 return LowerSETCC(Op, DAG);
121 case ISD::GlobalAddress:
122 return LowerGlobalAddress(Op, DAG);
123 }
124 }
125
126 const char *PTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
127 switch (Opcode) {
128 default:
129 llvm_unreachable("Unknown opcode");
130 case PTXISD::COPY_ADDRESS:
131 return "PTXISD::COPY_ADDRESS";
132 case PTXISD::LOAD_PARAM:
133 return "PTXISD::LOAD_PARAM";
134 case PTXISD::STORE_PARAM:
135 return "PTXISD::STORE_PARAM";
136 case PTXISD::READ_PARAM:
137 return "PTXISD::READ_PARAM";
138 case PTXISD::WRITE_PARAM:
139 return "PTXISD::WRITE_PARAM";
140 case PTXISD::EXIT:
141 return "PTXISD::EXIT";
142 case PTXISD::RET:
143 return "PTXISD::RET";
144 case PTXISD::CALL:
145 return "PTXISD::CALL";
146 }
147 }
148
149 //===----------------------------------------------------------------------===//
150 // Custom Lower Operation
151 //===----------------------------------------------------------------------===//
152
153 SDValue PTXTargetLowering::LowerSETCC(SDValue Op, SelectionDAG &DAG) const {
154 assert(Op.getValueType() == MVT::i1 && "SetCC type must be 1-bit integer");
155 SDValue Op0 = Op.getOperand(0);
156 SDValue Op1 = Op.getOperand(1);
157 SDValue Op2 = Op.getOperand(2);
158 DebugLoc dl = Op.getDebugLoc();
159 //ISD::CondCode CC = cast(Op.getOperand(2))->get();
160
161 // Look for X == 0, X == 1, X != 0, or X != 1
162 // We can simplify these to bitwise logic
163
164 //if (Op1.getOpcode() == ISD::Constant &&
165 // (cast(Op1)->getZExtValue() == 1 ||
166 // cast(Op1)->isNullValue()) &&
167 // (CC == ISD::SETEQ || CC == ISD::SETNE)) {
168 //
169 // return DAG.getNode(ISD::AND, dl, MVT::i1, Op0, Op1);
170 //}
171
172 //ConstantSDNode* COp1 = cast(Op1);
173 //if(COp1 && COp1->getZExtValue() == 1) {
174 // if(CC == ISD::SETNE) {
175 // return DAG.getNode(PTX::XORripreds, dl, MVT::i1, Op0);
176 // }
177 //}
178
179 llvm_unreachable("setcc was not matched by a pattern!");
180
181 return DAG.getNode(ISD::SETCC, dl, MVT::i1, Op0, Op1, Op2);
182 }
183
184 SDValue PTXTargetLowering::
185 LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const {
186 EVT PtrVT = getPointerTy();
187 DebugLoc dl = Op.getDebugLoc();
188 const GlobalValue *GV = cast(Op)->getGlobal();
189
190 assert(PtrVT.isSimple() && "Pointer must be to primitive type.");
191
192 SDValue targetGlobal = DAG.getTargetGlobalAddress(GV, dl, PtrVT);
193 SDValue movInstr = DAG.getNode(PTXISD::COPY_ADDRESS,
194 dl,
195 PtrVT.getSimpleVT(),
196 targetGlobal);
197
198 return movInstr;
199 }
200
201 //===----------------------------------------------------------------------===//
202 // Calling Convention Implementation
203 //===----------------------------------------------------------------------===//
204
205 SDValue PTXTargetLowering::
206 LowerFormalArguments(SDValue Chain,
207 CallingConv::ID CallConv,
208 bool isVarArg,
209 const SmallVectorImpl &Ins,
210 DebugLoc dl,
211 SelectionDAG &DAG,
212 SmallVectorImpl &InVals) const {
213 if (isVarArg) llvm_unreachable("PTX does not support varargs");
214
215 MachineFunction &MF = DAG.getMachineFunction();
216 const PTXSubtarget& ST = getTargetMachine().getSubtarget();
217 PTXMachineFunctionInfo *MFI = MF.getInfo();
218 PTXParamManager &PM = MFI->getParamManager();
219
220 switch (CallConv) {
221 default:
222 llvm_unreachable("Unsupported calling convention");
223 case CallingConv::PTX_Kernel:
224 MFI->setKernel(true);
225 break;
226 case CallingConv::PTX_Device:
227 MFI->setKernel(false);
228 break;
229 }
230
231 // We do one of two things here:
232 // IsKernel || SM >= 2.0 -> Use param space for arguments
233 // SM < 2.0 -> Use registers for arguments
234 if (MFI->isKernel() || ST.useParamSpaceForDeviceArgs()) {
235 // We just need to emit the proper LOAD_PARAM ISDs
236 for (unsigned i = 0, e = Ins.size(); i != e; ++i) {
237 assert((!MFI->isKernel() || Ins[i].VT != MVT::i1) &&
238 "Kernels cannot take pred operands");
239
240 unsigned ParamSize = Ins[i].VT.getStoreSizeInBits();
241 unsigned Param = PM.addArgumentParam(ParamSize);
242 const std::string &ParamName = PM.getParamName(Param);
243 SDValue ParamValue = DAG.getTargetExternalSymbol(ParamName.c_str(),
244 MVT::Other);
245 SDValue ArgValue = DAG.getNode(PTXISD::LOAD_PARAM, dl, Ins[i].VT, Chain,
246 ParamValue);
247 InVals.push_back(ArgValue);
248 }
249 }
250 else {
251 for (unsigned i = 0, e = Ins.size(); i != e; ++i) {
252 EVT RegVT = Ins[i].VT;
253 const TargetRegisterClass* TRC = getRegClassFor(RegVT);
254 unsigned RegType;
255
256 // Determine which register class we need
257 if (RegVT == MVT::i1)
258 RegType = PTXRegisterType::Pred;
259 else if (RegVT == MVT::i16)
260 RegType = PTXRegisterType::B16;
261 else if (RegVT == MVT::i32)
262 RegType = PTXRegisterType::B32;
263 else if (RegVT == MVT::i64)
264 RegType = PTXRegisterType::B64;
265 else if (RegVT == MVT::f32)
266 RegType = PTXRegisterType::F32;
267 else if (RegVT == MVT::f64)
268 RegType = PTXRegisterType::F64;
269 else
270 llvm_unreachable("Unknown parameter type");
271
272 // Use a unique index in the instruction to prevent instruction folding.
273 // Yes, this is a hack.
274 SDValue Index = DAG.getTargetConstant(i, MVT::i32);
275 unsigned Reg = MF.getRegInfo().createVirtualRegister(TRC);
276 SDValue ArgValue = DAG.getNode(PTXISD::READ_PARAM, dl, RegVT, Chain,
277 Index);
278
279 InVals.push_back(ArgValue);
280
281 MFI->addRegister(Reg, RegType, PTXRegisterSpace::Argument);
282 }
283 }
284
285 return Chain;
286 }
287
288 SDValue PTXTargetLowering::
289 LowerReturn(SDValue Chain,
290 CallingConv::ID CallConv,
291 bool isVarArg,
292 const SmallVectorImpl &Outs,
293 const SmallVectorImpl &OutVals,
294 DebugLoc dl,
295 SelectionDAG &DAG) const {
296 if (isVarArg) llvm_unreachable("PTX does not support varargs");
297
298 switch (CallConv) {
299 default:
300 llvm_unreachable("Unsupported calling convention.");
301 case CallingConv::PTX_Kernel:
302 assert(Outs.size() == 0 && "Kernel must return void.");
303 return DAG.getNode(PTXISD::EXIT, dl, MVT::Other, Chain);
304 case CallingConv::PTX_Device:
305 assert(Outs.size() <= 1 && "Can at most return one value.");
306 break;
307 }
308
309 MachineFunction& MF = DAG.getMachineFunction();
310 PTXMachineFunctionInfo *MFI = MF.getInfo();
311 PTXParamManager &PM = MFI->getParamManager();
312
313 SDValue Flag;
314 const PTXSubtarget& ST = getTargetMachine().getSubtarget();
315
316 if (ST.useParamSpaceForDeviceArgs()) {
317 assert(Outs.size() < 2 && "Device functions can return at most one value");
318
319 if (Outs.size() == 1) {
320 unsigned ParamSize = OutVals[0].getValueType().getSizeInBits();
321 unsigned Param = PM.addReturnParam(ParamSize);
322 const std::string &ParamName = PM.getParamName(Param);
323 SDValue ParamValue = DAG.getTargetExternalSymbol(ParamName.c_str(),
324 MVT::Other);
325 Chain = DAG.getNode(PTXISD::STORE_PARAM, dl, MVT::Other, Chain,
326 ParamValue, OutVals[0]);
327 }
328 } else {
329 for (unsigned i = 0, e = Outs.size(); i != e; ++i) {
330 EVT RegVT = Outs[i].VT;
331 const TargetRegisterClass* TRC;
332 unsigned RegType;
333
334 // Determine which register class we need
335 if (RegVT == MVT::i1) {
336 TRC = &PTX::RegPredRegClass;
337 RegType = PTXRegisterType::Pred;
338 } else if (RegVT == MVT::i16) {
339 TRC = &PTX::RegI16RegClass;
340 RegType = PTXRegisterType::B16;
341 } else if (RegVT == MVT::i32) {
342 TRC = &PTX::RegI32RegClass;
343 RegType = PTXRegisterType::B32;
344 } else if (RegVT == MVT::i64) {
345 TRC = &PTX::RegI64RegClass;
346 RegType = PTXRegisterType::B64;
347 } else if (RegVT == MVT::f32) {
348 TRC = &PTX::RegF32RegClass;
349 RegType = PTXRegisterType::F32;
350 } else if (RegVT == MVT::f64) {
351 TRC = &PTX::RegF64RegClass;
352 RegType = PTXRegisterType::F64;
353 } else {
354 llvm_unreachable("Unknown parameter type");
355 }
356
357 unsigned Reg = MF.getRegInfo().createVirtualRegister(TRC);
358
359 SDValue Copy = DAG.getCopyToReg(Chain, dl, Reg, OutVals[i]/*, Flag*/);
360 SDValue OutReg = DAG.getRegister(Reg, RegVT);
361
362 Chain = DAG.getNode(PTXISD::WRITE_PARAM, dl, MVT::Other, Copy, OutReg);
363
364 MFI->addRegister(Reg, RegType, PTXRegisterSpace::Return);
365 }
366 }
367
368 if (Flag.getNode() == 0) {
369 return DAG.getNode(PTXISD::RET, dl, MVT::Other, Chain);
370 }
371 else {
372 return DAG.getNode(PTXISD::RET, dl, MVT::Other, Chain, Flag);
373 }
374 }
375
376 SDValue
377 PTXTargetLowering::LowerCall(SDValue Chain, SDValue Callee,
378 CallingConv::ID CallConv, bool isVarArg,
379 bool doesNotRet, bool &isTailCall,
380 const SmallVectorImpl &Outs,
381 const SmallVectorImpl &OutVals,
382 const SmallVectorImpl &Ins,
383 DebugLoc dl, SelectionDAG &DAG,
384 SmallVectorImpl &InVals) const {
385
386 MachineFunction& MF = DAG.getMachineFunction();
387 PTXMachineFunctionInfo *PTXMFI = MF.getInfo();
388 PTXParamManager &PM = PTXMFI->getParamManager();
389 MachineFrameInfo *MFI = MF.getFrameInfo();
390
391 assert(getTargetMachine().getSubtarget().callsAreHandled() &&
392 "Calls are not handled for the target device");
393
394 // Identify the callee function
395 const GlobalValue *GV = cast(Callee)->getGlobal();
396 const Function *function = cast(GV);
397
398 // allow non-device calls only for printf
399 bool isPrintf = function->getName() == "printf" || function->getName() == "puts";
400
401 assert((isPrintf || function->getCallingConv() == CallingConv::PTX_Device) &&
402 "PTX function calls must be to PTX device functions");
403
404 unsigned outSize = isPrintf ? 2 : Outs.size();
405
406 std::vector Ops;
407 // The layout of the ops will be [Chain, #Ins, Ins, Callee, #Outs, Outs]
408 Ops.resize(outSize + Ins.size() + 4);
409
410 Ops[0] = Chain;
411
412 // Identify the callee function
413 Callee = DAG.getTargetGlobalAddress(GV, dl, getPointerTy());
414 Ops[Ins.size()+2] = Callee;
415
416 // #Outs
417 Ops[Ins.size()+3] = DAG.getTargetConstant(outSize, MVT::i32);
418
419 if (isPrintf) {
420 // first argument is the address of the global string variable in memory
421 unsigned Param0 = PM.addLocalParam(getPointerTy().getSizeInBits());
422 SDValue ParamValue0 = DAG.getTargetExternalSymbol(PM.getParamName(Param0).c_str(),
423 MVT::Other);
424 Chain = DAG.getNode(PTXISD::STORE_PARAM, dl, MVT::Other, Chain,
425 ParamValue0, OutVals[0]);
426 Ops[Ins.size()+4] = ParamValue0;
427
428 // alignment is the maximum size of all the arguments
429 unsigned alignment = 0;
430 for (unsigned i = 1; i < OutVals.size(); ++i) {
431 alignment = std::max(alignment,
432 OutVals[i].getValueType().getSizeInBits());
433 }
434
435 // size is the alignment multiplied by the number of arguments
436 unsigned size = alignment * (OutVals.size() - 1);
437
438 // second argument is the address of the stack object (unless no arguments)
439 unsigned Param1 = PM.addLocalParam(getPointerTy().getSizeInBits());
440 SDValue ParamValue1 = DAG.getTargetExternalSymbol(PM.getParamName(Param1).c_str(),
441 MVT::Other);
442 Ops[Ins.size()+5] = ParamValue1;
443
444 if (size > 0)
445 {
446 // create a local stack object to store the arguments
447 unsigned StackObject = MFI->CreateStackObject(size / 8, alignment / 8, false);
448 SDValue FrameIndex = DAG.getFrameIndex(StackObject, getPointerTy());
449
450 // store each of the arguments to the stack in turn
451 for (unsigned int i = 1; i != OutVals.size(); i++) {
452 SDValue FrameAddr = DAG.getNode(ISD::ADD, dl, getPointerTy(), FrameIndex, DAG.getTargetConstant((i - 1) * 8, getPointerTy()));
453 Chain = DAG.getStore(Chain, dl, OutVals[i], FrameAddr,
454 MachinePointerInfo(),
455 false, false, 0);
456 }
457
458 // copy the address of the local frame index to get the address in non-local space
459 SDValue genericAddr = DAG.getNode(PTXISD::COPY_ADDRESS, dl, getPointerTy(), FrameIndex);
460
461 // store this address in the second argument
462 Chain = DAG.getNode(PTXISD::STORE_PARAM, dl, MVT::Other, Chain, ParamValue1, genericAddr);
463 }
464 }
465 else
466 {
467 // Generate STORE_PARAM nodes for each function argument. In PTX, function
468 // arguments are explicitly stored into .param variables and passed as
469 // arguments. There is no register/stack-based calling convention in PTX.
470 for (unsigned i = 0; i != OutVals.size(); ++i) {
471 unsigned Size = OutVals[i].getValueType().getSizeInBits();
472 unsigned Param = PM.addLocalParam(Size);
473 const std::string &ParamName = PM.getParamName(Param);
474 SDValue ParamValue = DAG.getTargetExternalSymbol(ParamName.c_str(),
475 MVT::Other);
476 Chain = DAG.getNode(PTXISD::STORE_PARAM, dl, MVT::Other, Chain,
477 ParamValue, OutVals[i]);
478 Ops[i+Ins.size()+4] = ParamValue;
479 }
480 }
481
482 std::vector InParams;
483
484 // Generate list of .param variables to hold the return value(s).
485 Ops[1] = DAG.getTargetConstant(Ins.size(), MVT::i32);
486 for (unsigned i = 0; i < Ins.size(); ++i) {
487 unsigned Size = Ins[i].VT.getStoreSizeInBits();
488 unsigned Param = PM.addLocalParam(Size);
489 const std::string &ParamName = PM.getParamName(Param);
490 SDValue ParamValue = DAG.getTargetExternalSymbol(ParamName.c_str(),
491 MVT::Other);
492 Ops[i+2] = ParamValue;
493 InParams.push_back(ParamValue);
494 }
495
496 Ops[0] = Chain;
497
498 // Create the CALL node.
499 Chain = DAG.getNode(PTXISD::CALL, dl, MVT::Other, &Ops[0], Ops.size());
500
501 // Create the LOAD_PARAM nodes that retrieve the function return value(s).
502 for (unsigned i = 0; i < Ins.size(); ++i) {
503 SDValue Load = DAG.getNode(PTXISD::LOAD_PARAM, dl, Ins[i].VT, Chain,
504 InParams[i]);
505 InVals.push_back(Load);
506 }
507
508 return Chain;
509 }
510
511 unsigned PTXTargetLowering::getNumRegisters(LLVMContext &Context, EVT VT) {
512 // All arguments consist of one "register," regardless of the type.
513 return 1;
514 }
515
+0
-82
lib/Target/PTX/PTXISelLowering.h less more
None //===-- PTXISelLowering.h - PTX DAG Lowering Interface ----------*- C++ -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines the interfaces that PTX uses to lower LLVM code into a
10 // selection DAG.
11 //
12 //===----------------------------------------------------------------------===//
13
14 #ifndef PTX_ISEL_LOWERING_H
15 #define PTX_ISEL_LOWERING_H
16
17 #include "llvm/Target/TargetLowering.h"
18
19 namespace llvm {
20
21 namespace PTXISD {
22 enum NodeType {
23 FIRST_NUMBER = ISD::BUILTIN_OP_END,
24 LOAD_PARAM,
25 STORE_PARAM,
26 READ_PARAM,
27 WRITE_PARAM,
28 EXIT,
29 RET,
30 COPY_ADDRESS,
31 CALL
32 };
33 } // namespace PTXISD
34
35 class PTXTargetLowering : public TargetLowering {
36 public:
37 explicit PTXTargetLowering(TargetMachine &TM);
38
39 virtual const char *getTargetNodeName(unsigned Opcode) const;
40
41 virtual SDValue LowerOperation(SDValue Op, SelectionDAG &DAG) const;
42
43 virtual SDValue LowerSETCC(SDValue Op, SelectionDAG &DAG) const;
44
45 virtual SDValue
46 LowerFormalArguments(SDValue Chain,
47 CallingConv::ID CallConv,
48 bool isVarArg,
49 const SmallVectorImpl &Ins,
50 DebugLoc dl,
51 SelectionDAG &DAG,
52 SmallVectorImpl &InVals) const;
53
54 virtual SDValue
55 LowerReturn(SDValue Chain,
56 CallingConv::ID CallConv,
57 bool isVarArg,
58 const SmallVectorImpl &Outs,
59 const SmallVectorImpl &OutVals,
60 DebugLoc dl,
61 SelectionDAG &DAG) const;
62
63 virtual SDValue
64 LowerCall(SDValue Chain, SDValue Callee, CallingConv::ID CallConv,
65 bool isVarArg, bool doesNotRet, bool &isTailCall,
66 const SmallVectorImpl &Outs,
67 const SmallVectorImpl &OutVals,
68 const SmallVectorImpl &Ins,
69 DebugLoc dl, SelectionDAG &DAG,
70 SmallVectorImpl &InVals) const;
71
72 virtual EVT getSetCCResultType(EVT VT) const;
73
74 virtual unsigned getNumRegisters(LLVMContext &Context, EVT VT);
75
76 private:
77 SDValue LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const;
78 }; // class PTXTargetLowering
79 } // namespace llvm
80
81 #endif // PTX_ISEL_LOWERING_H
+0
-51
lib/Target/PTX/PTXInstrFormats.td less more
None //===-- PTXInstrFormats.td - PTX Instruction Formats -------*- tablegen -*-===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8
9
10 // Rounding Mode Specifier
11 /*class RoundingMode val> {
12 bits<3> Value = val;
13 }
14
15 def RndDefault : RoundingMode<0>;
16 def RndNearestEven : RoundingMode<1>;
17 def RndNearestZero : RoundingMode<2>;
18 def RndNegInf : RoundingMode<3>;
19 def RndPosInf : RoundingMode<4>;
20 def RndApprox : RoundingMode<5>;*/
21
22
23 // Rounding Mode Operand
24 def RndMode : Operand {
25 let PrintMethod = "printRoundingMode";
26 }
27
28 def RndDefault : PatLeaf<(i32 0)>;
29
30 // PTX Predicate operand, default to (0, 0) = (zero-reg, none).
31 // Leave PrintMethod empty; predicate printing is defined elsewhere.
32 def pred : PredicateOperand
33 (ops (i1 zero_reg), (i32 2))>;
34
35 def RndModeOperand : Operand {
36 let MIOperandInfo = (ops i32imm);
37 }
38
39 // Instruction Types
40 let Namespace = "PTX" in {
41
42 class InstPTX pattern>
43 : Instruction {
44 dag OutOperandList = oops;
45 dag InOperandList = !con(iops, (ins pred:$_p));
46 let AsmString = asmstr; // Predicate printing is defined elsewhere.
47 let Pattern = pattern;
48 let isPredicable = 1;
49 }
50 }
+0
-359
lib/Target/PTX/PTXInstrInfo.cpp less more
None //===-- PTXInstrInfo.cpp - PTX Instruction Information --------------------===//
1 //
2 // The LLVM Compiler Infrastructure
3 //
4 // This file is distributed under the University of Illinois Open Source
5 // License. See LICENSE.TXT for details.
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file contains the PTX implementation of the TargetInstrInfo class.
10 //
11 //===----------------------------------------------------------------------===//
12
13 #define DEBUG_TYPE "ptx-instrinfo"
14
15 #include "PTXInstrInfo.h"
16 #include "PTX.h"
17 #include "llvm/CodeGen/MachineInstrBuilder.h"
18 #include "llvm/CodeGen/MachineRegisterInfo.h"
19 #include "llvm/CodeGen/SelectionDAG.h"
20 #include "llvm/CodeGen/SelectionDAGNodes.h"
21 #include "llvm/Support/Debug.h"
22 #include "llvm/Support/TargetRegistry.h"
23 #include "llvm/Support/raw_ostream.h"
24
25 #define GET_INSTRINFO_CTOR
26 #include "PTXGenInstrInfo.inc"
27
28 using namespace llvm;
29
30 PTXInstrInfo::PTXInstrInfo(PTXTargetMachine &_TM)
31 : PTXGenInstrInfo(),
32 RI(_TM, *this), TM(_TM) {}
33
34 static const struct map_entry {
35 const TargetRegisterClass *cls;
36 const int opcode;
37 } map[] = {
38 { &PTX::RegI16RegClass, PTX::MOVU16rr },
39 { &PTX::RegI32RegClass, PTX::MOVU32rr },
40 { &PTX::RegI64RegClass, PTX::MOVU64rr },
41 { &PTX::RegF32RegClass, PTX::MOVF32rr },
42 { &PTX::RegF64RegClass, PTX::MOVF64rr },
43 { &PTX::RegPredRegClass, PTX::MOVPREDrr }
44 };
45
46 void PTXInstrInfo::copyPhysReg(MachineBasicBlock &MBB,
47 MachineBasicBlock::iterator I, DebugLoc DL,
48 unsigned DstReg, unsigned SrcReg,
49 bool KillSrc) const {
50
51 const MachineRegisterInfo& MRI = MBB.getParent()->getRegInfo();
52 //assert(MRI.getRegClass(SrcReg) == MRI.getRegClass(DstReg) &&
53 // "Invalid register copy between two register classes");
54
55 for (int i = 0, e = sizeof(map)/sizeof(map[0]); i != e; ++i) {
56 if (map[i].cls == MRI.getRegClass(DstReg)) {
57 const MCInstrDesc &MCID = get(map[i].opcode);
58 MachineInstr *MI = BuildMI(MBB, I, DL, MCID, DstReg).
59 addReg(SrcReg, getKillRegState(KillSrc));
60 AddDefaultPredicate(MI);
61 return;
62 }
63 }
64
65 llvm_unreachable("Impossible reg-to-reg copy");
66 }
67
68 bool PTXInstrInfo::copyRegToReg(MachineBasicBlock &MBB,
69 MachineBasicBlock::iterator I,
70 unsigned DstReg, unsigned SrcReg,
71 const TargetRegisterClass *DstRC,
72 const TargetRegisterClass *SrcRC,
73 DebugLoc DL) const {
74 if (DstRC != SrcRC)
75 return false;
76
77 for (int i = 0, e = sizeof(map)/sizeof(map[0]); i != e; ++ i)
78 if (DstRC == map[i].cls) {
79 const MCInstrDesc &MCID = get(map[i].opcode);
80 MachineInstr *MI = BuildMI(MBB, I, DL, MCID, DstReg).addReg(SrcReg);
81 AddDefaultPredicate(MI);
82 return true;
83 }
84
85 return false;
86 }
87
88 bool PTXInstrInfo::isMoveInstr(const MachineInstr& MI,
89 unsigned &SrcReg, unsigned &DstReg,
90 unsigned &SrcSubIdx, unsigned &DstSubIdx) const {
91 switch (MI.getOpcode()) {
92 default:
93 return false;
94 case PTX::MOVU16rr:
95 case PTX::MOVU32rr:
96 case PTX::MOVU64rr:
97 case PTX::MOVF32rr:
98 case PTX::MOVF64rr:
99 case PTX::MOVPREDrr:
100 assert(MI.getNumOperands() >= 2 &&
101 MI.getOperand(0).isReg() && MI.getOperand(1).isReg() &&
102 "Invalid register-register move instruction");
103 SrcSubIdx = DstSubIdx = 0; // No sub-registers
104 DstReg = MI.getOperand(0).getReg();
105 SrcReg = MI.getOperand(1).getReg();
106 return true;
107 }
108 }
109
110 // predicate support
111
112 bool PTXInstrInfo::isPredicated(const MachineInstr *MI) const {
113 int i = MI->findFirstPredOperandIdx();
114 return i != -1 && MI->getOperand(i).getReg() != PTX::NoRegister;
115 }
116
117 bool PTXInstrInfo::isUnpredicatedTerminator(const MachineInstr *MI) const {
118 return !isPredicated(MI) && MI->isTerminator();
119 }
120
121 bool PTXInstrInfo::
122 PredicateInstruction(MachineInstr *MI,
123 const SmallVectorImpl &Pred) const {
124 if (Pred.size() < 2)
125 llvm_unreachable("lesser than 2 predicate operands are provided");
126
127 int i = MI->findFirstPredOperandIdx();
128 if (i == -1)
129 llvm_unreachable("missing predicate operand");
130
131 MI->getOperand(i).setReg(Pred[0].getReg());
132 MI->getOperand(i+1).setImm(Pred[1].getImm());
133
134 return true;
135 }
136
137 bool PTXInstrInfo::
138 SubsumesPredicate(const SmallVectorImpl &Pred1,
139 const SmallVectorImpl &Pred2) const {
140 const MachineOperand &PredReg1 = Pred1[0];
141 const MachineOperand &PredReg2 = Pred2[0];
142 if (PredReg1.getReg() != PredReg2.getReg())
143 return false;
144
145 const MachineOperand &PredOp1 = Pred1[1];
146 const MachineOperand &PredOp2 = Pred2[1];
147 if (PredOp1.getImm() != PredOp2.getImm())
148 return false;
149
150 return true;
151 }
152
153 bool PTXInstrInfo::
154 DefinesPredicate(MachineInstr *MI,
155 std::vector &Pred) const {
156 // If an instruction sets a predicate register, it defines a predicate.
157
158 // TODO supprot 5-operand format of setp instruction
159
160 if (MI->getNumOperands() < 1)
161 return false;
162
163 const MachineOperand &MO = MI->getOperand(0);
164
165 if (!MO.isReg() || RI.getRegClass(MO.getReg()) != &PTX::RegPredRegClass)
166 return false;
167
168 Pred.push_back(MO);
169 Pred.push_back(MachineOperand::CreateImm(PTXPredicate::None));
170 return true;
171 }
172
173 // branch support
174
175 bool PTXInstrInfo::
176 AnalyzeBranch(MachineBasicBlock &MBB,
177 MachineBasicBlock *&TBB,
178 MachineBasicBlock *&FBB,
179 SmallVectorImpl &Cond,
180 bool AllowModify) const {
181 // TODO implement cases when AllowModify is true
182
183 if (MBB.empty())
184 return true;
185
186 MachineBasicBlock::iterator iter = MBB.end();
187 const MachineInstr& instLast1 = *--iter;
188 // for special case that MBB has only 1 instruction
189 const bool IsSizeOne = MBB.size() == 1;
190 // if IsSizeOne is true, *--iter and instLast2 are invalid
191 // we put a dummy value in instLast2 and desc2 since they are used
192 const MachineInstr& instLast2 = IsSizeOne ? instLast1 : *--iter;
193
194 DEBUG(dbgs() << "\n");
195 DEBUG(dbgs() << "AnalyzeBranch: opcode: " << instLast1.getOpcode() << "\n");
196 DEBUG(dbgs() << "AnalyzeBranch: MBB: " << MBB.getName().str() << "\n");
197 DEBUG(dbgs() << "AnalyzeBranch: TBB: " << TBB << "\n");
198 DEBUG(dbgs() << "AnalyzeBranch: FBB: " << FBB << "\n");
199
200 // this block ends with no branches
201 if (!IsAnyKindOfBranch(instLast1)) {
202 DEBUG(dbgs() << "AnalyzeBranch: ends with no branch\n");
203 return false;
204 }
205
206 // this block ends with only an unconditional branch
207 if (instLast1.isUnconditionalBranch() &&
208 // when IsSizeOne is true, it "absorbs" the evaluation of instLast2
209 (IsSizeOne || !IsAnyKindOfBranch(instLast2))) {
210 DEBUG(dbgs() << "AnalyzeBranch: ends with only uncond branch\n");
211 TBB = GetBranchTarget(instLast1);
212 return false;
213 }
214
215 // this block ends with a conditional branch and
216 // it falls through to a successor block
217 if (instLast1.isConditionalBranch() &&
218 IsAnySuccessorAlsoLayoutSuccessor(MBB)) {
219 DEBUG(dbgs() << "AnalyzeBranch: ends with cond branch and fall through\n");
220 TBB = GetBranchTarget(instLast1);
221 int i = instLast1.findFirstPredOperandIdx();
222 Cond.push_back(instLast1.getOperand(i));
223 Cond.push_back(instLast1.getOperand(i+1));
224 return false;
225 }
226
227 // when IsSizeOne is true, we are done
228 if (IsSizeOne)
229 return true;
230
231 // this block ends with a conditional branch
232 // followed by an unconditional branch
233 if (instLast2.isConditionalBranch() &&
234 instLast1.isUnconditionalBranch()) {
235 DEBUG(dbgs() << "AnalyzeBranch: ends with cond and uncond branch\n");
236 TBB = GetBranchTarget(instLast2);
237 FBB = GetBranchTarget(instLast1);
238 int i = instLast2.findFirstPredOperandIdx();
239 Cond.push_back(instLast2.getOperand(i));
240 Cond.push_back(instLast2.getOperand(i+1));
241 return false;
242 }
243
244 // branch cannot be understood
245 DEBUG(dbgs() << "AnalyzeBranch: cannot be understood\n");
246 return true;
247 }
248
249 unsigned PTXInstrInfo::RemoveBranch(MachineBasicBlock &MBB) const {
250 unsigned count = 0;
251 while (!MBB.empty())
252 if (IsAnyKindOfBranch(MBB.back())) {
253 MBB.pop_back();
254 ++count;
255 } else
256 break;
257 DEBUG(dbgs() << "RemoveBranch: MBB: " << MBB.getName().str() << "\n");
258 DEBUG(dbgs() << "RemoveBranch: remove " << count << " branch inst\n");
259 return count;
260 }
261
262 unsigned PTXInstrInfo::
263 InsertBranch(MachineBasicBlock &MBB,
264 MachineBasicBlock *TBB,
265 MachineBasicBlock *FBB,
266 const SmallVectorImpl &Cond,
267 DebugLoc DL) const {
268 DEBUG(dbgs() << "InsertBranch: MBB: " << MBB.getName().str() << "\n");
269 DEBUG(if (TBB) dbgs() << "InsertBranch: TBB: " << TBB->getName().str()
270 << "\n";
271 else dbgs() << "InsertBranch: TBB: (NULL)\n");
272 DEBUG(if (FBB) dbgs() << "InsertBranch: FBB: " << FBB->getName().str()
273 << "\n";
274 else dbgs() << "InsertBranch: FBB: (NULL)\n");
275 DEBUG(dbgs() << "InsertBranch: Cond size: " << Cond.size() << "\n");
276
277 assert(TBB && "TBB is NULL");
278
279 if (FBB) {
280 BuildMI(&MBB, DL, get(PTX::BRAdp))
281 .addMBB(TBB).addReg(Cond[0].getReg()).addImm(Cond[1].getImm());
282 BuildMI(&MBB, DL, get(PTX::BRAd))
283 .addMBB(FBB).addReg(PTX::NoRegister).addImm(PTXPredicate::None);
284 return 2;
285 } else if (Cond.size()) {
286 BuildMI(&MBB, DL, get(PTX::BRAdp))
287 .addMBB(TBB).addReg(Cond[0].getReg()).addImm(Cond[1].getImm());
288 return 1;
289 } else {
290 BuildMI(&MBB, DL, get(PTX::BRAd))
291 .addMBB(TBB).addReg(PTX::NoRegister).addImm(PTXPredicate::None);
292 return 1;
293 }
294 }
295
296 // Memory operand folding for spills
297 void PTXInstrInfo::storeRegToStackSlot(MachineBasicBlock &MBB,
298 MachineBasicBlock::iterator MII,
299 unsigned SrcReg, bool isKill, int FrameIdx,
300 const TargetRegisterClass *RC,
301 const TargetRegisterInfo *TRI) const {
302 llvm_unreachable("storeRegToStackSlot should not be called for PTX");
303 }
304
305 void PTXInstrInfo::loadRegFromStackSlot(MachineBasicBlock &MBB,
306 MachineBasicBlock::iterator MII,
307 unsigned DestReg, int FrameIdx,
308 const TargetRegisterClass *RC,
309 const TargetRegisterInfo *TRI) const {
310 llvm_unreachable("loadRegFromStackSlot should not be called for PTX");
311 }
312
313 // static helper routines
314
315 MachineSDNode *PTXInstrInfo::
316 GetPTXMachineNode(SelectionDAG *DAG, unsigned Opcode,
317 DebugLoc dl, EVT VT, SDValue Op1) {
318 SDValue predReg = DAG->getRegister(PTX::NoRegister, MVT::i1);
319 SDValue predOp = DAG->getTargetConstant(PTXPredicate::None, MVT::i32);
320 SDValue ops[] = { Op1, predReg, predOp };
321 return DAG->getMachineNode(Opcode, dl, VT, ops, array_lengthof(ops));
322 }
323
324