
/******************************************/
/* Begin Kernel                           */
/******************************************/
.amdgcn_target "amdgcn-amd-amdhsa--gfx936"
.text
.protected Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2
.globl Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2
.p2align 8
.type Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2,@function
.section .rodata,#alloc
.p2align 6
.amdhsa_kernel Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2
  .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= 8 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_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2
    .symbol: 'Cijk_Ailk_Bljk_HHS_BH_UserArgs_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2.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:    512
    .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_MT256x32x32_SN_K1_PGR6_SB1_TT8_2_WG16_16_2:
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 vgprValuA_X0_I0, 236
.set vgprValuA_X1_I0, 240
.set vgprValuB_X0_I0, 244
.set vgprValuB_X1_I0, 248

.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 vgprValuScalesF16, 180 
 

.set vgprGlobalReadOffsetScale, 196
.set vgprGlobalReadOffsetZero, 198


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

.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 vgprGlobalReadOffsetB, 214
.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 sgprWaveID, 74



.set sgprLDSWriteIter, 78
.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_32x32_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+2*2+0:vgprValuA_X2_I0+2*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*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+3*2+0:vgprValuA_X2_I0+3*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+3*4+0: vgprValuC+3*4+1: vgprValuC+3*4+2: vgprValuC+3*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+4*4+0:vgprValuC+4*4+1:vgprValuC+4*4+2:vgprValuC+4*4+3] v[vgprValuA_X2_I0+4*2+0:vgprValuA_X2_I0+4*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+4*4+0: vgprValuC+4*4+1: vgprValuC+4*4+2: vgprValuC+4*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+5*4+0:vgprValuC+5*4+1:vgprValuC+5*4+2:vgprValuC+5*4+3] v[vgprValuA_X2_I0+5*2+0:vgprValuA_X2_I0+5*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+5*4+0: vgprValuC+5*4+1: vgprValuC+5*4+2: vgprValuC+5*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+6*4+0:vgprValuC+6*4+1:vgprValuC+6*4+2:vgprValuC+6*4+3] v[vgprValuA_X2_I0+6*2+0:vgprValuA_X2_I0+6*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+6*4+0: vgprValuC+6*4+1: vgprValuC+6*4+2: vgprValuC+6*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+7*4+0:vgprValuC+7*4+1:vgprValuC+7*4+2:vgprValuC+7*4+3] v[vgprValuA_X2_I0+7*2+0:vgprValuA_X2_I0+7*2+1] v[vgprValuB_X0_I0+0*2+0:vgprValuB_X0_I0+0*2+1] v[vgprValuC+7*4+0: vgprValuC+7*4+1: vgprValuC+7*4+2: vgprValuC+7*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+8*4+0:vgprValuC+8*4+1:vgprValuC+8*4+2:vgprValuC+8*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+8*4+0: vgprValuC+8*4+1: vgprValuC+8*4+2: vgprValuC+8*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+9*4+0:vgprValuC+9*4+1:vgprValuC+9*4+2:vgprValuC+9*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+9*4+0: vgprValuC+9*4+1: vgprValuC+9*4+2: vgprValuC+9*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+10*4+0:vgprValuC+10*4+1:vgprValuC+10*4+2:vgprValuC+10*4+3] v[vgprValuA_X2_I0+2*2+0:vgprValuA_X2_I0+2*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+10*4+0: vgprValuC+10*4+1: vgprValuC+10*4+2: vgprValuC+10*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+11*4+0:vgprValuC+11*4+1:vgprValuC+11*4+2:vgprValuC+11*4+3] v[vgprValuA_X2_I0+3*2+0:vgprValuA_X2_I0+3*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+11*4+0: vgprValuC+11*4+1: vgprValuC+11*4+2: vgprValuC+11*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+12*4+0:vgprValuC+12*4+1:vgprValuC+12*4+2:vgprValuC+12*4+3] v[vgprValuA_X2_I0+4*2+0:vgprValuA_X2_I0+4*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+12*4+0: vgprValuC+12*4+1: vgprValuC+12*4+2: vgprValuC+12*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+13*4+0:vgprValuC+13*4+1:vgprValuC+13*4+2:vgprValuC+13*4+3] v[vgprValuA_X2_I0+5*2+0:vgprValuA_X2_I0+5*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+13*4+0: vgprValuC+13*4+1: vgprValuC+13*4+2: vgprValuC+13*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+14*4+0:vgprValuC+14*4+1:vgprValuC+14*4+2:vgprValuC+14*4+3] v[vgprValuA_X2_I0+6*2+0:vgprValuA_X2_I0+6*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+14*4+0: vgprValuC+14*4+1: vgprValuC+14*4+2: vgprValuC+14*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+15*4+0:vgprValuC+15*4+1:vgprValuC+15*4+2:vgprValuC+15*4+3] v[vgprValuA_X2_I0+7*2+0:vgprValuA_X2_I0+7*2+1] v[vgprValuB_X0_I0+1*2+0:vgprValuB_X0_I0+1*2+1] v[vgprValuC+15*4+0: vgprValuC+15*4+1: vgprValuC+15*4+2: vgprValuC+15*4+3] //
s_setprio 0
.endm

.macro MMAC_32x32_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+2*2+0:vgprValuA_X3_I0+2*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*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+3*2+0:vgprValuA_X3_I0+3*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+3*4+0: vgprValuC+3*4+1: vgprValuC+3*4+2: vgprValuC+3*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+4*4+0:vgprValuC+4*4+1:vgprValuC+4*4+2:vgprValuC+4*4+3] v[vgprValuA_X3_I0+4*2+0:vgprValuA_X3_I0+4*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+4*4+0: vgprValuC+4*4+1: vgprValuC+4*4+2: vgprValuC+4*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+5*4+0:vgprValuC+5*4+1:vgprValuC+5*4+2:vgprValuC+5*4+3] v[vgprValuA_X3_I0+5*2+0:vgprValuA_X3_I0+5*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+5*4+0: vgprValuC+5*4+1: vgprValuC+5*4+2: vgprValuC+5*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+6*4+0:vgprValuC+6*4+1:vgprValuC+6*4+2:vgprValuC+6*4+3] v[vgprValuA_X3_I0+6*2+0:vgprValuA_X3_I0+6*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+6*4+0: vgprValuC+6*4+1: vgprValuC+6*4+2: vgprValuC+6*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+7*4+0:vgprValuC+7*4+1:vgprValuC+7*4+2:vgprValuC+7*4+3] v[vgprValuA_X3_I0+7*2+0:vgprValuA_X3_I0+7*2+1] v[vgprValuB_X1_I0+0*2+0:vgprValuB_X1_I0+0*2+1] v[vgprValuC+7*4+0: vgprValuC+7*4+1: vgprValuC+7*4+2: vgprValuC+7*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+8*4+0:vgprValuC+8*4+1:vgprValuC+8*4+2:vgprValuC+8*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+8*4+0: vgprValuC+8*4+1: vgprValuC+8*4+2: vgprValuC+8*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+9*4+0:vgprValuC+9*4+1:vgprValuC+9*4+2:vgprValuC+9*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+9*4+0: vgprValuC+9*4+1: vgprValuC+9*4+2: vgprValuC+9*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+10*4+0:vgprValuC+10*4+1:vgprValuC+10*4+2:vgprValuC+10*4+3] v[vgprValuA_X3_I0+2*2+0:vgprValuA_X3_I0+2*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+10*4+0: vgprValuC+10*4+1: vgprValuC+10*4+2: vgprValuC+10*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+11*4+0:vgprValuC+11*4+1:vgprValuC+11*4+2:vgprValuC+11*4+3] v[vgprValuA_X3_I0+3*2+0:vgprValuA_X3_I0+3*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+11*4+0: vgprValuC+11*4+1: vgprValuC+11*4+2: vgprValuC+11*4+3] //

v_mmac_f32_16x16x16_f16 v[vgprValuC+12*4+0:vgprValuC+12*4+1:vgprValuC+12*4+2:vgprValuC+12*4+3] v[vgprValuA_X3_I0+4*2+0:vgprValuA_X3_I0+4*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+12*4+0: vgprValuC+12*4+1: vgprValuC+12*4+2: vgprValuC+12*4+3] // 
v_mmac_f32_16x16x16_f16 v[vgprValuC+13*4+0:vgprValuC+13*4+1:vgprValuC+13*4+2:vgprValuC+13*4+3] v[vgprValuA_X3_I0+5*2+0:vgprValuA_X3_I0+5*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+13*4+0: vgprValuC+13*4+1: vgprValuC+13*4+2: vgprValuC+13*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+14*4+0:vgprValuC+14*4+1:vgprValuC+14*4+2:vgprValuC+14*4+3] v[vgprValuA_X3_I0+6*2+0:vgprValuA_X3_I0+6*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+14*4+0: vgprValuC+14*4+1: vgprValuC+14*4+2: vgprValuC+14*4+3] //
v_mmac_f32_16x16x16_f16 v[vgprValuC+15*4+0:vgprValuC+15*4+1:vgprValuC+15*4+2:vgprValuC+15*4+3] v[vgprValuA_X3_I0+7*2+0:vgprValuA_X3_I0+7*2+1] v[vgprValuB_X1_I0+1*2+0:vgprValuB_X1_I0+1*2+1] v[vgprValuC+15*4+0: vgprValuC+15*4+1: vgprValuC+15*4+2: vgprValuC+15*4+3] //
s_setprio 0
.endm


