经过好多个月的零敲碎打，今天终于可以发布我写的CUDA SASS的汇编器CuAssembler！撒花~~~
其实核心的实现思想前年就有了，特别是指令编码的实现。不过当时的汇编还停留在几乎全手动的状态。现在我为了降低改汇编的门槛，把很多内容都尽量自动化了，这样更像一个真正的汇编器了。原理在前面的专栏文章已经有完整介绍，这里不重复。不过，因为是业余时间弄的，平时也会有很多别的事情，很多代码都是断断续续写的，做的很粗糙。不过功能的完整度还算可以，出错提示和log之类也算是比较靠谱。从用户门槛来讲，应该算是我知道的最低的了。大部分东西可以让CUDA C来做，看不懂的或是不关键的汇编部分可以直接保持不变，只改想改的部分就行。如果将来需要更细致的扩展，大部分的接口我都有留好，有兴趣的也可以在此基础上再改进。总之，供大家参考吧！
这里只是非常简单的介绍一下使用流程和注意事项，完整介绍大家可以到GitHub页面上看UserGuide和源码。虽然是英文的，但有基础知识的话大概流程也能猜个八九不离十。虽然我想尽量降低写汇编的门槛，但这显然不是CUDA初学者能掌握的。所以感兴趣的读者就去再深入研究，不感兴趣的大概看看热闹就行~
动手前先说基本功。首先CUDA知识要足够了解，关于微架构和汇编指令集的介绍可以看之前的一些文章。然后就是对汇编的结构和语法要有所了解。ELF的格式建议去了解一下。nvdisasm很多汇编directive与Gnu Assembler接近（其实只用了很少一部分），大家有兴趣的可以参考看看。还有就是编译流程要熟悉，NVCC、PTXAS，还有cuobjdump和nvdisasm都会用到一些。然后就是python基本操作了。这里牵扯到的东西很多，大家可能需要时不时的学点新东西，扩展一下知识面~
基本流程
尽管CuAssembler主要的目的是把nvdisasm的输出重新转回cubin，但它并不是从零开始写汇编。我的推荐是从CUDA C开始，除了kernel代码外，其他初始化代码都是尽量用Runtime API，这样使用和修改都最简单直接。其实大部分kernel代码也建议先写好初版，甚至把相关优化做得差不多，后面再做汇编微调，这样改动最小，最不容易出错。如果你只有cubin文件而且打算用driver API来运行，基本流程也类似，只是少了NVCC相关的dump和hack过程。从Runtime API开始的大致步骤如下：
- 新建一个CUDA工程，有若干
.cu文件，包含显式的kernel代码（如果调用了库里的kernel，当前改起来还比较麻烦）。这里最好先写一个大致能工作的版本，需要包含的信息包括：kernel的名字和参数，所有的全局变量，全局constant memory，还有texture/surface的reference，等等。这些会影响cubin文件中ELF的section和symbol相关信息，而这些当前不支持修改。 - 用NVCC的
--keep选项把.cu文件编译成.cubin，注意你可能需要选择相应的SM版本，否则可能只生成最低版本的PTX和cubin。 - 用NVCC的
--dryrun选项得到编译指令序列，把它存下来。后面hack cubin后需要用这个脚本把它直接生成可执行文件，这样程序运行起来和runtime API是一样的。 - 用CuAssembler的
CubinFile类将二进制的.cubin转为可编辑的文本格式.cuasm。 - 修改
.cuasm的内容，也就是做汇编级的优化。大部分的用户时间会耗在这一步…… - 用CuAssembler的
CuAsmParser类将修改后的.cuasm汇编成.cubin。 - 用前面保存好的
--dryrun得到的编译指令序列，把ptxas生成cubin前的指令都删掉，然后运行后面的指令。就可以得到相应的可执行文件。
如果是用driver API运行cubin的方式，那就只要4~6就可以了。如果你有一个很大的工程，有很多kernel但是只想优化其中几个，那这个步骤可能比较繁琐。一个比较妖的办法是在ptxas外面包一层，相当于做ptxas的劫持。因为ptxas每次处理一个文件，可以根据输入ptx路径和需求来判断是用原来的ptxas生成的cubin，还是copy一个hack过的cubin。具体实施大家可以自己想一想，我就不多说了……
一个简单的例子
比如你有一个这样的kernel：
__constant__ int C1[11]; // C1 will be stored in constant memory
__device__ int GlobalC1[7]; // GlobalC1 will be stored in device memory (RW), loaded with relocated address
__global__ void simpletest(const int4 VAL, int* v) // contents of VAL and address of v will be stored in constant memory
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
int a = v[idx]*VAL.x + GlobalC1[idx%16];
// SHFL is an instruction needs an associated .nv.info attribute.
a = __shfl_up_sync(0xffffffff, a, 1);
if (VAL.z > 0) // predicated statement
a += C1[VAL.y];
v[idx] = a;
}
生成的对应的cuasm大概是这个样子（省略其他section，这只是SM_75，其他版本会不一样）：
// --------------------- .text._Z10simpletest4int4Pi --------------------------
.section .text._Z10simpletest4int4Pi,"ax",@progbits
.__section_name 0x35b // offset in .shstrtab
.__section_type SHT_PROGBITS
.__section_flags 0x6
.__section_addr 0x0
.__section_offset 0x4500 // maybe updated by assembler
.__section_size 0x200 // maybe updated by assembler
.__section_link 3
.__section_info 0xc000030
.__section_entsize 0
.align 128 // equivalent to set sh_addralign
.sectioninfo @"SHI_REGISTERS=12"
.align 128
.global _Z10simpletest4int4Pi
.type _Z10simpletest4int4Pi,@function
.size _Z10simpletest4int4Pi,(.L_228 - _Z10simpletest4int4Pi)
.other _Z10simpletest4int4Pi,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z10simpletest4int4Pi:
.text._Z10simpletest4int4Pi:
[----:B------:R-:W-:Y:S08] /*0000*/ MOV R1, c[0x0][0x28] ;
[----:B------:R-:W0:-:S01] /*0010*/ S2R R2, SR_CTAID.X ;
[----:B------:R-:W-:-:S01] /*0020*/ UMOV UR4, 32@lo(GlobalC1) ;
[----:B------:R-:W-:-:S01] /*0030*/ MOV R9, 0x4 ;
[----:B------:R-:W-:-:S01] /*0040*/ UMOV UR5, 32@hi(GlobalC1) ;
[----:B------:R-:W0:-:S01] /*0050*/ S2R R3, SR_TID.X ;
[----:B------:R-:W-:-:S01] /*0060*/ MOV R4, UR4 ;
[----:B------:R-:W-:-:S02] /*0070*/ IMAD.U32 R5, RZ, RZ, UR5 ;
[----:B0-----:R-:W-:Y:S05] /*0080*/ IMAD R2, R2, c[0x0][0x0], R3 ;
[----:B------:R-:W-:Y:S04] /*0090*/ SHF.R.S32.HI R3, RZ, 0x1f, R2 ;
[----:B------:R-:W-:Y:S04] /*00a0*/ LEA.HI R3, R3, R2, RZ, 0x4 ;
[----:B------:R-:W-:Y:S05] /*00b0*/ LOP3.LUT R3, R3, 0xfffffff0, RZ, 0xc0, !PT ;
[R---:B------:R-:W-:-:S02] /*00c0*/ IMAD.IADD R7, R2.reuse, 0x1, -R3 ;
[----:B------:R-:W-:Y:S04] /*00d0*/ IMAD.WIDE R2, R2, R9, c[0x0][0x170] ;
[----:B------:R-:W-:Y:S04] /*00e0*/ IMAD.WIDE R4, R7, 0x4, R4 ;
[----:B------:R-:W2:-:S04] /*00f0*/ LDG.E.SYS R0, [R2] ;
[----:B------:R-:W2:-:S01] /*0100*/ LDG.E.SYS R5, [R4] ;
[----:B------:R-:W-:Y:S04] /*0110*/ MOV R6, c[0x0][0x168] ;
[----:B------:R-:W-:Y:S12] /*0120*/ ISETP.GE.AND P0, PT, R6, 0x1, PT ;
[----:B------:R-:W-:Y:S06] /*0130*/ @P0 IMAD R6, R9, c[0x0][0x164], RZ ;
[----:B------:R-:W0:-:S01] /*0140*/ @P0 LDC R6, c[0x3][R6] ;
[----:B--2---:R-:W-:Y:S08] /*0150*/ IMAD R0, R0, c[0x0][0x160], R5 ;
.CUASM_OFFSET_LABEL._Z10simpletest4int4Pi.EIATTR_COOP_GROUP_INSTR_OFFSETS.#:
[----:B------:R-:W0:-:S02] /*0160*/ SHFL.UP PT, R7, R0, 0x1, RZ ;
[----:B0-----:R-:W-:Y:S08] /*0170*/ @P0 IMAD.IADD R7, R7, 0x1, R6 ;
[----:B------:R-:W-:-:S01] /*0180*/ STG.E.SYS [R2], R7 ;
[----:B------:R-:W-:-:S05] /*0190*/ EXIT ;
.L_20:
[----:B------:R-:W-:Y:S00] /*01a0*/ BRA `(.L_20);
.L_228:这里大部分内容来自nvdisasm反汇编cubin的输出，但是补充了很多重要但是官方反汇编没有显示或是没有显示完整的东西，比如section的各个属性，control codes等等。具体介绍大家可以参考CuAssembler User Guide.
使用上需要注意的一些问题
- Cubin文件里有很多NVIDIA没有公开的格式信息，逆向这些东西既繁琐又无趣，还容易出错。所以我这里尽量把这些东西往官方编译路径上靠，只改那些比较有把握的东西。这也是CuAssembler为什么不支持改section和symbol相关内容的原因，毕竟这些东西可以从CUDA C实现，要靠谱得多。
- 当前指令编码的支持并不完整。Turing(SM_75)因为我以前测的比较多，所以可能相对完整些。除了
B2R有个modifier没显示，还有FSEL用QNAN做参数不能复现外（这个自己写的话可以用ISEL代替），其他都没有大问题。当然，肯定还有很多指令我没收集到。CuAssembler为用户提供了自动更新指令编码库的功能，也写在User Guide里。Ampere的LDG和STG也有modifier没显示，之前还有指令不显示的（报了bug，11.2修好了，不排除有我没发现的）。不过我这老打不开NV的bug report网站，算了，不管了，反正我也买不起新卡。Maxwell和Pascal的编码其实有很多坑，趟了很久还是在深水区，再说吧！反正就是懒。注意：之前有朋友介绍过AsKepler用了一个单bit翻转的方式来穷举编码，我试了试，问题非常多…… 严重不推荐。我只认官方ptxas从正常的ptx编译生成的指令编码，或者是官方库反汇编出来的编码。 - 当前没有完善的正确性检查，parser比较幼稚，特别是有些值越界也管不了。我打算借
nvdisasm的反汇编来做正确性检查。 - NVInfo里有很多属性现在不能自动设置，接口我虽然留了，但是多数功能还要自己一点点去逆向，每个SM版本还不一样，感觉有点烦……
- Control codes不能自动设置。这个也必须要自动收集，否则烦死了。有点思路，还没弄，再说吧……
- 当前这套流程其实不是特别顺手，比如也不支持变量，也不支持流控制之类。不过我觉得这些都可以用python包一层preprocess来解决，把CuAssembler当成后端汇编器就行了。有时间我会试着改改。
说了这么多，其实我自己还没怎么测过呢…… 手上没有卡，下班也比较懒，不想写程序。没错！我流下了贫穷和忙碌的泪水…… 我当前只是把一些CUDA sample的cubin反汇编后再汇编回去，对比了一下，大多数都没啥问题（有些Relocation的顺序会打乱，应该是没关系）。自己改汇编的话，可能一些NVInfo会有问题，当前还是需要用户自己控制。大家注意鉴别~
汇编器的介绍就先这么多，有功夫可以试着做一些比较深入的micro-benchmarking~ 不过下一次我会先分享一些架构和指令集设计中的一些理解和思考~
来源：知乎 www.zhihu.com
作者：cloudcore
【知乎日报】千万用户的选择，做朋友圈里的新鲜事分享大牛。
点击下载