
/******************************************/
/* Begin Kernel                           */
/******************************************/
.amdgcn_target "amdgcn-amd-amdhsa--gfx936"
.text
.protected Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq
.globl Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq
.p2align 8
.type Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq,@function
.section .rodata,#alloc
.p2align 6
.amdhsa_kernel Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq
  .amdhsa_user_sgpr_kernarg_segment_ptr 1
  .amdhsa_next_free_vgpr 256 // vgprs
  .amdhsa_next_free_sgpr 100 // sgprs
  .amdhsa_group_segment_fixed_size 65536 // lds bytes
  .amdhsa_private_segment_fixed_size 0
  .amdhsa_system_sgpr_workgroup_id_x 1
  .amdhsa_system_sgpr_workgroup_id_y 1
  .amdhsa_system_sgpr_workgroup_id_z 1
  .amdhsa_system_vgpr_workitem_id 0
  .amdhsa_float_denorm_mode_32 3
  .amdhsa_float_denorm_mode_16_64 3
.end_amdhsa_kernel
.text
/* Num VGPR   =256 */
/* Num AccVGPR=0 */
/* Num SGPR   =100 */

/******************************************/
/* Optimizations and Config:              */
/******************************************/
/* ThreadTile= 4 x 2 */
/* SubGroup= 16 x 16 */
/* VectorWidthA=-1 */
/* VectorWidthB=-1 */
/* GlobalReadVectorWidthA=1, GlobalReadVectorWidthB=1 */
/* DirectToLdsA=1 */
/* DirectToLdsB=1 */
/* UseSgprForGRO=0 */
.amdgpu_metadata
---
custom.config:
  InternalSupportParams:
    KernArgsVersion: 2
amdhsa.version:
  - 1
  - 1
amdhsa.kernels:
  - .name: Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq
    .symbol: 'Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq.kd'
    .language:                   OpenCL C
    .language_version:
      - 2
      - 0
    .args:
      - .name:            Gemm info
        .size:            4
        .offset:          0
        .value_kind:      by_value
        .value_type:      u32
      - .name:            kernel info0
        .size:            4
        .offset:          4
        .value_kind:      by_value
        .value_type:      u32
      - .name:            kernel info1
        .size:            4
        .offset:          8
        .value_kind:      by_value
        .value_type:      u32
      - .name:            numWG
        .size:            4
        .offset:          12
        .value_kind:      by_value
        .value_type:      u32
      - .name:            SizesFree0
        .size:            4
        .offset:          16
        .value_kind:      by_value
        .value_type:      u32
      - .name:            SizesFree1
        .size:            4
        .offset:          20
        .value_kind:      by_value
        .value_type:      u32
      - .name:            SizesFree2
        .size:            4
        .offset:          24
        .value_kind:      by_value
        .value_type:      u32
      - .name:            SizesSum0
        .size:            4
        .offset:          28
        .value_kind:      by_value
        .value_type:      u32
      - .name:            D
        .size:            8
        .offset:          32
        .value_kind:      global_buffer
        .value_type:      f16
        .address_space:   generic
      - .name:            C
        .size:            8
        .offset:          40
        .value_kind:      global_buffer
        .value_type:      f16
        .address_space:   generic
      - .name:            A
        .size:            8
        .offset:          48
        .value_kind:      global_buffer
        .value_type:      f16
        .address_space:   generic
      - .name:            B
        .size:            8
        .offset:          56
        .value_kind:      global_buffer
        .value_type:      f16
        .address_space:   generic
      - .name:            strideD0
        .size:            4
        .offset:          64
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideD1
        .size:            4
        .offset:          68
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideC0
        .size:            4
        .offset:          72
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideC1
        .size:            4
        .offset:          76
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideA0
        .size:            4
        .offset:          80
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideA1
        .size:            4
        .offset:          84
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideB0
        .size:            4
        .offset:          88
        .value_kind:      by_value
        .value_type:      u32
      - .name:            strideB1
        .size:            4
        .offset:          92
        .value_kind:      by_value
        .value_type:      u32
      - .name:            alpha
        .size:            4
        .offset:          96
        .value_kind:      by_value
        .value_type:      f32
      - .name:            beta
        .size:            4
        .offset:          100
        .value_kind:      by_value
        .value_type:      f32
      - .name:            AddressDbg
        .size:            8
        .offset:          104
        .value_kind:      global_buffer
        .value_type:      struct
        .address_space:   generic
      - .name:            dstD
        .size:            8
        .offset:          112
        .value_kind:      global_buffer
        .value_type:      f16
        .address_space:   generic
      - .name:            Synchronizer
        .size:            8
        .offset:          120
        .value_kind:      global_buffer
        .value_type:      f32
        .address_space:   generic
      - .name:            GSUSync
        .size:            4
        .offset:          128
        .value_kind:      by_value
        .value_type:      u32
    .group_segment_fixed_size:   65536
    .kernarg_segment_align:      8
    .kernarg_segment_size:       136
    .max_flat_workgroup_size:    768
    .private_segment_fixed_size: 0
    .sgpr_count:                 100
    .sgpr_spill_count:           0
    .vgpr_count:                 256
    .vgpr_spill_count:           0
    .wavefront_size:             64
...
.end_amdgpu_metadata
Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT64x32_dq:
label_ASM_Start:  /// Main body of the asm kernel

.set vgprValuC, 0
/*
.set vgprValuA_X0_I0, 16
.set vgprValuA_X1_I0, 20
.set vgprValuB_X0_I0, 24
.set vgprValuB_X1_I0, 28
*/
.set vgprValuB_X0_I0, 32
.set vgprValuB_X1_I0, 40

.set vgprValuA_X0_I0, 236
.set vgprValuA_X1_I0, 240

.set vgprValuA_X0_H0, 64
.set vgprValuA_X1_H0, 96
.set vgprValuA_X2_I0, 128
.set vgprValuA_X3_I0, 144
.set vgprValuZeros, 160
.set vgprValuZerosI32, 164
.set vgprValuScales, 176
.set vgprValuScalesF32, 180 
.set vgprGlobalReadOffsetScale, 196
.set vgprGlobalReadOffsetZero, 198

//user define
.set vgprGLA, 230
.set vgprGLB, 188
.set vgprLocalWriteA, 196
.set vgprLocalWriteB, 198

.set vgprLocalWriteC, 196
.set vgprLocalReadC, 230
.set vgprTmpValC, 0

.set vgprGlobalWriteD_Edge, 230

.set vgprAddressDbg, 200        //debugbuffer
.set vgprDebugTmp,   202        //debugbuffer
.set vgprSerial,     203
.set vgprTemp0,      204
.set vgprTemp1,      205
.set vgprTemp2,      206
.set vgprTemp3,      207
.set vgprGlobalWriteOffsetD,208
.set vgprGlobalReadOffsetA, 210
.set vgprGlobalReadOffsetA1, 212
.set vgprGlobalReadOffsetB, 214
.set vgprGlobalReadOffsetB1, 218
.set vgprLocalReadAddrA, 220
.set vgprLocalReadAddrB, 228
.set vgprLocalReadAddrB_ori, 228
.set vgprKeepSgprValue, 255



.set BufferLimit, 0   //0xffffffff
.set BufferOOB, 0x80000000
.set Srd127_96, 0x00020000
.set laneSrdA0, 0
.set laneSrdA1, 1
.set laneSrdA2, 3
.set laneSrdB0, 32
.set laneSrdB1, 33
.set laneSrdB2, 34


//Args define
.set sgprKernArgAddress, 0
.set sgprWorkGroup0, 2
.set sgprWorkGroup1, 3
.set sgprWorkGroup2, 4
.set sgprArgType, 5
.set sgprGSUSumIdx, 6
.set sgprNumWorkGroups1, 7
.set sgprGSULog2BpeC, 8
.set sgprGSULog2BpeD, 9
.set sgprStaggerU, 10
.set sgprLoopCounterL, 11
.set sgprGemmCount, 12
.set sgprGSU, 13
.set sgprWGM, 14
.set sgprNumWorkGroups0, 15
.set sgprSrdD, 16
.set sgprSrdC, 20