.set MT0, 128
.set MT1, 32
.set LDS_B_OFFSET, 8192
.set LDS_BLK_OFFSET, 10240
.set LDS_BLK_OFFSET_64Kmasked, 10240
.set LOG2BPE, 1
.set BPE, 2
.set DEPTHU, 32
.set LOG2DEPTHU, 5
.set PFTLOOPS, 6
.set GLWAVES, 4
.set LOG2GLWAVES, 2
.set MperWAVE, 32
.set NperWAVE, 32

/* 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_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], 0xf000


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
.endif


/******************************************/
/* Generate Global A parameters ...       */
/******************************************/
.set COALESCE_THREAD_A, 16     //x4 load
.set LOG2_COALESCE_THREAD_A, 4     //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, 16, v0                   
s_mul_i32 s[sgprTemp0], s[sgprStridesA], GLWAVES*BPE                               
v_mul_lo_u32 v1, v1, s[sgprTemp0]              
v_add_u32 v[vgprGlobalReadOffsetA], v0, v1
s_mul_i32 s[sgprTemp1], s[sgprStridesA], BPE
s_mul_i32 s[sgprTemp1], s[sgprTemp1], s[sgprGlWaveID]
v_add_u32 v[vgprGlobalReadOffsetA], v[vgprGlobalReadOffsetA], s[sgprTemp1]
s_mul_i32 s[sgprTemp0], s[sgprStridesA], 32
v_add_u32 v[vgprGlobalReadOffsetA+1], v[vgprGlobalReadOffsetA+0], s[sgprTemp0]

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], 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_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[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_cmp_eq_u64 s[sgprAddressA:sgprAddressA+1], 0 // s[sgprAddressA] == 0 ?
s_cbranch_scc1 label_SkipMmac


s_and_b32 s[sgprTemp1], s[sgprSizesFree+0], 1          // 
s_cmp_eq_u32 s[sgprTemp1], 0                     // 
s_cbranch_scc1 label_skiPadA
s_add_u32 s[sgprShadowLimitA+0], s[sgprShadowLimitA+0], 2 // extend limit for pre-pad
s_addc_u32 s[sgprShadowLimitA+1], s[sgprShadowLimitA+1], 0 // extend limit for pre-pad
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
label_skiPadA:

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

/******************************************/
/* Generate Global B parameters ...       */
/******************************************/
.set COALESCE_THREAD_B, 8     //x4 load
.set LOG2_COALESCE_THREAD_B, 3     //x4 load

v_and_b32 v[vgprTemp1], 255, v[vgprSerial] 
v_lshrrev_b32 v[vgprTemp1], 6, v[vgprTemp1]
v_lshlrev_b32 v[vgprTemp1], 2, v[vgprTemp1] 
v_and_b32 v[vgprTemp0], 63, v[vgprSerial] 
v_lshrrev_b32 v[vgprTemp0], 5, v[vgprTemp0]
v_lshlrev_b32 v[vgprTemp0], 4, v[vgprTemp0]  
v_add_u32 v[vgprTemp0], v[vgprTemp0], v[vgprTemp1]              
v_and_b32 v[vgprTemp1], 31, v[vgprSerial]     
v_lshrrev_b32 v[vgprTemp1], 3, v[vgprTemp1]  
v_add_u32 v[vgprTemp0], v[vgprTemp0], v[vgprTemp1]  
v_and_b32 v[vgprTemp1], 7, v[vgprSerial]    
v_lshlrev_b32 v[vgprTemp1], 2, v[vgprTemp1]    
s_lshr_b32 s[sgprTemp1], s[sgprStridesB], 0xd
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup1], MT1 // L624
s_sub_u32 s[sgprTemp0], s[sgprSizesFree+1], s[sgprTemp0]
s_sub_u32 s[sgprTemp0], s[sgprTemp0], 1 
v_mov_b32 v[vgprTemp3], s[sgprTemp0]                                 // 
v_min_i32 v[vgprTemp0], v[vgprTemp0], v[vgprTemp3]
v_lshlrev_b32 v[vgprGlobalReadOffsetB+0], s[sgprTemp1], v[vgprTemp0] // 
v_mov_b32 v[vgprGlobalReadOffsetB+3], v[vgprTemp1]
v_lshlrev_b32 v[vgprGlobalReadOffsetB+1], 0x1, v[vgprTemp1]  // offset *= bytes/element
s_lshr_b32 s[sgprStrideStruct], s[sgprStridesB], 0xd
s_cmp_gt_u32 s[sgprSizesSum+0], 512
s_cmov_b32 s[sgprStructNum], 512
s_cmov_b32 s[sgprStructBit], 26

s_cmp_gt_u32 s[sgprSizesSum+0], 1024
s_cmov_b32 s[sgprStructNum], 1024
s_cmov_b32 s[sgprStructBit], 27

s_cmp_gt_u32 s[sgprSizesSum+0], 2048
s_cmov_b32 s[sgprStructNum], 2048
s_cmov_b32 s[sgprStructBit], 28

s_cmp_gt_u32 s[sgprSizesSum+0], 4096
s_cmov_b32 s[sgprStructNum], 4096
s_cmov_b32 s[sgprStructBit], 29

s_mov_b32 s[sgprTemp7], 0
s_cmp_eq_u32 s[sgprSizesSum+0], 4096
s_cmov_b32 s[sgprTemp7], 1
s_cmp_eq_u32 s[sgprSizesSum+0], 2048
s_cmov_b32 s[sgprTemp7], 1
s_cmp_eq_u32 s[sgprSizesSum+0], 1024
s_cmov_b32 s[sgprTemp7], 1
s_cmp_eq_u32 s[sgprSizesSum+0], 512
s_cmov_b32 s[sgprTemp7], 1
s_cmp_eq_u32 s[sgprTemp7], 1
s_cbranch_scc1 Skip_K4096
s_cmp_le_u32 s[sgprSizesSum+0], 512
s_cbranch_scc1 Skip_K_Sub2048
s_sub_u32 s[sgprStrideStruct], s[sgprStridesB], s[sgprStructNum]
Skip_K_Sub2048:
Skip_K4096:

s_cmp_eq_u32 s[sgprTemp7], 1
s_cbranch_scc1 Skip2_K4096
s_cmp_le_u32 s[sgprSizesSum+0], 512
s_cbranch_scc1 Skip2_K_Sub2048

v_mul_lo_u32 v[vgprTemp2], s[sgprStrideStruct], v[vgprTemp0]
v_mov_b32 v[vgprGlobalReadOffsetB+0], v[vgprTemp0]
v_add_u32 v[vgprGlobalReadOffsetB+1], v[vgprTemp1], v[vgprTemp2]
v_lshlrev_b32 v[vgprGlobalReadOffsetB+1], 0x1, v[vgprGlobalReadOffsetB+1]
v_mov_b32 v[vgprGlobalReadOffsetB+3], v[vgprTemp1]
Skip2_K_Sub2048:
Skip2_K4096:
s_cmp_eq_u32 s[sgprTemp7], 1
s_cmov_b32 s[sgprStrideStruct], 0

s_mul_i32    s[sgprTensor2dSizeB+0], s[sgprSizesFree+1], s[sgprStridesB]
s_mul_hi_u32 s[sgprTensor2dSizeB+1], s[sgprSizesFree+1], s[sgprStridesB]
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup1], MT1 
s_mul_i32 s[sgprTemp0], s[sgprTemp0], s[sgprStridesB]

