文章目录
- 官方网站
- 常用指令
- 1. 基本写法
- 2. 一些基本指令用法
- 如何编译出ptx代码
- 为什么用ptx代码
官方网站
PTX ISA doc: 包括PTX的一些指令集的使用手册。
PTX compiler API: The PTX Compiler APIs are a set of APIs which can be used to compile a PTX program into GPU assembly code. 一个可能被使用的场景是在分离的ptx与主程序应用中，通过更新ptx模块代码，让主程序使用最新的ptx特性。（而不是将ptx嵌入到主程序中，然后每次都有编译所有的程序。）
PTX Writer’s Guide to Interoperability互操作性: This document defines the Application Binary Interface (ABI) for the CUDA® architecture when generating PTX. By following the ABI, external developers can generate compliant PTX code that can be linked with other code.
Inline PTX Assembly in CUDA: The reference guide for inlining PTX (parallel thread execution) assembly statements into CUDA. 也就是怎么在cuda中写PTX代码。
常用指令
1. 基本写法
PTX:
- 源模块是ASCII文本。行由换行字符（\n）分隔。
- PTX是区分大小写的，并使用小写字母作为关键字。
- PTX源模块具有汇编语言风格的语法，包括指令操作码和操作数。伪操作用于指定符号和地址管理。
- 每个PTX模块必须以.version指令开始，指定PTX语言版本，然后是.target指令，指定假定的目标架构。
- 指令关键字以点开头，因此不可能会与用户定义的标识符发生冲突
// 声明一个寄存器变量addr， 类型是u64.
.reg .u64 addr;
2. 一些基本指令用法
- 包括:
.pred, setp, @p, ld.global.v4.u32 ...以及cuda中如何嵌入asm代码
参考实例来源：https://github.com/mit-han-lab/torchsparse/blob/master/torchsparse/backend/utils/memory.cuh
template <int bytes>
struct global_load;template <>
struct global_load<16> {__device__ __inline__ global_load(uint4& D, void const* ptr, int pred_guard) {uint4& data = *reinterpret_cast<uint4*>(&D);// 应该也可以使用 __asm__ __volatile__.// 下面这段ptx指令表示为，但是更加有效：// if (static_cast<int>(pred_guard & (1 << ldg_idx)) != 0) {// data = *(ptr_ldg + ldg_idx);// },。asm volatile("{\n"// 下面一段话代表声明一个谓词变量， 谓词变量的应该可以等同于表达式或者一个行为变量。" .reg .pred p;\n"// .pred的变量经常是和setp一起使用的，这ptx的意思是将p = (int)(pred_guard & 1) != 0// setp: Comparison and Selection Instructions： Compare two numeric values with a relational operator, and (optionally) combine this result with a predicate value by applying a Boolean operator.// ne 表示不等于not equal.// %5 表示ascii后面的第五个参数，从0开始计数； 所以是(int)(pred_guard & 1)的值" setp.ne.b32 p, %5, 0;\n"// 注意下面这四行命令，表示将data.x = data.x, 因为D声明为 uint4 D = make_uint4(0,0,0,0);// 所以这四行命令其实是延迟执行这个初始化，也就是将D.x, D.y,D.z,D.w 初始化为0;" mov.b32 %0, %6;\n"" mov.b32 %1, %7;\n"" mov.b32 %2, %8;\n"" mov.b32 %3, %9;\n"// @p 表示if(p==true) / if(p)的意思； 如果p=true, 则执行ld.global.v4// 注意v4表示vector为4 elems。 写法固定，用大括号括起来目的地址，用中括号括起来源地址。" @p ld.global.v4.u32 {%0, %1, %2, %3}, [%4];\n""}\n"// =r 中有等号，表示目的地址； 而后面没有等号且用“：”分开的是原地址； 注意目的地址和原地址要用":"分开。: "=r"(data.x), "=r"(data.y), "=r"(data.z), "=r"(data.w): "l"(ptr), "r"((int)(pred_guard & 1)), "r"(data.x), "r"(data.y),"r"(data.z), "r"(data.w));}
};
如何编译出ptx代码
# 编译出ptx指令
nvcc -arch=sm_80 -ptx test.cu -o test.ptx
# 查看二进制文件中的汇编
nvcc -arch=sm_80 test.cu -o a.out
# 但是这个查看后的汇编和ptx不太一致，需要对比着看
objdump/cuobjdump -d a.out > log.txt
为什么用ptx代码
- 还是按照上面的例子来说，ptx代码比C++代码产生的ptx代码命令更少
// 首先，如上面例子中所示，两行if语句通过 nvcc编译成ptx后，的代码如下所示：
// C++:
if (static_cast<int>(pred_guard & 1) != 0) {*in_data = *reinterpret_cast<uint4*>(ptr);}// generated ptx： 12行指令
and.b32 %r17, %r16, 1;
setp.eq.b32 %p1, %r17, 1;
mov.pred %p2, 0;
xor.pred %p3, %p1, %p2;
not.pred %p4, %p3;
mov.u32 %r30, 0;
mov.u32 %r31, %r30;
mov.u32 %r32, %r30;
mov.u32 %r33, %r30;
@%p4 bra $L__BB2_2;cvta.to.global.u64 %rd3, %rd1;
ld.global.v4.u32 {%r33, %r32, %r31, %r30}, [%rd3];// 用户自定义的ptx，就少好一些复杂的指令
// 7行指令
asm volatile("{\n"" .reg .pred p;\n"" setp.ne.b32 p, %5, 0;\n"" mov.b32 %0, %6;\n"" mov.b32 %1, %7;\n"" mov.b32 %2, %8;\n"" mov.b32 %3, %9;\n"" @p ld.global.v4.u32 {%0, %1, %2, %3}, [%4];\n""}\n": "=r"(in_data->x), "=r"(in_data->y), "=r"(in_data->z),"=r"(in_data->w) : "l"(ptr), "r"(static_cast<int>(pred_guard & 1)),"r"(in_data->x),"r"(in_data->y), "r"(in_data->z), "r"(in_data->w));
