sgemm_pre_128.sass

# Kernel: sgemm_kernel_128
#
# SharedSize: 16384
# Params(8):
#   0:0x140:4:4 param_C,
#   1:0x144:4:0 param_m,
#   2:0x148:4:0 param_n,
#   3:0x14c:4:0 param_k,
#   4:0x150:4:0 param_lda,
#   5:0x154:4:0 param_ldb,
#   6:0x158:4:0 param_ldc
#   7:0x15c:4:0 param_alpha
#   8:0x160:4:4 param_D // for diagnostic printf output
#
# Globals:
#   c[0x0][0x164]: texA (the value is 1)
#   c[0x0][0x168]: texB (the value is 0)

<REGISTER_MAPPING>

    // Temporary registers to calculate the state registers. Reuse the C output registers.
    // These can be dynamically allocated (~) in the available registger space to elimiate any register bank conflicts.
    0-63    ~ blk, ldx, ldx2, ldx4, k, tid1, tid4, tid7, tid31_4, xmad_t0, xmad_end, bxOrig, byOrig, loy

    // Aliases for the C registers we use for initializing C (used as vectors)
    0-63    : cz<00-63>

    // The offset we store our zero value for initializing C. Reuse a register from the second blocking registers
    80      : zOffset

    // 64 C maxtrix output registers.
    // Use special mapping to avoid register bank conflicts between these registers and the blocking registers.
     3, 2,11,10,19,18,27,26 : cx00y<00-03|64-67>
     7, 6,15,14,23,22,31,30 : cx01y<00-03|64-67>
     1, 0, 9, 8,17,16,25,24 : cx02y<00-03|64-67>
     5, 4,13,12,21,20,29,28 : cx03y<00-03|64-67>
    35,34,43,42,51,50,59,58 : cx64y<00-03|64-67>
    39,38,47,46,55,54,63,62 : cx65y<00-03|64-67>
    33,32,41,40,49,48,57,56 : cx66y<00-03|64-67>
    37,36,45,44,53,52,61,60 : cx67y<00-03|64-67>

    // Double buffered register blocking used in vector loads.
    // Any bank conflicts that we can't avoid in these registers we can hide with .reuse flags
    64-79   : j0Ax<00-03|64-67>, j0By<00-03|64-67>
    80-95   : j1Ax<00-03|64-67>, j1By<00-03|64-67>

    // Registers to load A or B
    96-103  : loadX<0-7>

    // Key global state registers for main loop and some we reuse for outputing C.
    // Note, tweaking the register banks of track<0|4>, tex, writeS, readBs, readAs impacts performance because of
    // delayed bank conflicts between memory operations and ffmas.
    // The array index bracket notation can be used to request a bank in a dynamically allocated range.
    104-127 ~ track<0|4>[0], tex[2], readAs[2], readBs[3], writeS[3], end, ldx8, tid, bx, by, tid31, tid96, tid128 //, clock, smId, nSMs

    // Registers to store the results back to global memory. Reuse any register not needed after the main loop.
    // Statically allocate cs0-7 because they're vector registers.
    64-71   : cs<0-7>

    // dynamically allocated C output registers(~)
    72-103  ~ cy<00|04|08|12>, Cy<00|04|08|12>, ldc, ldc1, ldc4, ldc8, ldc60, writeCs, readCs, cx, ci, alpha, xmad_ci //, xmad_D, D, blckDimX, gridDimX

</REGISTER_MAPPING>

// Note the absense of the loading of the stack pointer into R1.
// No idea why ptxas does that anyway when it's not used for register spilling.
// Such a waste of a perfectly good register.

// Scheduler doesn't handle the dependency flags yet,
// so move these first instructions outside the block that's auto scheduled
//--:-:-:-:1      CS2R clock, SR_CLOCKLO;
//--:-:-:-:1      S2R smId, SR_VIRTID;
//--:-:-:-:1      S2R nSMs, SR_VIRTCFG;
--:-:1:-:1      S2R tid, SR_TID.X;   // Set Dep 1
--:-:2:-:1      S2R bx,  SR_CTAID.X; // Set Dep 2
--:-:3:-:1      S2R by,  SR_CTAID.Y; // Set Dep 3