s_sub_u32 s[sgprShadowLimitB+0], s[sgprTensor2dSizeB], s[sgprTemp0] 
s_subb_u32 s[sgprShadowLimitB+1], s[sgprTensor2dSizeB+1], 0 
// Set limit to use bytes fp16 = 0
s_lshl_b64 s[sgprShadowLimitB:sgprShadowLimitB+1], s[sgprShadowLimitB:sgprShadowLimitB+1], LOG2BPE 
s_cmp_eq_u32 s[sgprShadowLimitB+1], 0                           // are we within 2^32?
s_cselect_b32 s[sgprSrdB+2], s[sgprShadowLimitB+0], BufferLimit // Move shadow to real if we are within 2^32
s_mul_hi_u32 s[sgprTemp3], s[sgprStridesB+1], s[sgprWorkGroup2] // Stride*WG
s_mul_i32 s[sgprTemp2], s[sgprStridesB+1], s[sgprWorkGroup2]  // Stride*WG
s_add_u32 s[sgprTemp2], s[sgprTemp2], s[sgprTemp0]                            // accum wg term to tilestart
s_addc_u32 s[sgprTemp3], s[sgprTemp3], 0                           // accum wg term to tilestart
s_lshl_b64 s[sgprTemp2:sgprTemp3], s[sgprTemp2:sgprTemp3], 0x1   
s_add_u32 s[sgprSrdB+0], s[sgprAddressB+0], s[sgprTemp2]    // SRD base = Address+ tileStart0
s_addc_u32 s[sgprSrdB+1], s[sgprAddressB+1], s[sgprTemp3]   //
s_mov_b32 s[sgprSrdB+3], Srd127_96                 // Set bits 127_96 in SRD

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

s_lshr_b32 s[sgprTemp0], s[sgprStridesB], 13 //  
s_mul_i32 s[sgprTemp1], s[sgprStridesB], 2                 //  
s_lshr_b32 s[sgprTemp2], s[sgprTemp1], s[sgprTemp0]           //  
s_lshl_b32 s[sgprTemp0], s[sgprTemp2], 16                            //  
s_or_b32 s[sgprTemp1], s[sgprTemp0], 0x40000000                      // 

s_cmp_eq_u32 s[sgprTemp7], 1                     //  
s_cbranch_scc1 Skip1_K4096                         // 
s_cmp_le_u32 s[sgprSizesSum+0], 512                // 
s_cbranch_scc1 Skip1_K_Sub2048                     // 
s_mov_b32 s[sgprStructNum], 1                      //  
s_lshl_b32 s[sgprStructNum], s[sgprStructNum], s[sgprStructBit] // 
s_or_b32 s[sgprStructNum], s[sgprStructNum], 0x40000000 // 
s_mov_b32 s[sgprTemp1], s[sgprStructNum]                    //  
Skip1_K_Sub2048:                                   // 
Skip1_K4096: 

s_or_b32 s[sgprSrdB+1], s[sgprSrdB+1], s[sgprTemp1]         // struct buffer,

//Struct index Limit
s_and_b32 s[sgprTemp0], MT1-1, s[sgprSizesFree+1]
s_cmp_eq_u32 s[sgprTemp0], 0
s_cselect_b32 s[sgprTemp0], MT1, s[sgprTemp0]
s_sub_u32 s[sgprTemp1], s[sgprNumWorkGroups1], 1
s_cmp_eq_u32 s[sgprWorkGroup1], s[sgprTemp1]
s_cselect_b32 s[sgprTemp0], s[sgprTemp0], MT1
s_mov_b32 s[sgprSrdB+2], s[sgprTemp0]

s_mov_b32 s[sgprGlobalReadIncsB+0], DEPTHU*BPE

/******************************************/
/* Generate LDS A parameters ...          */
/******************************************/
.set WAVE_LDS_OFFSET_A, 64*16     //x4 load
.set WAVE_LDS_OFFSET, WAVE_LDS_OFFSET_A     //x4 load
.set LDS_SUB_M_OFFSET, BPE*32
.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]

//get lds read addrA
v_and_b32 v[vgprTemp0], v[vgprSerial], 63
v_and_b32 v0, 7, v[vgprTemp0]    
v_lshlrev_b32 v0, 4, v0
v_lshrrev_b32 v1, 2, v[vgprTemp0]

v_and_b32 v4, LOADxWAVES_K_A-1, v1
v_lshrrev_b32 v2, LOG2GLWAVES, v4
v_and_b32     v3, GLWAVES-1, v4
v_mul_u32_u24 v6, WAVE_LDS_OFFSET_A+0, v3                        //lds WaveOffset + pad
v_mul_u32_u24 v2, MT0*BPE, v2 
v_add_u32     v4, v2, v6
v_add_u32 v[vgprLocalReadAddrA], v4, v0
s_mul_i32 s[sgprTemp0], s[sgprWaveID], LDS_SUB_M_OFFSET
v_add_u32 v[vgprLocalReadAddrA], v[vgprLocalReadAddrA], s[sgprTemp0]
v_and_b32     v3, v3, 3
v_mul_u32_u24 v[vgprTemp0], 64, v3  
v_mov_b32 v[vgprTemp1], v6
v_add_u32 v[vgprTemp1], WAVE_LDS_OFFSET_A+0, v[vgprTemp1]                   //lds A/B offset
v_mov_b32 v[vgprTemp2], WAVE_LDS_OFFSET_A

v_add_u32 v[vgprLocalReadAddrA+0], 0, 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
ADDR_WRAP  v[vgprLocalReadAddrA+0]
.set WAVE_LDS_OFFSET, UNDEF     //x4 load
.set WAVE_LDS_OFFSET, UNDEF     //x4 load

/******************************************/
/* Generate LDS B parameters ...          */
/******************************************/
.set WAVE_LDS_OFFSET_B, 64*8     //x4 load
.set WAVE_LDS_OFFSET, WAVE_LDS_OFFSET_B     //x4 load
.set LDS_SUB_N_OFFSET, BPE*32
.set LOADxWAVES_K_B, 64/COALESCE_THREAD_B*GLWAVES
.set LOADxWAVES_K_B_LOG2, 5
.set LOADxWAVES_LDS_OFFSET_B, WAVE_LDS_OFFSET_B*GLWAVES

//Wrap Lds
s_mul_i32 s[sgprLocalWriteAddrB], s[sgprGlWaveID], WAVE_LDS_OFFSET_B+0                    // can add lds pad
s_and_b32 s[sgprTemp0], s[sgprGlWaveID], 3
//s_mul_i32 s[sgprTemp0], s[sgprTemp0], 2
s_lshl_b32 s[sgprTemp0], s[sgprTemp0], 16
s_or_b32 s[sgprLocalWriteAddrB], s[sgprLocalWriteAddrB], s[sgprTemp0]
s_add_u32 s[sgprLocalWriteAddrB], s[sgprLocalWriteAddrB], LDS_B_OFFSET
s_mov_b32 s[sgprLocalWriteAddrBori], s[sgprLocalWriteAddrB]

v_and_b32 v[vgprTemp0], v[vgprSerial], 63
v_and_b32 v[vgprTemp1], v[vgprTemp0], 3
v_lshrrev_b32 v[vgprTemp2], 4, v[vgprTemp0] 
v_mul_u32_u24 v[vgprTemp1], 64, v[vgprTemp1] 
v_lshlrev_b32 v[sgprTemp2], 3, v[vgprTemp2]
v_and_b32 v[vgprTemp0], v[vgprTemp0], 15
v_lshrrev_b32 v[vgprTemp0], 2, v[vgprTemp0]
s_mov_b32 s[sgprTemp1], 0x210 // L1477
v_mul_lo_u32 v[vgprTemp0], s[sgprTemp1], v[vgprTemp0]   
v_add_u32 v[vgprTemp1], v[vgprTemp1], v[vgprTemp0] 
v_add_u32 v[vgprLocalReadAddrB], v[vgprTemp1], v[sgprTemp2]
s_mov_b32 s[sgprTemp1], LDS_B_OFFSET
v_add_u32 v[vgprLocalReadAddrB], v[vgprLocalReadAddrB], s[sgprTemp1]
v_mov_b32 v[vgprLocalReadAddrB_ori], v[vgprLocalReadAddrB]
v_add_u32 v[vgprLocalReadAddrB+1], 0x100, v[vgprLocalReadAddrB]
v_add_u32 v[vgprLocalReadAddrB+2], 0x120, v[vgprLocalReadAddrB]

v_and_b32 v0, 63, v[vgprSerial]                    // 
v_and_b32 v1, 15, v0                               // 
v_lshrrev_b32 v2, 2, v1                            // 
v_lshlrev_b32 v3, 9, v2                           // 
v_add_u32 v3, LDS_B_OFFSET, v3                             // 
v_add_u32 v4, 512, v3                             // 
v_cmp_ge_u32  s[80:81], v[vgprLocalReadAddrB+1], v4
v_mov_b32 v5, 512                                 // 
v_cndmask_b32  v5, 0, v5, s[80:81]
v_sub_u32 v[vgprLocalReadAddrB+1], v[vgprLocalReadAddrB+1], v5 // 