.set sgprSizesFree, 24
.set sgprSizesSum, 27
.set sgprAddressD, 28
.set sgprAddressC, 30
.set sgprAddressA, 32
.set sgprAddressB, 34
.set sgprStridesD, 36
.set sgprStridesC, 38
.set sgprStridesA, 40
.set sgprStridesB, 42
.set sgprAlpha, 44
.set sgprBeta, 45
.set sgprWGMBuffer, 46
.set sgprAddressDbg, 48     //定义debug buffer address


//user define
.set sgprWorkGroup0Ori, 50
.set sgprLoopCntCommon, 51
.set sgprSrdA, 52
.set sgprSrdB, 56
.set sgprShadowLimitA, 60
.set sgprShadowLimitB, 30
.set sgprGlobalReadIncsA, 62
.set sgprGlobalReadIncsB, 63

.set sgprLocalWriteAddrA, 72
.set sgprLocalWriteAddrB, 73
.set sgprLDSMask, 75
//.set sgprLoopforPfIter, 76
//.set sgprLDSWriteIter, 78

.set sgprLocalWriteAddrA1, 76
.set sgprLocalWriteAddrB1, 77
.set sgprLocalWriteAddrA1ori, 78
.set sgprLocalWriteAddrB1ori, 79



.set sgprTemp0, 80
.set sgprTemp1, 81
.set sgprTemp2, 82
.set sgprTemp3, 83
.set sgprTensor2dSizeA, 84
.set sgprTensor2dSizeB, 84
.set sgprWaveID, 88
.set sgprLDSMask, 89

.set sgprLocalWriteAddrAori, 90
.set sgprLocalWriteAddrBori, 91
.set sgprGlWaveID, 92
.set sgprD_MEdge, 93
.set sgprTemp4, 96
.set sgprTemp5, 97
.set sgprTemp6, 98
.set sgprTemp7, 99
.set sgprStrideStruct, 96
.set sgprStructBit, 97
.set sgprStructNum, 98

.set sgprLocalWriteAddrScale, 74  
.set sgprLocalWriteAddrZero, 75
.set sgprZeroAddress, 68
.set sgprScaleAddress, 70
.set sgprZero, 64
.set sgprScale, 20



.macro MMAC_32x64_0
v_mmac_f32_16x16x16_f16 v[vgprValuC+0*4+0:vgprValuC+0*4+1:vgprValuC+0*4+2:vgprValuC+0*4+3] v[vgprValuA_X2_I0+0*2+0:vgprValuA_X2_I0+0*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+0*4+0: vgprValuC+0*4+1: vgprValuC+0*4+2: vgprValuC+0*4+3] // 
s_setprio 1
v_mmac_f32_16x16x16_f16 v[vgprValuC+1*4+0:vgprValuC+1*4+1:vgprValuC+1*4+2:vgprValuC+1*4+3] v[vgprValuA_X2_I0+1*2+0:vgprValuA_X2_I0+1*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+1*4+0: vgprValuC+1*4+1: vgprValuC+1*4+2: vgprValuC+1*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+2*4+0:vgprValuC+2*4+1:vgprValuC+2*4+2:vgprValuC+2*4+3] v[vgprValuA_X2_I0+0*2+0:vgprValuA_X2_I0+0*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+2*4+0: vgprValuC+2*4+1: vgprValuC+2*4+2: vgprValuC+2*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+3*4+0:vgprValuC+3*4+1:vgprValuC+3*4+2:vgprValuC+3*4+3] v[vgprValuA_X2_I0+1*2+0:vgprValuA_X2_I0+1*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+3*4+0: vgprValuC+3*4+1: vgprValuC+3*4+2: vgprValuC+3*4+3] // 
s_setprio 0
.endm

.macro MMAC_32x64_1
v_mmac_f32_16x16x16_f16 v[vgprValuC+0*4+0:vgprValuC+0*4+1:vgprValuC+0*4+2:vgprValuC+0*4+3] v[vgprValuA_X3_I0+0*2+0:vgprValuA_X3_I0+0*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+0*4+0: vgprValuC+0*4+1: vgprValuC+0*4+2: vgprValuC+0*4+3] // 
s_setprio 1
v_mmac_f32_16x16x16_f16 v[vgprValuC+1*4+0:vgprValuC+1*4+1:vgprValuC+1*4+2:vgprValuC+1*4+3] v[vgprValuA_X3_I0+1*2+0:vgprValuA_X3_I0+1*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+1*4+0: vgprValuC+1*4+1: vgprValuC+1*4+2: vgprValuC+1*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+2*4+0:vgprValuC+2*4+1:vgprValuC+2*4+2:vgprValuC+2*4+3] v[vgprValuA_X3_I0+0*2+0:vgprValuA_X3_I0+0*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+2*4+0: vgprValuC+2*4+1: vgprValuC+2*4+2: vgprValuC+2*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+3*4+0:vgprValuC+3*4+1:vgprValuC+3*4+2:vgprValuC+3*4+3] v[vgprValuA_X3_I0+1*2+0:vgprValuA_X3_I0+1*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+3*4+0: vgprValuC+3*4+1: vgprValuC+3*4+2: vgprValuC+3*4+3] // 
s_setprio 0
.endm

.set MT0, 32
.set MT1, 32

.set LDS_B_OFFSET, 2048
.set LDS_BLK_OFFSET, 4096
.set LDS_BLK_OFFSET_64Kmasked, 4096
.set LOG2BPE, 1
//.set BPE, 2
.set BPE, 1
.set DEPTHU, 32
.set LOG2DEPTHU, 5
.set PFTLOOPS, 6
.set GLWAVES, 4
.set LOG2GLWAVES, 2
.set MperWAVE, 32
.set NperWAVE, 64

/* Load num of Gemms */
s_load_dword s[sgprGemmCount], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x0

/* Load GSU data */
s_load_dword s[sgprGSU], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x4

/* Load WGM data */
s_load_dword s[sgprWGM], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x8

s_add_u32 s[sgprKernArgAddress], s[sgprKernArgAddress], 16 // Shift common args
s_addc_u32 s[sgprKernArgAddress+1], s[sgprKernArgAddress+1], 0x0
s_load_dwordx16 s[24:39], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x0
s_load_dwordx4 s[40:43], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x40
s_load_dwordx2 s[44:45], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x50
s_load_dwordx2 s[sgprAddressDbg:sgprAddressDbg+1], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x58
s_load_dwordx2 s[sgprZeroAddress:sgprZeroAddress+1], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x60
s_load_dwordx2 s[sgprScaleAddress:sgprScaleAddress+1], s[sgprKernArgAddress:sgprKernArgAddress+1], 0x68
s_waitcnt lgkmcnt(0)
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], 8
//s_load_dwordx2 s[sgprTemp2:sgprTemp3], s[sgprWGMBuffer:sgprWGMBuffer+1], s[sgprTemp0]
s_waitcnt lgkmcnt(0)
s_mov_b32 s[sgprWorkGroup0Ori], s[sgprWorkGroup0]
// A uint8 to fp16 mcc
s_lshr_b32 s[sgprSizesFree+0], s[sgprSizesFree+0], 1
//s_lshr_b32 s[sgprStridesA], s[sgprStridesA], 1

v_mov_b32 v[vgprSerial], v0
v_readfirstlane_b32 s[sgprWaveID], v[vgprSerial]
s_lshr_b32 s[sgprWaveID], s[sgprWaveID], 6
s_and_b32 s[sgprGlWaveID], s[sgprWaveID], GLWAVES-1
s_mov_b32 s[sgprLDSMask], 0x4000


