PTX系列—语法 Nvidia cuda zone中文站
2014-07-18 09:26阅读：
PTX程序是文本源文件的集合。PTX源文件拥有配套的语言风格语法，具有指令操作代码和操作数。Pseudo操作指定符号和地址管理。Ptxas程序组合PTX源文件来生成相应的二进制目标文件。英文原文链接：http://www.nvidia.com/object/io_1195170102263.html
PTX程序是文本源文件的集合。PTX源文件拥有配套的语言风格语法，具有指令操作代码和操作数。Pseudo操作指定符号和地址管理。Ptxas程序组合PTX源文件来生成相应的二进制目标文件。
4.1 源格式
源文件为ASCII文本。使用换行符(‘’)来分行。
所有空白字符都相等；忽略空白，除非空白在语言中用作分隔标记。
C预处理程序cpp可以用于处理PTX源文件。以#开头的行为预处理程序。以下是常见的预处理程序命令：
#include, #define, #if, #ifdef, #else, #endif, #line, #file
C：Harbison 和 Steele提供的《参考手册》很好地描述了C预处理器。
PTX区分大小写，并且使用
每个PTX文件必须以版本指令开头，指定PTX语言的版本，后面跟着目标指令，指定预期的目标架构。有关这些指令的详细信息，请参阅第9章。
4.2 注释
PTX中的注释遵循C/C++语法，使用非嵌套的进行可能跨越多行的注释，在延伸到当前行末尾的注释前使用//。
PTX中的注释都当作空白来处理。
4.3
语句
PTX语句可以是命令(directive)或指令(instruction)。语句以一个可选标签开始，以分号结束。
示例：
.reg .b32 r1, r2;
.global .f32 array[N];
start: mov.b32
r1, %tid.0;
shl.b32 r1, r1, 2; // shift thread id by 2 bits
ld.b32 r2, array[r1]; // thread[tid] gets array[tid]
add.f32 r2, r2, 0.5; // add 1/2
4.3.1 命令语句
命令关键字以点开头，所以不会与用户定义的标识符混淆。PTX中的
表1
.align .func .section .tex
.byte .global .shared .union
.const .local .sreg version
.entry .loc .struct .visible
.extern .param .surf
.file .reg
.target
4.3.2 指令语句
指令的构成如下：指令操作数，后跟以逗号分隔的操作数列表，以分号结束。操作数可以是寄存器变量、常量表达式、地址表达式或标签名称。指令有一个可选的保护谓语，用于控制条件执行。保护谓语位于操作标签之后，操作数之前，写作@p，其中p是谓语寄存器。保护谓语可以忽略，写作@!p。
目标操作数在最前面，后跟源操作数。
指令关键字在表2中列出。所有指令关键字都是PTX中的保留标记。
表2 保留的指令关键字
abs ex2 not sin
add exit or slct
and ld rcp sqrt
atom lg2 rem st
bar mad ret sub
bra mad24 rsqrt tex
brkpt max sad trap
call min selp vote
cnot mov set xor
cos mul setp
cvt mul24 shl
div neg shr
4.4 标识符
用户定义的标识符遵循扩展的C++规则：以字母字符、下划线、美元符号或百分号([ A-Za-z_$%
])开头，并且后跟0个或多个字母数字、下划线或美元符号([ A-Za-z_$% ])。
很多高级语言（比如C和C++）遵循类似的标识符名称规则，只是不允许百分号。PTX允许百分号作为标识符的第一个字符。百分号可用于避免名称冲突，例如在用户定义的变量名称和编译器生成的名称之间的冲突。
PTX预定义了少量以百分号开头的特殊寄存器，如表3所列。
表3 预定义的标识符
%clock %ctaid %ntid
%gridid %nctaid %tid
4.5 立即常量
PTX中的立即常量仅限于整型和浮点类型。
4.5.1 整数立即常量
整数立即常量可以按十进制、十六进制、八进制和二进制符号来书写。
十进制常量以非零数字开头，后跟0个或多个数字（0-9）。
十六进制常量以0x和0X开头，后跟一个和多个十六进制数字（来自集合[0-9a-fA-F]）。
八进制常量以零开头，后跟0个或多个八进制数字（0-7）。
二进制常量以0b和0B开头，后跟一个或多个二进制数字（01）。
4.5.2 浮点立即常量
浮点立即常量可以使用可选小数点和可选有符号指数来书写。不同于C和C++，没有后缀字符用于指定大小（例如，float和double）。
PTX包含浮点常量的二次表示，精确的机器表示给作十六进制常量。对于64位浮点值，该常量以0d或0D开头，后跟十六进制数字。对于32位浮点值，该常量以0f或0F开头，后跟八进制数字。
4.5.3
谓语立即常量
布尔值TRUE和FALSE的谓语立即常量分别写作二进制数字1和0。
4.5.4常量表达式
常量表达式在编译时进行评估，以便构建简单的值供立即操作数和地址表达式使用。整数和浮点常量表达式均受支持，但是请注意，浮点常量表达式可能评估为不同的值，而不是在目标架构上计算，因为编译器可能使用比目标架构更高的精度来对表达式求值。
常量表达式从词汇常量、基本算法运算符（加、减、乘和除）和插入语中构建。整数常量表达式可以包含余数（%）、shift运算符（<<
和 >>）以及逻辑运算符（&、|和^）。
PTX中运算符的意义和C和C++中的是一样的。
引文来源
Nvidia cuda zone中文站