v_add_u32 v4, 512, v3                             // 
v_cmp_ge_u32  s[80:81], v[vgprLocalReadAddrB+2], v4
v_mov_b32 v5, 512                                 // 
v_cndmask_b32  v5, 0, v5, s[80:81]
v_sub_u32 v[vgprLocalReadAddrB+2], v[vgprLocalReadAddrB+2], v5 // 

/******************************************/
/* Keep Sgpr Values for use later ...     */
/******************************************/

//store sgprs to keep value
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdA+0], laneSrdA0
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdA+1], laneSrdA1
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdA+2], laneSrdA2
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdB+0], laneSrdB0
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdB+1], laneSrdB1
v_writelane_b32 v[vgprKeepSgprValue], s[sgprSrdB+2], laneSrdB2

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

v_and_b32 v[vgprGlobalReadOffsetZero], v[vgprSerial], 15   
v_lshlrev_b32 v[vgprGlobalReadOffsetZero], 0x1, v[vgprGlobalReadOffsetZero]
s_and_b32 s[sgprTemp0], s[sgprWaveID], 3
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 64
v_add_u32 v[vgprGlobalReadOffsetZero], s[sgprTemp0], v[vgprGlobalReadOffsetZero]

s_mul_hi_u32 s[sgprTemp1], s[sgprWorkGroup0], MT0                       // WorkGroup[01] * MT
s_mul_i32 s[sgprTemp0], s[sgprWorkGroup0], MT0                          // WorkGroup[01] * MT
s_mul_hi_u32 s[sgprTemp1], s[sgprTemp1], 8              // tlu=0, scaled tile-offset by stride
s_mul_i32 s[sgprTemp0], s[sgprTemp0], 8                 // tlu=0, scaled tile-offset by stride

s_lshr_b64 s[sgprTemp2:sgprTemp3], s[sgprShadowLimitA:sgprShadowLimitA+1], 4
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_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_lshr_b64 s[sgprTemp2:sgprTemp3], s[sgprTemp0:sgprTemp1], 2
s_add_u32 s[sgprZero+0], s[sgprZeroAddress+0], s[sgprTemp2]        // SRD base = Address+ tileStart0
s_addc_u32 s[sgprZero+1], s[sgprZeroAddress+1], s[sgprTemp3]       // 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_dwordx4 v[vgprGlobalReadOffsetA+0], s[sgprSrdA:sgprSrdA+3], 0 offen offset:0, lds
s_add_u32 m0, m0, WAVE_LDS_OFFSET_A*4  
buffer_load_dwordx4 v[vgprGlobalReadOffsetA+1], s[sgprSrdA:sgprSrdA+3], 0 offen offset:0, lds

s_add_u32 m0, s[sgprLocalWriteAddrB], \offset
buffer_load_dwordx2 v[vgprGlobalReadOffsetB:vgprGlobalReadOffsetB+1], s[sgprSrdB:sgprSrdB+3], 0, idxen offen offset:0,  lds
.endm

.macro GLOBAL_LOAD_Scale_Zero
buffer_load_short_d16  v[vgprValuZeros+0], v[vgprGlobalReadOffsetZero+0], s[sgprZero:sgprZero+3], 0 offen offset:0
buffer_load_short_d16  v[vgprValuZeros+1], v[vgprGlobalReadOffsetZero+0], s[sgprZero:sgprZero+3], 0 offen offset:32

buffer_load_dwordx2  v[vgprValuScales+0:vgprValuScales+1], v[vgprGlobalReadOffsetScale+0], s[sgprScale:sgprScale+3], 0 offen offset:0
buffer_load_dwordx2  v[vgprValuScales+2:vgprValuScales+3], v[vgprGlobalReadOffsetScale+0], s[sgprScale:sgprScale+3], 0 offen offset:128
.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_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], s[sgprTemp1]
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_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_lshr_b64 s[sgprTemp2:sgprTemp3], s[sgprShadowLimitA:sgprShadowLimitA+1], 6
s_sub_u32 s[sgprTemp2], s[sgprTemp2],  s[sgprTemp0]
s_subb_u32 s[sgprTemp3], s[sgprTemp3], s[sgprTemp1]
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

.endm


.macro I32ToF16 vgprIn:req vgprZero:req vgprScale:req
v_sub_i32 v[\vgprIn], v[\vgprIn], v[\vgprZero]
v_cvt_f32_i32 v[\vgprIn], v[\vgprIn]
v_cvt_f16_f32 v[\vgprIn], v[\vgprIn]
v_mul_f16 v[\vgprIn], v[\vgprIn], v[\vgprScale]
.endm

.macro UnPackB32ToTwoF16 vgprScale:req vgprOut:req
v_lshlrev_b32 v[\vgprOut+4], 16, v[\vgprScale+2]
v_lshrrev_b32 v[\vgprOut+4], 16, v[\vgprOut+4]
v_lshrrev_b32 v[\vgprOut+5], 16, v[\vgprScale+2]

v_lshlrev_b32 v[\vgprOut+6], 16, v[\vgprScale+3]
v_lshrrev_b32 v[\vgprOut+6], 16, v[\vgprOut+6]
v_lshrrev_b32 v[\vgprOut+7], 16, v[\vgprScale+3]

v_lshlrev_b32 v[\vgprOut+2], 16, v[\vgprScale+1]
v_lshrrev_b32 v[\vgprOut+2], 16, v[\vgprOut+2]
v_lshrrev_b32 v[\vgprOut+3], 16, v[\vgprScale+1]

v_lshlrev_b32 v[\vgprOut+0], 16, v[\vgprScale+0]
v_lshrrev_b32 v[\vgprOut+0], 16, v[\vgprOut+0]
v_lshrrev_b32 v[\vgprOut+1], 16, v[\vgprScale+0]
.endm

.macro UnPackB32To8B4 vgprZero:req vgprOut:req
v_lshrrev_b32 v[\vgprOut+3], 12, v[\vgprZero]
v_lshrrev_b32 v[\vgprOut+2], 8, v[\vgprZero]
v_lshrrev_b32 v[\vgprOut+1], 4, v[\vgprZero]
v_lshrrev_b32 v[\vgprOut+0], 0, v[\vgprZero]

v_and_b32 v[\vgprOut+0], v[\vgprOut+0], 0xf
v_and_b32 v[\vgprOut+1], v[\vgprOut+1], 0xf
v_and_b32 v[\vgprOut+2], v[\vgprOut+2], 0xf
v_and_b32 v[\vgprOut+3], v[\vgprOut+3], 0xf
.endm

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

v_lshrrev_b32 v[vgprValuA_X0_H0+7], 28, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+6], 24, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+5], 20, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+4], 16, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+3], 12, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+2], 8, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+1], 4, v[\vgprIn]
v_lshrrev_b32 v[vgprValuA_X0_H0+0], 0, v[\vgprIn]
v_and_b32 v[vgprValuA_X0_H0+0], v[vgprValuA_X0_H0+0], 0xf
v_and_b32 v[vgprValuA_X0_H0+1], v[vgprValuA_X0_H0+1], 0xf
v_and_b32 v[vgprValuA_X0_H0+2], v[vgprValuA_X0_H0+2], 0xf
v_and_b32 v[vgprValuA_X0_H0+3], v[vgprValuA_X0_H0+3], 0xf
v_and_b32 v[vgprValuA_X0_H0+4], v[vgprValuA_X0_H0+4], 0xf
v_and_b32 v[vgprValuA_X0_H0+5], v[vgprValuA_X0_H0+5], 0xf
v_and_b32 v[vgprValuA_X0_H0+6], v[vgprValuA_X0_H0+6], 0xf
v_and_b32 v[vgprValuA_X0_H0+7], v[vgprValuA_X0_H0+7], 0xf

I32ToF16 vgprValuA_X0_H0+0 \vgprZero+0 \vgprScale+0
I32ToF16 vgprValuA_X0_H0+1 \vgprZero+1 \vgprScale+1
I32ToF16 vgprValuA_X0_H0+2 \vgprZero+2 \vgprScale+2
I32ToF16 vgprValuA_X0_H0+3 \vgprZero+3 \vgprScale+3
I32ToF16 vgprValuA_X0_H0+4 \vgprZero+0 \vgprScale+0
I32ToF16 vgprValuA_X0_H0+5 \vgprZero+1 \vgprScale+1
I32ToF16 vgprValuA_X0_H0+6 \vgprZero+2 \vgprScale+2
I32ToF16 vgprValuA_X0_H0+7 \vgprZero+3 \vgprScale+3