v_mov_b32 v8, MT0                                  // set MT0 into sgpr
v_mov_b32 v7, s[sgprSizesFree+0]                   // set Free0 size
v_cvt_f32_u32 v6, v8                               // v6 = ceil(v7 / v8)
v_rcp_iflag_f32 v6, v6                             // v6 = ceil(v7 / v8)
v_cvt_f32_u32 v9, v7                               // v6 = ceil(v7 / v8)
v_mul_f32 v6, v6, v9                               // v6 = ceil(v7 / v8)
v_cvt_u32_f32 v6, v6                               // v6 = ceil(v7 / v8)
v_mul_u32_u24 v9, v6, v8                           // v6 = ceil(v7 / v8)
v_sub_u32 v9, v7, v9                               // v6 = ceil(v7 / v8)
v_cmp_ne_u32 vcc, v9, 0                            // v6 = ceil(v7 / v8)
v_addc_co_u32 v6, vcc, v6, 0, vcc                  // ceil
v_mov_b32 v8, MT1                                  // set MT1 into sgpr
v_mov_b32 v7, s[sgprSizesFree+1]                   // set Free1 size
v_readfirstlane_b32 s[sgprNumWorkGroups0], v6      // set back to numWorkGroup0
v_cvt_f32_u32 v6, v8                               // v6 = ceil(v7 / v8)
v_rcp_iflag_f32 v6, v6                             // v6 = ceil(v7 / v8)
v_cvt_f32_u32 v9, v7                               // v6 = ceil(v7 / v8)
v_mul_f32 v6, v6, v9                               // v6 = ceil(v7 / v8)
v_cvt_u32_f32 v6, v6                               // v6 = ceil(v7 / v8)
v_mul_u32_u24 v9, v6, v8                           // v6 = ceil(v7 / v8)
v_sub_u32 v9, v7, v9                               // v6 = ceil(v7 / v8)
v_cmp_ne_u32 vcc, v9, 0                            // v6 = ceil(v7 / v8)
v_addc_co_u32 v6, vcc, v6, 0, vcc                  // ceil
v_readfirstlane_b32 s[sgprNumWorkGroups1], v6      // set back to numWorkGroup1

/* remap wg from 1D(idxWG012) to 3D(wg2,wg1,wg0) */
/* wg2 = idxWG012 * smallMagicNumber(1/(numWG0*numWG1)) */
s_mul_i32 s78, s[sgprNumWorkGroups0], s[sgprNumWorkGroups1]
s_and_b32 s79, s[sgprGSU], 0x3fff                  // Restore GSU
s_mul_i32 s78, s78, s79
v_cvt_f32_u32 v6, s78                              // s78 = s[sgprWorkGroup0] / s78
v_rcp_iflag_f32 v6, v6                             // s78 = s[sgprWorkGroup0] / s78
v_cvt_f32_u32 v7, s[sgprWorkGroup0]                // s78 = s[sgprWorkGroup0] / s78
v_mul_f32 v6, v6, v7                               // s78 = s[sgprWorkGroup0] / s78
v_cvt_u32_f32 v6, v6                               // s78 = s[sgprWorkGroup0] / s78
v_mul_u32_u24 v7, v6, s78                          // s78 = s[sgprWorkGroup0] / s78
v_sub_u32 v7, s[sgprWorkGroup0], v7                // s78 = s[sgprWorkGroup0] / s78
v_cmpx_eq_u32 exec, v7, s78                        // s78 = s[sgprWorkGroup0] / s78
v_add_u32 v6, 1, v6                                // s78 = s[sgprWorkGroup0] / s78
s_mov_b64 exec, -1                                 // s78 = s[sgprWorkGroup0] / s78
v_readfirstlane_b32 s78, v6                        // quotient
s_mov_b32 s[sgprWorkGroup2], s78
/* idxWG01 = idxWG012 - wg2 * numWG0 * numWG1 */
s_mul_i32 s78, s[sgprNumWorkGroups1], s[sgprNumWorkGroups0]
s_mul_i32 s78, s78, s[sgprWorkGroup2]
s_mul_i32 s78, s78, s79
s_sub_u32 s[sgprWorkGroup0], s[sgprWorkGroup0], s78
/* wg1 = idxWG01 * smallMagicNumber(1/numWG0) */
v_cvt_f32_u32 v6, s[sgprNumWorkGroups0]            // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_rcp_iflag_f32 v6, v6                             // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_cvt_f32_u32 v7, s[sgprWorkGroup0]                // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_mul_f32 v6, v6, v7                               // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_cvt_u32_f32 v6, v6                               // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_mul_u32_u24 v7, v6, s[sgprNumWorkGroups0]        // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_sub_u32 v7, s[sgprWorkGroup0], v7                // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_cmpx_eq_u32 exec, v7, s[sgprNumWorkGroups0]      // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_add_u32 v6, 1, v6                                // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
s_mov_b64 exec, -1                                 // s78 = s[sgprWorkGroup0] / s[sgprNumWorkGroups0]
v_readfirstlane_b32 s78, v6                        // quotient
s_mov_b32 s[sgprWorkGroup1], s78
/* wg0 = idxWG01 - wg1 * numWG0 */
s_mul_i32 s78, s[sgprWorkGroup1], s[sgprNumWorkGroups0]
s_sub_u32 s[sgprWorkGroup0], s[sgprWorkGroup0], s78

s_cmp_le_i32 s[sgprWGM], 1
s_cbranch_scc1 label_WGM
v_cvt_f32_u32 v6, s[sgprWGM]                       // WGM
v_rcp_iflag_f32 v6, v6                             // WGM
v_cvt_f32_u32 v7, s[sgprWorkGroup1]                // WGM
v_mul_f32 v6, v6, v7                               // WGM
v_cvt_u32_f32 v6, v6                               // WGM
v_mul_u32_u24 v7, v6, s[sgprWGM]                   // WGM
v_sub_u32 v7, s[sgprWorkGroup1], v7                // WGM
v_cmpx_eq_u32 exec, v7, s[sgprWGM]                 // WGM
v_add_u32 v6, 1, v6                                // WGM
s_mov_b64 exec, -1                                 // WGM
v_readfirstlane_b32 s76, v6                        // quotient
s_mul_i32 s77, s76, s[sgprWGM]                     // quotient * non-magic divisor
s_sub_u32 s77, s[sgprWorkGroup1], s77              // WorkGroup1=remainder
s_mul_i32 s77, s77, s[sgprNumWorkGroups0]          // (wg1 % WGM)*NumWorkGroups0
s_add_u32 s77, s77, s[sgprWorkGroup0]              // wgSerial = wg0 + (wg1 % WGM)*NumWorkGroups0
v_cvt_f32_u32 v6, s[sgprWGM]                       // WGM
v_rcp_iflag_f32 v6, v6                             // WGM
v_cvt_f32_u32 v7, s[sgprNumWorkGroups1]            // WGM
v_mul_f32 v6, v6, v7                               // WGM
v_cvt_u32_f32 v6, v6                               // WGM
v_mul_u32_u24 v7, v6, s[sgprWGM]                   // WGM
v_sub_u32 v7, s[sgprNumWorkGroups1], v7            // WGM
v_cmpx_eq_u32 exec, v7, s[sgprWGM]                 // WGM
v_add_u32 v6, 1, v6                                // WGM
s_mov_b64 exec, -1                                 // WGM
v_readfirstlane_b32 s78, v6                        // quotient
s_mul_i32 s79, s[sgprWGM], s78                     // quotient * non-magic divisor
s_sub_u32 s79, s[sgprNumWorkGroups1], s79          // NumWorkGroups1=remainder
s_cmp_eq_u32 s79, 0                                // remainder == 0 ?
s_cmov_b32 s79, s[sgprWGM]                         // remainder = WGM if remainder == 0
s_cmp_ge_u32 s76, s78                              // blockId >= numFullBlocks ?
s_cselect_b32 s78, s79, s[sgprWGM]
v_cvt_f32_u32 v6, s78                              // s[sgprWorkGroup0] = s77 / s78
v_rcp_iflag_f32 v6, v6                             // s[sgprWorkGroup0] = s77 / s78
v_cvt_f32_u32 v7, s77                              // s[sgprWorkGroup0] = s77 / s78
v_mul_f32 v6, v6, v7                               // s[sgprWorkGroup0] = s77 / s78
v_cvt_u32_f32 v6, v6                               // s[sgprWorkGroup0] = s77 / s78
v_mul_u32_u24 v7, v6, s78                          // s[sgprWorkGroup0] = s77 / s78
v_sub_u32 v7, s77, v7                              // s[sgprWorkGroup0] = s77 / s78
v_cmpx_eq_u32 exec, v7, s78                        // s[sgprWorkGroup0] = s77 / s78
v_add_u32 v6, 1, v6                                // s[sgprWorkGroup0] = s77 / s78
v_mov_b32 v7, 0                                    // s[sgprWorkGroup1] = s77 % s78
s_mov_b64 exec, -1                                 // s[sgprWorkGroup0] = s77 / s78
v_readfirstlane_b32 s[sgprWorkGroup0], v6          // quotient
v_readfirstlane_b32 s[sgprWorkGroup1], v7          // remainder
s_mul_i32 s[sgprWorkGroup1], s[sgprWorkGroup0], s78 // quotient * non-magic divisor
s_sub_u32 s[sgprWorkGroup1], s77, s[sgprWorkGroup1] // WorkGroup1=remainder
s_mul_i32 s76, s76, s[sgprWGM]                     // blockId * WGM
s_add_u32 s[sgprWorkGroup1], s[sgprWorkGroup1], s76 // wg1 += blockId * WGM
label_WGM:


