Introduction
本文系翻译文章Demystifying PTX Code。
在我最近的文章里,我展示了怎样从CUDA和OpenCL代码生成PTX文件,本文则将重点解读PTX文件里的复杂指令。
我们还是使用向量加法的代码
代码片段:
https://gist.github.com/FindHao/394b2f069788e5a4c80a069638a47e1c
原作者的代码项目:https://github.com/pentschev/ptxtract
本文中,我们将关注OpenCL的PTX文件,未来我会写下OpenCL和CUDA代码生成的PTX文件的区别。
1. 向量加法的PTX代码
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sun May 18 04:44:51 2014 (1400399091)
// Driver 331.79
//
.version 3.0
.target sm_21, texmode_independent
.address_size 32
.entry add_vectors(
.param .u32 .ptr .global .align 4 add_vectors_param_0,
.param .u32 .ptr .global .align 4 add_vectors_param_1,
.param .u32 .ptr .global .align 4 add_vectors_param_2,
.param .u32 add_vectors_param_3
)
{
.reg .pred %p<2>;
.reg .s32 %r<21>;
ld.param.u32 %r9, [add_vectors_param_3];
mov.u32 %r5, %envreg3;
mov.u32 %r6, %ntid.x;
mov.u32 %r7, %ctaid.x;
mov.u32 %r8, %tid.x;
add.s32 %r10, %r8, %r5;
mad.lo.s32 %r4, %r7, %r6, %r10;
setp.lt.s32 %p1, %r4, %r9;
@%p1 bra BB0_2;
ret;
BB0_2:
shl.b32 %r11, %r4, 2;
ld.param.u32 %r18, [add_vectors_param_0];
add.s32 %r12, %r18, %r11;
ld.param.u32 %r19, [add_vectors_param_1];
add.s32 %r13, %r19, %r11;
ld.global.u32 %r14, [%r13];
ld.global.u32 %r15, [%r12];
add.s32 %r16, %r14, %r15;
ld.param.u32 %r20, [add_vectors_param_2];
add.s32 %r17, %r20, %r11;
st.global.u32 [%r17], %r16;
ret;
}
其原始c代码:
__kernel
void add_vectors(
__global const int* a,
__global const int* b,
__global int* c,
int n)
{
int idx = get_global_id(0);
if (idx >= n) return;
c[idx] = a[idx] + b[idx];
}
1.1 文件头
文件以一些编译器的信息注释作为开头,紧接着跟着三行:
- PTX ISA版本
- 目标架构,计算能力
- 使用的地址模式
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sun May 18 04:44:51 2014 (1400399091)
// Driver 331.79
//
.version 3.0
.target sm_21, texmode_independent
.address_size 32
1.2 entry Kernel函数入口
紧接着是.entry
指令开头的kernel函数:
.entry add_vectors(
...
)
1.3 parameters Kernel函数的参数
这个Kernel函数的参数是三个指向32位整型global数组的指针(注意,常量和非常量指针在这里没有区别)和一个32位整型变量(输入和输出向量的长度)。
每个参数以.param
指令(这里的翻译欠妥当,现在仍旧不知道directive和instruction对应的中文是什么,暂时都翻译为指令)开头,跟着它的数据类型,比如.u32
。在我们的样例中,由于都是.u32
,表示编译器将所有函数参数转换成了32位有符号整数。
指针参数后面跟着
- 一个
.ptr
的指令 - 他们指向了什么地址空间(
.global
) - 数据对齐的方式,本例中是4字节对齐:
.align 4
.param .u32 .ptr .global .align 4 add_vectors_param_0,
.param .u32 .ptr .global .align 4 add_vectors_param_1,
.param .u32 .ptr .global .align 4 add_vectors_param_2,
.param .u32 add_vectors_param_3
此时函数原型结束。
1.4 寄存器
1.4.1 寄存器定义
函数的代码区正式开始,一开始是寄存器的定义,以.reg
开头:
# .reg 数据类型 寄存器名字
.reg .pred %p<2>;
.reg .s32 %r<21>;
寄存器名字是以%
作为前缀。
有.pred
指令的用来条件分配(@todo 修正),比如分支指令。由于PTX是中间语言,因为寄存器的定义是虚拟的,不一定完全和硬件寄存器是一对一的关系。
一组包含N个虚拟寄存器的寄存器组可以用<N>
的形式来定义,并且可以通过r0, r1, .. , rN-1,r是通过%r
给寄存器组赋的名字。
1.4.2 拷贝参数到寄存器
mov.u32 %r5, %envreg3;
mov.u32 %r6, %ntid.x;
mov.u32 %r7, %ctaid.x;
mov.u32 %r8, %tid.x
其后,ld.param
指令则将函数参数拷贝给了寄存器。因为第二个参数传递的是地址,因此需要添加[ ]
中括号来获取其数据。拷贝参数到寄存器是必须的,因为绝大多数PTX指令都不能直接操作函数参数。
ld.param.u32 %r9, [add_vectors_param_3];
随后,一些特殊寄存器的数值被拷贝给了GPU寄存器:
envreg3
: 只读的驱动定义的特殊寄存器ntid.x
:每个CTA的x维度的线程数量,相当于get_local_sizeo(0)
ctaid.x
:grid里的CTA标识符,相当于get_group_id(0)
tid.x
:CTA x维度的线程号,相当于get_local_id(0)
1.5 加法
add.s32 %r10, %r8, %r5;
现在开始才是真正的代码部分。envreg3
和tid.x
的数值做加法,将结果存储到寄存器r10
,既然寄存器envreg3
是驱动定义的计算而不是代码的直接部分,这里不再深挖它的细节(注意:envreg
没有出现在cuda的PTX代码里)。
1.6 乘加指令
mad.lo.s32 %r4, %r7, %r6, %r10;
mad
指令是multiply add 的缩写,这里计算了全局的线程id(函数get_global_id(0)
)。上面的操作相当于r4 = r7 * r6 + r10
,r6和r7分别是ntid.x
和ctaid.x
,他们乘积的低16位(.lo
后缀表明的)和r10想家,存储计算结果到r4寄存器。
1.7 分支跳转
setp.lt.s32 %p1, %r4, %r9;
@%p1 bra BB0_2;
ret;
接下来就是本例中唯一的一个条件指令。setp
指令是指比较r4(全局线程id)是否比r9(输入输出的长度)小(lower than,.lt
指令)来设定谓词p1。@
指令则判断p1,如果p1是true,执行分支BB0_2(bra
指令,注意bra
指令的目标一定要是label或者指向label的寄存器),执行完BB0_2,回来继续执行后面的代码。如果p1是false,则直接执行后面的代码。本例中,后面就一条ret
指令表示当前分支的结束。
1.8 计算,存储,结束
BB0_2:
shl.b32 %r11, %r4, 2;
ld.param.u32 %r18, [add_vectors_param_0];
add.s32 %r12, %r18, %r11;
ld.param.u32 %r19, [add_vectors_param_1];
add.s32 %r13, %r19, %r11;
真正的数据处理是在分支BB0_2中做的。由于我们是处理的32位的数据指针,第一条指令shl
将32位的无类型二进制寄存器左移了两位(相当于乘以4)后放到了寄存器r11里,这实际上是根据线程的全局id计算从0开始的数据的位置。(@todo 这里的翻译明显不对)。其后,r12 = r18 + r11 ; r13 = r19 + r11,是指传入的矩阵参数a和b分别加上当前线程的全局id,即当前要加的对应的a[i]和b[i]的地址。
ld.global.u32 %r14, [%r13];
ld.global.u32 %r15, [%r12];
add.s32 %r16, %r14, %r15;
ld.param.u32 %r20, [add_vectors_param_2];
add.s32 %r17, %r20, %r11;
st.global.u32 [%r17], %r16;
ret;
ld.global
表示从global memory load(ld即load的缩写) 数据,而r13则是前面计算的&b[i]
,因此需要用[%r13]
来提取r13的内容(即b[i]的地址)对应的内容(即b[i])。r12同理。继而a[i] + b[i] 保存到了r16里。而[add_vectors_param_2]
则是指c,加上r11(即线程全局id左移两位以后的数据)即表示c[i]的地址。将计算的到的r16存储到r17里存的地址指向的地方,使用的是st.global.u32
指令。最后调用ret
指令,程序结束。
2. 结束语
尽管PTX代码不如C/C++代码那么直白,但是在我看来,对于大多数架构来说,相对纯汇编代码还是比较容易理解的,因为我们不用去处理各种中断。比如读懂PTX的代码可以很大程度上帮你理解编译器到底对你的C/C++代码做了什么,以及程序如果在硬件上流动。这有时对你优化你的代码,达到GPU的最大加速效果有很大帮助。
译者按
在最近阅读PTX手册时,Directive和Instruction一直没搞明白到底该怎样翻译。如果读者有了解这方面的欢迎邮件交流,或者在本篇文章下留言。
同时我也在尽力翻译nVIDIA官网的PTX手册,不长,但是还有略吃力。翻译结束以后会挂在github or somewhere,欢迎star和提pull requests。
Comments