CUDA PTX ISA阅读笔记(一)

不知道这是个啥的看这里:Parallel Thread Execution ISA Version 5.0. 简要来说,PTX就是.cu代码编译出来的一种东西,然后再由PTX编译生成执行代码。如果不想看网页版,cuda的安装目录下的doc文件夹里有pdf版本,看起来也很舒服。 ps:因为文档是英文的(而且有二百多页= =),鉴于博主英语水平有限并且时间也有限(主要是懒),因此只意译了一些自以为重点的内容,如想要深入学习,还是乖乖看文档去吧

第一章 介绍

1.1. 使用GPU进行可扩展数据并行计算

介绍了一波并行计算的知识。

1.2. PTX的目标

PTX为提供了一个稳定的编程模型和指令集,这个ISA能够跨越多种GPU,并且能够优化代码的编译等等。

1.3. PTX ISA 5.0版本

就是PTX ISA5.0的一些新特性

1.4. 文档结构

  • 编程模型:编程模型的概要
  • PTX 机器模型:大致介绍PTX虚拟机
  • 语法:描述PTX语言的基础语法
  • 状态空间、类型和变量:就是描述这些玩意
  • 指令操作数
  • 应用二进制接口:描述了函数定义和调用的语法,以及PTX支持的应用二进制接口
  • 指令集
  • 特殊的寄存器
  • 版本更新介绍

第二章 编程模型

2.1. 一个高并行度的协处理器

继续科普GPU。

2.2. 线程层级

2.2.1 合作线程阵列

2.2.2 线程阵列网格

上边这两节主要就是讲一些基本的GPU的block啊grid啊之类的东西,想了解的可以看我的另一篇文章:《GPU高性能编程 CUDA实战》(CUDA By Example)读书笔记-第五章。这里的图就用了这个手册里的。

2.3. 内存层级

这个图实在是太好了:

第三章 PTX机器模型

3.1. 一组带有片上共享内存的SIMT多处理器

主要讲一下硬件层级结构,果然图还是最好的:

第四章 语法

PTX语言是由操作指令和操作数组成。

4.1. 代码格式

使用\n换行,空格木有意义,#这个符号和C差不多,就是预编译指令,而且大小写敏感,每个PTX代码都是由.version打头,表示PTX的版本。

4.2. 注释

和C一样

4.3. 语句

以一个可选的标记开始,以分号结束,就像这样: start: mov.b32 r1, %tid.x;

4.3.1. 指示

提供了PTX的指示

4.3.2. 指令

提供了PTX的指令:

ps:关于directive和 instruction这两个词的区别涉及一些汇编上的知识,前者这里翻译为指示,后者这里翻译成指令,因为一般directive并不会产生代码而是指示编译器的一些行为,而instruction则会产生实际的代码,想了解的可以看这里:What-is-the-difference-between-an-instruction-and-a-directive-in-assembly-language

4.4. 标识符

这个大概就是变量名的命名规则吧,基本就和C一样啦,然后系统预定义的变量都是以%开头的大佬变量。

4.5. 常量

这个,我猜,大概是是标号标错了,应该是包含下面各种常量的大标题才对。

4.6. 整型常量

每个整型常量都是64哒,分为有符号和无符号,由.s64和.u64定义,其中各个进制的数是如下定义的:

X进制

表示方式

十六进制

0[xX]{十六进制数}+U?

十进制

0{octal 十进制数}+U?

二进制

0[bB]{0/1}+U?

小数

{非零数}{十进制数}*U?

4.6.1. 浮点常量

浮点数都是64位的,除了用一个32位十六进制去精确表达一个单精度浮点数(黑人问号脸???),具体表达方式如下:

精度

表达方式

单精度

0[fF]{十六进制数}{8}

双精度

0[dD]{十六进制数}{16}

4.6.2. 判断值常量

0就是false,非零就是true

4.6.3. 常量表达式

这个大概是可以对常量能够使用的表达式,也和C基本一致啦:

4.6.4 整型常量表达式求值

和C语言一样一样的

4.6.5 表达式求值规则总结

C语言+1

第五章 状态空间、类型和变量

5.1. 状态空间

这个状态空间就我理解吧,就是在哪块内存上操作。

5.1.1. 寄存器状态空间

利用.reg来声明寄存器状态空间,该空间可以使用几乎形式的数据,但是不同于其他状态空间的是寄存器是没有地址的。

5.1.2. 特殊寄存器状态空间