.set debug_buffer, 0
.if debug_buffer     //计算每个 workgroup 的 debugbuffer 的地址偏移
v_mov_b32 v4, s[sgprWorkGroup0Ori]                    // v1=wg1*nwg0+wg0
v_lshlrev_b32 v5, 0x8, v4                          // v1 = v1 * 256  //这里是 thread 总数，随着wave数量自行更改
v_mul_lo_u32 v6, 2, v5
v_add_u32 v7, v6, v[vgprSerial]                    // v1=tid+NT*(wg1*nwg0+wg0)=serial
v_mul_lo_u32 v8, 0x40, v7                          // v1=serial*nipt*4
v_mov_b32 v2, 0                                    // 
v_mov_b32 v3, s[sgprAddressDbg+1]                  // v3=AddressD1
v_add_co_u32 v[vgprAddressDbg], vcc, s[sgprAddressDbg], v8 // v[vgprAddressDbg]=AddrD* + serial*nipt*4
v_addc_co_u32 v[vgprAddressDbg+1], vcc, v3, v2, vcc // v[vgprAddressDbg]=AddrD* + serial*nipt*4
.endif

.if debug_buffer
v_mov_b32 v[vgprDebugTmp], v[vgprSerial]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc
.endif

.if debug_buffer
v_mov_b32 v[vgprDebugTmp], s[sgprWorkGroup0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprWorkGroup1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprWaveID]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprStridesA]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

.endif

/******************************************/
/* Generate Global A parameters ...       */
/******************************************/
.set COALESCE_THREAD_A, 8     //x4 load
.set LOG2_COALESCE_THREAD_A, 3     //x4 load

v_and_b32 v[vgprTemp0], v[vgprSerial], 63
v_lshrrev_b32 v1, LOG2_COALESCE_THREAD_A, v[vgprTemp0]               
v_and_b32 v0, COALESCE_THREAD_A-1, v[vgprTemp0]  
v_mul_lo_u32 v0, 8, v0                   
//s_mul_i32 s[sgprTemp0], s[sgprStridesA], GLWAVES*BPE  
s_mul_i32 s[sgprTemp0], s[sgprStridesA], GLWAVES                             
v_mul_lo_u32 v1, v1, s[sgprTemp0]              
v_add_u32 v[vgprGlobalReadOffsetA], v0, v1
//s_mul_i32 s[sgprTemp1], s[sgprStridesA], BPE
s_mov_b32 s[sgprTemp1], s[sgprStridesA]
s_mul_i32 s[sgprTemp1], s[sgprTemp1], s[sgprGlWaveID]
v_add_u32 v[vgprGlobalReadOffsetA], v[vgprGlobalReadOffsetA], s[sgprTemp1]


//s_lshr_b32 s[sgprTemp1], s[sgprSizesSum+0], 1 //
s_mul_i32 s[sgprTemp1], s[sgprSizesSum+0], s[sgprStridesA] // notice
s_lshr_b32 s[sgprTemp1], s[sgprTemp1], 1 
v_add_u32 v[vgprGlobalReadOffsetA1], v[vgprGlobalReadOffsetA], s[sgprTemp1]







s_mul_i32    s[sgprTensor2dSizeA+0], s[sgprStridesA+0], s[sgprSizesSum]
s_mul_hi_u32 s[sgprTensor2dSizeA+1], s[sgprStridesA+0], s[sgprSizesSum]
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], 2*MT0                      
s_sub_u32 s[sgprShadowLimitA+0], s[sgprTensor2dSizeA], s[sgprTemp0] 
s_subb_u32 s[sgprShadowLimitA+1], s[sgprTensor2dSizeA+1], 0 
// Set limit to use bytes fp16 = 0
//s_lshl_b64 s[sgprShadowLimitA:sgprShadowLimitA+1], s[sgprShadowLimitA:sgprShadowLimitA+1], LOG2BPE 
s_cmp_eq_u32 s[sgprShadowLimitA+1], 0                           // are we within 2^32?
s_cselect_b32 s[sgprSrdA+2], s[sgprShadowLimitA+0], BufferLimit // Move shadow to real if we are within 2^32
s_mul_hi_u32 s[sgprTemp3], s[sgprStridesA+1], s[sgprWorkGroup2] // Stride*WG
s_mul_i32 s[sgprTemp2], s[sgprStridesA+1], s[sgprWorkGroup2]  // Stride*WG

s_lshl_b64 s[sgprTemp2:sgprTemp3], s[sgprTemp2:sgprTemp3], 0x1   

s_add_u32 s[sgprTemp2], s[sgprTemp2], s[sgprTemp0]                           
s_addc_u32 s[sgprTemp3], s[sgprTemp3], 0
s_add_u32 s[sgprSrdA+0], s[sgprAddressA+0], s[sgprTemp2]   
s_addc_u32 s[sgprSrdA+1], s[sgprAddressA+1], s[sgprTemp3]  
s_mov_b32 s[sgprSrdA+3], Srd127_96                 // Set bits 127_96 in SRD

s_and_b32 s[sgprTemp0], s[sgprStridesA], 15
s_sub_u32 s[sgprTemp1], 16, s[sgprTemp0] 
s_add_u32 s[sgprTemp1], s[sgprShadowLimitA+0], s[sgprTemp1]
s_cmp_gt_u32 s[sgprTemp0], 0
s_cmov_b32 s[sgprShadowLimitA+0], s[sgprTemp1]

//s_cmp_eq_u64 s[sgprAddressA:sgprAddressA+1], 0 // s[sgprAddressA] == 0 ?
//s_cbranch_scc1 label_SkipMmac

//s_mul_i32 s[sgprGlobalReadIncsA+0], DEPTHU*BPE, s[sgprStridesA]  //depthU*PEB
s_mul_i32 s[sgprGlobalReadIncsA+0], DEPTHU, s[sgprStridesA]  //depthU*PEB

/******************************************/
/* Generate LDS A parameters ...          */
/******************************************/
.set WAVE_LDS_OFFSET_A, 64*8     //x4 load
.set WAVE_LDS_OFFSET, WAVE_LDS_OFFSET_A     //x4 load
.set LDS_SUB_M_OFFSET, 16
.set LOADxWAVES_K_A, 64/COALESCE_THREAD_A*GLWAVES
.set LOADxWAVES_K_A_LOG2, 4
.set LOADxWAVES_LDS_OFFSET_A, WAVE_LDS_OFFSET_A*GLWAVES

//Wrap Lds
s_mul_i32 s[sgprLocalWriteAddrA], s[sgprGlWaveID], WAVE_LDS_OFFSET_A+0                    // can add lds pad
s_and_b32 s[sgprTemp0], s[sgprGlWaveID], 1
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 4
s_lshl_b32 s[sgprTemp0], s[sgprTemp0], 16
s_or_b32 s[sgprLocalWriteAddrA], s[sgprLocalWriteAddrA], s[sgprTemp0]
s_mov_b32 s[sgprLocalWriteAddrAori], s[sgprLocalWriteAddrA]

s_add_u32 s[sgprLocalWriteAddrA1], 0x8000, s[sgprLocalWriteAddrA]

//get lds read addrA
v_and_b32 v[vgprTemp0], v[vgprSerial], 63
v_and_b32 v0, 15, v[vgprTemp0]  
v_lshrrev_b32 v1, 4, v[vgprTemp0]
v_mul_u32_u24 v2, 0x40, v1

v_add_u32 v[vgprLocalReadAddrA], v2, v0
s_and_b32 s[sgprTemp0], s[sgprWaveID], 3
s_mul_i32 s[sgprTemp0], s[sgprTemp0], LDS_SUB_M_OFFSET
v_add_u32 v[vgprLocalReadAddrA], v[vgprLocalReadAddrA], s[sgprTemp0]

