Date Modified Tags CUDA / ptx

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 文件头

文件以一些编译器的信息注释作为开头,紧接着跟着三行:

  1. PTX ISA版本
  2. 目标架构,计算能力
  3. 使用的地址模式
//
// 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;

现在开始才是真正的代码部分。envreg3tid.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.xctaid.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。

Reference

原文 Demystifying PTX Code


文章版权归 FindHao 所有丨本站默认采用CC-BY-NC-SA 4.0协议进行授权|
转载必须包含本声明,并以超链接形式注明作者 FindHao 和本文原始地址:
https://findhao.net/easycoding/2064.html

Comments