PTX (Parallel Thread Execution)
是英伟达设计的一种中间语言，可以看作GPU的一种汇编语言。
以分号作为指令的分隔符。
状态空间
明确变量存储在GPU的哪种内存中。
| 状态空间 | 关键字 | 描述 | 生命周期 |
|---|---|---|---|
| 寄存器 | .reg | 单个线程私有 | 线程 |
| 特殊寄存器 | .sreg | 预定义的只读寄存器，如线程/块ID | 线程/块 |
| 共享内存 | .shared | 同线程块共享 | 线程块 |
| 全局内存 | .global | 所有线程块可读写 | 应用程序 |
| 本地内存 | .local | 单个线程私有，物理上位于全局内存 | 线程 |
| 常量内存 | .const | 所有线程块只读，由主机初始化 | 应用程序 |
| 参数内存 | .param | 用于向内核传递参数 | 内核调用 |
变量和操作数
- 立即数：普通数值。
- 标识符：变量名以
%开头。 - 地址：用
[]括起来。 - 向量：支持
.v2、.v4等，写作{ %f1, %f2, %f3, %f4 }
格式
opcode{.type}{.modifier} d, a, b, c;
| 部分 | 描述 | 示例 |
|---|---|---|
opcode | 操作码 | add |
.type | 数据类型 | .u32 |
.modifier | 修饰符 | .lo |
d | 目标操作数 | 通常是寄存器 |
a, b, c | 源操作数 | 可以是寄存器、立即数或地址 |
基础
寄存器声明
.reg .b32 %r<10>; // 声明了 10 个 32 位位寄存器，名字是 %r0 到 %r9
.reg .f32 %f<5>; // 声明了 5 个 32 位浮点寄存器，名字是 %f0 到 %f4
.reg .pred %p<3>; // 声明了 3 个谓词（布尔）寄存器，用于条件判断
数据搬运
ld.global.f32 %f1, [%r1]; // 从 gmem 加载一个 fp32 到 %f1
st.shared.v2.b32 [%r2], {%r3, %r4}; // 将两个 32 位值存储到 smem
谓词执行
@%p1 add.f32 %f1, %f2, %f3; // 如果布尔寄存器 %p1 为 true，则执行
@!%p1 add.f32 %f1, %f2, %f4; // 如果 %p1 为 false，则执行