## 1 需求 在 Z100/K100 上分析用 Triton 写的 kernel 对于向量/标量寄存器的压力 ## 2 现有方法 hipprof 工具提供了寄存器压力分析的功能，可以定位到每条汇编指令对向量/标量寄存器的使用情况，使用方法如下： 对于一个用 hip 编写的 kernel 而言，先用 hipcc 编译成 elf 文件： ``` root@h01r4n18:/workspace# hipcc -o vector_add vector_add.cpp root@h01r4n18:/workspace# file vector_add vector_add: ELF 64-bit LSB executable, x86-64, version 1 (SYSV), dynamically linked, interpreter /lib64/ld-linux-x86-64.so.2, for GNU/Linux 3.2.0, not stripped root@h01r4n18:/workspace# cat vector_add.cpp #include __global__ void vector_add(const float* A, const float* B, float* C, int N) { int idx = blockDim.x * blockIdx.x + threadIdx.x; if (idx < N) C[idx] = A[idx] + B[idx]; return; } int main(){ return 0; } ``` 然后借助 hipprof 的寄存器压力分析功能，来进行分析： ``` root@h01r4n18:/workspace# hipprof --codeobj-analyze ./vector_add 1:ELF file : vector_add-hipv4-amdgcn-amd-amdhsa--gfx906.out 2:ELF file : vector_add-hipv4-amdgcn-amd-amdhsa--gfx926.out 3:ELF file : vector_add-hipv4-amdgcn-amd-amdhsa--gfx928.out 4:ELF file : vector_add-hipv4-amdgcn-amd-amdhsa--gfx936.out input and select elf file index: 1 1:kernalName: _Z10vector_addPKfS0_Pfi input and select kernel index: root@h01r4n18:/workspace# ls vector_add-hipv4-amdgcn-amd-amdhsa--gfx906.out.obj.html vector_add-hipv4-amdgcn-amd-amdhsa--gfx906.out.obj.html ``` 最后打开生成的 html 文件即可查看每条指令对于向量/标量寄存器的使用情况：  ## 3 问题 Triton 编写的 kernel 经过 Triton 编译器会生成elf文件，位于~/.triton/cache下，如下面所示的gemv_INT_revsplitK_kernel.hsaco： ```language root@h01r4n18:/workspace# tree ~/.triton/cache/ /root/.triton/cache/ ├── 1708943e4228e69a0defc2ff9233a07461d75e9aa7578c7f654ef90f2f14bc1d │   ├── gemv_INT_revsplitK_kernel.amdgcn │   ├── gemv_INT_revsplitK_kernel.hsaco │   ├── gemv_INT_revsplitK_kernel.json │   ├── gemv_INT_revsplitK_kernel.llir │   ├── gemv_INT_revsplitK_kernel.ttgir │   ├── gemv_INT_revsplitK_kernel.ttir │   └── __grp__gemv_INT_revsplitK_kernel.json ├── 59f5f5022c214acc7a65c351366f27d49e16a7cafc7173cbd955202e635c8d0f │   └── hip_utils.so └── a5132e2ae49b39161ba2431c5769840a3e679b1b2f8c42f361d5ca461ad71c7d └── __triton_launcher.so 3 directories, 9 files root@h01r4n18:/workspace# file ~/.triton/cache/1708943e4228e69a0defc2ff9233a07461d75e9aa7578c7f654ef90f2f14bc1d/gemv_INT_revsplitK_kernel.hsaco /root/.triton/cache/1708943e4228e69a0defc2ff9233a07461d75e9aa7578c7f654ef90f2f14bc1d/gemv_INT_revsplitK_kernel.hsaco: ELF 64-bit LSB shared object, *unknown arch 0xe0* version 1, dynamically linked, with debug_info, not stripped ``` 而当用 hipprof 对其进行寄存器压力分析时，会出现如图所示的情况：