s_lshr_b32 s[sgprTemp0], s[sgprWaveID], 2
s_mul_i32 s[sgprTemp2], s[sgprTemp0], 0x8000
v_add_u32 v[vgprLocalReadAddrA], v[vgprLocalReadAddrA], s[sgprTemp2]

v_add_u32 v[vgprLocalReadAddrA+1], 0x240, v[vgprLocalReadAddrA]
v_add_u32 v[vgprLocalReadAddrA+2], 0x400, v[vgprLocalReadAddrA]
v_add_u32 v[vgprLocalReadAddrA+3], 0x640, v[vgprLocalReadAddrA]

v_add_u32 v[vgprLocalReadAddrA+4], 0x100, v[vgprLocalReadAddrA]
v_add_u32 v[vgprLocalReadAddrA+5], 0x340, v[vgprLocalReadAddrA]
v_add_u32 v[vgprLocalReadAddrA+6], 0x500, v[vgprLocalReadAddrA]
v_add_u32 v[vgprLocalReadAddrA+7], 0x740, v[vgprLocalReadAddrA]


//fixed input: v[vgprTemp0]->added address;
//             v[vgprTemp1]->max clips; v[vgprTemp2]->reducer
.macro ADDR_WRAP vaddr:req 
v_cmp_ge_u32 s[sgprTemp0:sgprTemp1], \vaddr, v[vgprTemp1]
v_cndmask_b32 v[vgprTemp3], 0, v[vgprTemp2], s[sgprTemp0:sgprTemp1]
v_sub_u32 \vaddr, \vaddr, v[vgprTemp3]
.endm
v_mov_b32 v[vgprTemp2], 0x200
v_mov_b32 v[vgprTemp1], 0x400
v_add_u32 v[vgprTemp1], v[vgprTemp1], s[sgprTemp2]
ADDR_WRAP  v[vgprLocalReadAddrA+5]
v_mov_b32 v[vgprTemp1], 0x800
v_add_u32 v[vgprTemp1], v[vgprTemp1], s[sgprTemp2]
ADDR_WRAP  v[vgprLocalReadAddrA+7]

.set WAVE_LDS_OFFSET, UNDEF     //x4 load
.set WAVE_LDS_OFFSET, UNDEF     //x4 load

/******************************************/
/* Generate Scale Zeros                      */
/******************************************/
v_and_b32 v[vgprGlobalReadOffsetScale], v[vgprSerial], 15   
v_lshlrev_b32 v[vgprGlobalReadOffsetScale], 0x2, v[vgprGlobalReadOffsetScale]
s_and_b32 s[sgprTemp0], s[sgprWaveID], 3  
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 64
v_add_u32 v[vgprGlobalReadOffsetScale], s[sgprTemp0], v[vgprGlobalReadOffsetScale]

s_lshr_b32 s[sgprTemp2], s[sgprSizesSum+0], 6 //
//s_lshl_b32 s[sgprTemp1], s[sgprStridesA], 2
s_lshl_b32 s[sgprTemp1], s[sgprStridesA], 1
s_mul_i32 s[sgprTemp1], s[sgprTemp2], s[sgprTemp1]
s_lshr_b32 s[sgprTemp0], s[sgprWaveID], 2
s_mul_i32 s[sgprTemp0], s[sgprTemp0], s[sgprTemp1]
v_add_u32 v[vgprGlobalReadOffsetScale], s[sgprTemp0], v[vgprGlobalReadOffsetScale]


v_and_b32 v[vgprGlobalReadOffsetZero], v[vgprSerial], 15 
s_and_b32 s[sgprTemp0], s[sgprWaveID], 3  
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 16
v_add_u32 v[vgprGlobalReadOffsetZero], s[sgprTemp0], v[vgprGlobalReadOffsetZero]


s_lshr_b32 s[sgprTemp2], s[sgprSizesSum+0], 7 //
//s_lshl_b32 s[sgprTemp1], s[sgprStridesA], 2
s_mul_i32 s[sgprTemp1], s[sgprTemp2], s[sgprStridesA]
s_lshr_b32 s[sgprTemp0], s[sgprWaveID], 2
s_mul_i32 s[sgprTemp0], s[sgprTemp0], s[sgprTemp1]
v_add_u32 v[vgprGlobalReadOffsetZero], s[sgprTemp0], v[vgprGlobalReadOffsetZero]


s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], MT0*2 
s_add_u32 s[sgprTemp2], s[sgprShadowLimitA], s[sgprTemp0] 
s_addc_u32 s[sgprTemp3], s[sgprShadowLimitA+1], 0 // extend limit for pre-pad
s_cmp_lt_i32 s[sgprWaveID], 8
s_cmov_b32 s[sgprShadowLimitA], s[sgprTemp2]
s_cmov_b32 s[sgprShadowLimitA+1], s[sgprTemp3]

s_mul_hi_u32 s[sgprTemp1], s[sgprWorkGroup0], 8*MT0                       // WorkGroup[01] * MT
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], 8*MT0                          // WorkGroup[01] * MT

s_lshr_b64 s[sgprTemp2:sgprTemp3], s[sgprShadowLimitA:sgprShadowLimitA+1], 4

s_sub_u32 s[sgprTemp2], s[sgprTemp2], s[sgprTemp0] 
s_subb_u32 s[sgprTemp3], s[sgprTemp3], 0 

s_cmp_eq_u32 s[sgprTemp3], 0                           // are we within 2^32?
s_cselect_b32 s[sgprScale+2], s[sgprTemp2], BufferLimit // Move shadow to real if we are within 2^32

s_add_u32 s[sgprScale+0], s[sgprScaleAddress+0], s[sgprTemp0]        // SRD base = Address+ tileStart0
s_addc_u32 s[sgprScale+1], s[sgprScaleAddress+1], s[sgprTemp1]       //

s_mov_b32 s[sgprScale+3], Srd127_96                              // Set bits 127_96 in SRD
s_lshr_b64 s[sgprTemp2:sgprTemp3], s[sgprShadowLimitA:sgprShadowLimitA+1], 6

s_lshr_b64 s[sgprTemp0:sgprTemp1], s[sgprTemp0:sgprTemp1], 2
s_sub_u32 s[sgprTemp2], s[sgprTemp2], s[sgprTemp0] 
s_subb_u32 s[sgprTemp3], s[sgprTemp3], 0 

s_cmp_eq_u32 s[sgprTemp3], 0                           // are we within 2^32?
s_cselect_b32 s[sgprZero+2], s[sgprTemp2], BufferLimit // Move shadow to real if we are within 2^32
s_add_u32 s[sgprZero+0], s[sgprZeroAddress+0], s[sgprTemp0]        // SRD base = Address+ tileStart0
s_addc_u32 s[sgprZero+1], s[sgprZeroAddress+1], s[sgprTemp1]       // SRD base = Address+ tileStart1
s_mov_b32 s[sgprZero+3], Srd127_96   

/******************************************/
/* Define Global Load...                  */
/******************************************/

.macro GLOBAL_LOADAB offset:req

s_add_u32 m0, s[sgprLocalWriteAddrA], \offset            
buffer_load_dwordx2 v[vgprGlobalReadOffsetA+0], s[sgprSrdA:sgprSrdA+3], 0 offen offset:0, lds

s_add_u32 m0, s[sgprLocalWriteAddrA1], \offset            
buffer_load_dwordx2 v[vgprGlobalReadOffsetA1+0], s[sgprSrdA:sgprSrdA+3], 0 offen offset:0, lds
.endm

.macro GLOBAL_LOAD_Scale_Zero
buffer_load_ubyte  v[vgprValuZeros+0], v[vgprGlobalReadOffsetZero+0], s[sgprZero:sgprZero+3], 0 offen offset:0
buffer_load_dword  v[vgprValuScales+0], v[vgprGlobalReadOffsetScale+0], s[sgprScale:sgprScale+3], 0 offen offset:0
.endm