// Instructions in a SCHEDULE_BLOCK are automatically reordered and appropriately stalled for simple dependancies
// Memory dependencies are left up to the auther to deal with manually for now.
01:-:-:Y:1      ISETP.GE.AND P0, PT, tid, 128, PT; // Wait Dep 1
--:-:-:-:1      LOP.AND tid31, tid, 31;
--:-:-:-:1      BFE.U32 tid4, tid, 0x205; // 2 bits at position 5
--:-:-:-:1      MOV k, c[0x0][0x14c];
--:-:-:-:1      BFE.U32 tid7, tid, 0x301; // 3 bits at position 1
--:-:-:-:1      LOP.AND tid128, tid, 128;
--:-:-:-:1      LOP.AND readBs, tid, 0x70;
--:-:-:-:1      SHL tid31_4, tid31, 4;
--:-:-:-:1      LOP.AND tid1, tid, 1;
--:-:-:-:1      IADD k, k, -8;
--:-:-:-:1      LOP.AND zOffset, tid, -32;
--:-:-:-:1      SHR.U32 readAs, tid128, 4;
--:-:-:-:1      LOP.AND tid96, tid, 96;
--:-:-:-:1      SHR.U32 readBs, readBs, 3;
--:-:-:-:0 @!P0 MOV ldx4, c[0x0][0x150];
--:-:-:-:1      STS.128 [zOffset + 4x<16*128>], RZ;
--:-:-:-:1  @P0 MOV ldx4, c[0x0][0x154];
--:-:-:-:1      ISCADD writeS, tid4, tid31_4, 9;
06:-:-:-:1      SEL blk, by, bx, P0;               // Wait Dep 2 & 3
--:-:-:-:1 @!P0 MOV32I tex, 0x80000001; // texA
--:-:-:-:1  @P0 MOV32I tex, 0x80000000; // texB
--:-:-:-:1      LOP.OR readAs, readAs, tid7;
--:-:-:-:1      SHR.U32 ldx, ldx4, 2;
--:-:-:-:1      LOP.OR readBs, readBs, tid1;
--:-:-:-:1      ISCADD track0, blk, tid31, 5;
--:-:-:-:1      IADD ldx8, ldx4, ldx4;
--:-:-:-:1  @P0 IADD writeS, writeS, 4x<8*128>;
--:-:-:-:1      SHL readAs, readAs, 4;
--:-:-:-:1      XMAD.MRG xmad_t0, ldx, tid4.H1, RZ; // XMAD.LO is a macro that is expanded out into the 3 XMADs
--:-:-:-:1      ISCADD readBs, readBs, 4x<8*128>, 4;
--:-:-:-:1      XMAD track0, ldx, tid4, track0;
--:-:-:Y:5      XMAD.MRG xmad_end, k, ldx.H1, RZ;
--:-:-:-:2      XMAD.PSL.CBCC track0, ldx.H1, xmad_t0.H1, track0;
--:-:1:-:4      TLD.B.LZ.P loadX0, track0, tex, 0x0, 1D, 0xf; // Set Dep 1
--:-:-:-:1      IADD track4, track0, ldx4;
--:-:-:-:1      XMAD end, k, ldx, track0;
--:-:2:Y:5      TLD.B.LZ.P loadX4, track4, tex, 0x0, 1D, 0xf; // Set Dep 2
--:-:-:-:1      XMAD.PSL.CBCC end, k.H1, xmad_end.H1, end;

// Initialize C registeres to zero
// Using LDS.U.128 is a neat trick to save a few clock cyles
// (when you have enough warps to hide the latency.)
--:-:3:-:1      LDS.U.128 cz00, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz04, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz08, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz12, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz16, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz20, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz24, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz28, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz32, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz36, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz40, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz44, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz48, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz52, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz56, [zOffset + 4x<16*128>];
--:-:3:-:1      LDS.U.128 cz60, [zOffset + 4x<16*128>];

// These instuctions need to occur after the textures load so put them in a new block
// that starts with a dependency barrier wait.
01:-:-:-:1      STS.128 [writeS + 4x<0*128>], loadX0; // Wait Dep 1
--:-:-:-:0      IADD track0, track0, ldx8;
02:-:-:-:1      STS.128 [writeS + 4x<4*128>], loadX4; // Wait Dep 2
--:-:-:-:0      IADD track4, track4, ldx8;
04:-:-:-:5      BAR.SYNC 0;

// The next store to shared goes to high area.
// Having 2 share buffers allows us to eliminate a bar.sync in the main loop.
// This way we don't have to wait for all threads to arrive before writing fresh data to shared.
// Other threads can continue reading from the last batch while the new data is being written.
--:-:-:-:0      LOP.XOR writeS, writeS, 4x<16*128>;

// Preload the fist lines of A and B from shared
--:-:-:-:1      LDS.U.128 j0Ax00, [readAs + 4x<0*128 + 00>];
--:-:-:-:1      LDS.U.128 j0By00, [readBs + 4x<0*128 + 00>];
--:-:-:-:1      LDS.U.128 j0Ax64, [readAs + 4x<0*128 + 64>];
--:-:1:-:1      LDS.U.128 j0By64, [readBs + 4x<0*128 + 64>]; // Set Dep 1


// The main loop
// While calculating the first line, load in the next line from shared.
// Shared memory stores enough to do this 8 times per loop.
// Also pull in the next block of memory from global and store it to shared.

// Efficiency:
// ffma: 512
// lds:  32 dual issued
// sts:  2  dual issued
// tex:  2  dual issued
// add:  2
// xor:  3
// setp: 1
// bar:  1  dual issued
// bra:  1  dual issued
// Total: 524 (512/518 = 98.8% FFMA)

// Memory Throughput Upper Bound:
// 2 * 4 * 4 bytes per thread per 518 clocks
// 128 threads per SM
// 16 SM's (GM204)
// 1640Mhz (boost overclock)
// .931 GiB/GB  (1000^3 / 1024^3)
// 193 GiB/sec
// Available: 224 GiB/sec (or 256 GiB/sec overclocked at 8GHz)

LOOP:

// Loop end condition
--:-:-:-:1      ISETP.LE.AND P0, PT, track0, end, PT;

