typedef unsigned int uint; #define INFINITY (__int_as_float(0x7f800000)) #define NAN (__int_as_float(0x7fffffff)) template __device__ __forceinline__ T tg_bitcast(F v) { union U { F f; T t; }; U u; u.f = v; return u.t; } #include struct __align__(16) half8 { half x, y, z, w, a, b, c, d; }; __device__ half8 make_half8(half x, half y, half z, half w, half a, half b, half c, half d) { half8 r={x, y, z, w, a, b, c, d}; return r; } struct __align__(8) half4 { half x, y, z, w; }; __device__ half4 make_half4(half x, half y, half z, half w) { half4 r={x, y, z, w}; return r; } __device__ float4 __WMMA_8_16_16_f16_f32(half8 a, half4 b, float4 c){ int *a_pk = (int *)(&a), *b_pk = (int *)(&b), *c_pk = (int *)(&c); asm("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32" "{%0, %1, %2, %3}, {%4, %5, %6, %7}," "{%8, %9}, {%0, %1, %2, %3};" : "+r"(c_pk[0]), "+r"(c_pk[1]), "+r"(c_pk[2]), "+r"(c_pk[3]) : "r"(a_pk[0]), "r"(a_pk[1]), "r"(a_pk[2]), "r"(a_pk[3]), "r"(b_pk[0]), "r"(b_pk[1])); return c; } extern "C" __global__ void __launch_bounds__(64) r_8_32_2_32_2_2_2_4_4_24_2(half* data0_1048576, half* data1_786432, half* data2_786432, half* data3_1024) { float buf0[64]; int gidx0 = blockIdx.x; /* 32 */ int gidx1 = blockIdx.y; /* 8 */ int lidx0 = threadIdx.x; /* 32 */ int lidx1 = threadIdx.y; /* 2 */ float4 cast0 = make_float4(0.0f,0.0f,0.0f,0.0f); int alu0 = (lidx0>>2); int alu1 = (alu0*768); int alu2 = ((lidx0&3)<<1); *(buf0+0) = 0.0f; *(buf0+1) = 0.0f; *(buf0+10) = 0.0f; *(buf0+11) = 0.0f; *(buf0+12) = 0.0f; *(buf0+13) = 0.0f; *(buf0+14) = 0.0f; *(buf0+15) = 0.0f; *(buf0+16) = 0.0f; *(buf0+17) = 0.0f; *(buf0+18) = 0.0f; *(buf0+19) = 0.0f; *(buf0+2) = 0.0f; *(buf0+20) = 0.0f; *(buf0+21) = 0.0f; *(buf0+22) = 0.0f; *(buf0+23) = 0.0f; *(buf0+24) = 0.0f; *(buf0+25) = 0.0f; *(buf0+26) = 0.0f; *(buf0+27) = 0.0f; *(buf0+28) = 0.0f; *(buf0+29) = 0.0f; *(buf0+3) = 0.0f; *(buf0+30) = 0.0f; *(buf0+31) = 0.0f; *(buf0+32) = 0.0f; *(buf0+33) = 0.0f; *(buf0+34) = 0.0f; *(buf0+35) = 0.0f; *(buf0+36) = 0.0f; *(buf0+37) = 0.0f; *(buf0+38) = 0.0f; *(buf0+39) = 0.0f; *(buf0+4) = 0.0f; *(buf0+40) = 0.0f; *(buf0+41) = 0.0f; *(buf0+42) = 0.0f; *(buf0+43) = 0.0f; *(buf0+44) = 0.0f; *(buf0+45) = 0.0f; *(buf0+46) = 0.0f; *(buf0+47) = 0.0f; *(buf0+48) = 0.0f; *(buf0+49) = 0.0f; *(buf0+5) = 0.0f; *(buf0+50) = 0.0f; *(buf0+51) = 0.0f; *(buf0+52) = 0.0f; *(buf0+53) = 0.0f; *(buf0+54) = 0.0f; *(buf0+55) = 0.0f; *(buf0+56) = 0.0f; *(buf0+57) = 0.0f; *(buf0+58) = 0.0f; *(buf0+59) = 0.0f; *(buf0+6) = 0.0f; *(buf0+60) = 0.0f; *(buf0+61) = 0.0f; *(buf0+62) = 0.0f; *(buf0+63) = 0.0f; *(buf0+7) = 0.0f; *(buf0+8) = 0.0f; *(buf0+9) = 0.0f; for (int Lidx0 = 0; Lidx0 < 24; Lidx0++) { int alu67 = (alu2+(Lidx0<<5)); int alu68 = (alu1+(gidx1*98304)+(lidx1*12288)+alu67); half2 val0 = (*((half2*)((data1_786432+(alu68+16))))); half2 val1 = (*((half2*)((data1_786432+(alu68+24))))); half2 val2 = (*((half2*)((data1_786432+(alu68+24576))))); half2 val3 = (*((half2*)((data1_786432+(alu68+24584))))); half2 val4 = (*((half2*)((data1_786432+(alu68+24592))))); half2 val5 = (*((half2*)((data1_786432+(alu68+24600))))); half2 val6 = (*((half2*)((data1_786432+(alu68+30720))))); half2 val7 = (*((half2*)((data1_786432+(alu68+30728))))); half2 val8 = (*((half2*)((data1_786432+(alu68+30736))))); half2 val9 = (*((half2*)((data1_786432+(alu68+30744))))); half2 val10 = (*((half2*)((data1_786432+(alu68+49152))))); half2 val11 = (*((half2*)((data1_786432+(alu68+49160))))); half2 val12 = (*((half2*)((data1_786432+(alu68+49168))))); half2 val13 = (*((half2*)((data1_786432+(alu68+49176))))); half2 val14 = (*((half2*)((data1_786432+(alu68+55296))))); half2 val15 = (*((half2*)((data1_786432+(alu68+55304))))); half2 val16 = (*((half2*)((data1_786432+(alu68+55312))))); half2 val17 = (*((half2*)((data1_786432+(alu68+55320))))); half2 val18 = (*((half2*)((data1_786432+(alu68+6144))))); half2 val19 = (*((half2*)((data1_786432+(alu68+6152))))); half2 val20 = (*((half2*)((data1_786432+(alu68+6160))))); half2 val21 = (*((half2*)((data1_786432+(alu68+6168))))); half2 val22 = (*((half2*)((data1_786432+(alu68+73728))))); half2 val23 = (*((half2*)((data1_786432+(alu68+73736))))); half2 val24 = (*((half2*)((data1_786432+(alu68+73744))))); half2 val25 = (*((half2*)((data1_786432+(alu68+73752))))); half2 val26 = (*((half2*)((data1_786432+(alu68+79872))))); half2 val27 = (*((half2*)((data1_786432+(alu68+79880))))); half2 val28 = (*((half2*)((data1_786432+(alu68+79888))))); half2 val29 = (*((half2*)((data1_786432+(alu68+79896))))); half2 val30 = (*((half2*)((data1_786432+(alu68+8))))); half2 val31 = (*((half2*)((data1_786432+alu68)))); int alu69 = (alu1+(gidx0*24576)+alu67); half2 val32 = (*((half2*)((data2_786432+(alu69+12288))))); half2 val33 = (*((half2*)((data2_786432+(alu69+12296))))); half2 val34 = (*((half2*)((data2_786432+(alu69+12304))))); half2 val35 = (*((half2*)((data2_786432+(alu69+12312))))); half2 val36 = (*((half2*)((data2_786432+(alu69+16))))); half2 val37 = (*((half2*)((data2_786432+(alu69+18432))))); half2 val38 = (*((half2*)((data2_786432+(alu69+18440))))); half2 val39 = (*((half2*)((data2_786432+(alu69+18448))))); half2 val40 = (*((half2*)((data2_786432+(alu69+18456))))); half2 val41 = (*((half2*)((data2_786432+(alu69+24))))); half2 val42 = (*((half2*)((data2_786432+(alu69+6144))))); half2 val43 = (*((half2*)((data2_786432+(alu69+6152))))); half2 val44 = (*((half2*)((data2_786432+(alu69+6160))))); half2 val45 = (*((half2*)((data2_786432+(alu69+6168))))); half2 val46 = (*((half2*)((data2_786432+(alu69+8))))); half2 val47 = (*((half2*)((data2_786432+alu69)))); half8 cast1 = make_half8(val0.x,val0.y,val20.x,val20.y,val1.x,val1.y,val21.x,val21.y); half8 cast2 = make_half8(val2.x,val2.y,val6.x,val6.y,val3.x,val3.y,val7.x,val7.y); half8 cast3 = make_half8(val4.x,val4.y,val8.x,val8.y,val5.x,val5.y,val9.x,val9.y); half8 cast4 = make_half8(val10.x,val10.y,val14.x,val14.y,val11.x,val11.y,val15.x,val15.y); half8 cast5 = make_half8(val12.x,val12.y,val16.x,val16.y,val13.x,val13.y,val17.x,val17.y); half8 cast6 = make_half8(val22.x,val22.y,val26.x,val26.y,val23.x,val23.y,val27.x,val27.y); half8 cast7 = make_half8(val24.x,val24.y,val28.x,val28.y,val25.x,val25.y,val29.x,val29.y); half8 cast8 = make_half8(val31.x,val31.y,val18.x,val18.y,val30.x,val30.y,val19.x,val19.y); half4 cast9 = make_half4(val32.x,val32.y,val33.x,val33.y); half4 cast10 = make_half4(val34.x,val34.y,val35.x,val35.y); half4 cast11 = make_half4(val36.x,val36.y,val41.x,val41.y); half4 cast12 = make_half4(val37.x,val37.y,val38.x,val38.y); half4 cast13 = make_half4(val39.x,val39.y,val40.x,val40.y); half4 cast14 = make_half4(val42.x,val42.y,val43.x,val43.y); half4 cast15 = make_half4(val44.x,val44.y,val45.x,val45.y); half4 cast16 = make_half4(val47.x,val47.y,val46.x,val46.y); float4 wmma0 = __WMMA_8_16_16_f16_f32(cast1, cast10, cast0); float4 wmma1 = __WMMA_8_16_16_f16_f32(cast1, cast11, cast0); float4 wmma2 = __WMMA_8_16_16_f16_f32(cast1, cast13, cast0); float4 wmma3 = __WMMA_8_16_16_f16_f32(cast1, cast15, cast0); float4 wmma4 = __WMMA_8_16_16_f16_f32(cast2, cast9, cast0); float4 wmma5 = __WMMA_8_16_16_f16_f32(cast2, cast12, cast0); float4 wmma6 = __WMMA_8_16_16_f16_f32(cast2, cast14, cast0); float4 wmma7 = __WMMA_8_16_16_f16_f32(cast2, cast16, cast0); float4 wmma8 = __WMMA_8_16_16_f16_f32(cast3, cast10, cast0); float4 wmma9 = __WMMA_8_16_16_f16_f32(cast3, cast11, cast0); float4 wmma10 = __WMMA_8_16_16_f16_f32(cast3, cast13, cast0); float4 wmma11 = __WMMA_8_16_16_f16_f32(cast3, cast15, cast0); float4 wmma12 = __WMMA_8_16_16_f16_f32(cast4, cast9, cast0); float4 wmma13 = __WMMA_8_16_16_f16_f32(cast4, cast12, cast0); float4 wmma14 = __WMMA_8_16_16_f16_f32(cast4, cast14, cast0); float4 wmma15 = __WMMA_8_16_16_f16_f32(cast4, cast16, cast0); float4 wmma16 = __WMMA_8_16_16_f16_f32(cast5, cast10, cast0); float4 wmma17 = __WMMA_8_16_16_f16_f32(cast5, cast11, cast0); float4 wmma18 = __WMMA_8_16_16_f16_f32(cast5, cast13, cast0); float4 wmma19 = __WMMA_8_16_16_f16_f32(cast5, cast15, cast0); float4 wmma20 = __WMMA_8_16_16_f16_f32(cast6, cast9, cast0); float4 wmma21 = __WMMA_8_16_16_f16_f32(cast6, cast12, cast0); float4 wmma22 = __WMMA_8_16_16_f16_f32(cast6, cast14, cast0); float4 wmma23 = __WMMA_8_16_16_f16_f32(cast6, cast16, cast0); float4 wmma24 = __WMMA_8_16_16_f16_f32(cast7, cast10, cast0); float4 wmma25 = __WMMA_8_16_16_f16_f32(cast7, cast11, cast0); float4 wmma26 = __WMMA_8_16_16_f16_f32(cast7, cast13, cast0); float4 wmma27 = __WMMA_8_16_16_f16_f32(cast7, cast15, cast0); float4 wmma28 = __WMMA_8_16_16_f16_f32(cast8, cast9, cast0); float4 wmma29 = __WMMA_8_16_16_f16_f32(cast8, cast12, cast0); float4 wmma30 = __WMMA_8_16_16_f16_f32(cast8, cast14, cast0); float4 wmma31 = __WMMA_8_16_16_f16_f32(cast8, cast16, cast0); *(buf0+0) = ((*(buf0+0))+wmma31.x+wmma1.x); *(buf0+1) = ((*(buf0+1))+wmma31.y+wmma1.y); *(buf0+10) = ((*(buf0+10))+wmma30.z+wmma3.z); *(buf0+11) = ((*(buf0+11))+wmma30.w+wmma3.w); *(buf0+12) = ((*(buf0+12))+wmma28.z+wmma0.z); *(buf0+13) = ((*(buf0+13))+wmma28.w+wmma0.w); *(buf0+14) = ((*(buf0+14))+wmma29.z+wmma2.z); *(buf0+15) = ((*(buf0+15))+wmma29.w+wmma2.w); *(buf0+16) = ((*(buf0+16))+wmma7.x+wmma9.x); *(buf0+17) = ((*(buf0+17))+wmma7.y+wmma9.y); *(buf0+18) = ((*(buf0+18))+wmma6.x+wmma11.x); *(buf0+19) = ((*(buf0+19))+wmma6.y+wmma11.y); *(buf0+2) = ((*(buf0+2))+wmma30.x+wmma3.x); *(buf0+20) = ((*(buf0+20))+wmma4.x+wmma8.x); *(buf0+21) = ((*(buf0+21))+wmma4.y+wmma8.y); *(buf0+22) = ((*(buf0+22))+wmma5.x+wmma10.x); *(buf0+23) = ((*(buf0+23))+wmma5.y+wmma10.y); *(buf0+24) = ((*(buf0+24))+wmma7.z+wmma9.z); *(buf0+25) = ((*(buf0+25))+wmma7.w+wmma9.w); *(buf0+26) = ((*(buf0+26))+wmma6.z+wmma11.z); *(buf0+27) = ((*(buf0+27))+wmma6.w+wmma11.w); *(buf0+28) = ((*(buf0+28))+wmma4.z+wmma8.z); *(buf0+29) = ((*(buf0+29))+wmma4.w+wmma8.w); *(buf0+3) = ((*(buf0+3))+wmma30.y+wmma3.y); *(buf0+30) = ((*(buf0+30))+wmma5.z+wmma10.z); *(buf0+31) = ((*(buf0+31))+wmma5.w+wmma10.w); *(buf0+32) = ((*(buf0+32))+wmma15.x+wmma17.x); *(buf0+33) = ((*(buf0+33))+wmma15.y+wmma17.y); *(buf0+34) = ((*(buf0+34))+wmma14.x+wmma19.x); *(buf0+35) = ((*(buf0+35))+wmma14.y+wmma19.y); *(buf0+36) = ((*(buf0+36))+wmma12.x+wmma16.x); *(buf0+37) = ((*(buf0+37))+wmma12.y+wmma16.y); *(buf0+38) = ((*(buf0+38))+wmma13.x+wmma18.x); *(buf0+39) = ((*(buf0+39))+wmma13.y+wmma18.y); *(buf0+4) = ((*(buf0+4))+wmma28.x+wmma0.x); *(buf0+40) = ((*(buf0+40))+wmma15.z+wmma17.z); *(buf0+41) = ((*(buf0+41))+wmma15.w+wmma17.w); *(buf0+42) = ((*(buf0+42))+wmma14.z+wmma19.z); *(buf0+43) = ((*(buf0+43))+wmma14.w+wmma19.w); *(buf0+44) = ((*(buf0+44))+wmma12.z+wmma16.z); *(buf0+45) = ((*(buf0+45))+wmma12.w+wmma16.w); *(buf0+46) = ((*(buf0+46))+wmma13.z+wmma18.z); *(buf0+47) = ((*(buf0+47))+wmma13.w+wmma18.w); *(buf0+48) = ((*(buf0+48))+wmma23.x+wmma25.x); *(buf0+49) = ((*(buf0+49))+wmma23.y+wmma25.y); *(buf0+5) = ((*(buf0+5))+wmma28.y+wmma0.y); *(buf0+50) = ((*(buf0+50))+wmma22.x+wmma27.x); *(buf0+51) = ((*(buf0+51))+wmma22.y+wmma27.y); *(buf0+52) = ((*(buf0+52))+wmma20.x+wmma24.x); *(buf0+53) = ((*(buf0+53))+wmma20.y+wmma24.y); *(buf0+54) = ((*(buf0+54))+wmma21.x+wmma26.x); *(buf0+55) = ((*(buf0+55))+wmma21.y+wmma26.y); *(buf0+56) = ((*(buf0+56))+wmma23.z+wmma25.z); *(buf0+57) = ((*(buf0+57))+wmma23.w+wmma25.w); *(buf0+58) = ((*(buf0+58))+wmma22.z+wmma27.z); *(buf0+59) = ((*(buf0+59))+wmma22.w+wmma27.w); *(buf0+6) = ((*(buf0+6))+wmma29.x+wmma2.x); *(buf0+60) = ((*(buf0+60))+wmma20.z+wmma24.z); *(buf0+61) = ((*(buf0+61))+wmma20.w+wmma24.w); *(buf0+62) = ((*(buf0+62))+wmma21.z+wmma26.z); *(buf0+63) = ((*(buf0+63))+wmma21.w+wmma26.w); *(buf0+7) = ((*(buf0+7))+wmma29.y+wmma2.y); *(buf0+8) = ((*(buf0+8))+wmma31.z+wmma1.z); *(buf0+9) = ((*(buf0+9))+wmma31.w+wmma1.w); } int alu135 = (alu2+(gidx0<<5)); half2 val48 = (*((half2*)((data3_1024+(alu135+16))))); half2 val49 = (*((half2*)((data3_1024+(alu135+24))))); half2 val50 = (*((half2*)((data3_1024+(alu135+8))))); half2 val51 = (*((half2*)((data3_1024+alu135)))); int alu136 = ((alu0<<10)+(gidx1<<17)+(lidx1<<14)+alu135); *((half2*)((data0_1048576+(alu136+106496)))) = make_half2((((half)((*(buf0+56))))+val51.x),(((half)((*(buf0+57))))+val51.y)); *((half2*)((data0_1048576+(alu136+106504)))) = make_half2((((half)((*(buf0+58))))+val50.x),(((half)((*(buf0+59))))+val50.y)); *((half2*)((data0_1048576+(alu136+106512)))) = make_half2((((half)((*(buf0+60))))+val48.x),(((half)((*(buf0+61))))+val48.y)); *((half2*)((data0_1048576+(alu136+106520)))) = make_half2((((half)((*(buf0+62))))+val49.x),(((half)((*(buf0+63))))+val49.y)); *((half2*)((data0_1048576+(alu136+16)))) = make_half2((((half)((*(buf0+4))))+val48.x),(((half)((*(buf0+5))))+val48.y)); *((half2*)((data0_1048576+(alu136+24)))) = make_half2((((half)((*(buf0+6))))+val49.x),(((half)((*(buf0+7))))+val49.y)); *((half2*)((data0_1048576+(alu136+32768)))) = make_half2((((half)((*(buf0+16))))+val51.x),(((half)((*(buf0+17))))+val51.y)); *((half2*)((data0_1048576+(alu136+32776)))) = make_half2((((half)((*(buf0+18))))+val50.x),(((half)((*(buf0+19))))+val50.y)); *((half2*)((data0_1048576+(alu136+32784)))) = make_half2((((half)((*(buf0+20))))+val48.x),(((half)((*(buf0+21))))+val48.y)); *((half2*)((data0_1048576+(alu136+32792)))) = make_half2((((half)((*(buf0+22))))+val49.x),(((half)((*(buf0+23))))+val49.y)); *((half2*)((data0_1048576+(alu136+40960)))) = make_half2((((half)((*(buf0+24))))+val51.x),(((half)((*(buf0+25))))+val51.y)); *((half2*)((data0_1048576+(alu136+40968)))) = make_half2((((half)((*(buf0+26))))+val50.x),(((half)((*(buf0+27))))+val50.y)); *((half2*)((data0_1048576+(alu136+40976)))) = make_half2((((half)((*(buf0+28))))+val48.x),(((half)((*(buf0+29))))+val48.y)); *((half2*)((data0_1048576+(alu136+40984)))) = make_half2((((half)((*(buf0+30))))+val49.x),(((half)((*(buf0+31))))+val49.y)); *((half2*)((data0_1048576+(alu136+65536)))) = make_half2((((half)((*(buf0+32))))+val51.x),(((half)((*(buf0+33))))+val51.y)); *((half2*)((data0_1048576+(alu136+65544)))) = make_half2((((half)((*(buf0+34))))+val50.x),(((half)((*(buf0+35))))+val50.y)); *((half2*)((data0_1048576+(alu136+65552)))) = make_half2((((half)((*(buf0+36))))+val48.x),(((half)((*(buf0+37))))+val48.y)); *((half2*)((data0_1048576+(alu136+65560)))) = make_half2((((half)((*(buf0+38))))+val49.x),(((half)((*(buf0+39))))+val49.y)); *((half2*)((data0_1048576+(alu136+73728)))) = make_half2((((half)((*(buf0+40))))+val51.x),(((half)((*(buf0+41))))+val51.y)); *((half2*)((data0_1048576+(alu136+73736)))) = make_half2((((half)((*(buf0+42))))+val50.x),(((half)((*(buf0+43))))+val50.y)); *((half2*)((data0_1048576+(alu136+73744)))) = make_half2((((half)((*(buf0+44))))+val48.x),(((half)((*(buf0+45))))+val48.y)); *((half2*)((data0_1048576+(alu136+73752)))) = make_half2((((half)((*(buf0+46))))+val49.x),(((half)((*(buf0+47))))+val49.y)); *((half2*)((data0_1048576+(alu136+8)))) = make_half2((((half)((*(buf0+2))))+val50.x),(((half)((*(buf0+3))))+val50.y)); *((half2*)((data0_1048576+(alu136+8192)))) = make_half2((((half)((*(buf0+8))))+val51.x),(((half)((*(buf0+9))))+val51.y)); *((half2*)((data0_1048576+(alu136+8200)))) = make_half2((((half)((*(buf0+10))))+val50.x),(((half)((*(buf0+11))))+val50.y)); *((half2*)((data0_1048576+(alu136+8208)))) = make_half2((((half)((*(buf0+12))))+val48.x),(((half)((*(buf0+13))))+val48.y)); *((half2*)((data0_1048576+(alu136+8216)))) = make_half2((((half)((*(buf0+14))))+val49.x),(((half)((*(buf0+15))))+val49.y)); *((half2*)((data0_1048576+(alu136+98304)))) = make_half2((((half)((*(buf0+48))))+val51.x),(((half)((*(buf0+49))))+val51.y)); *((half2*)((data0_1048576+(alu136+98312)))) = make_half2((((half)((*(buf0+50))))+val50.x),(((half)((*(buf0+51))))+val50.y)); *((half2*)((data0_1048576+(alu136+98320)))) = make_half2((((half)((*(buf0+52))))+val48.x),(((half)((*(buf0+53))))+val48.y)); *((half2*)((data0_1048576+(alu136+98328)))) = make_half2((((half)((*(buf0+54))))+val49.x),(((half)((*(buf0+55))))+val49.y)); *((half2*)((data0_1048576+alu136))) = make_half2((((half)((*(buf0+0))))+val51.x),(((half)((*(buf0+1))))+val51.y)); }