用.sreg来声明,存的主要是系统预定义的一些变量,比如grid的维数之类的数据。

5.1.3. 常量状态空间

常量状态空间使用.const来表示,被限制在64KB之内。并且被组织成10个区域,驱动要在这十个区域中申请空间,然后可以将这些申请到的空间用指针传递给核函数。

5.1.3.1. 存储体常量寄存器(弃用)

以前这种是需要指定确定的区域号才可以的,就像这样: .extern .const[2] .b32 const_buffer[];

5.1.4. 全局状态空间

使用ld.global,st.globle和atom.global来访问全局状态空间。而且,访问全局变量空间是没有顺序的,是需要使用bar.sync来同步的。

5.1.5. 本地状态空间

.local声明本地状态空间,而且只能在线程内部使用。

5.1.6. 参数状态空间

参数状态空间被用于1.将输入的参数从主机传递给核函数。2.为在核函数内调用的设备函数声明形式化输入和返回参数。3.声明作为函数调用参数的本地数组,特别是用来传递大的结构体给函数。

5.1.6.1. 核函数参数

.entry foo ( .param .b32 N, .param .align 8 .b8 buffer[64]) 
{ 
    .reg .u32 %n; 
    .reg .f64 %d; 
    ld.param.u32 %n, [N]; 
    ld.param.f64 %d, [buffer]; 
    ...

5.1.6.2. 核函数参数属性

5.1.6.3. 核函数参数属性: .ptr

使用这个相当于一个指针,还可以指定内存对齐的大小。

 .entry foo ( 
    .param .u32 param1, 
    .param .u32 .ptr.const.align 8 param3, 
    .param .u32 .ptr.align 16 param4 
    ) { .. }

5.1.6.4. 设备函数参数

这个最常用于传递大小和寄存器大小不一样的变量,比如结构体。

 .func foo ( .reg .b32 N, .param .align 8 .b8 buffer[12] ) {
    .reg .s32 %y; 
    ld.param.f64 %d, [buffer]; 
    ld.param.s32 %y, [buffer+8]; 
    ... 
}

5.1.7. 共享状态空间

用.shared定义,共享内存有一个特点是可以广播,并且能够顺序访问(有某种一致性机制?)

5.1.8. 纹理状态空间(弃用)

纹理内存也是全局内存的一部分,被上下文的所有线程共享并且是只读的。使用.tex应该被.global里的.texref来代替。就像:

  .tex .u32 tex_a;
  //转换成下面这样
  .global .texref tex_a;

5.2. 类型

5.2.1. 基本类型

这些基本类型就好像C语言中的int,float之类的,用来定义变量的:

5.2.2. 使用子字段的尺寸限制

像.u8, .s8,和.b8这种类型仅限于在ld,st和cvt中使用。.f16只能被转换成并且只能从.f32,.f64类型。.f16×2这种浮点类型只允许被用在浮点数算法指令和纹理获取指令上。

5.3. 纹理采集器和表面类型

下面这段话是从专家手册里摘录的关于表面引用的解释:

读写纹理和表面的指令相对于其他指令涉及了更多隐秘状态。参数,例如基地址、维度、格式和纹理内容的解释方式,都包含在一个header头结构中。header是一个中间数据结构,它的软件抽象被称为纹理引用(texture reference)或表面引用(surface reference)。 这里有个表用来讲专门为纹理状态空间提供的不透明类型:

5.3.1. 纹理和表面设置

像上表中所提到的width, height, 和 depth都用来说明纹理内存的大小之类的特性。

5.3.2. 采集器设置

它有各种模式,看CUDA C Programming Guide获取更多细节。

5.3.3. 频道数据类型和频道指令字段

以前之后OpenCL能用,现在都能用了。 讲真,由于对纹理内存了解太少,这节看得很勉强。

5.4. 变量

5.4.1. 变量声明

变量声明需要同时声明状态空间和数据类型比如:

.global .u32 loc; 
.reg .s32 i; 
.const .f32 bias[] = {-1.0, 1.0}; 
.global .u8 bg[4] = {0, 0, 0, 0}; 
.reg .v4 .f32 accel; 
.reg .pred p, q, r;

5.4.2. 向量

这里的向量的长度是被ptx固定的,只能是2或者4,也不能是判断值(true of false),定义同普通变量:global .v4 .f32 V;

5.4.3. 数组声明

数组的定义和C差不多,可以指定长度也可以不指定然后初始化:

      .local  .u16 kernel[19][19];
      .shared .u8  mailbox[128];
      .global .u32 index[] = { 0, 1, 2, 3, 4, 5, 6, 7 };
      .global .s32 offset[][2] = { {-1, 0}, {0, -1}, {1, 0}, {0, 1} };

5.4.4. 初始化器

对于初始化,是这样的:

.const .f32 vals[8] = { 0.33, 0.25, 0.125 };
.global .s32 x[3][2] = { {1,2}, {3} };
//相当于
.const .f32 vals[4] = { 0.33, 0.25, 0.125, 0.0, 0.0 }; 
.global .s32 x[3][2] = { {1,2}, {3,0}, {0,0} };

当前,变量的初始化只对常量和global状态空间支持,默认的初始化值是0。对于数组,还可以采用以下神奇的方法来初始化:

.const .u32 foo = 42; 
.global .u32 p1 = foo; // offset of foo in .const space .global .u32 p2 = generic(foo); // generic address of foo 

// array of generic-address pointers to elements of bar .global .u32 parr[] = { generic(bar), generic(bar)+4, generic(bar)+8 };

5.4.5. 内存对齐

就是可以在定义数组什么的时候指定内存对齐的大小:

// allocate array at 4-byte aligned address. Elements are bytes. .const .align 4 .b8 bar[8] = {0,0,0,0,2,0,0,0};

5.4.6. 参数化的变量名称

这里提供了一种快捷声明变量的方法:.reg .b32 %r<100>; //就相当于声明了 %r0, %r1, ..., %r99

5.4.7. 变量属性

参见下一节

5.4.8. 变量属性指示: .attribute

变量有个.manage属性,这个属性只能在.global状态空间上使用,使用了这个属性之后能召唤神龙可以将变量放置在一个虚拟空间上,这个空间主机和设备都能够访问。具体是这样使用的:.global .attribute(.managed) .s32 g;

第六章 指令操作数

6.1. 操作数类型信息

每个指令里的操作数都要声明其类型,而且类型必须符合指令的模板,并没有自动的类型转换。

6.2. 源操作数

PTX描述的是一个存储读取机,因此对于ALU指令的操作数必须在.reg寄存器状态空间。cvt指令可以的参数有多种类型和大小,可以转换一种类型(或者大小)到另一种类型(或大小)。ld,st,mov和cvt指令从一个地址拷贝数据到另一个地址。ld,st将内容拷贝到寄存器或者从寄存器中拷贝出来,mov指令把数据从一个寄存器换到另一个寄存器。大多数指令有个可选的判断操作,一些指令有附加的判断类型的源操作数,这些经常被定义名为p,q,r,s.

6.3. 目的操作数

用来得到一个结果,一般都在寄存器中。

6.4. 使用地址,数组和向量

6.4.1. 地址作为操作数

就类似各种类型的定义:

.shared .u16 x; 
.reg .u16 r0; 
.global .v4 .f32 V; 
.reg .v4 .f32 W; 
.const .s32 tbl[256];
.reg .b32 p; .reg .s32 q; 

ld.shared.u16 r0,[x]; 
ld.global.v4.f32 W, [V]; 
ld.const.s32 q, [tbl+12]; 
mov.u32 p, tbl;

6.4.2. 数组作为操作数

数组的使用也基本上和C语言是一样的:

ld.global.u32 s, a[0]; 
ld.global.u32 s, a[N-1]; 
mov.u32 s, a[1]; // move address of a[1] into s

6.4.3. 向量作为操作数

向量的感觉更像是一个结构体或者数组,使用向量可以快速地给多个数复制,很强:

.reg .v4 .f32 V; 
.reg .f32 a, b, c, d; 
mov.v4.f32 {a,b,c,d}, V;
//也可以反过来
ld.global.v4.f32  {a,b,c,d}, [addr+offset];
ld.global.v2.u32  V2, [addr+offset2];

6.4.4. 标记和函数名作为操作数

这个主要是用来获得标记或者函数名,在分支语句中做跳转使用。

6.5. 类型转换

6.5.1. 标量转换

sext:符号扩展。zext:零扩展。chop:只保留低位。s是有符号整数,f是浮点数,u是无符号整数。2就是转换成

6.5.2. 取整修改器

这里是表示取整的标志,有什么向下取证向上取整之类的。(最低有效位(英语:Least Significant Bit,lsb)是指一个二进制数字中的第0位(即最低位),权值为2^0,可以用它来检测数的奇偶性。)

6.6. 操作数耗时

不同状态空空间的操作数会影响一个操作的速度。寄存器最快,全局变量最慢,而多线程可以掩盖这种延迟,或者让取值指令越简单越好。下面是从这些地方取值的延迟:

第七章 抽象ABI

ABI是Application Binary Interface的缩写,翻译过来是二进制程序接口。直白点讲就是系统提供的一系列函数。

7.1. 函数的声明和定义

话不多说就看代码好了

//定义了一个结构体
struct {
    double dbl;
    char   c[4];
};
//有返回值和传入参数
.func (.reg .s32 out) bar (.reg .s32 x, .param .align 8 .b8 y[12]) 
{ 
    .reg .f64 f1; 
    .reg .b32 c1, c2, c3, c4; 
    ... 
    ld.param.f64 f1, [y+0]; 
    ld.param.b8 c1, [y+8];
    ld.param.b8 c2, [y+9];
    ld.param.b8 c3, [y+10]; 
    ld.param.b8 c4, [y+11]; 
    ... ... // computation using x,f1,c1,c2,c3,c4; 
} 
{
    .param .b8 .align 8 py[12]; 
    ...
    //通过位移来使用参数
    st.param.b64 [py+ 0], %rd; 
    st.param.b8 [py+ 8], %rc1; 
    st.param.b8 [py+ 9], %rc2; 
    st.param.b8 [py+10], %rc1; 
    st.param.b8 [py+11], %rc2; 
    // scalar args in .reg space, byte array in .param space 
    call (%out), bar, (%x, py); 
    ...

要注意,对于参数的st.param和对返回值的ld.out都必须紧跟着函数调用call。这样才能让编译器优化是的.param不占用多余的空间。并且这个.param允许简单的映射将有多个地址的结构映射到能够传给函数的变量上。

7.1.1. PTX ISA Version 1.x的改变

1.x只支持.reg,后来开始支持.param。

7.2. 列表函数

现在的ptx并不支持列表函数。(不支持说个毛,下一位!)

7.3. Alloca

同上同上

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏java思维导图

JAVA容器-自问自答学ArrayList

前言 这次我和大家一起学习HashMap,HashMap我们在工作中经常会使用,而且面试中也很频繁会问到,因为它里面蕴含着很多知识点,可以很好的考察个人基础。但...

37490
来自专栏一个会写诗的程序员的博客

Sting str = "aaaa" 的形式定义一个字符串最大长度只能有 65534 个。

为什么呢?因为在class文件的规范中, CONSTANT_Utf8_info表中使用一个16位的无符号整数来记录字符串的长度的,最多能表示 65536个字节,...

8920
来自专栏yl 成长笔记

UML 类图基础

定义:两个类之间的强依赖关系, 可以为单向,亦可为双向。常见表现形式 为 A 类中有 B 类型的成员变量。

11740
来自专栏企鹅号快讯

Python算法分享系列-查找,排序,递归

iTesting,爱测试,爱分享 沉寂了一段时间,继续学习。 算法这个系列我想分享很久了,奈何本身对算法不是特别了解,又找不到合适的载体来分享。 最近看了本有趣...

46560
来自专栏java一日一条

Python开发的10个小贴士

下面是十个Python中很有用的贴士和技巧。其中一些是初学这门语言常常会犯的错误。

8220
来自专栏测试开发架构之路

C/C++常用头文件及函数汇总

C/C++头文件一览 C #include <assert.h>    //设定插入点 #include <ctype.h>     //字符处理 #inclu...

51050
来自专栏互联网开发者交流社区

HashMap相关(二)

12850
来自专栏mathor

LeetCode75.颜色分类

 这道题两种做法,一种是用计数排序,因为告诉了你只有3种数字,所以直接创建一个长度为3的数组arr,然后遍历一遍原数组,每出现一次某个数,arr对应位置的值...

7120
来自专栏机器学习从入门到成神

牛客网刷题汇总(一)附解析

版权声明:本文为博主原创文章,未经博主允许不得转载。 https://blog.csdn.net/sinat_35512245/articl...

87920
来自专栏Java技术栈

递归算法介绍及Java应用实战

什么是递归算法 递归算法是把问题转化为规模缩小了的同类问题的子问题,然后递归调用函数(或过程)来表示问题的解。一个过程(或函数)直接或间接调用自己本身,这种过程...

409100

扫码关注云+社区

领取腾讯云代金券