<html><head></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space; ">Dan,<div>It looks like cvt.ll and step.ll are failing. Can you please take a look when you have a moment?</div><div><br></div><div> Regards,</div><div> Chad</div><div><br></div><div><span class="Apple-style-span" style="font-family: Times; "><pre><span class="stdout" style="font-family: 'Courier New', courier, monotype; ">******************** TEST 'LLVM :: CodeGen/PTX/cvt.ll' FAILED ********************
Script:
--
/home/baldrick/osuosl/slave/llvm-x86_64/llvm/Debug+Asserts/bin/llc < /home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/cvt.ll -march=ptx32 | /home/baldrick/osuosl/slave/llvm-x86_64/llvm/Debug+Asserts/bin/FileCheck /home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/cvt.ll
--
Exit Code: 1
Command Output (stderr):
--
/home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/cvt.ll:7:10: error: expected string not found in input
; CHECK: setp.gt.b16 p[[P0:[0-9]+]], rh{{[0-9]+}}, 0
^
<stdin>:1:2: note: scanning from here
.version 2.0
^
<stdin>:9:2: note: possible intended match here
setp.gt.u16 p0, rh12, 0;
^
--
********************</span></pre></span><div><span class="Apple-style-span" style="font-family: Times; "><pre><span class="stdout" style="font-family: 'Courier New', courier, monotype; ">******************** TEST 'LLVM :: CodeGen/PTX/setp.ll' FAILED ********************
Script:
--
/home/baldrick/osuosl/slave/llvm-x86_64/llvm/Debug+Asserts/bin/llc < /home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/setp.ll -march=ptx32 | /home/baldrick/osuosl/slave/llvm-x86_64/llvm/Debug+Asserts/bin/FileCheck /home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/setp.ll
--
Exit Code: 1
Command Output (stderr):
--
/home/baldrick/osuosl/slave/llvm-x86_64/llvm/test/CodeGen/PTX/setp.ll:196:10: error: expected string not found in input
; CHECK: setp.gt.b32 p[[P0:[0-9]+]], r{{[0-9]+}}, 0;
^
<stdin>:193:1: note: scanning from here
}
^
<stdin>:199:2: note: possible intended match here
setp.gt.u32 p0, r14, 0;
^
--
********************</span></pre></span><div><br></div></div></div><div><br></div><div><br><div><div>On Jun 25, 2011, at 11:16 AM, Dan Bailey wrote:</div><br class="Apple-interchange-newline"><blockquote type="cite"><div>Author: drb<br>Date: Sat Jun 25 13:16:28 2011<br>New Revision: 133873<br><br>URL: <a href="http://llvm.org/viewvc/llvm-project?rev=133873&view=rev">http://llvm.org/viewvc/llvm-project?rev=133873&view=rev</a><br>Log:<br>PTX: Reverting implementation of i8.<br><br>The .b8 operations in PTX are far more limiting than I first thought. The mov operation isn't even supported, so there's no way of converting a .pred value into a .b8 without going via .b16, which is<br>not sensible. An improved implementation needs to use the fact that loads and stores automatically extend and truncate to implement support for EXTLOAD and TRUNCSTORE in order to correctly support<br>boolean values.<br><br><br>Modified:<br> llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp<br> llvm/trunk/lib/Target/PTX/PTXCallingConv.td<br> llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp<br> llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp<br> llvm/trunk/lib/Target/PTX/PTXInstrInfo.td<br> llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td<br> llvm/trunk/lib/Target/PTX/generate-register-td.py<br> llvm/trunk/test/CodeGen/PTX/cvt.ll<br> llvm/trunk/test/CodeGen/PTX/ld.ll<br> llvm/trunk/test/CodeGen/PTX/mov.ll<br> llvm/trunk/test/CodeGen/PTX/st.ll<br><br>Modified: llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXAsmPrinter.cpp Sat Jun 25 13:16:28 2011<br>@@ -92,7 +92,6 @@<br> #define TEST_REGCLS(cls, clsstr) \<br> if (PTX::cls ## RegisterClass->contains(RegNo)) return # clsstr;<br> TEST_REGCLS(RegPred, pred);<br>- TEST_REGCLS(RegI8, b8);<br> TEST_REGCLS(RegI16, b16);<br> TEST_REGCLS(RegI32, b32);<br> TEST_REGCLS(RegI64, b64);<br>@@ -125,7 +124,6 @@<br> case Type::IntegerTyID:<br> switch (type->getPrimitiveSizeInBits()) {<br> default: llvm_unreachable("Unknown integer bit-width");<br>- case 8: return ".u8";<br> case 16: return ".u16";<br> case 32: return ".u32";<br> case 64: return ".u64";<br><br>Modified: llvm/trunk/lib/Target/PTX/PTXCallingConv.td<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXCallingConv.td?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXCallingConv.td?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXCallingConv.td (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXCallingConv.td Sat Jun 25 13:16:28 2011<br>@@ -15,7 +15,6 @@<br> // PTX Formal Parameter Calling Convention<br> def CC_PTX : CallingConv<[<br> CCIfType<[i1], CCAssignToReg<[P12, P13, P14, P15, P16, P17, P18, P19, P20, P21, P22, P23, P24, P25, P26, P27, P28, P29, P30, P31, P32, P33, P34, P35, P36, P37, P38, P39, P40, P41, P42, P43, P44, P45, P46, P47, P48, P49, P50, P51, P52, P53, P54, P55, P56, P57, P58, P59, P60, P61, P62, P63, P64, P65, P66, P67, P68, P69, P70, P71, P72, P73, P74, P75, P76, P77, P78, P79, P80, P81, P82, P83, P84, P85, P86, P87, P88, P89, P90, P91, P92, P93, P94, P95, P96, P97, P98, P99, P100, P101, P102, P103, P104, P105, P106, P107, P108, P109, P110, P111, P112, P113, P114, P115, P116, P117, P118, P119, P120, P121, P122, P123, P124, P125, P126, P127]>>,<br>- CCIfType<[i8], CCAssignToReg<[RQ12, RQ13, RQ14, RQ15, RQ16, RQ17, RQ18, RQ19, RQ20, RQ21, RQ22, RQ23, RQ24, RQ25, RQ26, RQ27, RQ28, RQ29, RQ30, RQ31, RQ32, RQ33, RQ34, RQ35, RQ36, RQ37, RQ38, RQ39, RQ40, RQ41, RQ42, RQ43, RQ44, RQ45, RQ46, RQ47, RQ48, RQ49, RQ50, RQ51, RQ52, RQ53, RQ54, RQ55, RQ56, RQ57, RQ58, RQ59, RQ60, RQ61, RQ62, RQ63, RQ64, RQ65, RQ66, RQ67, RQ68, RQ69, RQ70, RQ71, RQ72, RQ73, RQ74, RQ75, RQ76, RQ77, RQ78, RQ79, RQ80, RQ81, RQ82, RQ83, RQ84, RQ85, RQ86, RQ87, RQ88, RQ89, RQ90, RQ91, RQ92, RQ93, RQ94, RQ95, RQ96, RQ97, RQ98, RQ99, RQ100, RQ101, RQ102, RQ103, RQ104, RQ105, RQ106, RQ107, RQ108, RQ109, RQ110, RQ111, RQ112, RQ113, RQ114, RQ115, RQ116, RQ117, RQ118, RQ119, RQ120, RQ121, RQ122, RQ123, RQ124, RQ125, RQ126, RQ127]>>,<br> CCIfType<[i16], CCAssignToReg<[RH12, RH13, RH14, RH15, RH16, RH17, RH18, RH19, RH20, RH21, RH22, RH23, RH24, RH25, RH26, RH27, RH28, RH29, RH30, RH31, RH32, RH33, RH34, RH35, RH36, RH37, RH38, RH39, RH40, RH41, RH42, RH43, RH44, RH45, RH46, RH47, RH48, RH49, RH50, RH51, RH52, RH53, RH54, RH55, RH56, RH57, RH58, RH59, RH60, RH61, RH62, RH63, RH64, RH65, RH66, RH67, RH68, RH69, RH70, RH71, RH72, RH73, RH74, RH75, RH76, RH77, RH78, RH79, RH80, RH81, RH82, RH83, RH84, RH85, RH86, RH87, RH88, RH89, RH90, RH91, RH92, RH93, RH94, RH95, RH96, RH97, RH98, RH99, RH100, RH101, RH102, RH103, RH104, RH105, RH106, RH107, RH108, RH109, RH110, RH111, RH112, RH113, RH114, RH115, RH116, RH117, RH118, RH119, RH120, RH121, RH122, RH123, RH124, RH125, RH126, RH127]>>,<br> CCIfType<[i32,f32], CCAssignToReg<[R12, R13, R14, R15, R16, R17, R18, R19, R20, R21, R22, R23, R24, R25, R26, R27, R28, R29, R30, R31, R32, R33, R34, R35, R36, R37, R38, R39, R40, R41, R42, R43, R44, R45, R46, R47, R48, R49, R50, R51, R52, R53, R54, R55, R56, R57, R58, R59, R60, R61, R62, R63, R64, R65, R66, R67, R68, R69, R70, R71, R72, R73, R74, R75, R76, R77, R78, R79, R80, R81, R82, R83, R84, R85, R86, R87, R88, R89, R90, R91, R92, R93, R94, R95, R96, R97, R98, R99, R100, R101, R102, R103, R104, R105, R106, R107, R108, R109, R110, R111, R112, R113, R114, R115, R116, R117, R118, R119, R120, R121, R122, R123, R124, R125, R126, R127]>>,<br> CCIfType<[i64,f64], CCAssignToReg<[RD12, RD13, RD14, RD15, RD16, RD17, RD18, RD19, RD20, RD21, RD22, RD23, RD24, RD25, RD26, RD27, RD28, RD29, RD30, RD31, RD32, RD33, RD34, RD35, RD36, RD37, RD38, RD39, RD40, RD41, RD42, RD43, RD44, RD45, RD46, RD47, RD48, RD49, RD50, RD51, RD52, RD53, RD54, RD55, RD56, RD57, RD58, RD59, RD60, RD61, RD62, RD63, RD64, RD65, RD66, RD67, RD68, RD69, RD70, RD71, RD72, RD73, RD74, RD75, RD76, RD77, RD78, RD79, RD80, RD81, RD82, RD83, RD84, RD85, RD86, RD87, RD88, RD89, RD90, RD91, RD92, RD93, RD94, RD95, RD96, RD97, RD98, RD99, RD100, RD101, RD102, RD103, RD104, RD105, RD106, RD107, RD108, RD109, RD110, RD111, RD112, RD113, RD114, RD115, RD116, RD117, RD118, RD119, RD120, RD121, RD122, RD123, RD124, RD125, RD126, RD127]>><br>@@ -24,7 +23,6 @@<br> // PTX Return Value Calling Convention<br> def RetCC_PTX : CallingConv<[<br> CCIfType<[i1], CCAssignToReg<[P0, P1, P2, P3, P4, P5, P6, P7, P8, P9, P10, P11]>>,<br>- CCIfType<[i8], CCAssignToReg<[RQ0, RQ1, RQ2, RQ3, RQ4, RQ5, RQ6, RQ7, RQ8, RQ9, RQ10, RQ11]>>,<br> CCIfType<[i16], CCAssignToReg<[RH0, RH1, RH2, RH3, RH4, RH5, RH6, RH7, RH8, RH9, RH10, RH11]>>,<br> CCIfType<[i32,f32], CCAssignToReg<[R0, R1, R2, R3, R4, R5, R6, R7, R8, R9, R10, R11]>>,<br> CCIfType<[i64,f64], CCAssignToReg<[RD0, RD1, RD2, RD3, RD4, RD5, RD6, RD7, RD8, RD9, RD10, RD11]>><br><br>Modified: llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXISelLowering.cpp Sat Jun 25 13:16:28 2011<br>@@ -40,7 +40,6 @@<br> : TargetLowering(TM, new TargetLoweringObjectFileELF()) {<br> // Set up the register classes.<br> addRegisterClass(MVT::i1, PTX::RegPredRegisterClass);<br>- addRegisterClass(MVT::i8, PTX::RegI8RegisterClass);<br> addRegisterClass(MVT::i16, PTX::RegI16RegisterClass);<br> addRegisterClass(MVT::i32, PTX::RegI32RegisterClass);<br> addRegisterClass(MVT::i64, PTX::RegI64RegisterClass);<br>@@ -48,48 +47,59 @@<br> addRegisterClass(MVT::f64, PTX::RegF64RegisterClass);<br><br> setBooleanContents(ZeroOrOneBooleanContent);<br>-<br>- setOperationAction(ISD::EXCEPTIONADDR, MVT::i32, Expand);<br>-<br>- setOperationAction(ISD::ConstantFP, MVT::f32, Legal);<br>- setOperationAction(ISD::ConstantFP, MVT::f64, Legal);<br>- <br>- // Promote i1 type<br>- setLoadExtAction(ISD::EXTLOAD, MVT::i1, Promote);<br>- setLoadExtAction(ISD::ZEXTLOAD, MVT::i1, Promote);<br>- setLoadExtAction(ISD::SEXTLOAD, MVT::i1, Promote);<br>+ setMinFunctionAlignment(2);<br><br>- setTruncStoreAction(MVT::i8, MVT::i1, Promote);<br>+ ////////////////////////////////////<br>+ /////////// Expansion //////////////<br>+ ////////////////////////////////////<br><br>- setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i1, Expand);<br>+ // (any/zero/sign) extload => load + (any/zero/sign) extend<br><br>- // Turn i16 (z)extload into load + (z)extend<br> setLoadExtAction(ISD::EXTLOAD, MVT::i16, Expand);<br> setLoadExtAction(ISD::ZEXTLOAD, MVT::i16, Expand);<br> setLoadExtAction(ISD::SEXTLOAD, MVT::i16, Expand);<br>-<br>- // Turn f32 extload into load + fextend<br>- setLoadExtAction(ISD::EXTLOAD, MVT::f32, Expand);<br>-<br>- // Turn f64 truncstore into trunc + store.<br>- setTruncStoreAction(MVT::f64, MVT::f32, Expand);<br>-<br>- // Customize translation of memory addresses<br>- setOperationAction(ISD::GlobalAddress, MVT::i32, Custom);<br>- setOperationAction(ISD::GlobalAddress, MVT::i64, Custom);<br>-<br>- // Expand BR_CC into BRCOND<br>+ <br>+ // f32 extload => load + fextend<br>+ <br>+ setLoadExtAction(ISD::EXTLOAD, MVT::f32, Expand); <br>+ <br>+ // f64 truncstore => trunc + store<br>+ <br>+ setTruncStoreAction(MVT::f64, MVT::f32, Expand); <br>+ <br>+ // sign_extend_inreg => sign_extend<br>+ <br>+ setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::i1, Expand);<br>+ <br>+ // br_cc => brcond<br>+ <br> setOperationAction(ISD::BR_CC, MVT::Other, Expand);<br><br>- // Expand SELECT_CC into SETCC<br>+ // select_cc => setcc<br>+ <br> setOperationAction(ISD::SELECT_CC, MVT::Other, Expand);<br> setOperationAction(ISD::SELECT_CC, MVT::f32, Expand);<br> setOperationAction(ISD::SELECT_CC, MVT::f64, Expand);<br>-<br>- // need to lower SETCC of RegPred into bitwise logic<br>+ <br>+ ////////////////////////////////////<br>+ //////////// Legal /////////////////<br>+ ////////////////////////////////////<br>+ <br>+ setOperationAction(ISD::ConstantFP, MVT::f32, Legal);<br>+ setOperationAction(ISD::ConstantFP, MVT::f64, Legal);<br>+ <br>+ ////////////////////////////////////<br>+ //////////// Custom ////////////////<br>+ ////////////////////////////////////<br>+ <br>+ // customise setcc to use bitwise logic if possible<br>+ <br> setOperationAction(ISD::SETCC, MVT::i1, Custom);<br><br>- setMinFunctionAlignment(2);<br>+ // customize translation of memory addresses<br>+ <br>+ setOperationAction(ISD::GlobalAddress, MVT::i32, Custom);<br>+ setOperationAction(ISD::GlobalAddress, MVT::i64, Custom);<br><br> // Compute derived properties from the register classes<br> computeRegisterProperties();<br>@@ -187,7 +197,6 @@<br> bool operator==(MVT::SimpleValueType _VT) const { return VT == _VT; }<br> } argmap[] = {<br> argmap_entry(MVT::i1, PTX::RegPredRegisterClass),<br>- argmap_entry(MVT::i8, PTX::RegI8RegisterClass),<br> argmap_entry(MVT::i16, PTX::RegI16RegisterClass),<br> argmap_entry(MVT::i32, PTX::RegI32RegisterClass),<br> argmap_entry(MVT::i64, PTX::RegI64RegisterClass),<br>@@ -264,9 +273,6 @@<br> if (RegVT == MVT::i1) {<br> TRC = PTX::RegPredRegisterClass;<br> }<br>- else if (RegVT == MVT::i8) {<br>- TRC = PTX::RegI8RegisterClass;<br>- }<br> else if (RegVT == MVT::i16) {<br> TRC = PTX::RegI16RegisterClass;<br> }<br><br>Modified: llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXInstrInfo.cpp Sat Jun 25 13:16:28 2011<br>@@ -33,7 +33,6 @@<br> const TargetRegisterClass *cls;<br> const int opcode;<br> } map[] = {<br>- { &PTX::RegI8RegClass, PTX::MOVU8rr },<br> { &PTX::RegI16RegClass, PTX::MOVU16rr },<br> { &PTX::RegI32RegClass, PTX::MOVU32rr },<br> { &PTX::RegI64RegClass, PTX::MOVU64rr },<br>@@ -303,9 +302,7 @@<br> int OpCode;<br><br> // Select the appropriate opcode based on the register class<br>- if (RC == PTX::RegI8RegisterClass) {<br>- OpCode = PTX::STACKSTOREI8;<br>- } else if (RC == PTX::RegI16RegisterClass) {<br>+ if (RC == PTX::RegI16RegisterClass) {<br> OpCode = PTX::STACKSTOREI16;<br> } else if (RC == PTX::RegI32RegisterClass) {<br> OpCode = PTX::STACKSTOREI32;<br>@@ -340,9 +337,7 @@<br> int OpCode;<br><br> // Select the appropriate opcode based on the register class<br>- if (RC == PTX::RegI8RegisterClass) {<br>- OpCode = PTX::STACKLOADI8;<br>- } else if (RC == PTX::RegI16RegisterClass) {<br>+ if (RC == PTX::RegI16RegisterClass) {<br> OpCode = PTX::STACKLOADI16;<br> } else if (RC == PTX::RegI32RegisterClass) {<br> OpCode = PTX::STACKLOADI32;<br><br>Modified: llvm/trunk/lib/Target/PTX/PTXInstrInfo.td<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.td?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXInstrInfo.td?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXInstrInfo.td (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXInstrInfo.td Sat Jun 25 13:16:28 2011<br>@@ -537,7 +537,6 @@<br> }<br><br> multiclass PTX_LD_ALL<string opstr, PatFrag pat_load> {<br>- defm u8 : PTX_LD<opstr, ".u8", RegI8, pat_load>;<br> defm u16 : PTX_LD<opstr, ".u16", RegI16, pat_load>;<br> defm u32 : PTX_LD<opstr, ".u32", RegI32, pat_load>;<br> defm u64 : PTX_LD<opstr, ".u64", RegI64, pat_load>;<br>@@ -573,7 +572,6 @@<br> }<br><br> multiclass PTX_ST_ALL<string opstr, PatFrag pat_store> {<br>- defm u8 : PTX_ST<opstr, ".u8", RegI8, pat_store>;<br> defm u16 : PTX_ST<opstr, ".u16", RegI16, pat_store>;<br> defm u32 : PTX_ST<opstr, ".u32", RegI32, pat_store>;<br> defm u64 : PTX_ST<opstr, ".u64", RegI64, pat_store>;<br>@@ -785,27 +783,22 @@<br> let neverHasSideEffects = 1 in {<br> def MOVPREDrr<br> : InstPTX<(outs RegPred:$d), (ins RegPred:$a), "mov.pred\t$d, $a", []>;<br>- def MOVU8rr<br>- : InstPTX<(outs RegI8:$d), (ins RegI8:$a), "mov.u8\t$d, $a", []>;<br> def MOVU16rr<br>- : InstPTX<(outs RegI16:$d), (ins RegI16:$a), "mov.u16\t$d, $a", []>;<br>+ : InstPTX<(outs RegI16:$d), (ins RegI16:$a), "mov.u16\t$d, $a", []>;<br> def MOVU32rr<br>- : InstPTX<(outs RegI32:$d), (ins RegI32:$a), "mov.u32\t$d, $a", []>;<br>+ : InstPTX<(outs RegI32:$d), (ins RegI32:$a), "mov.u32\t$d, $a", []>;<br> def MOVU64rr<br>- : InstPTX<(outs RegI64:$d), (ins RegI64:$a), "mov.u64\t$d, $a", []>;<br>+ : InstPTX<(outs RegI64:$d), (ins RegI64:$a), "mov.u64\t$d, $a", []>;<br> def MOVF32rr<br>- : InstPTX<(outs RegF32:$d), (ins RegF32:$a), "mov.f32\t$d, $a", []>;<br>+ : InstPTX<(outs RegF32:$d), (ins RegF32:$a), "mov.f32\t$d, $a", []>;<br> def MOVF64rr<br>- : InstPTX<(outs RegF64:$d), (ins RegF64:$a), "mov.f64\t$d, $a", []>;<br>+ : InstPTX<(outs RegF64:$d), (ins RegF64:$a), "mov.f64\t$d, $a", []>;<br> }<br><br> let isReMaterializable = 1, isAsCheapAsAMove = 1 in {<br> def MOVPREDri<br> : InstPTX<(outs RegPred:$d), (ins i1imm:$a), "mov.pred\t$d, $a",<br> [(set RegPred:$d, imm:$a)]>;<br>- def MOVU8ri<br>- : InstPTX<(outs RegI8:$d), (ins i8imm:$a), "mov.u8\t$d, $a",<br>- [(set RegI8:$d, imm:$a)]>;<br> def MOVU16ri<br> : InstPTX<(outs RegI16:$d), (ins i16imm:$a), "mov.u16\t$d, $a",<br> [(set RegI16:$d, imm:$a)]>;<br>@@ -845,9 +838,6 @@<br> def LDpiPred : InstPTX<(outs RegPred:$d), (ins MEMpi:$a),<br> "ld.param.pred\t$d, [$a]",<br> [(set RegPred:$d, (PTXloadparam timm:$a))]>;<br>- def LDpiU8 : InstPTX<(outs RegI8:$d), (ins MEMpi:$a),<br>- "ld.param.u8\t$d, [$a]",<br>- [(set RegI8:$d, (PTXloadparam timm:$a))]>;<br> def LDpiU16 : InstPTX<(outs RegI16:$d), (ins MEMpi:$a),<br> "ld.param.u16\t$d, [$a]",<br> [(set RegI16:$d, (PTXloadparam timm:$a))]>;<br>@@ -867,9 +857,6 @@<br> def STpiPred : InstPTX<(outs), (ins MEMret:$d, RegPred:$a),<br> "st.param.pred\t[$d], $a",<br> [(PTXstoreparam timm:$d, RegPred:$a)]>;<br>- def STpiU8 : InstPTX<(outs), (ins MEMret:$d, RegI8:$a),<br>- "st.param.u8\t[$d], $a",<br>- [(PTXstoreparam timm:$d, RegI8:$a)]>;<br> def STpiU16 : InstPTX<(outs), (ins MEMret:$d, RegI16:$a),<br> "st.param.u16\t[$d], $a",<br> [(PTXstoreparam timm:$d, RegI16:$a)]>;<br>@@ -900,62 +887,34 @@<br> // PTX does not directly support converting to a predicate type, so we fake it<br> // by performing a greater-than test between the value and zero. This follows<br> // the C convention that any non-zero value is equivalent to 'true'.<br>-def CVT_pred_u8<br>- : InstPTX<(outs RegPred:$d), (ins RegI8:$a), "setp.gt.b8\t$d, $a, 0",<br>- [(set RegPred:$d, (trunc RegI8:$a))]>;<br>-<br> def CVT_pred_u16<br>- : InstPTX<(outs RegPred:$d), (ins RegI16:$a), "setp.gt.b16\t$d, $a, 0",<br>+ : InstPTX<(outs RegPred:$d), (ins RegI16:$a), "setp.gt.u16\t$d, $a, 0",<br> [(set RegPred:$d, (trunc RegI16:$a))]>;<br><br> def CVT_pred_u32<br>- : InstPTX<(outs RegPred:$d), (ins RegI32:$a), "setp.gt.b32\t$d, $a, 0",<br>+ : InstPTX<(outs RegPred:$d), (ins RegI32:$a), "setp.gt.u32\t$d, $a, 0",<br> [(set RegPred:$d, (trunc RegI32:$a))]>;<br><br> def CVT_pred_u64<br>- : InstPTX<(outs RegPred:$d), (ins RegI64:$a), "setp.gt.b64\t$d, $a, 0",<br>+ : InstPTX<(outs RegPred:$d), (ins RegI64:$a), "setp.gt.u64\t$d, $a, 0",<br> [(set RegPred:$d, (trunc RegI64:$a))]>;<br><br> def CVT_pred_f32<br>- : InstPTX<(outs RegPred:$d), (ins RegF32:$a), "setp.gt.b32\t$d, $a, 0",<br>+ : InstPTX<(outs RegPred:$d), (ins RegF32:$a), "setp.gt.f32\t$d, $a, 0",<br> [(set RegPred:$d, (fp_to_uint RegF32:$a))]>;<br><br> def CVT_pred_f64<br>- : InstPTX<(outs RegPred:$d), (ins RegF64:$a), "setp.gt.b64\t$d, $a, 0",<br>+ : InstPTX<(outs RegPred:$d), (ins RegF64:$a), "setp.gt.f64\t$d, $a, 0",<br> [(set RegPred:$d, (fp_to_uint RegF64:$a))]>;<br><br>-// Conversion to u8<br>-// PTX does not directly support converting a predicate to a value, so we<br>-// use a select instruction to select either 0 or 1 (integer or fp) based<br>-// on the truth value of the predicate.<br>-def CVT_u8_pred<br>- : InstPTX<(outs RegI8:$d), (ins RegPred:$a), "selp.u8\t$d, 1, 0, $a",<br>- [(set RegI8:$d, (zext RegPred:$a))]>;<br>-<br>-def CVT_u8_preds<br>- : InstPTX<(outs RegI8:$d), (ins RegPred:$a), "selp.u8\t$d, 1, 0, $a",<br>- [(set RegI8:$d, (sext RegPred:$a))]>;<br>-<br>-def CVT_u8_u32<br>- : InstPTX<(outs RegI8:$d), (ins RegI32:$a), "cvt.u8.u32\t$d, $a",<br>- [(set RegI8:$d, (trunc RegI32:$a))]>;<br>-<br>-def CVT_u8_u64<br>- : InstPTX<(outs RegI8:$d), (ins RegI64:$a), "cvt.u8.u64\t$d, $a",<br>- [(set RegI8:$d, (trunc RegI64:$a))]>;<br>-<br>-def CVT_u8_f32<br>- : InstPTX<(outs RegI8:$d), (ins RegF32:$a), "cvt.rzi.u8.f32\t$d, $a",<br>- [(set RegI8:$d, (fp_to_uint RegF32:$a))]>;<br>-<br>-def CVT_u8_f64<br>- : InstPTX<(outs RegI8:$d), (ins RegF64:$a), "cvt.rzi.u8.f64\t$d, $a",<br>- [(set RegI8:$d, (fp_to_uint RegF64:$a))]>;<br>-<br> // Conversion to u16<br> // PTX does not directly support converting a predicate to a value, so we<br> // use a select instruction to select either 0 or 1 (integer or fp) based<br> // on the truth value of the predicate.<br>+def CVT_u16_preda<br>+ : InstPTX<(outs RegI16:$d), (ins RegPred:$a), "selp.u16\t$d, 1, 0, $a",<br>+ [(set RegI16:$d, (anyext RegPred:$a))]>;<br>+<br> def CVT_u16_pred<br> : InstPTX<(outs RegI16:$d), (ins RegPred:$a), "selp.u16\t$d, 1, 0, $a",<br> [(set RegI16:$d, (zext RegPred:$a))]>;<br>@@ -964,14 +923,6 @@<br> : InstPTX<(outs RegI16:$d), (ins RegPred:$a), "selp.u16\t$d, 1, 0, $a",<br> [(set RegI16:$d, (sext RegPred:$a))]>;<br><br>-def CVT_u16_u8<br>- : InstPTX<(outs RegI16:$d), (ins RegI8:$a), "cvt.u16.u8\t$d, $a",<br>- [(set RegI16:$d, (zext RegI8:$a))]>;<br>-<br>-def CVT_u16_s8<br>- : InstPTX<(outs RegI16:$d), (ins RegI8:$a), "cvt.u16.s8\t$d, $a",<br>- [(set RegI16:$d, (sext RegI8:$a))]>;<br>-<br> def CVT_u16_u32<br> : InstPTX<(outs RegI16:$d), (ins RegI32:$a), "cvt.u16.u32\t$d, $a",<br> [(set RegI16:$d, (trunc RegI32:$a))]>;<br>@@ -994,9 +945,9 @@<br> : InstPTX<(outs RegI32:$d), (ins RegPred:$a), "selp.u32\t$d, 1, 0, $a",<br> [(set RegI32:$d, (zext RegPred:$a))]>;<br><br>-def CVT_u32_u8<br>- : InstPTX<(outs RegI32:$d), (ins RegI8:$a), "cvt.u32.u8\t$d, $a",<br>- [(set RegI32:$d, (zext RegI8:$a))]>;<br>+def CVT_u32_b16<br>+ : InstPTX<(outs RegI32:$d), (ins RegI16:$a), "cvt.u32.u16\t$d, $a",<br>+ [(set RegI32:$d, (anyext RegI16:$a))]>;<br><br> def CVT_u32_u16<br> : InstPTX<(outs RegI32:$d), (ins RegI16:$a), "cvt.u32.u16\t$d, $a",<br>@@ -1006,10 +957,6 @@<br> : InstPTX<(outs RegI32:$d), (ins RegPred:$a), "selp.u32\t$d, 1, 0, $a",<br> [(set RegI32:$d, (sext RegPred:$a))]>;<br><br>-def CVT_u32_s8<br>- : InstPTX<(outs RegI32:$d), (ins RegI8:$a), "cvt.u32.s8\t$d, $a",<br>- [(set RegI32:$d, (zext RegI8:$a))]>;<br>-<br> def CVT_u32_s16<br> : InstPTX<(outs RegI32:$d), (ins RegI16:$a), "cvt.u32.s16\t$d, $a",<br> [(set RegI32:$d, (sext RegI16:$a))]>;<br>@@ -1032,30 +979,22 @@<br> : InstPTX<(outs RegI64:$d), (ins RegPred:$a), "selp.u64\t$d, 1, 0, $a",<br> [(set RegI64:$d, (zext RegPred:$a))]>;<br><br>-def CVT_u64_u8<br>- : InstPTX<(outs RegI64:$d), (ins RegI8:$a), "cvt.u64.u8\t$d, $a",<br>- [(set RegI64:$d, (zext RegI8:$a))]>;<br>-<br>-def CVT_u64_u16<br>- : InstPTX<(outs RegI64:$d), (ins RegI16:$a), "cvt.u64.u16\t$d, $a",<br>- [(set RegI64:$d, (zext RegI16:$a))]>;<br>-<br>-def CVT_u64_u32<br>- : InstPTX<(outs RegI64:$d), (ins RegI32:$a), "cvt.u64.u32\t$d, $a",<br>- [(set RegI64:$d, (zext RegI32:$a))]>;<br>-<br> def CVT_u64_preds<br> : InstPTX<(outs RegI64:$d), (ins RegPred:$a), "selp.u64\t$d, 1, 0, $a",<br> [(set RegI64:$d, (sext RegPred:$a))]>;<br><br>-def CVT_u64_s8<br>- : InstPTX<(outs RegI64:$d), (ins RegI8:$a), "cvt.u64.s8\t$d, $a",<br>- [(set RegI64:$d, (zext RegI8:$a))]>;<br>+def CVT_u64_u16<br>+ : InstPTX<(outs RegI64:$d), (ins RegI16:$a), "cvt.u64.u16\t$d, $a",<br>+ [(set RegI64:$d, (zext RegI16:$a))]>;<br><br> def CVT_u64_s16<br> : InstPTX<(outs RegI64:$d), (ins RegI16:$a), "cvt.u64.s16\t$d, $a",<br> [(set RegI64:$d, (sext RegI16:$a))]>;<br><br>+def CVT_u64_u32<br>+ : InstPTX<(outs RegI64:$d), (ins RegI32:$a), "cvt.u64.u32\t$d, $a",<br>+ [(set RegI64:$d, (zext RegI32:$a))]>;<br>+<br> def CVT_u64_s32<br> : InstPTX<(outs RegI64:$d), (ins RegI32:$a), "cvt.u64.s32\t$d, $a",<br> [(set RegI64:$d, (sext RegI32:$a))]>;<br>@@ -1075,10 +1014,6 @@<br> "selp.f32\t$d, 0F3F800000, 0F00000000, $a", // 1.0<br> [(set RegF32:$d, (uint_to_fp RegPred:$a))]>;<br><br>-def CVT_f32_u8<br>- : InstPTX<(outs RegF32:$d), (ins RegI8:$a), "cvt.rn.f32.u8\t$d, $a",<br>- [(set RegF32:$d, (uint_to_fp RegI8:$a))]>;<br>-<br> def CVT_f32_u16<br> : InstPTX<(outs RegF32:$d), (ins RegI16:$a), "cvt.rn.f32.u16\t$d, $a",<br> [(set RegF32:$d, (uint_to_fp RegI16:$a))]>;<br>@@ -1102,10 +1037,6 @@<br> "selp.f64\t$d, 0D3F80000000000000, 0D0000000000000000, $a", // 1.0<br> [(set RegF64:$d, (uint_to_fp RegPred:$a))]>;<br><br>-def CVT_f64_u8<br>- : InstPTX<(outs RegF64:$d), (ins RegI8:$a), "cvt.rn.f64.u8\t$d, $a",<br>- [(set RegF64:$d, (uint_to_fp RegI8:$a))]>;<br>-<br> def CVT_f64_u16<br> : InstPTX<(outs RegF64:$d), (ins RegI16:$a), "cvt.rn.f64.u16\t$d, $a",<br> [(set RegF64:$d, (uint_to_fp RegI16:$a))]>;<br>@@ -1144,8 +1075,6 @@<br><br> ///===- Spill Instructions ------------------------------------------------===//<br> // Special instructions used for stack spilling<br>-def STACKSTOREI8 : InstPTX<(outs), (ins i32imm:$d, RegI8:$a),<br>- "mov.u8\ts$d, $a", []>;<br> def STACKSTOREI16 : InstPTX<(outs), (ins i32imm:$d, RegI16:$a),<br> "mov.u16\ts$d, $a", []>;<br> def STACKSTOREI32 : InstPTX<(outs), (ins i32imm:$d, RegI32:$a),<br>@@ -1157,8 +1086,6 @@<br> def STACKSTOREF64 : InstPTX<(outs), (ins i32imm:$d, RegF64:$a),<br> "mov.f64\ts$d, $a", []>;<br><br>-def STACKLOADI8 : InstPTX<(outs), (ins RegI8:$d, i32imm:$a),<br>- "mov.u8\t$d, s$a", []>;<br> def STACKLOADI16 : InstPTX<(outs), (ins RegI16:$d, i32imm:$a),<br> "mov.u16\t$d, s$a", []>;<br> def STACKLOADI32 : InstPTX<(outs), (ins RegI32:$d, i32imm:$a),<br><br>Modified: llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td (original)<br>+++ llvm/trunk/lib/Target/PTX/PTXRegisterInfo.td Sat Jun 25 13:16:28 2011<br>@@ -151,137 +151,6 @@<br> def P126 : PTXReg<"p126">;<br> def P127 : PTXReg<"p127">;<br><br>-///===- 8-Bit Registers --------------------------------------------------===//<br>-<br>-def RQ0 : PTXReg<"rq0">;<br>-def RQ1 : PTXReg<"rq1">;<br>-def RQ2 : PTXReg<"rq2">;<br>-def RQ3 : PTXReg<"rq3">;<br>-def RQ4 : PTXReg<"rq4">;<br>-def RQ5 : PTXReg<"rq5">;<br>-def RQ6 : PTXReg<"rq6">;<br>-def RQ7 : PTXReg<"rq7">;<br>-def RQ8 : PTXReg<"rq8">;<br>-def RQ9 : PTXReg<"rq9">;<br>-def RQ10 : PTXReg<"rq10">;<br>-def RQ11 : PTXReg<"rq11">;<br>-def RQ12 : PTXReg<"rq12">;<br>-def RQ13 : PTXReg<"rq13">;<br>-def RQ14 : PTXReg<"rq14">;<br>-def RQ15 : PTXReg<"rq15">;<br>-def RQ16 : PTXReg<"rq16">;<br>-def RQ17 : PTXReg<"rq17">;<br>-def RQ18 : PTXReg<"rq18">;<br>-def RQ19 : PTXReg<"rq19">;<br>-def RQ20 : PTXReg<"rq20">;<br>-def RQ21 : PTXReg<"rq21">;<br>-def RQ22 : PTXReg<"rq22">;<br>-def RQ23 : PTXReg<"rq23">;<br>-def RQ24 : PTXReg<"rq24">;<br>-def RQ25 : PTXReg<"rq25">;<br>-def RQ26 : PTXReg<"rq26">;<br>-def RQ27 : PTXReg<"rq27">;<br>-def RQ28 : PTXReg<"rq28">;<br>-def RQ29 : PTXReg<"rq29">;<br>-def RQ30 : PTXReg<"rq30">;<br>-def RQ31 : PTXReg<"rq31">;<br>-def RQ32 : PTXReg<"rq32">;<br>-def RQ33 : PTXReg<"rq33">;<br>-def RQ34 : PTXReg<"rq34">;<br>-def RQ35 : PTXReg<"rq35">;<br>-def RQ36 : PTXReg<"rq36">;<br>-def RQ37 : PTXReg<"rq37">;<br>-def RQ38 : PTXReg<"rq38">;<br>-def RQ39 : PTXReg<"rq39">;<br>-def RQ40 : PTXReg<"rq40">;<br>-def RQ41 : PTXReg<"rq41">;<br>-def RQ42 : PTXReg<"rq42">;<br>-def RQ43 : PTXReg<"rq43">;<br>-def RQ44 : PTXReg<"rq44">;<br>-def RQ45 : PTXReg<"rq45">;<br>-def RQ46 : PTXReg<"rq46">;<br>-def RQ47 : PTXReg<"rq47">;<br>-def RQ48 : PTXReg<"rq48">;<br>-def RQ49 : PTXReg<"rq49">;<br>-def RQ50 : PTXReg<"rq50">;<br>-def RQ51 : PTXReg<"rq51">;<br>-def RQ52 : PTXReg<"rq52">;<br>-def RQ53 : PTXReg<"rq53">;<br>-def RQ54 : PTXReg<"rq54">;<br>-def RQ55 : PTXReg<"rq55">;<br>-def RQ56 : PTXReg<"rq56">;<br>-def RQ57 : PTXReg<"rq57">;<br>-def RQ58 : PTXReg<"rq58">;<br>-def RQ59 : PTXReg<"rq59">;<br>-def RQ60 : PTXReg<"rq60">;<br>-def RQ61 : PTXReg<"rq61">;<br>-def RQ62 : PTXReg<"rq62">;<br>-def RQ63 : PTXReg<"rq63">;<br>-def RQ64 : PTXReg<"rq64">;<br>-def RQ65 : PTXReg<"rq65">;<br>-def RQ66 : PTXReg<"rq66">;<br>-def RQ67 : PTXReg<"rq67">;<br>-def RQ68 : PTXReg<"rq68">;<br>-def RQ69 : PTXReg<"rq69">;<br>-def RQ70 : PTXReg<"rq70">;<br>-def RQ71 : PTXReg<"rq71">;<br>-def RQ72 : PTXReg<"rq72">;<br>-def RQ73 : PTXReg<"rq73">;<br>-def RQ74 : PTXReg<"rq74">;<br>-def RQ75 : PTXReg<"rq75">;<br>-def RQ76 : PTXReg<"rq76">;<br>-def RQ77 : PTXReg<"rq77">;<br>-def RQ78 : PTXReg<"rq78">;<br>-def RQ79 : PTXReg<"rq79">;<br>-def RQ80 : PTXReg<"rq80">;<br>-def RQ81 : PTXReg<"rq81">;<br>-def RQ82 : PTXReg<"rq82">;<br>-def RQ83 : PTXReg<"rq83">;<br>-def RQ84 : PTXReg<"rq84">;<br>-def RQ85 : PTXReg<"rq85">;<br>-def RQ86 : PTXReg<"rq86">;<br>-def RQ87 : PTXReg<"rq87">;<br>-def RQ88 : PTXReg<"rq88">;<br>-def RQ89 : PTXReg<"rq89">;<br>-def RQ90 : PTXReg<"rq90">;<br>-def RQ91 : PTXReg<"rq91">;<br>-def RQ92 : PTXReg<"rq92">;<br>-def RQ93 : PTXReg<"rq93">;<br>-def RQ94 : PTXReg<"rq94">;<br>-def RQ95 : PTXReg<"rq95">;<br>-def RQ96 : PTXReg<"rq96">;<br>-def RQ97 : PTXReg<"rq97">;<br>-def RQ98 : PTXReg<"rq98">;<br>-def RQ99 : PTXReg<"rq99">;<br>-def RQ100 : PTXReg<"rq100">;<br>-def RQ101 : PTXReg<"rq101">;<br>-def RQ102 : PTXReg<"rq102">;<br>-def RQ103 : PTXReg<"rq103">;<br>-def RQ104 : PTXReg<"rq104">;<br>-def RQ105 : PTXReg<"rq105">;<br>-def RQ106 : PTXReg<"rq106">;<br>-def RQ107 : PTXReg<"rq107">;<br>-def RQ108 : PTXReg<"rq108">;<br>-def RQ109 : PTXReg<"rq109">;<br>-def RQ110 : PTXReg<"rq110">;<br>-def RQ111 : PTXReg<"rq111">;<br>-def RQ112 : PTXReg<"rq112">;<br>-def RQ113 : PTXReg<"rq113">;<br>-def RQ114 : PTXReg<"rq114">;<br>-def RQ115 : PTXReg<"rq115">;<br>-def RQ116 : PTXReg<"rq116">;<br>-def RQ117 : PTXReg<"rq117">;<br>-def RQ118 : PTXReg<"rq118">;<br>-def RQ119 : PTXReg<"rq119">;<br>-def RQ120 : PTXReg<"rq120">;<br>-def RQ121 : PTXReg<"rq121">;<br>-def RQ122 : PTXReg<"rq122">;<br>-def RQ123 : PTXReg<"rq123">;<br>-def RQ124 : PTXReg<"rq124">;<br>-def RQ125 : PTXReg<"rq125">;<br>-def RQ126 : PTXReg<"rq126">;<br>-def RQ127 : PTXReg<"rq127">;<br>-<br> ///===- 16-Bit Registers --------------------------------------------------===//<br><br> def RH0 : PTXReg<"rh0">;<br>@@ -679,7 +548,6 @@<br> // Register classes<br> //===----------------------------------------------------------------------===//<br> def RegPred : RegisterClass<"PTX", [i1], 8, (sequence "P%u", 0, 127)>;<br>-def RegI8 : RegisterClass<"PTX", [i8], 8, (sequence "RQ%u", 0, 127)>;<br> def RegI16 : RegisterClass<"PTX", [i16], 16, (sequence "RH%u", 0, 127)>;<br> def RegI32 : RegisterClass<"PTX", [i32], 32, (sequence "R%u", 0, 127)>;<br> def RegI64 : RegisterClass<"PTX", [i64], 64, (sequence "RD%u", 0, 127)>;<br><br>Modified: llvm/trunk/lib/Target/PTX/generate-register-td.py<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/generate-register-td.py?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/PTX/generate-register-td.py?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/lib/Target/PTX/generate-register-td.py (original)<br>+++ llvm/trunk/lib/Target/PTX/generate-register-td.py Sat Jun 25 13:16:28 2011<br>@@ -15,16 +15,15 @@<br> from sys import argv, exit, stdout<br><br><br>-if len(argv) != 6:<br>- print('Usage: generate-register-td.py <num_preds> <num_8> <num_16> <num_32> <num_64>')<br>+if len(argv) != 5:<br>+ print('Usage: generate-register-td.py <num_preds> <num_16> <num_32> <num_64>')<br> exit(1)<br><br> try:<br> num_pred = int(argv[1])<br>- num_8bit = int(argv[2])<br>- num_16bit = int(argv[3])<br>- num_32bit = int(argv[4])<br>- num_64bit = int(argv[5])<br>+ num_16bit = int(argv[2])<br>+ num_32bit = int(argv[3])<br>+ num_64bit = int(argv[4])<br> except:<br> print('ERROR: Invalid integer parameter')<br> exit(1)<br>@@ -61,11 +60,6 @@<br> for r in range(0, num_pred):<br> td_file.write('def P%d : PTXReg<"p%d">;\n' % (r, r))<br><br>-# Print 8-bit registers<br>-td_file.write('\n///===- 8-Bit Registers --------------------------------------------------===//\n\n')<br>-for r in range(0, num_8bit):<br>- td_file.write('def RQ%d : PTXReg<"rq%d">;\n' % (r, r))<br>-<br> # Print 16-bit registers<br> td_file.write('\n///===- 16-Bit Registers --------------------------------------------------===//\n\n')<br> for r in range(0, num_16bit):<br>@@ -92,7 +86,6 @@<br> # Print register classes<br><br> td_file.write('def RegPred : RegisterClass<"PTX", [i1], 8, (sequence "P%%u", 0, %d)>;\n' % (num_pred-1))<br>-td_file.write('def RegI8 : RegisterClass<"PTX", [i8], 8, (sequence "RQ%%u", 0, %d)>;\n' % (num_8bit-1))<br> td_file.write('def RegI16 : RegisterClass<"PTX", [i16], 16, (sequence "RH%%u", 0, %d)>;\n' % (num_16bit-1))<br> td_file.write('def RegI32 : RegisterClass<"PTX", [i32], 32, (sequence "R%%u", 0, %d)>;\n' % (num_32bit-1))<br> td_file.write('def RegI64 : RegisterClass<"PTX", [i64], 64, (sequence "RD%%u", 0, %d)>;\n' % (num_64bit-1))<br>@@ -108,20 +101,16 @@<br> # Reserve 10% of the available registers for return values, and the other 90%<br> # for parameters<br> num_ret_pred = int(0.1 * num_pred)<br>-num_ret_8bit = int(0.1 * num_8bit)<br> num_ret_16bit = int(0.1 * num_16bit)<br> num_ret_32bit = int(0.1 * num_32bit)<br> num_ret_64bit = int(0.1 * num_64bit)<br> num_param_pred = num_pred - num_ret_pred<br>-num_param_8bit = num_8bit - num_ret_8bit<br> num_param_16bit = num_16bit - num_ret_16bit<br> num_param_32bit = num_32bit - num_ret_32bit<br> num_param_64bit = num_64bit - num_ret_64bit<br><br> param_regs_pred = [('P%d' % (i+num_ret_pred)) for i in range(0, num_param_pred)]<br> ret_regs_pred = ['P%d' % i for i in range(0, num_ret_pred)]<br>-param_regs_8bit = [('RQ%d' % (i+num_ret_8bit)) for i in range(0, num_param_8bit)]<br>-ret_regs_8bit = ['RQ%d' % i for i in range(0, num_ret_8bit)]<br> param_regs_16bit = [('RH%d' % (i+num_ret_16bit)) for i in range(0, num_param_16bit)]<br> ret_regs_16bit = ['RH%d' % i for i in range(0, num_ret_16bit)]<br> param_regs_32bit = [('R%d' % (i+num_ret_32bit)) for i in range(0, num_param_32bit)]<br>@@ -131,8 +120,6 @@<br><br> param_list_pred = reduce(lambda x, y: '%s, %s' % (x, y), param_regs_pred)<br> ret_list_pred = reduce(lambda x, y: '%s, %s' % (x, y), ret_regs_pred)<br>-param_list_8bit = reduce(lambda x, y: '%s, %s' % (x, y), param_regs_8bit)<br>-ret_list_8bit = reduce(lambda x, y: '%s, %s' % (x, y), ret_regs_8bit)<br> param_list_16bit = reduce(lambda x, y: '%s, %s' % (x, y), param_regs_16bit)<br> ret_list_16bit = reduce(lambda x, y: '%s, %s' % (x, y), ret_regs_16bit)<br> param_list_32bit = reduce(lambda x, y: '%s, %s' % (x, y), param_regs_32bit)<br>@@ -157,7 +144,6 @@<br> // PTX Formal Parameter Calling Convention<br> def CC_PTX : CallingConv<[<br> CCIfType<[i1], CCAssignToReg<[%s]>>,<br>- CCIfType<[i8], CCAssignToReg<[%s]>>,<br> CCIfType<[i16], CCAssignToReg<[%s]>>,<br> CCIfType<[i32,f32], CCAssignToReg<[%s]>>,<br> CCIfType<[i64,f64], CCAssignToReg<[%s]>><br>@@ -166,13 +152,12 @@<br> // PTX Return Value Calling Convention<br> def RetCC_PTX : CallingConv<[<br> CCIfType<[i1], CCAssignToReg<[%s]>>,<br>- CCIfType<[i8], CCAssignToReg<[%s]>>,<br> CCIfType<[i16], CCAssignToReg<[%s]>>,<br> CCIfType<[i32,f32], CCAssignToReg<[%s]>>,<br> CCIfType<[i64,f64], CCAssignToReg<[%s]>><br> ]>;<br>-''' % (param_list_pred, param_list_8bit, param_list_16bit, param_list_32bit, param_list_64bit,<br>- ret_list_pred, ret_list_8bit, ret_list_16bit, ret_list_32bit, ret_list_64bit))<br>+''' % (param_list_pred, param_list_16bit, param_list_32bit, param_list_64bit,<br>+ ret_list_pred, ret_list_16bit, ret_list_32bit, ret_list_64bit))<br><br><br> td_file.close()<br><br>Modified: llvm/trunk/test/CodeGen/PTX/cvt.ll<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/cvt.ll?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/cvt.ll?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/test/CodeGen/PTX/cvt.ll (original)<br>+++ llvm/trunk/test/CodeGen/PTX/cvt.ll Sat Jun 25 13:16:28 2011<br>@@ -3,17 +3,6 @@<br> ; preds <br> ; (note: we convert back to i32 to return)<br><br>-define ptx_device i32 @cvt_pred_i8(i8 %x, i1 %y) {<br>-; CHECK: setp.gt.b8 p[[P0:[0-9]+]], rq{{[0-9]+}}, 0<br>-; CHECK-NEXT: and.pred p0, p[[P0:[0-9]+]], p{{[0-9]+}};<br>-; CHECK-NEXT: selp.u32 r{{[0-9]+}}, 1, 0, p[[P0:[0-9]+]];<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = trunc i8 %x to i1<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%b = and i1 %a, %y<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%c = zext i1 %b to i32<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i32 %c<br>-}<br>-<br> define ptx_device i32 @cvt_pred_i16(i16 %x, i1 %y) {<br> ; CHECK: setp.gt.b16 p[[P0:[0-9]+]], rh{{[0-9]+}}, 0<br> ; CHECK-NEXT: and.pred p0, p[[P0:[0-9]+]], p{{[0-9]+}};<br>@@ -69,43 +58,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret i32 %c<br> }<br><br>-; i8<br>-<br>-define ptx_device i8 @cvt_i8_preds(i1 %x) {<br>-; CHECK: selp.u8 rq{{[0-9]+}}, 1, 0, p{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = zext i1 %x to i8<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %a<br>-}<br>-<br>-define ptx_device i8 @cvt_i8_i32(i32 %x) {<br>-; CHECK: cvt.u8.u32 rq{{[0-9]+}}, r{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = trunc i32 %x to i8<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %a<br>-}<br>-<br>-define ptx_device i8 @cvt_i8_i64(i64 %x) {<br>-; CHECK: cvt.u8.u64 rq{{[0-9]+}}, rd{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = trunc i64 %x to i8<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %a<br>-}<br>-<br>-define ptx_device i8 @cvt_i8_f32(float %x) {<br>-; CHECK: cvt.rzi.u8.f32 rq{{[0-9]+}}, r{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = fptoui float %x to i8<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %a<br>-}<br>-<br>-define ptx_device i8 @cvt_i8_f64(double %x) {<br>-; CHECK: cvt.rzi.u8.f64 rq{{[0-9]+}}, rd{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = fptoui double %x to i8<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %a<br>-}<br>-<br> ; i16<br><br> define ptx_device i16 @cvt_i16_preds(i1 %x) {<br>@@ -115,13 +67,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret i16 %a<br> }<br><br>-define ptx_device i16 @cvt_i16_i8(i8 %x) {<br>-; CHECK: cvt.u16.u8 rh{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = zext i8 %x to i16<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i16 %a<br>-}<br>-<br> define ptx_device i16 @cvt_i16_i32(i32 %x) {<br> ; CHECK: cvt.u16.u32 rh{{[0-9]+}}, r{{[0-9]+}};<br> ; CHECK-NEXT: ret;<br>@@ -159,13 +104,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret i32 %a<br> }<br><br>-define ptx_device i32 @cvt_i32_i8(i8 %x) {<br>-; CHECK: cvt.u32.u8 r{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = zext i8 %x to i32<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i32 %a<br>-}<br>-<br> define ptx_device i32 @cvt_i32_i16(i16 %x) {<br> ; CHECK: cvt.u32.u16 r{{[0-9]+}}, rh{{[0-9]+}};<br> ; CHECK-NEXT: ret;<br>@@ -203,13 +141,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret i64 %a<br> }<br><br>-define ptx_device i64 @cvt_i64_i8(i8 %x) {<br>-; CHECK: cvt.u64.u8 rd{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = zext i8 %x to i64<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i64 %a<br>-}<br>-<br> define ptx_device i64 @cvt_i64_i16(i16 %x) {<br> ; CHECK: cvt.u64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};<br> ; CHECK-NEXT: ret;<br>@@ -247,13 +178,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret float %a<br> }<br><br>-define ptx_device float @cvt_f32_i8(i8 %x) {<br>-; CHECK: cvt.rn.f32.u8 r{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = uitofp i8 %x to float<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret float %a<br>-}<br>-<br> define ptx_device float @cvt_f32_i16(i16 %x) {<br> ; CHECK: cvt.rn.f32.u16 r{{[0-9]+}}, rh{{[0-9]+}};<br> ; CHECK-NEXT: ret;<br>@@ -291,13 +215,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret double %a<br> }<br><br>-define ptx_device double @cvt_f64_i8(i8 %x) {<br>-; CHECK: cvt.rn.f64.u8 rd{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK-NEXT: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>%a = uitofp i8 %x to double<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret double %a<br>-}<br>-<br> define ptx_device double @cvt_f64_i16(i16 %x) {<br> ; CHECK: cvt.rn.f64.u16 rd{{[0-9]+}}, rh{{[0-9]+}};<br> ; CHECK-NEXT: ret;<br><br>Modified: llvm/trunk/test/CodeGen/PTX/ld.ll<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/ld.ll?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/ld.ll?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/test/CodeGen/PTX/ld.ll (original)<br>+++ llvm/trunk/test/CodeGen/PTX/ld.ll Sat Jun 25 13:16:28 2011<br>@@ -1,17 +1,5 @@<br> ; RUN: llc < %s -march=ptx32 | FileCheck %s<br><br>-;CHECK: .extern .global .b8 array_i8[10];<br>-@array_i8 = external global [10 x i8]<br>-<br>-;CHECK: .extern .const .b8 array_constant_i8[10];<br>-@array_constant_i8 = external addrspace(1) constant [10 x i8]<br>-<br>-;CHECK: .extern .local .b8 array_local_i8[10];<br>-@array_local_i8 = external addrspace(2) global [10 x i8]<br>-<br>-;CHECK: .extern .shared .b8 array_shared_i8[10];<br>-@array_shared_i8 = external addrspace(4) global [10 x i8]<br>-<br> ;CHECK: .extern .global .b8 array_i16[20];<br> @array_i16 = external global [10 x i16]<br><br>@@ -72,13 +60,6 @@<br> ;CHECK: .extern .shared .b8 array_shared_double[80];<br> @array_shared_double = external addrspace(4) global [10 x double]<br><br>-define ptx_device i8 @t1_u8(i8* %p) {<br>-entry:<br>-;CHECK: ld.global.u8 rq{{[0-9]+}}, [r{{[0-9]+}}];<br>-;CHECK-NEXT: ret;<br>- %x = load i8* %p<br>- ret i8 %x<br>-}<br><br> define ptx_device i16 @t1_u16(i16* %p) {<br> entry:<br>@@ -120,15 +101,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t2_u8(i8* %p) {<br>-entry:<br>-;CHECK: ld.global.u8 rq{{[0-9]+}}, [r{{[0-9]+}}+1];<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr i8* %p, i32 1<br>- %x = load i8* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t2_u16(i16* %p) {<br> entry:<br> ;CHECK: ld.global.u16 rh{{[0-9]+}}, [r{{[0-9]+}}+2];<br>@@ -174,15 +146,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t3_u8(i8* %p, i32 %q) {<br>-entry:<br>-;CHECK: add.u32 r[[R0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};<br>-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]];<br>- %i = getelementptr i8* %p, i32 %q<br>- %x = load i8* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t3_u16(i16* %p, i32 %q) {<br> entry:<br> ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;<br>@@ -233,16 +196,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t4_global_u8() {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;<br>-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]];<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 0<br>- %x = load i8* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t4_global_u16() {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;<br>@@ -343,16 +296,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t4_local_u8() {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i8;<br>-;CHECK-NEXT: ld.local.u8 rq{{[0-9]+}}, [r[[R0]]];<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8] addrspace(2)* @array_local_i8, i32 0, i32 0<br>- %x = load i8 addrspace(2)* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t4_local_u16() {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;<br>@@ -403,16 +346,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t4_shared_u8() {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i8;<br>-;CHECK-NEXT: ld.shared.u8 rq{{[0-9]+}}, [r[[R0]]];<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8] addrspace(4)* @array_shared_i8, i32 0, i32 0<br>- %x = load i8 addrspace(4)* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t4_shared_u16() {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;<br>@@ -463,16 +396,6 @@<br> ret double %x<br> }<br><br>-define ptx_device i8 @t5_u8() {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;<br>-;CHECK-NEXT: ld.global.u8 rq{{[0-9]+}}, [r[[R0]]+1];<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 1<br>- %x = load i8* %i<br>- ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t5_u16() {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;<br><br>Modified: llvm/trunk/test/CodeGen/PTX/mov.ll<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mov.ll?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/mov.ll?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/test/CodeGen/PTX/mov.ll (original)<br>+++ llvm/trunk/test/CodeGen/PTX/mov.ll Sat Jun 25 13:16:28 2011<br>@@ -1,11 +1,5 @@<br> ; RUN: llc < %s -march=ptx32 | FileCheck %s<br><br>-define ptx_device i8 @t1_u8() {<br>-; CHECK: mov.u8 rq{{[0-9]+}}, 0;<br>-; CHECK: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 0<br>-}<br>-<br> define ptx_device i16 @t1_u16() {<br> ; CHECK: mov.u16 rh{{[0-9]+}}, 0;<br> ; CHECK: ret;<br>@@ -36,12 +30,6 @@<br> <span class="Apple-tab-span" style="white-space:pre"> </span>ret double 0.0<br> }<br><br>-define ptx_device i8 @t2_u8(i8 %x) {<br>-; CHECK: mov.u8 rq{{[0-9]+}}, rq{{[0-9]+}};<br>-; CHECK: ret;<br>-<span class="Apple-tab-span" style="white-space:pre"> </span>ret i8 %x<br>-}<br>-<br> define ptx_device i16 @t2_u16(i16 %x) {<br> ; CHECK: mov.u16 rh{{[0-9]+}}, rh{{[0-9]+}};<br> ; CHECK: ret;<br><br>Modified: llvm/trunk/test/CodeGen/PTX/st.ll<br>URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/st.ll?rev=133873&r1=133872&r2=133873&view=diff">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/PTX/st.ll?rev=133873&r1=133872&r2=133873&view=diff</a><br>==============================================================================<br>--- llvm/trunk/test/CodeGen/PTX/st.ll (original)<br>+++ llvm/trunk/test/CodeGen/PTX/st.ll Sat Jun 25 13:16:28 2011<br>@@ -1,17 +1,5 @@<br> ; RUN: llc < %s -march=ptx32 | FileCheck %s<br><br>-;CHECK: .extern .global .b8 array_i8[10];<br>-@array_i8 = external global [10 x i8]<br>-<br>-;CHECK: .extern .const .b8 array_constant_i8[10];<br>-@array_constant_i8 = external addrspace(1) constant [10 x i8]<br>-<br>-;CHECK: .extern .local .b8 array_local_i8[10];<br>-@array_local_i8 = external addrspace(2) global [10 x i8]<br>-<br>-;CHECK: .extern .shared .b8 array_shared_i8[10];<br>-@array_shared_i8 = external addrspace(4) global [10 x i8]<br>-<br> ;CHECK: .extern .global .b8 array_i16[20];<br> @array_i16 = external global [10 x i16]<br><br>@@ -72,13 +60,6 @@<br> ;CHECK: .extern .shared .b8 array_shared_double[80];<br> @array_shared_double = external addrspace(4) global [10 x double]<br><br>-define ptx_device void @t1_u8(i8* %p, i8 %x) {<br>-entry:<br>-;CHECK: st.global.u8 [r{{[0-9]+}}], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- store i8 %x, i8* %p<br>- ret void<br>-}<br><br> define ptx_device void @t1_u16(i16* %p, i16 %x) {<br> entry:<br>@@ -120,15 +101,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t2_u8(i8* %p, i8 %x) {<br>-entry:<br>-;CHECK: st.global.u8 [r{{[0-9]+}}+1], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr i8* %p, i32 1<br>- store i8 %x, i8* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t2_u16(i16* %p, i16 %x) {<br> entry:<br> ;CHECK: st.global.u16 [r{{[0-9]+}}+2], rh{{[0-9]+}};<br>@@ -174,16 +146,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t3_u8(i8* %p, i32 %q, i8 %x) {<br>-entry:<br>-;CHECK: add.u32 r[[R0:[0-9]+]], r{{[0-9]+}}, r{{[0-9]+}};<br>-;CHECK-NEXT: st.global.u8 [r[[R0]]], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr i8* %p, i32 %q<br>- store i8 %x, i8* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t3_u16(i16* %p, i32 %q, i16 %x) {<br> entry:<br> ;CHECK: shl.b32 r[[R0:[0-9]+]], r{{[0-9]+}}, 1;<br>@@ -239,16 +201,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t4_global_u8(i8 %x) {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;<br>-;CHECK-NEXT: st.global.u8 [r[[R0]]], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8]* @array_i8, i8 0, i8 0<br>- store i8 %x, i8* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t4_global_u16(i16 %x) {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;<br>@@ -299,16 +251,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t4_local_u8(i8 %x) {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i8;<br>-;CHECK-NEXT: st.local.u8 [r[[R0]]], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8] addrspace(2)* @array_local_i8, i32 0, i32 0<br>- store i8 %x, i8 addrspace(2)* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t4_local_u16(i16 %x) {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_local_i16;<br>@@ -359,16 +301,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t4_shared_u8(i8 %x) {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i8;<br>-;CHECK-NEXT: st.shared.u8 [r[[R0]]], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8] addrspace(4)* @array_shared_i8, i32 0, i32 0<br>- store i8 %x, i8 addrspace(4)* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t4_shared_u16(i16 %x) {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_shared_i16;<br>@@ -419,16 +351,6 @@<br> ret void<br> }<br><br>-define ptx_device void @t5_u8(i8 %x) {<br>-entry:<br>-;CHECK: mov.u32 r[[R0:[0-9]+]], array_i8;<br>-;CHECK-NEXT: st.global.u8 [r[[R0]]+1], rq{{[0-9]+}};<br>-;CHECK-NEXT: ret;<br>- %i = getelementptr [10 x i8]* @array_i8, i32 0, i32 1<br>- store i8 %x, i8* %i<br>- ret void<br>-}<br>-<br> define ptx_device void @t5_u16(i16 %x) {<br> entry:<br> ;CHECK: mov.u32 r[[R0:[0-9]+]], array_i16;<br><br><br>_______________________________________________<br>llvm-commits mailing list<br><a href="mailto:llvm-commits@cs.uiuc.edu">llvm-commits@cs.uiuc.edu</a><br>http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits<br></div></blockquote></div><br></div></body></html>