.macro GLOBAL_INC_Scale_Zero
s_lshr_b32 s[sgprTemp0], s[sgprGlobalReadIncsA+0], 3
s_mov_b32 s[sgprTemp1], 0
s_add_u32 s[sgprScale+0], s[sgprScale+0],  s[sgprTemp0]
s_addc_u32 s[sgprScale+1], s[sgprScale+1], s[sgprTemp1]
s_sub_u32 s[sgprTemp2], s[sgprScale+2],  s[sgprTemp0]
s_cmp_lt_u32 s[sgprScale+2], s[sgprTemp0]                            // are we within 2^32?
s_cmov_b32 s[sgprScale+2], 0
s_cmp_ge_u32 s[sgprScale+2], s[sgprTemp0]                            // are we within 2^32?
s_cmov_b32 s[sgprScale+2], s[sgprTemp2] 

s_lshr_b32 s[sgprTemp0], s[sgprGlobalReadIncsA+0], 5
s_mov_b32 s[sgprTemp1], 0
s_add_u32 s[sgprZero+0], s[sgprZero+0],  s[sgprTemp0]
s_addc_u32 s[sgprZero+1], s[sgprZero+1], s[sgprTemp1]
s_sub_u32 s[sgprTemp2], s[sgprZero+2],  s[sgprTemp0]
s_cmp_lt_u32 s[sgprZero+2], s[sgprTemp0]                            // are we within 2^32?
s_cmov_b32 s[sgprZero+2], 0
s_cmp_ge_u32 s[sgprZero+2], s[sgprTemp0]                            // are we within 2^32?
s_cmov_b32 s[sgprZero+2], s[sgprTemp2] 

.endm

.macro UnPackB32ToTwoF32 vgprScale:req vgprOut:req
v_lshlrev_b32 v[\vgprOut+0], 16, v[\vgprScale]
v_lshrrev_b32 v[\vgprOut+0], 16, v[\vgprOut+0]
v_lshrrev_b32 v[\vgprOut+1], 16, v[\vgprScale+0]
v_cvt_f32_f16 v[\vgprOut+0], v[\vgprOut+0]
v_cvt_f32_f16 v[\vgprOut+1], v[\vgprOut+1]
.endm

.macro UnPackB8To2F32 vgprZero:req vgprOut:req
v_lshrrev_b32 v[\vgprOut+1], 4, v[\vgprZero]
v_and_b32 v[\vgprOut+0], v[\vgprZero], 0xf
v_and_b32 v[\vgprOut+1], v[\vgprOut+1], 0xf
v_cvt_f32_ubyte0   v[\vgprOut+0], v[\vgprOut+0]
v_cvt_f32_ubyte0   v[\vgprOut+1], v[\vgprOut+1]
.endm

.macro I4ToFp16 vgprIn:req vgprZero:req vgprScale:req vgprPack:req

v_lshrrev_b32 v[vgprValuA_X0_H0+1], 4, v[\vgprIn]
v_and_b32 v[vgprValuA_X0_H0+0], v[\vgprIn], 0xf

v_lshrrev_b32 v[vgprValuA_X0_H0+3], 4, v[\vgprIn+1]
v_and_b32 v[vgprValuA_X0_H0+2], v[\vgprIn+1], 0xf

v_lshrrev_b32 v[vgprValuA_X0_H0+5], 4, v[\vgprIn+2]
v_and_b32 v[vgprValuA_X0_H0+4], v[\vgprIn+2], 0xf

v_lshrrev_b32 v[vgprValuA_X0_H0+7], 4, v[\vgprIn+3]
v_and_b32 v[vgprValuA_X0_H0+6], v[\vgprIn+3], 0xf

v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+0], v[vgprValuA_X0_H0+0]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+1], v[vgprValuA_X0_H0+1]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+2], v[vgprValuA_X0_H0+2]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+3], v[vgprValuA_X0_H0+3]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+4], v[vgprValuA_X0_H0+4]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+5], v[vgprValuA_X0_H0+5]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+6], v[vgprValuA_X0_H0+6]
v_cvt_f32_ubyte0   v[vgprValuA_X0_H0+7], v[vgprValuA_X0_H0+7]

v_pk_fma_f32 v[vgprValuA_X0_H0+0:vgprValuA_X0_H0+1], v[vgprValuA_X0_H0+0:vgprValuA_X0_H0+1], v[\vgprScale:\vgprScale+1], v[\vgprZero:\vgprZero+1]
v_pk_fma_f32 v[vgprValuA_X0_H0+2:vgprValuA_X0_H0+3], v[vgprValuA_X0_H0+2:vgprValuA_X0_H0+3], v[\vgprScale:\vgprScale+1], v[\vgprZero:\vgprZero+1]
v_pk_fma_f32 v[vgprValuA_X0_H0+4:vgprValuA_X0_H0+5], v[vgprValuA_X0_H0+4:vgprValuA_X0_H0+5], v[\vgprScale:\vgprScale+1], v[\vgprZero:\vgprZero+1]
v_pk_fma_f32 v[vgprValuA_X0_H0+6:vgprValuA_X0_H0+7], v[vgprValuA_X0_H0+6:vgprValuA_X0_H0+7], v[\vgprScale:\vgprScale+1], v[\vgprZero:\vgprZero+1]

v_cvt_f16_f32 v[\vgprPack+0], v[vgprValuA_X0_H0+0]
v_cvt_f16_f32 v[\vgprPack+1], v[vgprValuA_X0_H0+1]
v_cvt_f16_f32 v[\vgprPack+2], v[vgprValuA_X0_H0+2]
v_cvt_f16_f32 v[\vgprPack+3], v[vgprValuA_X0_H0+3]

v_cvt_f16_f32 v[\vgprPack+4], v[vgprValuA_X0_H0+4]
v_cvt_f16_f32 v[\vgprPack+5], v[vgprValuA_X0_H0+5]
v_cvt_f16_f32 v[\vgprPack+6], v[vgprValuA_X0_H0+6]
v_cvt_f16_f32 v[\vgprPack+7], v[vgprValuA_X0_H0+7]
.endm



.macro WriteFp16ToGlobal vgprOut:req
buffer_store_short v[\vgprOut+0],  v[vgprGlobalWriteOffsetD+0], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0 
buffer_store_short v[\vgprOut+1],  v[vgprGlobalWriteOffsetD+1], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0

s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*1               
s_add_u32  s[sgprSrdD+0], s[sgprSrdD+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdD+1], s[sgprSrdD+1], 0        
s_subb_u32 s[sgprSrdD+2], s[sgprSrdD+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdD+2], 0

buffer_store_short v[\vgprOut+2],  v[vgprGlobalWriteOffsetD+0], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0 
buffer_store_short v[\vgprOut+3],  v[vgprGlobalWriteOffsetD+1], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0

s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*1               
s_add_u32  s[sgprSrdD+0], s[sgprSrdD+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdD+1], s[sgprSrdD+1], 0        
s_subb_u32 s[sgprSrdD+2], s[sgprSrdD+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdD+2], 0

buffer_store_short v[\vgprOut+4],  v[vgprGlobalWriteOffsetD+0], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0 
buffer_store_short v[\vgprOut+5],  v[vgprGlobalWriteOffsetD+1], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0

s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*1               
s_add_u32  s[sgprSrdD+0], s[sgprSrdD+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdD+1], s[sgprSrdD+1], 0        
s_subb_u32 s[sgprSrdD+2], s[sgprSrdD+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdD+2], 0

buffer_store_short v[\vgprOut+6],  v[vgprGlobalWriteOffsetD+0], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0 
buffer_store_short v[\vgprOut+7],  v[vgprGlobalWriteOffsetD+1], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0

s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*13               
s_add_u32  s[sgprSrdD+0], s[sgprSrdD+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdD+1], s[sgprSrdD+1], 0        
s_subb_u32 s[sgprSrdD+2], s[sgprSrdD+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdD+2], 0

.endm




/******************************************/
/* Define Global Load adress Increase...  */
/******************************************/

.macro GLOBAL_INC 

s_mov_b32 s[sgprTemp0], s[sgprGlobalReadIncsA+0]
s_mov_b32 s[sgprTemp1], 0
s_add_u32 s[sgprSrdA+0], s[sgprSrdA+0],  s[sgprTemp0]
s_addc_u32 s[sgprSrdA+1], s[sgprSrdA+1], s[sgprTemp1]
s_sub_u32 s[sgprShadowLimitA+0], s[sgprShadowLimitA+0],  s[sgprTemp0]
s_subb_u32 s[sgprShadowLimitA+1], s[sgprShadowLimitA+1], s[sgprTemp1]
s_cmp_eq_u32 s[sgprShadowLimitA+1], 0                            // are we within 2^32?
s_cselect_b32 s[sgprSrdA+2], s[sgprShadowLimitA+0], BufferLimit // Move shadow to real if we are within 2^32
.endm

