从hip cuda kernel 的汇编语言来感受 AMD GPU内部工作方式

0, 无参数 kernel 汇编语言示例

./param_00.hip

__global__ void WWWWW()
{((int*)0x8888888)[3] = 0x77777;
}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_00.hip -o param_00.s

_Z5WWWWWv:                              ; @_Z5WWWWWv
; %bb.0:                                ; %entry
   v_mov_b32_e32 v0, 0x8888894
   v_mov_b32_e32 v1, 0
   v_mov_b32_e32 v2, 0x77777
   flat_store_dword v[0:1], v2
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5WWWWWv

1，一个参数的 kernel 的汇编语言示例

param_01.hip

__global__ void MMMMM(int* A)
{A[10] = 0x77777;}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_01.hip -o param_01.s

_Z5MMMMMPi:                             ; @_Z5MMMMMPi
; %bb.0:                                ; %entry
   s_load_dwordx2 s[0:1], s[4:5], 0x0
   v_mov_b32_e32 v0, 0
   v_mov_b32_e32 v1, 0x77777
   s_waitcnt lgkmcnt(0)
   global_store_dword v0, v1, s[0:1] offset:40
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5MMMMMPi

2, 两个参数的kernel 的汇编语言

param_02.hip

__global__ void MMMMM(int* AA, int *BB)
{AA[10] = 0x77777;BB[20] = 0x33333;}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_02.hip -o param_02.s

_Z5MMMMMPiS_:                           ; @_Z5MMMMMPiS_
; %bb.0:                                ; %entry
   s_load_dwordx4 s[0:3], s[4:5], 0x0
   v_mov_b32_e32 v0, 0
   v_mov_b32_e32 v1, 0x77777
   v_mov_b32_e32 v2, 0x33333
   s_waitcnt lgkmcnt(0)
   global_store_dword v0, v1, s[0:1] offset:40
   global_store_dword v0, v2, s[2:3] offset:80
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5MMMMMPiS_

3, 三个参数的kernel 的汇编语言

param_03.hip

__global__ void MMMMM(int* AA, int* BB, int* CC)
{AA[10] = 0x77777;BB[20] = 0x33333;CC[30] = 0x12121;}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_03.hip -o param_03.s

_Z5MMMMMPiS_S_:                         ; @_Z5MMMMMPiS_S_
; %bb.0:                                ; %entry
   s_load_dwordx4 s[0:3], s[4:5], 0x0
   v_mov_b32_e32 v0, 0
   s_load_dwordx2 s[4:5], s[4:5], 0x10
   v_mov_b32_e32 v1, 0x77777
   v_mov_b32_e32 v2, 0x33333
   s_waitcnt lgkmcnt(0)
   global_store_dword v0, v1, s[0:1] offset:40
   global_store_dword v0, v2, s[2:3] offset:80
   v_mov_b32_e32 v1, 0x12121
   global_store_dword v0, v1, s[4:5] offset:120
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5MMMMMPiS_S_

4，四个参数的kernel 的汇编语言

param_04.hip

__global__ void MMMMM(int* AA, int* BB, int* CC, int* DD)
{AA[10] = 0x77777;BB[20] = 0x33333;CC[30] = 0x22222;DD[40] = 0x44444;}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_04.hip -o param_04.s

_Z5MMMMMPiS_S_S_:                       ; @_Z5MMMMMPiS_S_S_
; %bb.0:                                ; %entry
   s_load_dwordx8 s[0:7], s[4:5], 0x0
   v_mov_b32_e32 v0, 0
   v_mov_b32_e32 v1, 0x77777
   v_mov_b32_e32 v2, 0x33333
   s_waitcnt lgkmcnt(0)
   global_store_dword v0, v1, s[0:1] offset:40
   global_store_dword v0, v2, s[2:3] offset:80
   v_mov_b32_e32 v1, 0x22222
   global_store_dword v0, v1, s[4:5] offset:120
   v_mov_b32_e32 v1, 0x44444
   global_store_dword v0, v1, s[6:7] offset:160
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5MMMMMPiS_S_S_

5，带线程索引的 kernel 的汇编语言

param_threadid.hip

#include <hip/hip_runtime.h>__global__ void MMMMM(int* AA)
{AA[threadIdx.x + 10] = 0x77777;}

../../../local_amdgpu/bin/clang++ -O1 --cuda-device-only --offload-arch=gfx906 -S ./param_threadid.hip -o param_threadid.s

_Z5MMMMMPi:                             ; @_Z5MMMMMPi
; %bb.0:                                ; %entry
   s_load_dwordx2 s[0:1], s[4:5], 0x0
   v_lshlrev_b32_e32 v0, 2, v0
   v_mov_b32_e32 v1, 0x77777
   s_waitcnt lgkmcnt(0)
   global_store_dword v0, v1, s[0:1] offset:40
   s_endpgm
   .section   .rodata,"a",@progbits
   .p2align   6, 0x0
   .amdhsa_kernel _Z5MMMMMPi