v_pack_b32_f16 v[\vgprPack+0], v[vgprValuA_X0_H0+0], v[vgprValuA_X0_H0+4]
v_pack_b32_f16 v[\vgprPack+2], v[vgprValuA_X0_H0+1], v[vgprValuA_X0_H0+5]
v_pack_b32_f16 v[\vgprPack+4], v[vgprValuA_X0_H0+2], v[vgprValuA_X0_H0+6]
v_pack_b32_f16 v[\vgprPack+6], v[vgprValuA_X0_H0+3], v[vgprValuA_X0_H0+7]

v_lshrrev_b32 v[vgprValuA_X0_H0+15], 28, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+14], 24, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+13], 20, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+12], 16, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+11], 12, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+10], 8, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+9], 4, v[\vgprIn+1]
v_lshrrev_b32 v[vgprValuA_X0_H0+8], 0, v[\vgprIn+1]

v_and_b32 v[vgprValuA_X0_H0+8], v[vgprValuA_X0_H0+8], 0xf
v_and_b32 v[vgprValuA_X0_H0+9], v[vgprValuA_X0_H0+9], 0xf
v_and_b32 v[vgprValuA_X0_H0+10], v[vgprValuA_X0_H0+10], 0xf
v_and_b32 v[vgprValuA_X0_H0+11], v[vgprValuA_X0_H0+11], 0xf
v_and_b32 v[vgprValuA_X0_H0+12], v[vgprValuA_X0_H0+12], 0xf
v_and_b32 v[vgprValuA_X0_H0+13], v[vgprValuA_X0_H0+13], 0xf
v_and_b32 v[vgprValuA_X0_H0+14], v[vgprValuA_X0_H0+14], 0xf
v_and_b32 v[vgprValuA_X0_H0+15], v[vgprValuA_X0_H0+15], 0xf

I32ToF16 vgprValuA_X0_H0+8 \vgprZero+0 \vgprScale+0
I32ToF16 vgprValuA_X0_H0+9 \vgprZero+1 \vgprScale+1
I32ToF16 vgprValuA_X0_H0+10 \vgprZero+2 \vgprScale+2
I32ToF16 vgprValuA_X0_H0+11 \vgprZero+3 \vgprScale+3
I32ToF16 vgprValuA_X0_H0+12 \vgprZero+0 \vgprScale+0
I32ToF16 vgprValuA_X0_H0+13 \vgprZero+1 \vgprScale+1
I32ToF16 vgprValuA_X0_H0+14 \vgprZero+2 \vgprScale+2
I32ToF16 vgprValuA_X0_H0+15 \vgprZero+3 \vgprScale+3

v_pack_b32_f16 v[\vgprPack+1], v[vgprValuA_X0_H0+8], v[vgprValuA_X0_H0+12]
v_pack_b32_f16 v[\vgprPack+3], v[vgprValuA_X0_H0+9], v[vgprValuA_X0_H0+13]
v_pack_b32_f16 v[\vgprPack+5], v[vgprValuA_X0_H0+10], v[vgprValuA_X0_H0+14]
v_pack_b32_f16 v[\vgprPack+7], v[vgprValuA_X0_H0+11], v[vgprValuA_X0_H0+15]

v_lshrrev_b32 v[vgprValuA_X0_H0+23], 28, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+22], 24, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+21], 20, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+20], 16, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+19], 12, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+18], 8, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+17], 4, v[\vgprIn+2]
v_lshrrev_b32 v[vgprValuA_X0_H0+16], 0, v[\vgprIn+2]

v_and_b32 v[vgprValuA_X0_H0+16], v[vgprValuA_X0_H0+16], 0xf
v_and_b32 v[vgprValuA_X0_H0+17], v[vgprValuA_X0_H0+17], 0xf
v_and_b32 v[vgprValuA_X0_H0+18], v[vgprValuA_X0_H0+18], 0xf
v_and_b32 v[vgprValuA_X0_H0+19], v[vgprValuA_X0_H0+19], 0xf
v_and_b32 v[vgprValuA_X0_H0+20], v[vgprValuA_X0_H0+20], 0xf
v_and_b32 v[vgprValuA_X0_H0+21], v[vgprValuA_X0_H0+21], 0xf
v_and_b32 v[vgprValuA_X0_H0+22], v[vgprValuA_X0_H0+22], 0xf
v_and_b32 v[vgprValuA_X0_H0+23], v[vgprValuA_X0_H0+23], 0xf

I32ToF16 vgprValuA_X0_H0+16 \vgprZero+4 \vgprScale+4
I32ToF16 vgprValuA_X0_H0+17 \vgprZero+5 \vgprScale+5
I32ToF16 vgprValuA_X0_H0+18 \vgprZero+6 \vgprScale+6
I32ToF16 vgprValuA_X0_H0+19 \vgprZero+7 \vgprScale+7
I32ToF16 vgprValuA_X0_H0+20 \vgprZero+4 \vgprScale+4
I32ToF16 vgprValuA_X0_H0+21 \vgprZero+5 \vgprScale+5
I32ToF16 vgprValuA_X0_H0+22 \vgprZero+6 \vgprScale+6
I32ToF16 vgprValuA_X0_H0+23 \vgprZero+7 \vgprScale+7

v_pack_b32_f16 v[\vgprPack+8], v[vgprValuA_X0_H0+16], v[vgprValuA_X0_H0+20]
v_pack_b32_f16 v[\vgprPack+10], v[vgprValuA_X0_H0+17], v[vgprValuA_X0_H0+21]
v_pack_b32_f16 v[\vgprPack+12], v[vgprValuA_X0_H0+18], v[vgprValuA_X0_H0+22]
v_pack_b32_f16 v[\vgprPack+14], v[vgprValuA_X0_H0+19], v[vgprValuA_X0_H0+23]

v_lshrrev_b32 v[vgprValuA_X0_H0+31], 28, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+30], 24, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+29], 20, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+28], 16, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+27], 12, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+26], 8, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+25], 4, v[\vgprIn+3]
v_lshrrev_b32 v[vgprValuA_X0_H0+24], 0, v[\vgprIn+3]

v_and_b32 v[vgprValuA_X0_H0+24], v[vgprValuA_X0_H0+24], 0xf
v_and_b32 v[vgprValuA_X0_H0+25], v[vgprValuA_X0_H0+25], 0xf
v_and_b32 v[vgprValuA_X0_H0+26], v[vgprValuA_X0_H0+26], 0xf
v_and_b32 v[vgprValuA_X0_H0+27], v[vgprValuA_X0_H0+27], 0xf
v_and_b32 v[vgprValuA_X0_H0+28], v[vgprValuA_X0_H0+28], 0xf
v_and_b32 v[vgprValuA_X0_H0+29], v[vgprValuA_X0_H0+29], 0xf
v_and_b32 v[vgprValuA_X0_H0+30], v[vgprValuA_X0_H0+30], 0xf
v_and_b32 v[vgprValuA_X0_H0+31], v[vgprValuA_X0_H0+31], 0xf

I32ToF16 vgprValuA_X0_H0+24 \vgprZero+4 \vgprScale+4
I32ToF16 vgprValuA_X0_H0+25 \vgprZero+5 \vgprScale+5
I32ToF16 vgprValuA_X0_H0+26 \vgprZero+6 \vgprScale+6
I32ToF16 vgprValuA_X0_H0+27 \vgprZero+7 \vgprScale+7
I32ToF16 vgprValuA_X0_H0+28 \vgprZero+4 \vgprScale+4
I32ToF16 vgprValuA_X0_H0+29 \vgprZero+5 \vgprScale+5
I32ToF16 vgprValuA_X0_H0+30 \vgprZero+6 \vgprScale+6
I32ToF16 vgprValuA_X0_H0+31 \vgprZero+7 \vgprScale+7

v_pack_b32_f16 v[\vgprPack+9], v[vgprValuA_X0_H0+24], v[vgprValuA_X0_H0+28]
v_pack_b32_f16 v[\vgprPack+11], v[vgprValuA_X0_H0+25], v[vgprValuA_X0_H0+29]
v_pack_b32_f16 v[\vgprPack+13], v[vgprValuA_X0_H0+26], v[vgprValuA_X0_H0+30]
v_pack_b32_f16 v[\vgprPack+15], v[vgprValuA_X0_H0+27], v[vgprValuA_X0_H0+31]

.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

s_mov_b32 s[sgprTemp0], s[sgprGlobalReadIncsB+0]
s_mov_b32 s[sgprTemp1], 0
s_add_u32 s[sgprSrdB+0], s[sgprSrdB+0],  s[sgprTemp0]
s_addc_u32 s[sgprSrdB+1], s[sgprSrdB+1], s[sgprTemp1]
.endm

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