/******************************************/
/* Define LDS Load...                     */
/******************************************/

.macro LDS_LOADAB off:req

ds_read_u8 v[vgprValuA_X0_I0+ 0], v[vgprLocalReadAddrA+0] offset:\off 
ds_read_u8 v[vgprValuA_X0_I0+ 1], v[vgprLocalReadAddrA+1] offset:\off 
ds_read_u8 v[vgprValuA_X0_I0+ 2], v[vgprLocalReadAddrA+2] offset:\off 
ds_read_u8 v[vgprValuA_X0_I0+ 3], v[vgprLocalReadAddrA+3] offset:\off 
.endm

.macro LDS_LOADAB1 off:req

ds_read_u8 v[vgprValuA_X1_I0+ 0], v[vgprLocalReadAddrA+4] offset:\off 
ds_read_u8 v[vgprValuA_X1_I0+ 1], v[vgprLocalReadAddrA+5] offset:\off 
ds_read_u8 v[vgprValuA_X1_I0+ 2], v[vgprLocalReadAddrA+6] offset:\off 
ds_read_u8 v[vgprValuA_X1_I0+ 3], v[vgprLocalReadAddrA+7] offset:\off 
.endm

/******************************************/
/* Use Global Load Wave process ...       */
/******************************************/

//s_lshr_b32 s[sgprLoopCounterL], s[sgprSizesSum], 5
s_lshr_b32 s[sgprLoopCounterL], s[sgprSizesSum], 6


s_min_u32 s[sgprLoopCntCommon], 4, s[sgprLoopCounterL] 
s_and_b32 s[sgprTemp0], s[sgprSizesSum], 31   
s_add_u32 s[sgprLoopCntCommon], s[sgprLoopCounterL], scc
s_sub_u32 s[sgprLoopCounterL], s[sgprLoopCntCommon], 1   // -1 for tail

s_cmp_lt_i32 s[sgprWaveID], 8
s_cbranch_scc1 SkipGL


.if debug_buffer
v_mov_b32 v[vgprDebugTmp], s[sgprLoopCounterL]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprLocalWriteAddrA]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprLocalWriteAddrA1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], v[vgprGlobalReadOffsetA]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], v[vgprGlobalReadOffsetA1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprSrdA+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp], s[sgprSrdA+1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc



.endif



s_cmp_eq_i32 s[sgprLoopCntCommon], 1
s_cbranch_scc1 SkipToLastLoad
s_mov_b32 s[sgprTemp3], 0
PreFetchBegin:
GLOBAL_LOADAB LDS_BLK_OFFSET*0
GLOBAL_INC

s_addk_i32 s[sgprLocalWriteAddrA], LDS_BLK_OFFSET_64Kmasked 
s_add_u32 s[sgprTemp0], s[sgprLocalWriteAddrAori], s[sgprLDSMask]      
s_cmp_ge_u32 s[sgprLocalWriteAddrA], s[sgprTemp0]      
s_cmov_b32  s[sgprLocalWriteAddrA], s[sgprLocalWriteAddrAori] 


s_add_u32 s[sgprLocalWriteAddrA1], 0x8000, s[sgprLocalWriteAddrA]

s_cmp_lt_i32 s[sgprTemp3], 2
s_cbranch_scc1 SkipWait
s_waitcnt vmcnt(4)
s_barrier
SkipWait:
s_add_u32 s[sgprTemp3], s[sgprTemp3], 1
s_sub_u32 s[sgprLoopCntCommon], s[sgprLoopCntCommon], 1
s_cmp_gt_i32 s[sgprLoopCntCommon], 1
s_cbranch_scc1 PreFetchBegin
s_cmp_eq_i32 s[sgprTemp3], 1
s_cbranch_scc1 Last1
s_waitcnt vmcnt(2)
s_barrier
Last1:
s_waitcnt vmcnt(0)
s_barrier
s_barrier

SkipToLastLoad:
s_mov_b32  s[sgprLocalWriteAddrA], s[sgprLocalWriteAddrAori]  
s_add_u32 s[sgprLocalWriteAddrA1], 0x8000, s[sgprLocalWriteAddrA]

GLOBAL_LOADAB LDS_BLK_OFFSET*0
s_waitcnt vmcnt(0)
s_barrier
s_endpgm
SkipGL:

/******************************************/
/* Generate SrcD ...                      */
/******************************************/

s_mov_b32 s[sgprSrdD+0], s[sgprAddressD+0]         // init SRD base address (lower)
s_mov_b32 s[sgprSrdD+1], s[sgprAddressD+1]         // init SRD base address (upper) + other fields
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], 128
s_mul_i32 s[sgprTemp1], s[sgprWorkGroup1], MT1
s_mul_i32 s[sgprTemp1], s[sgprTemp1], s[sgprStridesD]
s_add_u32 s[sgprTemp0], s[sgprTemp0], s[sgprTemp1]
s_mul_hi_u32 s[sgprTemp3], s[sgprStridesC+1], s[sgprWorkGroup2] 
s_mul_i32 s[sgprTemp2], s[sgprStridesC+1], s[sgprWorkGroup2] 
s_add_u32 s[sgprTemp2], s[sgprTemp2], s[sgprTemp0]
s_addc_u32 s[sgprTemp3], s[sgprTemp3], 0
s_lshl_b64 s[sgprTemp2:sgprTemp3], s[sgprTemp2:sgprTemp3], 0x1   
s_add_u32 s[sgprSrdD+0], s[sgprSrdD+0], s[sgprTemp2] 
s_addc_u32 s[sgprSrdD+1], s[sgprSrdD+1], s[sgprTemp3]   
//s_mul_i32 s[sgprTemp0], s[sgprStridesD+0], s[sgprSizesFree+1]

s_mul_i32 s[sgprTemp0], s[sgprStridesD+0], s[sgprSizesSum]
s_sub_u32 s[sgprTemp1], s[sgprTemp0], s[sgprTemp1]
s_lshl_b32 s[sgprSrdD+2], s[sgprTemp1], 1
s_mov_b32 s[sgprSrdD+3], Srd127_96                 // Set bits 127_96 in post-loop SRD

v_and_b32 v[vgprTemp0], v[vgprSerial], 63
v_lshrrev_b32 v[vgprTemp1], 4, v[vgprTemp0]  
v_lshlrev_b32 v[vgprTemp1], 2, v[vgprTemp1] 
s_lshr_b32 s[sgprTemp0], s[sgprWaveID], 0x2
//s_lshl_b32 s[sgprTemp0], s[sgprTemp0], 0x4 

s_lshr_b32 s[sgprTemp1], s[sgprSizesSum], 0x1 
s_mul_i32 s[sgprTemp0], s[sgprTemp0], s[sgprTemp1]
v_add_u32 v[vgprTemp1], s[sgprTemp0], v[vgprTemp1]
v_and_b32 v[vgprTemp2], 15, v[vgprTemp0]                    
v_lshlrev_b32 v[vgprTemp2], 0x2, v[vgprTemp2] 
v_mul_lo_u32 v[vgprTemp1], v[vgprTemp1], s[sgprStridesD]
v_lshlrev_b32 v[vgprTemp1], 0x1, v[vgprTemp1]        
v_add_u32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], v[vgprTemp2]
s_and_b32 s[sgprTemp0], s[sgprWaveID], 3
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 32
s_lshl_b32 s[sgprTemp0], s[sgprTemp0], 0x1 
v_add_u32 v[vgprGlobalWriteOffsetD], v[vgprGlobalWriteOffsetD], s[sgprTemp0]

v_add_u32 v[vgprGlobalWriteOffsetD+1], v[vgprGlobalWriteOffsetD], 2

/******************************************/
/* Init ValueC ...                        */
/******************************************/
GLOBAL_LOAD_Scale_Zero

GLOBAL_INC_Scale_Zero