01:-:-:-:0      FFMA cx02y00, j0Ax02, j0By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j1Ax00, [readAs + 4x<1*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j0Ax02, j0By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j0Ax00, j0By01, cx00y01;
--:-:-:-:1      LDS.U.128 j1By00, [readBs + 4x<1*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j0Ax00, j0By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j0Ax03, j0By00, cx03y00;
--:-:-:-:1      LDS.U.128 j1Ax64, [readAs + 4x<1*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j0Ax03, j0By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j0Ax01, j0By01, cx01y01;
--:-:1:-:1      LDS.U.128 j1By64, [readBs + 4x<1*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j0Ax01, j0By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j0Ax66, j0By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j0Ax66, j0By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j0Ax64, j0By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j0Ax64, j0By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j0Ax67, j0By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j0Ax67, j0By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j0Ax65, j0By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j0Ax65, j0By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j0Ax67, j0By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j0Ax67, j0By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j0Ax65, j0By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j0Ax65, j0By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j0Ax66, j0By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j0Ax66, j0By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j0Ax64, j0By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j0Ax64, j0By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j0Ax03, j0By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j0Ax03, j0By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j0Ax01, j0By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j0Ax01, j0By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j0Ax02, j0By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j0Ax02, j0By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j0Ax00, j0By03, cx00y03;
--:-:-:-:0      FFMA cx00y02, j0Ax00, j0By02, cx00y02;
--:-:2:-:1  @P0 TLD.B.LZ.P loadX0, track0, tex, 0x0, 1D, 0xf; // Set Dep 2
--:-:-:Y:1      FFMA cx02y64, j0Ax02, j0By64, cx02y64;
--:-:-:-:0      FFMA cx02y65, j0Ax02, j0By65, cx02y65;
--:-:3:-:1  @P0 TLD.B.LZ.P loadX4, track4, tex, 0x0, 1D, 0xf; // Set Dep 3
--:-:-:-:1      FFMA cx00y65, j0Ax00, j0By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j0Ax00, j0By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j0Ax03, j0By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j0Ax03, j0By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j0Ax01, j0By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j0Ax01, j0By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j0Ax66, j0By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j0Ax66, j0By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j0Ax64, j0By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j0Ax64, j0By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j0Ax67, j0By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j0Ax67, j0By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j0Ax65, j0By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j0Ax65, j0By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j0Ax67, j0By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j0Ax67, j0By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j0Ax65, j0By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j0Ax65, j0By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j0Ax66, j0By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j0Ax66, j0By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j0Ax64, j0By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j0Ax64, j0By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j0Ax03, j0By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j0Ax03, j0By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j0Ax01, j0By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j0Ax01, j0By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j0Ax02, j0By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j0Ax02, j0By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j0Ax00, j0By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j0Ax00, j0By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j1Ax02, j1By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j0Ax00, [readAs + 4x<2*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j1Ax02, j1By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j1Ax00, j1By01, cx00y01;
--:-:-:-:1      LDS.U.128 j0By00, [readBs + 4x<2*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j1Ax00, j1By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j1Ax03, j1By00, cx03y00;
--:-:-:-:1      LDS.U.128 j0Ax64, [readAs + 4x<2*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j1Ax03, j1By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j1Ax01, j1By01, cx01y01;
--:-:1:-:1      LDS.U.128 j0By64, [readBs + 4x<2*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j1Ax01, j1By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j1Ax66, j1By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j1Ax66, j1By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j1Ax64, j1By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j1Ax64, j1By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j1Ax67, j1By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j1Ax67, j1By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j1Ax65, j1By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j1Ax65, j1By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j1Ax67, j1By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j1Ax67, j1By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j1Ax65, j1By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j1Ax65, j1By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j1Ax66, j1By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j1Ax66, j1By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j1Ax64, j1By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j1Ax64, j1By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j1Ax03, j1By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j1Ax03, j1By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j1Ax01, j1By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j1Ax01, j1By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j1Ax02, j1By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j1Ax02, j1By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j1Ax00, j1By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j1Ax00, j1By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j1Ax02, j1By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j1Ax02, j1By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j1Ax00, j1By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j1Ax00, j1By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j1Ax03, j1By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j1Ax03, j1By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j1Ax01, j1By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j1Ax01, j1By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j1Ax66, j1By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j1Ax66, j1By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j1Ax64, j1By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j1Ax64, j1By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j1Ax67, j1By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j1Ax67, j1By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j1Ax65, j1By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j1Ax65, j1By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j1Ax67, j1By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j1Ax67, j1By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j1Ax65, j1By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j1Ax65, j1By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j1Ax66, j1By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j1Ax66, j1By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j1Ax64, j1By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j1Ax64, j1By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j1Ax03, j1By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j1Ax03, j1By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j1Ax01, j1By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j1Ax01, j1By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j1Ax02, j1By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j1Ax02, j1By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j1Ax00, j1By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j1Ax00, j1By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j0Ax02, j0By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j1Ax00, [readAs + 4x<3*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j0Ax02, j0By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j0Ax00, j0By01, cx00y01;
--:-:-:-:1      LDS.U.128 j1By00, [readBs + 4x<3*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j0Ax00, j0By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j0Ax03, j0By00, cx03y00;
--:-:-:-:1      LDS.U.128 j1Ax64, [readAs + 4x<3*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j0Ax03, j0By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j0Ax01, j0By01, cx01y01;
--:-:1:-:1      LDS.U.128 j1By64, [readBs + 4x<3*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j0Ax01, j0By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j0Ax66, j0By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j0Ax66, j0By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j0Ax64, j0By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j0Ax64, j0By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j0Ax67, j0By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j0Ax67, j0By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j0Ax65, j0By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j0Ax65, j0By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j0Ax67, j0By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j0Ax67, j0By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j0Ax65, j0By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j0Ax65, j0By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j0Ax66, j0By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j0Ax66, j0By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j0Ax64, j0By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j0Ax64, j0By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j0Ax03, j0By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j0Ax03, j0By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j0Ax01, j0By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j0Ax01, j0By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j0Ax02, j0By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j0Ax02, j0By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j0Ax00, j0By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j0Ax00, j0By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j0Ax02, j0By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j0Ax02, j0By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j0Ax00, j0By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j0Ax00, j0By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j0Ax03, j0By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j0Ax03, j0By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j0Ax01, j0By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j0Ax01, j0By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j0Ax66, j0By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j0Ax66, j0By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j0Ax64, j0By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j0Ax64, j0By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j0Ax67, j0By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j0Ax67, j0By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j0Ax65, j0By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j0Ax65, j0By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j0Ax67, j0By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j0Ax67, j0By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j0Ax65, j0By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j0Ax65, j0By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j0Ax66, j0By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j0Ax66, j0By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j0Ax64, j0By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j0Ax64, j0By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j0Ax03, j0By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j0Ax03, j0By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j0Ax01, j0By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j0Ax01, j0By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j0Ax02, j0By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j0Ax02, j0By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j0Ax00, j0By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j0Ax00, j0By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j1Ax02, j1By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j0Ax00, [readAs + 4x<4*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j1Ax02, j1By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j1Ax00, j1By01, cx00y01;
--:-:-:-:1      LDS.U.128 j0By00, [readBs + 4x<4*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j1Ax00, j1By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j1Ax03, j1By00, cx03y00;
--:-:-:-:1      LDS.U.128 j0Ax64, [readAs + 4x<4*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j1Ax03, j1By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j1Ax01, j1By01, cx01y01;
--:-:1:-:1      LDS.U.128 j0By64, [readBs + 4x<4*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j1Ax01, j1By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j1Ax66, j1By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j1Ax66, j1By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j1Ax64, j1By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j1Ax64, j1By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j1Ax67, j1By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j1Ax67, j1By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j1Ax65, j1By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j1Ax65, j1By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j1Ax67, j1By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j1Ax67, j1By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j1Ax65, j1By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j1Ax65, j1By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j1Ax66, j1By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j1Ax66, j1By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j1Ax64, j1By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j1Ax64, j1By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j1Ax03, j1By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j1Ax03, j1By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j1Ax01, j1By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j1Ax01, j1By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j1Ax02, j1By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j1Ax02, j1By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j1Ax00, j1By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j1Ax00, j1By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j1Ax02, j1By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j1Ax02, j1By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j1Ax00, j1By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j1Ax00, j1By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j1Ax03, j1By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j1Ax03, j1By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j1Ax01, j1By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j1Ax01, j1By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j1Ax66, j1By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j1Ax66, j1By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j1Ax64, j1By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j1Ax64, j1By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j1Ax67, j1By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j1Ax67, j1By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j1Ax65, j1By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j1Ax65, j1By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j1Ax67, j1By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j1Ax67, j1By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j1Ax65, j1By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j1Ax65, j1By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j1Ax66, j1By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j1Ax66, j1By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j1Ax64, j1By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j1Ax64, j1By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j1Ax03, j1By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j1Ax03, j1By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j1Ax01, j1By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j1Ax01, j1By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j1Ax02, j1By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j1Ax02, j1By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j1Ax00, j1By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j1Ax00, j1By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j0Ax02, j0By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j1Ax00, [readAs + 4x<5*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j0Ax02, j0By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j0Ax00, j0By01, cx00y01;
--:-:-:-:1      LDS.U.128 j1By00, [readBs + 4x<5*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j0Ax00, j0By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j0Ax03, j0By00, cx03y00;
--:-:-:-:1      LDS.U.128 j1Ax64, [readAs + 4x<5*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j0Ax03, j0By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j0Ax01, j0By01, cx01y01;
--:-:1:-:1      LDS.U.128 j1By64, [readBs + 4x<5*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j0Ax01, j0By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j0Ax66, j0By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j0Ax66, j0By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j0Ax64, j0By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j0Ax64, j0By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j0Ax67, j0By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j0Ax67, j0By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j0Ax65, j0By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j0Ax65, j0By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j0Ax67, j0By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j0Ax67, j0By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j0Ax65, j0By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j0Ax65, j0By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j0Ax66, j0By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j0Ax66, j0By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j0Ax64, j0By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j0Ax64, j0By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j0Ax03, j0By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j0Ax03, j0By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j0Ax01, j0By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j0Ax01, j0By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j0Ax02, j0By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j0Ax02, j0By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j0Ax00, j0By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j0Ax00, j0By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j0Ax02, j0By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j0Ax02, j0By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j0Ax00, j0By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j0Ax00, j0By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j0Ax03, j0By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j0Ax03, j0By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j0Ax01, j0By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j0Ax01, j0By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j0Ax66, j0By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j0Ax66, j0By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j0Ax64, j0By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j0Ax64, j0By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j0Ax67, j0By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j0Ax67, j0By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j0Ax65, j0By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j0Ax65, j0By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j0Ax67, j0By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j0Ax67, j0By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j0Ax65, j0By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j0Ax65, j0By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j0Ax66, j0By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j0Ax66, j0By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j0Ax64, j0By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j0Ax64, j0By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j0Ax03, j0By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j0Ax03, j0By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j0Ax01, j0By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j0Ax01, j0By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j0Ax02, j0By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j0Ax02, j0By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j0Ax00, j0By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j0Ax00, j0By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j1Ax02, j1By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j0Ax00, [readAs + 4x<6*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j1Ax02, j1By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j1Ax00, j1By01, cx00y01;
--:-:-:-:1      LDS.U.128 j0By00, [readBs + 4x<6*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j1Ax00, j1By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j1Ax03, j1By00, cx03y00;
--:-:-:-:1      LDS.U.128 j0Ax64, [readAs + 4x<6*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j1Ax03, j1By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j1Ax01, j1By01, cx01y01;
--:-:1:-:1      LDS.U.128 j0By64, [readBs + 4x<6*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j1Ax01, j1By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j1Ax66, j1By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j1Ax66, j1By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j1Ax64, j1By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j1Ax64, j1By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j1Ax67, j1By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j1Ax67, j1By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j1Ax65, j1By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j1Ax65, j1By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j1Ax67, j1By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j1Ax67, j1By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j1Ax65, j1By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j1Ax65, j1By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j1Ax66, j1By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j1Ax66, j1By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j1Ax64, j1By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j1Ax64, j1By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j1Ax03, j1By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j1Ax03, j1By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j1Ax01, j1By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j1Ax01, j1By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j1Ax02, j1By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j1Ax02, j1By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j1Ax00, j1By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j1Ax00, j1By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j1Ax02, j1By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j1Ax02, j1By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j1Ax00, j1By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j1Ax00, j1By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j1Ax03, j1By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j1Ax03, j1By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j1Ax01, j1By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j1Ax01, j1By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j1Ax66, j1By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j1Ax66, j1By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j1Ax64, j1By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j1Ax64, j1By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j1Ax67, j1By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j1Ax67, j1By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j1Ax65, j1By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j1Ax65, j1By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j1Ax67, j1By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j1Ax67, j1By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j1Ax65, j1By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j1Ax65, j1By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j1Ax66, j1By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j1Ax66, j1By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j1Ax64, j1By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j1Ax64, j1By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j1Ax03, j1By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j1Ax03, j1By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j1Ax01, j1By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j1Ax01, j1By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j1Ax02, j1By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j1Ax02, j1By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j1Ax00, j1By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j1Ax00, j1By66, cx00y66;
01:-:-:-:0      FFMA cx02y00, j0Ax02, j0By00, cx02y00; // Wait Dep 1
--:-:-:-:1      LDS.U.128 j1Ax00, [readAs + 4x<7*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j0Ax02, j0By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j0Ax00, j0By01, cx00y01;
--:-:-:-:1      LDS.U.128 j1By00, [readBs + 4x<7*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j0Ax00, j0By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j0Ax03, j0By00, cx03y00;
--:-:-:-:1      LDS.U.128 j1Ax64, [readAs + 4x<7*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j0Ax03, j0By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j0Ax01, j0By01, cx01y01;
--:-:1:-:1      LDS.U.128 j1By64, [readBs + 4x<7*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j0Ax01, j0By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j0Ax66, j0By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j0Ax66, j0By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j0Ax64, j0By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j0Ax64, j0By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j0Ax67, j0By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j0Ax67, j0By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j0Ax65, j0By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j0Ax65, j0By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j0Ax67, j0By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j0Ax67, j0By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j0Ax65, j0By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j0Ax65, j0By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j0Ax66, j0By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j0Ax66, j0By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j0Ax64, j0By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j0Ax64, j0By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j0Ax03, j0By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j0Ax03, j0By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j0Ax01, j0By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j0Ax01, j0By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j0Ax02, j0By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j0Ax02, j0By03, cx02y03;
--:-:-:-:0      FFMA cx00y03, j0Ax00, j0By03, cx00y03;
02:-:-:-:1  @P0 STS.128 [writeS + 4x<0*128>], loadX0; // Wait Dep 2
--:-:-:-:1      FFMA cx00y02, j0Ax00, j0By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j0Ax02, j0By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j0Ax02, j0By65, cx02y65;
--:-:-:-:0      FFMA cx00y65, j0Ax00, j0By65, cx00y65;
04:-:-:-:1  @P0 STS.128 [writeS + 4x<4*128>], loadX4; // Wait Dep 3
--:-:-:-:1      FFMA cx00y64, j0Ax00, j0By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j0Ax03, j0By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j0Ax03, j0By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j0Ax01, j0By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j0Ax01, j0By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j0Ax66, j0By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j0Ax66, j0By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j0Ax64, j0By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j0Ax64, j0By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j0Ax67, j0By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j0Ax67, j0By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j0Ax65, j0By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j0Ax65, j0By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j0Ax67, j0By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j0Ax67, j0By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j0Ax65, j0By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j0Ax65, j0By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j0Ax66, j0By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j0Ax66, j0By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j0Ax64, j0By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j0Ax64, j0By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j0Ax03, j0By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j0Ax03, j0By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j0Ax01, j0By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j0Ax01, j0By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j0Ax02, j0By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j0Ax02, j0By67, cx02y67;
--:-:-:-:0      FFMA cx00y67, j0Ax00, j0By67, cx00y67;
01:-:-:-:5      BAR.SYNC 0;                            // Wait Dep 1
--:-:-:-:1  @P0 LOP.XOR readAs, readAs, 4x<16*128>;
--:-:-:-:1  @P0 LOP.XOR readBs, readBs, 4x<16*128>;
--:-:-:-:1  @P0 LOP.XOR writeS, writeS, 4x<16*128>;
--:-:-:-:1      FFMA cx00y66, j0Ax00, j0By66, cx00y66;
--:-:-:-:0      FFMA cx02y00, j1Ax02, j1By00, cx02y00;
--:-:-:-:1  @P0 LDS.U.128 j0Ax00, [readAs + 4x<0*128 + 00>];
--:-:-:-:1      FFMA cx02y01, j1Ax02, j1By01, cx02y01;
--:-:-:-:0      FFMA cx00y01, j1Ax00, j1By01, cx00y01;
--:-:-:-:1  @P0 LDS.U.128 j0By00, [readBs + 4x<0*128 + 00>];
--:-:-:-:1      FFMA cx00y00, j1Ax00, j1By00, cx00y00;
--:-:-:-:0      FFMA cx03y00, j1Ax03, j1By00, cx03y00;
--:-:-:-:1  @P0 LDS.U.128 j0Ax64, [readAs + 4x<0*128 + 64>];
--:-:-:-:1      FFMA cx03y01, j1Ax03, j1By01, cx03y01;
--:-:-:-:0      FFMA cx01y01, j1Ax01, j1By01, cx01y01;
--:-:1:-:1  @P0 LDS.U.128 j0By64, [readBs + 4x<0*128 + 64>]; // Set Dep 1
--:-:-:-:1      FFMA cx01y00, j1Ax01, j1By00, cx01y00;
--:-:-:-:1      FFMA cx66y00, j1Ax66, j1By00, cx66y00;
--:-:-:-:1      FFMA cx66y01, j1Ax66, j1By01, cx66y01;
--:-:-:-:1      FFMA cx64y01, j1Ax64, j1By01, cx64y01;
--:-:-:-:1      FFMA cx64y00, j1Ax64, j1By00, cx64y00;
--:-:-:-:1      FFMA cx67y00, j1Ax67, j1By00, cx67y00;
--:-:-:-:1      FFMA cx67y01, j1Ax67, j1By01, cx67y01;
--:-:-:-:1      FFMA cx65y01, j1Ax65, j1By01, cx65y01;
--:-:-:-:1      FFMA cx65y00, j1Ax65, j1By00, cx65y00;
--:-:-:-:1      FFMA cx67y02, j1Ax67, j1By02, cx67y02;
--:-:-:-:1      FFMA cx67y03, j1Ax67, j1By03, cx67y03;
--:-:-:-:1      FFMA cx65y03, j1Ax65, j1By03, cx65y03;
--:-:-:-:1      FFMA cx65y02, j1Ax65, j1By02, cx65y02;
--:-:-:-:1      FFMA cx66y02, j1Ax66, j1By02, cx66y02;
--:-:-:-:1      FFMA cx66y03, j1Ax66, j1By03, cx66y03;
--:-:-:-:1      FFMA cx64y03, j1Ax64, j1By03, cx64y03;
--:-:-:-:1      FFMA cx64y02, j1Ax64, j1By02, cx64y02;
--:-:-:-:1      FFMA cx03y02, j1Ax03, j1By02, cx03y02;
--:-:-:-:1      FFMA cx03y03, j1Ax03, j1By03, cx03y03;
--:-:-:-:1      FFMA cx01y03, j1Ax01, j1By03, cx01y03;
--:-:-:-:1      FFMA cx01y02, j1Ax01, j1By02, cx01y02;
--:-:-:-:1      FFMA cx02y02, j1Ax02, j1By02, cx02y02;
--:-:-:-:1      FFMA cx02y03, j1Ax02, j1By03, cx02y03;
--:-:-:-:1      FFMA cx00y03, j1Ax00, j1By03, cx00y03;
--:-:-:-:1      FFMA cx00y02, j1Ax00, j1By02, cx00y02;
--:-:-:Y:1      FFMA cx02y64, j1Ax02, j1By64, cx02y64;
--:-:-:-:1      FFMA cx02y65, j1Ax02, j1By65, cx02y65;
--:-:-:-:1      FFMA cx00y65, j1Ax00, j1By65, cx00y65;
--:-:-:-:1      FFMA cx00y64, j1Ax00, j1By64, cx00y64;
--:-:-:-:1      FFMA cx03y64, j1Ax03, j1By64, cx03y64;
--:-:-:-:1      FFMA cx03y65, j1Ax03, j1By65, cx03y65;
--:-:-:-:1      FFMA cx01y65, j1Ax01, j1By65, cx01y65;
--:-:-:-:1      FFMA cx01y64, j1Ax01, j1By64, cx01y64;
--:-:-:-:1      FFMA cx66y64, j1Ax66, j1By64, cx66y64;
--:-:-:-:1      FFMA cx66y65, j1Ax66, j1By65, cx66y65;
--:-:-:-:1      FFMA cx64y65, j1Ax64, j1By65, cx64y65;
--:-:-:-:1      FFMA cx64y64, j1Ax64, j1By64, cx64y64;
--:-:-:-:1      FFMA cx67y64, j1Ax67, j1By64, cx67y64;
--:-:-:-:1      FFMA cx67y65, j1Ax67, j1By65, cx67y65;
--:-:-:-:1      FFMA cx65y65, j1Ax65, j1By65, cx65y65;
--:-:-:-:1      FFMA cx65y64, j1Ax65, j1By64, cx65y64;
--:-:-:-:1      FFMA cx67y66, j1Ax67, j1By66, cx67y66;
--:-:-:-:1      FFMA cx67y67, j1Ax67, j1By67, cx67y67;
--:-:-:-:1      FFMA cx65y67, j1Ax65, j1By67, cx65y67;
--:-:-:-:1      FFMA cx65y66, j1Ax65, j1By66, cx65y66;
--:-:-:-:1      FFMA cx66y66, j1Ax66, j1By66, cx66y66;
--:-:-:-:1      FFMA cx66y67, j1Ax66, j1By67, cx66y67;
--:-:-:-:1      FFMA cx64y67, j1Ax64, j1By67, cx64y67;
--:-:-:-:1      FFMA cx64y66, j1Ax64, j1By66, cx64y66;
--:-:-:-:1      FFMA cx03y66, j1Ax03, j1By66, cx03y66;
--:-:-:-:1      FFMA cx03y67, j1Ax03, j1By67, cx03y67;
--:-:-:-:1      FFMA cx01y67, j1Ax01, j1By67, cx01y67;
--:-:-:-:1      FFMA cx01y66, j1Ax01, j1By66, cx01y66;
--:-:-:-:1      FFMA cx02y66, j1Ax02, j1By66, cx02y66;
--:-:-:-:1      FFMA cx02y67, j1Ax02, j1By67, cx02y67;
--:-:-:-:1      FFMA cx00y67, j1Ax00, j1By67, cx00y67;
--:-:-:-:1      FFMA cx00y66, j1Ax00, j1By66, cx00y66;
--:-:-:-:1  @P0 IADD track0, track0, ldx8;
--:-:-:-:0  @P0 IADD track4, track4, ldx8;
--:-:-:Y:5  @P0 BRA LOOP;

// Main loop is done, time to write C to global memory.
--:-:-:-:1      SHR.U32 cx, tid128, 2;
--:-:-:-:1      MOV ldc, c[0x0][0x158];
--:-:-:-:1      SHR.U32 cy00, tid96, 1;
--:-:-:-:1      MOV alpha, c[0x0][0x15c];
--:-:-:-:1      SHL readCs, tid96, 4;
--:-:-:-:1      LOP.AND readAs, readAs, 0xfff;
--:-:-:-:1      LOP.OR cx, tid31, cx;
--:-:-:-:1      SHL ldc1, ldc, 2;
--:-:-:-:1      LOP.AND readBs, readBs, 0xfff;
--:-:-:-:1      ISCADD cy00, by, cy00, 7;
--:-:-:-:1      FMUL cs0, cx00y00, alpha;
--:-:-:-:1      SHL ldc4, ldc, 4;
--:-:-:-:1      LOP.OR readCs, readCs, cx;
--:-:-:-:1      ISCADD cx, bx, cx, 7;
--:-:-:-:1      FMUL cs1, cx01y00, alpha;
--:-:-:-:1      SHL ldc8, ldc, 5;
--:-:-:-:1      XMAD.MRG xmad_ci, cy00, ldc.H1, RZ;
--:-:-:-:1      ISCADD writeCs, readBs, readAs, 5;
--:-:-:-:1      FMUL cs2, cx02y00, alpha;
--:-:-:-:1      SHL readCs, readCs, 2;
--:-:-:-:1      XMAD ci, cy00, ldc, cx;
--:-:-:-:1      ISETP.LT.AND P5, PT, cx, c[0x0][0x144], PT; // cx +  0 < m
--:-:-:-:1      IADD cx, cx, 64;
--:-:-:-:1      ISCADD ldc60, ldc, -ldc4, 8;
--:-:-:-:1      FMUL cs3, cx03y00, alpha;
--:-:-:-:1      FMUL cs4, cx64y00, alpha;
--:-:-:-:1      XMAD.PSL.CBCC ci, cy00.H1, xmad_ci.H1, ci;
--:-:-:-:1      IADD cy00, cy00, -1;
--:-:-:-:1      ISETP.LT.AND P6, PT, cx, c[0x0][0x144], PT; // cx + 64 < m
--:-:-:-:1      FMUL cs5, cx65y00, alpha;
--:-:-:-:1      FMUL cs6, cx66y00, alpha;
--:-:-:-:1      FMUL cs7, cx67y00, alpha;
--:-:-:-:1      ISCADD Cy00, ci, c[0x0][0x140], 2;
--:-:-:-:1      IADD cy04, cy00, 4;
--:-:-:-:1      IADD cy08, cy00, 8;
--:-:-:-:3      IADD cy12, cy00, 12;
--:-:-:Y:6      IADD Cy00, Cy00, -ldc1;
--:-:-:-:1      IADD Cy04, Cy00, ldc4;
--:-:-:Y:5      IADD Cy08, Cy00, ldc8;
--:-:-:-:0      IADD Cy12, Cy04, ldc8; // Dual Issue (last instruction after reordering)

// There's nothing yet in place to handle dependecies with subroutines.
// So don't schedule this block.
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y01, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y01, alpha;
--:-:-:-:1      FMUL cs2, cx02y01, alpha;
--:-:-:-:1      FMUL cs3, cx03y01, alpha;
--:-:-:-:1      FMUL cs4, cx64y01, alpha;
--:-:-:-:1      FMUL cs5, cx65y01, alpha;
--:-:-:-:1      FMUL cs6, cx66y01, alpha;
--:-:-:-:0      FMUL cs7, cx67y01, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y02, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y02, alpha;
--:-:-:-:1      FMUL cs2, cx02y02, alpha;
--:-:-:-:1      FMUL cs3, cx03y02, alpha;
--:-:-:-:1      FMUL cs4, cx64y02, alpha;
--:-:-:-:1      FMUL cs5, cx65y02, alpha;
--:-:-:-:1      FMUL cs6, cx66y02, alpha;
--:-:-:-:0      FMUL cs7, cx67y02, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y03, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y03, alpha;
--:-:-:-:1      FMUL cs2, cx02y03, alpha;
--:-:-:-:1      FMUL cs3, cx03y03, alpha;
--:-:-:-:1      FMUL cs4, cx64y03, alpha;
--:-:-:-:1      FMUL cs5, cx65y03, alpha;
--:-:-:-:1      FMUL cs6, cx66y03, alpha;
--:-:-:-:0      FMUL cs7, cx67y03, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

--:-:-:-:1      IADD cy00, cy00, 60;
--:-:-:-:1      IADD cy04, cy04, 60;
--:-:-:-:1      IADD cy08, cy08, 60;
--:-:-:-:1      IADD cy12, cy12, 60;

02:-:-:-:1      IADD Cy00, Cy00, ldc60; // Wait Dep 2
--:-:-:-:1      IADD Cy04, Cy04, ldc60;
--:-:-:-:1      IADD Cy08, Cy08, ldc60;
--:-:-:-:1      IADD Cy12, Cy12, ldc60;

--:-:-:-:1      FMUL cs0, cx00y64, alpha;
--:-:-:-:1      FMUL cs1, cx01y64, alpha;
--:-:-:-:1      FMUL cs2, cx02y64, alpha;
--:-:-:-:1      FMUL cs3, cx03y64, alpha;
--:-:-:-:1      FMUL cs4, cx64y64, alpha;
--:-:-:-:1      FMUL cs5, cx65y64, alpha;
--:-:-:-:1      FMUL cs6, cx66y64, alpha;
--:-:-:-:0      FMUL cs7, cx67y64, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y65, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y65, alpha;
--:-:-:-:1      FMUL cs2, cx02y65, alpha;
--:-:-:-:1      FMUL cs3, cx03y65, alpha;
--:-:-:-:1      FMUL cs4, cx64y65, alpha;
--:-:-:-:1      FMUL cs5, cx65y65, alpha;
--:-:-:-:1      FMUL cs6, cx66y65, alpha;
--:-:-:-:0      FMUL cs7, cx67y65, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y66, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y66, alpha;
--:-:-:-:1      FMUL cs2, cx02y66, alpha;
--:-:-:-:1      FMUL cs3, cx03y66, alpha;
--:-:-:-:1      FMUL cs4, cx64y66, alpha;
--:-:-:-:1      FMUL cs5, cx65y66, alpha;
--:-:-:-:1      FMUL cs6, cx66y66, alpha;
--:-:-:-:0      FMUL cs7, cx67y66, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;

02:-:-:-:1      FMUL cs0, cx00y67, alpha; // Wait Dep 2
--:-:-:-:1      FMUL cs1, cx01y67, alpha;
--:-:-:-:1      FMUL cs2, cx02y67, alpha;
--:-:-:-:1      FMUL cs3, cx03y67, alpha;
--:-:-:-:1      FMUL cs4, cx64y67, alpha;
--:-:-:-:1      FMUL cs5, cx65y67, alpha;
--:-:-:-:1      FMUL cs6, cx66y67, alpha;
--:-:-:-:0      FMUL cs7, cx67y67, alpha; // Dual Issue
--:-:-:-:5      CAL STORE_C;


// And we'd done.  The remainder is the STORE_C subroutine that's defined at the end of the kernel.
--:-:-:-:5      EXIT;

// This routine does warp synchronous shuffling of our output data so as to be able
// to have coalesced writes to global memory.  This is actually faster because the shared
// memory latencies can be hidden by other warps and we're only adding a few extra clocks
// to this thread.  Global memory here is the bottleneck and being able to half the needed
// bandwidth at the expense of a few clocks is a modest win.  This also keeps power lower
// and our chip running faster.

// Note, the SHFL instruction doesn't help us here because we're swaping different registers
// from different threads.
STORE_C:

--:-:-:-:0      IADD cy00, cy00, 1;
--:-:-:-:1      STS.128 [writeCs+4x<00>], cs0;
--:-:-:-:0      IADD cy04, cy04, 1;
--:-:-:-:1      STS.128 [writeCs+4x<64>], cs4;
--:-:-:-:0      IADD cy08, cy08, 1;
--:-:-:-:1      LDS cs0, [readCs + 4x<0*128 + 00>];
--:-:-:-:0      IADD cy12, cy12, 1;
--:-:-:-:1      LDS cs1, [readCs + 4x<0*128 + 64>];
--:-:-:-:0      IADD Cy00, Cy00, ldc1;
--:-:-:-:1      LDS cs2, [readCs + 4x<1*128 + 00>];
--:-:-:-:0      IADD Cy04, Cy04, ldc1;
--:-:-:-:1      LDS cs3, [readCs + 4x<1*128 + 64>];
--:-:-:-:0      IADD Cy08, Cy08, ldc1;
--:-:-:-:1      LDS cs4, [readCs + 4x<2*128 + 00>];
--:-:-:-:0      IADD Cy12, Cy12, ldc1;
--:-:-:-:1      LDS cs5, [readCs + 4x<2*128 + 64>];
--:-:-:-:1      ISETP.LT.AND P0, PT, cy00, c[0x0][0x148], P5; // cy00 < n && cx +  0 < m
--:-:-:-:1      LDS cs6, [readCs + 4x<3*128 + 00>];
--:-:-:-:1      ISETP.LT.AND P1, PT, cy00, c[0x0][0x148], P6; // cy00 < n && cx + 64 < m
--:-:1:-:1      LDS cs7, [readCs + 4x<3*128 + 64>]; // Set Dep 1
--:-:-:-:2      ISETP.LT.AND P2, PT, cy04, c[0x0][0x148], P5; // cy04 < n && cx +  0 < m
--:-:-:Y:7      ISETP.LT.AND P3, PT, cy04, c[0x0][0x148], P6; // cy04 < n && cx + 64 < m
01:-:-:-:1  @P0 STG.CG [Cy00 + 4x<00>], cs0; // Wait Dep 1
--:-:-:-:1      ISETP.LT.AND P0, PT, cy08, c[0x0][0x148], P5; // cy08 < n && cx +  0 < m
--:-:-:-:1  @P1 STG.CG [Cy00 + 4x<64>], cs1;
--:-:-:-:1      ISETP.LT.AND P1, PT, cy08, c[0x0][0x148], P6; // cy08 < n && cx + 64 < m
--:-:-:-:1  @P2 STG.CG [Cy04 + 4x<00>], cs2;
--:-:-:-:1      ISETP.LT.AND P2, PT, cy12, c[0x0][0x148], P5; // cy12 < n && cx +  0 < m
--:-:-:-:1  @P3 STG.CG [Cy04 + 4x<64>], cs3;
--:-:-:Y:7      ISETP.LT.AND P3, PT, cy12, c[0x0][0x148], P6; // cy12 < n && cx + 64 < m
--:-:-:-:2  @P0 STG.CG [Cy08 + 4x<00>], cs4;
--:-:-:-:2  @P1 STG.CG [Cy08 + 4x<64>], cs5;
--:-:-:-:2  @P2 STG.CG [Cy12 + 4x<00>], cs6;
--:2:-:-:1  @P3 STG.CG [Cy12 + 4x<64>], cs7; // Set Dep 2

--:-:-:-:5      RET;