.macro LDS_LOADAB off:req

ds_read_m32x16_b16 v[vgprValuA_X0_I0+ 0:vgprValuA_X0_I0+ 3], v[vgprLocalReadAddrA+0] offset:\off 
ds_read_b64 v[vgprValuB_X0_I0+ 0:vgprValuB_X0_I0+ 1], v[vgprLocalReadAddrB] offset:\off+0
ds_read_b64 v[vgprValuB_X0_I0+ 2:vgprValuB_X0_I0+ 3], v[vgprLocalReadAddrB+1] offset:\off+0
.endm

.macro LDS_LOADAB1 off:req

ds_read_m32x16_b16 v[vgprValuA_X1_I0+ 0:vgprValuA_X1_I0+ 3], v[vgprLocalReadAddrA+0] offset:\off+4096
ds_read_b64 v[vgprValuB_X1_I0+ 0:vgprValuB_X1_I0+ 1], v[vgprLocalReadAddrB] offset:\off+32
ds_read_b64 v[vgprValuB_X1_I0+ 2:vgprValuB_X1_I0+ 3], v[vgprLocalReadAddrB+2] offset:\off+0

.endm

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

s_lshr_b32 s[sgprLoopCounterL], s[sgprSizesSum], 5
s_min_u32 s[sgprLoopCntCommon], 6, 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], 4
s_cbranch_scc1 SkipGL
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_addk_i32 s[sgprLocalWriteAddrB], LDS_BLK_OFFSET_64Kmasked      
s_add_u32 s[sgprTemp0], s[sgprLocalWriteAddrBori], s[sgprLDSMask]    
s_cmp_ge_u32 s[sgprLocalWriteAddrB], s[sgprTemp0]                           
s_cmov_b32  s[sgprLocalWriteAddrB], s[sgprLocalWriteAddrBori]                  

s_cmp_lt_i32 s[sgprTemp3], 4
s_cbranch_scc1 SkipWait
s_waitcnt vmcnt(12)
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_cmp_eq_i32 s[sgprTemp3], 2
s_cbranch_scc1 Last2
s_cmp_eq_i32 s[sgprTemp3], 3
s_cbranch_scc1 Last3
s_waitcnt vmcnt(9)
s_barrier
Last3:
s_waitcnt vmcnt(6)
s_barrier
Last2:
s_waitcnt vmcnt(3)
s_barrier
Last1:
s_waitcnt vmcnt(0)
s_barrier
s_barrier

SkipToLastLoad:
s_mov_b32  s[sgprLocalWriteAddrA], s[sgprLocalWriteAddrAori]  
s_mov_b32  s[sgprLocalWriteAddrB], s[sgprLocalWriteAddrBori]  
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], 512

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_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_and_b32 v[vgprTemp2], 15, v[vgprTemp0]                    
v_lshlrev_b32 v[vgprTemp2], 0x3, 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_mul_i32 s[sgprTemp0], s[sgprWaveID], 128
s_lshl_b32 s[sgprTemp0], s[sgprTemp0], 0x1 
v_add_u32 v[vgprGlobalWriteOffsetD], v[vgprGlobalWriteOffsetD], s[sgprTemp0]


/******************************************/
/* Init ValueC ...                        */
/******************************************/
GLOBAL_LOAD_Scale_Zero
v_mov_b32 v[vgprValuC+0], 0x0
v_mov_b32 v[vgprValuC+1], 0x0
v_mov_b32 v[vgprValuC+2], 0x0
v_mov_b32 v[vgprValuC+3], 0x0
v_mov_b32 v[vgprValuC+4], 0x0
v_mov_b32 v[vgprValuC+5], 0x0
v_mov_b32 v[vgprValuC+6], 0x0
v_mov_b32 v[vgprValuC+7], 0x0
v_mov_b32 v[vgprValuC+8], 0x0
v_mov_b32 v[vgprValuC+9], 0x0
v_mov_b32 v[vgprValuC+10], 0x0
v_mov_b32 v[vgprValuC+11], 0x0
v_mov_b32 v[vgprValuC+12], 0x0
v_mov_b32 v[vgprValuC+13], 0x0
v_mov_b32 v[vgprValuC+14], 0x0
v_mov_b32 v[vgprValuC+15], 0x0
v_mov_b32 v[vgprValuC+16], 0x0                     // initC
v_mov_b32 v[vgprValuC+17], 0x0                     // initC
v_mov_b32 v[vgprValuC+18], 0x0                     // initC
v_mov_b32 v[vgprValuC+19], 0x0                     // initC
v_mov_b32 v[vgprValuC+20], 0x0                     // initC
v_mov_b32 v[vgprValuC+21], 0x0                     // initC
v_mov_b32 v[vgprValuC+22], 0x0                     // initC
v_mov_b32 v[vgprValuC+23], 0x0                     // initC
v_mov_b32 v[vgprValuC+24], 0x0                     // initC
v_mov_b32 v[vgprValuC+25], 0x0                     // initC
v_mov_b32 v[vgprValuC+26], 0x0                     // initC
v_mov_b32 v[vgprValuC+27], 0x0                     // initC
v_mov_b32 v[vgprValuC+28], 0x0                     // initC
v_mov_b32 v[vgprValuC+29], 0x0                     // initC
v_mov_b32 v[vgprValuC+30], 0x0                     // initC
v_mov_b32 v[vgprValuC+31], 0x0                     // initC
v_mov_b32 v[vgprValuC+32], 0x0                     // initC
v_mov_b32 v[vgprValuC+33], 0x0                     // initC
v_mov_b32 v[vgprValuC+34], 0x0                     // initC
v_mov_b32 v[vgprValuC+35], 0x0                     // initC
v_mov_b32 v[vgprValuC+36], 0x0                     // initC
v_mov_b32 v[vgprValuC+37], 0x0                     // initC
v_mov_b32 v[vgprValuC+38], 0x0                     // initC
v_mov_b32 v[vgprValuC+39], 0x0                     // initC
v_mov_b32 v[vgprValuC+40], 0x0                     // initC
v_mov_b32 v[vgprValuC+41], 0x0                     // initC
v_mov_b32 v[vgprValuC+42], 0x0                     // initC
v_mov_b32 v[vgprValuC+43], 0x0                     // initC
v_mov_b32 v[vgprValuC+44], 0x0                     // initC
v_mov_b32 v[vgprValuC+45], 0x0                     // initC
v_mov_b32 v[vgprValuC+46], 0x0                     // initC
v_mov_b32 v[vgprValuC+47], 0x0                     // initC
v_mov_b32 v[vgprValuC+48], 0x0                     // initC
v_mov_b32 v[vgprValuC+49], 0x0                     // initC
v_mov_b32 v[vgprValuC+50], 0x0                     // initC
v_mov_b32 v[vgprValuC+51], 0x0                     // initC
v_mov_b32 v[vgprValuC+52], 0x0                     // initC
v_mov_b32 v[vgprValuC+53], 0x0                     // initC
v_mov_b32 v[vgprValuC+54], 0x0                     // initC
v_mov_b32 v[vgprValuC+55], 0x0                     // initC
v_mov_b32 v[vgprValuC+56], 0x0                     // initC
v_mov_b32 v[vgprValuC+57], 0x0                     // initC
v_mov_b32 v[vgprValuC+58], 0x0                     // initC
v_mov_b32 v[vgprValuC+59], 0x0                     // initC
v_mov_b32 v[vgprValuC+60], 0x0                     // initC
v_mov_b32 v[vgprValuC+61], 0x0                     // initC
v_mov_b32 v[vgprValuC+62], 0x0                     // initC
v_mov_b32 v[vgprValuC+63], 0x0                     // initC
GLOBAL_INC_Scale_Zero