s_lshl_b32 s[sgprSizesFree+0], s[sgprSizesFree+0], 2
s_mul_i32 s[sgprD_MEdge], s[sgprWorkGroup0], 128
s_sub_u32 s[sgprD_MEdge], s[sgprSizesFree+0], s[sgprD_MEdge]
s_lshl_b32 s[sgprD_MEdge], s[sgprD_MEdge], 1
s_min_u32 s[sgprD_MEdge], s[sgprD_MEdge], 128*2

v_and_b32 v[vgprGlobalWriteD_Edge+2], v[vgprSerial], 15 
v_lshlrev_b32  v[vgprGlobalWriteD_Edge+2], 1, v[vgprGlobalWriteD_Edge+2] 
s_and_b32 s[sgprTemp1], s[sgprWaveID], 3
s_mul_i32 s[sgprTemp1], s[sgprTemp1], 32
v_add_u32 v[vgprGlobalWriteD_Edge+2], v[vgprGlobalWriteD_Edge+2], s[sgprTemp1]         
v_lshlrev_b32  v[vgprGlobalWriteD_Edge+2], 1, v[vgprGlobalWriteD_Edge+2]                 
//v_mov_b32 v[vgprGlobalWriteD_Edge+3], v[vgprGlobalWriteD_Edge+2]               //store inittial addr
//v_mov_b32 v[vgprGlobalWriteD_Edge+0], v[vgprGlobalWriteOffsetD]  //store inittial addr

//v_mov_b32 v[vgprGlobalWriteD_Edge+1], v[vgprGlobalWriteD_Edge+0]  
//v_mov_b32 v[vgprGlobalWriteD_Edge+2], v[vgprGlobalWriteD_Edge+3]   
//v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprGlobalWriteD_Edge+2], s[sgprD_MEdge]
v_cmp_gt_u32 s[sgprTemp2:sgprTemp2+1], v[vgprGlobalWriteD_Edge+2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprGlobalWriteOffsetD], -1, s[sgprTemp2:sgprTemp2+1]
v_cndmask_b32 v[vgprGlobalWriteOffsetD+1], v[vgprGlobalWriteOffsetD+1], -1, s[sgprTemp2:sgprTemp2+1]

.if debug_buffer
v_mov_b32 v[vgprDebugTmp],  v[vgprGlobalWriteOffsetD]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc
.endif


/******************************************/
/* LoopCounter == 0, Skip to tail/last loop ... */
/******************************************/
s_cmp_le_i32 s[sgprLoopCounterL], 0
s_cbranch_scc1 TAIL_LOOP
s_waitcnt vmcnt(0)
s_barrier

LDS_LOADAB LDS_BLK_OFFSET*0
UnPackB32ToTwoF32 vgprValuScales+0 vgprValuScalesF32+0
UnPackB8To2F32 vgprValuZeros+0 vgprValuZerosI32+0

.if debug_buffer

v_mov_b32 v[vgprDebugTmp],  v[vgprValuScales+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuZeros+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

.endif



v_mul_f32 v[vgprValuZerosI32+0], -v[vgprValuZerosI32+0], v[vgprValuScalesF32+0]
v_mul_f32 v[vgprValuZerosI32+1], -v[vgprValuZerosI32+1], v[vgprValuScalesF32+1]

/******************************************/
/* Main Loop Process ...                  */
/******************************************/

s_mov_b32 s[sgprLoopCntCommon], s[sgprLoopCounterL]
MainLoopBeginW0_3:

LDS_LOADAB1 LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X2_I0



.if debug_buffer

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X0_I0+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X0_I0+1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc


//v_mov_b32 v[vgprDebugTmp],  v[vgprLocalReadAddrA+0]
//flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
//v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

//v_mov_b32 v[vgprDebugTmp],  v[vgprLocalReadAddrA+1]
//flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
//v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+2]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+3]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

//v_mov_b32 v[vgprDebugTmp],  v[vgprGlobalWriteOffsetD+0]
//flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
//v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

//v_mov_b32 v[vgprDebugTmp],  v[vgprGlobalWriteOffsetD+1]
//flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
//v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc





.endif





WriteFp16ToGlobal vgprValuA_X2_I0

//s_waitcnt vmcnt(0)
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*1
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X3_I0

//s_waitcnt vmcnt(0)
WriteFp16ToGlobal vgprValuA_X3_I0



s_sub_u32 s[sgprLoopCntCommon], s[sgprLoopCntCommon], 1
s_cmp_le_i32 s[sgprLoopCntCommon], 0
s_cbranch_scc1 TAIL_LOOP

GLOBAL_LOAD_Scale_Zero
LDS_LOADAB1 LDS_BLK_OFFSET*1
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X2_I0
WriteFp16ToGlobal vgprValuA_X2_I0
s_waitcnt vmcnt(8)
s_barrier
GLOBAL_INC_Scale_Zero
LDS_LOADAB LDS_BLK_OFFSET*2
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X3_I0
WriteFp16ToGlobal vgprValuA_X3_I0
UnPackB32ToTwoF32 vgprValuScales+0 vgprValuScalesF32+0
UnPackB8To2F32 vgprValuZeros+0 vgprValuZerosI32+0
v_mul_f32 v[vgprValuZerosI32+0], -v[vgprValuZerosI32+0], v[vgprValuScalesF32+0]
v_mul_f32 v[vgprValuZerosI32+1], -v[vgprValuZerosI32+1], v[vgprValuScalesF32+1]

s_sub_u32 s[sgprLoopCntCommon], s[sgprLoopCntCommon], 1
s_cmp_le_i32 s[sgprLoopCntCommon], 0
s_cbranch_scc1 TAIL_LOOP

LDS_LOADAB1 LDS_BLK_OFFSET*2
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X2_I0
WriteFp16ToGlobal vgprValuA_X2_I0
//s_waitcnt vmcnt(0)
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*3
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X3_I0
WriteFp16ToGlobal vgprValuA_X3_I0

s_sub_u32 s[sgprLoopCntCommon], s[sgprLoopCntCommon], 1
s_cmp_le_i32 s[sgprLoopCntCommon], 0
s_cbranch_scc1 TAIL_LOOP

GLOBAL_LOAD_Scale_Zero
LDS_LOADAB1 LDS_BLK_OFFSET*3
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X2_I0
WriteFp16ToGlobal vgprValuA_X2_I0
s_waitcnt vmcnt(8)
s_barrier
GLOBAL_INC_Scale_Zero
LDS_LOADAB LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(6)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X3_I0
WriteFp16ToGlobal vgprValuA_X3_I0
UnPackB32ToTwoF32 vgprValuScales+0 vgprValuScalesF32+0
UnPackB8To2F32 vgprValuZeros+0 vgprValuZerosI32+0
v_mul_f32 v[vgprValuZerosI32+0], -v[vgprValuZerosI32+0], v[vgprValuScalesF32+0]
v_mul_f32 v[vgprValuZerosI32+1], -v[vgprValuZerosI32+1], v[vgprValuScalesF32+1]

s_sub_u32 s[sgprLoopCntCommon], s[sgprLoopCntCommon], 1
s_cmp_gt_i32 s[sgprLoopCntCommon], 0
s_cbranch_scc1 MainLoopBeginW0_3

s_cmp_le_i32 s[sgprLoopCntCommon], 0
s_cbranch_scc1 TAIL_LOOP

/******************************************/
/* Tail Loop Process ...                  */
/******************************************/
TAIL_LOOP:

s_waitcnt vmcnt(0)
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(0)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X2_I0

/*
.if debug_buffer
v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X0_I0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  v[vgprValuA_X2_I0+1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  s[sgprSrdD+0]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  s[sgprSrdD+1]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  s[sgprSrdD+2]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc

v_mov_b32 v[vgprDebugTmp],  s[sgprSrdD+3]
flat_store_dword v[vgprAddressDbg:vgprAddressDbg+1], v[vgprDebugTmp] // debug dump store
v_add_u32 v[vgprAddressDbg], v[vgprAddressDbg], 0x4 // debug dump inc
.endif
*/



WriteFp16ToGlobal vgprValuA_X2_I0
s_waitcnt vmcnt(0)
LDS_LOADAB1 LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(0)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF32+0 vgprValuA_X3_I0
WriteFp16ToGlobal vgprValuA_X3_I0


s_endpgm