/******************************************/
/* 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

UnPackB32ToTwoF16 vgprValuScales+0 vgprValuScalesF16+0
UnPackB32To8B4 vgprValuZeros+0 vgprValuZerosI32+0
UnPackB32To8B4 vgprValuZeros+1 vgprValuZerosI32+4

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

s_cmp_ge_u32 s[sgprWaveID], 4
s_cbranch_scc1 WaveID_gecase

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

LDS_LOADAB1 LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*1
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_1

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(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_waitcnt vmcnt(0)
s_barrier

GLOBAL_INC_Scale_Zero
LDS_LOADAB LDS_BLK_OFFSET*2
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
UnPackB32ToTwoF16 vgprValuScales+0 vgprValuScalesF16+0
UnPackB32To8B4 vgprValuZeros+0 vgprValuZerosI32+0
UnPackB32To8B4 vgprValuZeros+1 vgprValuZerosI32+4
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_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(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*3
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_1
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(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_waitcnt vmcnt(0)
s_barrier
GLOBAL_INC_Scale_Zero
LDS_LOADAB LDS_BLK_OFFSET*4
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
UnPackB32ToTwoF16 vgprValuScales+0 vgprValuScalesF16+0
UnPackB32To8B4 vgprValuZeros+0 vgprValuZerosI32+0
UnPackB32To8B4 vgprValuZeros+1 vgprValuZerosI32+4
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_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*4
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*5
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_1
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*5
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
s_waitcnt vmcnt(0)
s_barrier
GLOBAL_INC_Scale_Zero
LDS_LOADAB LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(3)
I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
UnPackB32ToTwoF16 vgprValuScales+0 vgprValuScalesF16+0
UnPackB32To8B4 vgprValuZeros+0 vgprValuZerosI32+0
UnPackB32To8B4 vgprValuZeros+1 vgprValuZerosI32+4
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_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
s_branch WaveID_EndSwitch
WaveID_gecase:

WaveID_EndSwitch:

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


//reload sgprs value from vgprs
v_readlane_b32 s[sgprSrdA+0], v[vgprKeepSgprValue], laneSrdA0
v_readlane_b32 s[sgprSrdA+1], v[vgprKeepSgprValue], laneSrdA1
v_readlane_b32 s[sgprSrdA+2], v[vgprKeepSgprValue], laneSrdA2
v_readlane_b32 s[sgprSrdB+0], v[vgprKeepSgprValue], laneSrdB0
v_readlane_b32 s[sgprSrdB+1], v[vgprKeepSgprValue], laneSrdB1
v_readlane_b32 s[sgprSrdB+2], v[vgprKeepSgprValue], laneSrdB2
s_mov_b32 s[sgprGlobalReadIncsB+0], DEPTHU*BPE

s_waitcnt vmcnt(0)
s_waitcnt lgkmcnt(0)
s_barrier
s_mov_b32 s[sgprLocalWriteAddrA], s[sgprLocalWriteAddrAori]
s_mov_b32 s[sgprLocalWriteAddrB], s[sgprLocalWriteAddrBori]

s_lshr_b32 s[sgprTemp0], s[sgprSizesSum], LOG2DEPTHU // 
s_and_b32 s[sgprTemp1], s[sgprSizesSum], DEPTHU-1
s_cselect_b32 s[sgprTemp1], 0, 1                     //if has tail noneed -1
s_sub_u32 s[sgprTemp0], s[sgprTemp0], s[sgprTemp1]   //if has tail noneed -1, get increase blocks

s_mul_i32    s[sgprTemp2], s[sgprTemp0], s[sgprGlobalReadIncsA]
s_mul_hi_i32 s[sgprTemp3], s[sgprTemp0], 0
s_add_u32  s[sgprSrdA+0], s[sgprSrdA+0], s[sgprTemp2]
s_addc_u32 s[sgprSrdA+1], s[sgprSrdA+1], s[sgprTemp3]
s_sub_u32  s[sgprSrdA+2], s[sgprSrdA+2], s[sgprTemp2]

s_mul_i32    s[sgprTemp2], s[sgprTemp0], s[sgprGlobalReadIncsB]
s_mul_hi_i32 s[sgprTemp3], s[sgprTemp0], 0
s_add_u32  s[sgprSrdB+0], s[sgprSrdB+0], s[sgprTemp2]
s_addc_u32 s[sgprSrdB+1], s[sgprSrdB+1], s[sgprTemp3]
//s_sub_u32  s[sgprSrdB+2], s[sgprSrdB+2], s[sgprTemp2]

s_and_b32 s[sgprLoopCounterL], s[sgprSizesSum], DEPTHU-1
s_cmp_eq_i32 s[sgprLoopCounterL], 0
s_cmov_b32 s[sgprLoopCounterL], DEPTHU             //如果没有多余k就设为整数k

s_mul_i32 s[sgprTemp3], s[sgprWorkGroup0], MT0
s_sub_u32 s[sgprTemp3], s[sgprSizesFree+0], s[sgprTemp3]
s_cmp_ge_u32 s[sgprTemp3], MT0
s_cbranch_scc1 NOT_EDGE_A

s_lshl_b32 s[sgprTemp3], s[sgprTemp3], LOG2BPE
s_lshr_b32 s[sgprTemp3], s[sgprTemp3], 2
s_lshl_b32 s[sgprTemp3], s[sgprTemp3], 2                    //计算16byte整数偏移


//fp16 情况只会出现少2个byte情况，当m>=1时只是最后一列k会出现读不进的情况
//当m=1时则会出现最后1列k读不进数据情况
//fp16 由于是每个数是2 byte所以只会出现最后一列读不到数的情况，所以只用刷新最后一列K的数据

s_mov_b32  s[sgprTemp0], 1
s_cmp_eq_i32 s[sgprSizesFree+0], 1
s_cmov_b32 s[sgprTemp0], 1            //这个条件只有fp16需要

s_sub_u32  s[sgprTemp0], s[sgprLoopCounterL], s[sgprTemp0]             //move to last K
s_mul_i32  s[sgprTemp1], s[sgprTemp0], s[sgprStridesA]
s_lshl_b32 s[sgprTemp1], s[sgprTemp1], LOG2BPE               //乘上每个点  byte数
v_and_b32  v[vgprTemp0], v[vgprSerial], 15     
v_lshlrev_b32 v[vgprTemp0], LOG2BPE, v[vgprTemp0]
v_add_u32  v[vgprTemp0], v[vgprTemp0], s[sgprTemp1]          //计算出最后一列k的global偏移

s_mul_i32  s[sgprTemp2], s[sgprTemp0], MT0
s_lshl_b32 s[sgprTemp2], s[sgprTemp2], LOG2BPE              //乘上每个点  byte数
v_and_b32  v[vgprTemp2], v[vgprSerial], 15
v_lshlrev_b32 v[vgprTemp2], LOG2BPE, v[vgprTemp2]
v_add_u32  v[vgprTemp2], v[vgprTemp2], s[sgprTemp2]         //计算出最后一列k的lds偏移

v_add_u32  v[vgprTemp0], v[vgprTemp0], s[sgprTemp3]
v_add_u32  v[vgprTemp2], v[vgprTemp2], s[sgprTemp3]         //gl 和lds都加上N方向偏移

v_and_b32  v[vgprTemp3], v[vgprSerial], 63
v_lshrrev_b32 v[vgprTemp3], 4, v[vgprTemp3]
v_lshlrev_b32 v[vgprTemp3], LOG2BPE, v[vgprTemp3]
v_mul_lo_u32  v[vgprTemp1], v[vgprTemp3], s[sgprStridesA]
v_add_u32  v[vgprTemp0], v[vgprTemp0], v[vgprTemp1]     //向后读取4k列 用limit保证
v_mul_u32_u24  v[vgprTemp1], MT0, v[vgprTemp3]
v_add_u32  v[vgprTemp2], v[vgprTemp2], v[vgprTemp1]     //向后写入4k列 用mask保证不越界


buffer_load_ushort v[vgprTemp1], v[vgprTemp0], s[sgprSrdA:sgprSrdA+3], 0 offen offset:0
v_cmp_lt_u32 s[sgprTemp0:sgprTemp1], v[vgprTemp0], s[sgprSrdA+2]    //计算超出范围的地址，作为写入lds的mask
s_waitcnt vmcnt(0)
s_mov_b64 exec, s[sgprTemp0:sgprTemp1]                      //写多了会覆盖后面的数据，特别是A会覆盖B的数据
ds_write_b16 v[vgprTemp2], v[vgprTemp1]
s_mov_b64 exec, 0xffffffffffffffff
s_waitcnt lgkmcnt(0)
NOT_EDGE_A:


s_mul_i32 s[sgprTemp3], s[sgprWorkGroup1], MT1
s_sub_u32 s[sgprTemp3], s[sgprSizesFree+1], s[sgprTemp3]
//s_cmp_ge_u32 s[sgprTemp3], MT1
s_cmp_ge_u32 s[sgprTemp3], 0
s_cbranch_scc1 NOT_EDGE_B

s_lshl_b32 s[sgprTemp3], s[sgprTemp3], LOG2BPE
s_lshr_b32 s[sgprTemp3], s[sgprTemp3], 4
s_lshl_b32 s[sgprTemp3], s[sgprTemp3], 4                    //计算16byte整数偏移

//fp16 情况只会出现少2个byte情况，当m>=1时只是最后一列k会出现读不进的情况
//当m=1时则会出现最后1列k读不进数据情况
//fp16 由于是每个数是2 byte所以只会出现最后一列读不到数的情况，所以只用刷新最后一列K的数据

s_mov_b32  s[sgprTemp0], 1
s_cmp_eq_i32 s[sgprSizesFree+1], 1
s_cmov_b32 s[sgprTemp0], 1            //这个条件只有fp16需要

s_sub_u32  s[sgprTemp0], s[sgprLoopCounterL], s[sgprTemp0]              //move to last K
s_mul_i32  s[sgprTemp1], s[sgprTemp0], s[sgprStridesB]
s_lshl_b32 s[sgprTemp1], s[sgprTemp1], LOG2BPE               //乘上每个点  byte数
v_and_b32  v[vgprTemp0], v[vgprSerial], 15     
v_lshlrev_b32 v[vgprTemp0], LOG2BPE, v[vgprTemp0]
v_add_u32  v[vgprTemp0], v[vgprTemp0], s[sgprTemp1]          //计算出最后一列k的global偏移

s_mul_i32  s[sgprTemp2], s[sgprTemp0], MT1
s_lshl_b32 s[sgprTemp2], s[sgprTemp2], LOG2BPE              //乘上每个点  byte数
v_and_b32  v[vgprTemp2], v[vgprSerial], 15
v_lshlrev_b32 v[vgprTemp2], LOG2BPE, v[vgprTemp2]
v_add_u32  v[vgprTemp2], v[vgprTemp2], s[sgprTemp2]
v_add_u32  v[vgprTemp2], LDS_B_OFFSET, v[vgprTemp2]         //计算出最后一列k的lds偏移

v_add_u32  v[vgprTemp0], v[vgprTemp0], s[sgprTemp3]
v_add_u32  v[vgprTemp2], v[vgprTemp2], s[sgprTemp3]         //gl 和lds都加上N方向偏移

v_and_b32  v[vgprTemp3], v[vgprSerial], 63
v_lshrrev_b32 v[vgprTemp3], 4, v[vgprTemp3]
v_lshlrev_b32 v[vgprTemp3], LOG2BPE, v[vgprTemp3]
v_mul_lo_u32  v[vgprTemp1], v[vgprTemp3], s[sgprStridesB]
v_add_u32  v[vgprTemp0], v[vgprTemp0], v[vgprTemp1]     //向后读取4k列 用limit保证
v_mul_u32_u24  v[vgprTemp1], MT1, v[vgprTemp3]
v_add_u32  v[vgprTemp2], v[vgprTemp2], v[vgprTemp1]     //向后写入4k列 用mask保证不越界

buffer_load_ushort v[vgprTemp1], v[vgprTemp0], s[sgprSrdB:sgprSrdB+3], 0 offen offset:0
v_cmp_lt_u32 s[sgprTemp0:sgprTemp1], v[vgprTemp0], s[sgprSrdB+2]    //计算超出范围的地址，作为写入lds的mask
s_waitcnt vmcnt(0)
s_mov_b64 exec, s[sgprTemp0:sgprTemp1]                      //写多了会覆盖后面的数据，特别是A会覆盖B的数据
ds_write_b16 v[vgprTemp2], v[vgprTemp1]
s_mov_b64 exec, 0xffffffffffffffff
s_waitcnt lgkmcnt(0)

NOT_EDGE_B:

//GLOBAL_LOAD_Scale_Zero
s_waitcnt vmcnt(0)
s_barrier
LDS_LOADAB LDS_BLK_OFFSET*0

s_waitcnt lgkmcnt(0)
I4ToFp16 vgprValuA_X0_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X2_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_0
LDS_LOADAB1 LDS_BLK_OFFSET*0
s_waitcnt lgkmcnt(0)

I4ToFp16 vgprValuA_X1_I0+0 vgprValuZerosI32+0 vgprValuScalesF16+0 vgprValuA_X3_I0
.align32 8, 0xbf800001
s_nop (1)
MMAC_32x32_1

/******************************************/
/* Global Write Process ...               */
/******************************************/

label_SkipMmac:

s_lshl_b32 s[sgprSizesFree+0], s[sgprSizesFree+0], 2
s_mul_i32 s[sgprD_MEdge], s[sgprWorkGroup0], 512

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], 512*2
v_and_b32 v[vgprTemp2], v[vgprSerial], 15  
s_mov_b32 s[sgprTemp1], s[sgprWaveID]
s_mul_i32 s[sgprTemp1], s[sgprTemp1], 128
v_add_u32 v[vgprTemp2], v[vgprTemp2], s[sgprTemp1]         
v_lshlrev_b32  v[vgprTemp2], 1, v[vgprTemp2]                 
v_mov_b32 v[vgprTemp3], v[vgprTemp2]               //store inittial addr
v_mov_b32 v[vgprTemp0], v[vgprGlobalWriteOffsetD]  //store inittial addr

.set Nvoff, 0
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]

.if debug_buffer
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[vgprValuC+Nvoff+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


buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]


v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 1
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 2
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 3
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 32
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 33
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 34
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0

.set Nvoff, 35
v_mov_b32 v[vgprTemp1], v[vgprTemp0]   //v[vgprTemp1] <- v[vgprGlobalWriteOffsetD] with strideD
v_mov_b32 v[vgprTemp2], v[vgprTemp3]   //v[vgprTemp2] is just M offset without strideD for compare edge 
v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+0], s[sgprAlpha], v[vgprValuC+Nvoff+0]
v_cvt_f16_f32 v[vgprValuC+Nvoff+0], v[vgprValuC+Nvoff+0]
buffer_store_short v[vgprValuC+Nvoff+0],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+4], s[sgprAlpha], v[vgprValuC+Nvoff+4]
v_cvt_f16_f32 v[vgprValuC+Nvoff+4], v[vgprValuC+Nvoff+4]
buffer_store_short v[vgprValuC+Nvoff+4],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+8], s[sgprAlpha], v[vgprValuC+Nvoff+8]
v_cvt_f16_f32 v[vgprValuC+Nvoff+8], v[vgprValuC+Nvoff+8]
buffer_store_short v[vgprValuC+Nvoff+8],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+12], s[sgprAlpha], v[vgprValuC+Nvoff+12]
v_cvt_f16_f32 v[vgprValuC+Nvoff+12], v[vgprValuC+Nvoff+12]
buffer_store_short v[vgprValuC+Nvoff+12],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 122, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 122, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+16], s[sgprAlpha], v[vgprValuC+Nvoff+16]
v_cvt_f16_f32 v[vgprValuC+Nvoff+16], v[vgprValuC+Nvoff+16]
buffer_store_short v[vgprValuC+Nvoff+16],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+20], s[sgprAlpha], v[vgprValuC+Nvoff+20]
v_cvt_f16_f32 v[vgprValuC+Nvoff+20], v[vgprValuC+Nvoff+20]
buffer_store_short v[vgprValuC+Nvoff+20],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+24], s[sgprAlpha], v[vgprValuC+Nvoff+24]
v_cvt_f16_f32 v[vgprValuC+Nvoff+24], v[vgprValuC+Nvoff+24]
buffer_store_short v[vgprValuC+Nvoff+24],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 
v_add_u32 v[vgprTemp1], 2, v[vgprTemp1]
v_add_u32 v[vgprTemp2], 2, v[vgprTemp2]

v_cmp_ge_u32 s[sgprTemp2:sgprTemp2+1], v[vgprTemp2], s[sgprD_MEdge]
v_cndmask_b32 v[vgprGlobalWriteOffsetD], v[vgprTemp1], -1, s[sgprTemp2:sgprTemp2+1]
v_mul_f32 v[vgprValuC+Nvoff+28], s[sgprAlpha], v[vgprValuC+Nvoff+28]
v_cvt_f16_f32 v[vgprValuC+Nvoff+28], v[vgprValuC+Nvoff+28]
buffer_store_short v[vgprValuC+Nvoff+28],  v[vgprGlobalWriteOffsetD], s[sgprSrdD:sgprSrdD+3], 0 offen offset:0   // store D 

.set Nvoff, UNDEF
s_mul_i32  s[sgprTemp0], s[sgprStridesD], 2*4                
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

s_add_u32  s[sgprSrdC+0], s[sgprSrdC+0], s[sgprTemp0]        
s_addc_u32 s[sgprSrdC+1], s[sgprSrdC+1], 0        
s_subb_u32 s[sgprSrdC+2], s[sgprSrdC+2],  s[sgprTemp0]
s_cmov_b32 s[sgprSrdC+2], 0


Beta_EndSwitch:
s_endpgm
