CUDA 中的内联 PTX 汇编

将 PTX（并行线程执行）汇编语句内联到 CUDA 中的参考指南。

1. 在 CUDA 中使用内联 PTX 汇编

NVIDIA^® CUDA^® 编程环境提供了一种并行线程执行 (PTX) 指令集架构 (ISA)，用于将 GPU 用作数据并行计算设备。有关 PTX ISA 的更多信息，请参阅最新版本的 PTX ISA 参考文档。

本应用笔记介绍了如何将 PTX 汇编语言语句内联到 CUDA 代码中。

1.1. 汇编器 (ASM) 语句

汇编器语句 asm() 提供了一种将任意 PTX 代码插入到 CUDA 程序中的方法。一个简单的例子是

asm("membar.gl;");

这会将 PTX membar.gl 插入到生成的 PTX 代码中 asm() 语句的位置。

1.1.1. 参数

当我们在 asm 中传入和传出值时，asm() 语句变得更加复杂和有用。基本语法如下

asm("template-string" : "constraint"(output) : "constraint"(input));

其中您可以有多个输入或输出操作数，以逗号分隔。模板字符串包含对操作数的引用 PTX 指令。可以通过用分号分隔来给出多个 PTX 指令。

一个简单的例子如下

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));

模板字符串中的每个 %n 都是指向以下操作数列表的索引，按文本顺序排列。因此，%0 指的是第一个操作数，%1 指的是第二个操作数，依此类推。由于输出操作数始终列在输入操作数之前，因此它们被分配了最小的索引。此示例在概念上等效于以下内容

add.s32 i, j, k;

请注意，字符串中的编号引用可以按任意顺序排列。以下内容与上面的示例等效

asm("add.s32 %0, %2, %1;" : "=r"(i) : "r"(k), "r"(j));

您也可以重复引用，例如

asm("add.s32 %0, %1, %1;" : "=r"(i) : "r"(k));

在概念上是

add.s32 i, k, k;

如果没有输入操作数，则可以删除最后一个冒号，例如

asm("mov.s32 %0, 2;" : "=r"(i));

如果没有输出操作数，则冒号分隔符是相邻的，例如

asm("mov.s32 r1, %0;" :: "r"(i));

如果您想要 ptx 指令中的 %，则应使用双 %% 对其进行转义，例如

asm("mov.u32 %0, %%clock;" : "=r"(x));

上面进行了简化，以解释字符串 % 引用的顺序。实际上，操作数值是通过约束指定的任何机制传递的。约束的完整列表将在稍后解释，但“r”约束指的是 32 位整数寄存器。因此，之前的示例 asm() 语句

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));

在编译器生成的输出中生成以下代码序列

ld.s32 r1, [j];
ld.s32 r2, [k];
add.s32 r3, r1, r2;
st.s32 [i], r3;

这就是输入操作数和输出操作数之间的区别变得重要的地方。输入操作数在 asm() 语句之前加载到寄存器中，然后结果寄存器存储到输出操作数中。“=r” 中的 “=” 修饰符指定寄存器是写入的。还有一个 “+” 修饰符可用，它指定寄存器既被读取又被写入，例如

asm("add.s32 %0, %0, %1;" : "+r"(i) : "r" (j));

可以将多个指令组合成单个 asm() 语句；基本上，任何合法的东西都可以放入 asm 字符串中。可以通过利用 C/C++ 的隐式字符串连接将多个指令拆分到多行。C++ 样式行尾注释 “//” 和经典 C 样式注释 “/**/” 可以与这些字符串交错。为了在 PTX 中间文件中生成可读的输出，最佳实践是使用 “nt” 终止除最后一个指令字符串之外的每个指令字符串。

例如，立方体例程可以编写为

__device__ int cube (int x)
{
  int y;
  asm(".reg .u32 t1;\n\t"              // temp reg t1
      " mul.lo.u32 t1, %1, %1;\n\t"    // t1 = x * x
      " mul.lo.u32 %0, t1, %1;"        // y = t1 * x
      : "=r"(y) : "r" (x));
  return y;
}

如果输出操作数由 asm 指令有条件地更新，则应使用 “+” 修饰符。在这种情况下，输出操作数有一个隐式用法。例如，

__device__ int cond (int x)
{
  int y = 0;
  asm("{\n\t"
      " .reg .pred %p;\n\t"
      " setp.eq.s32 %p, %1, 34;\n\t" // x == 34?
      " @%p mov.s32 %0, 1;\n\t"      // set y to 1 if true
      "}"                            // conceptually y = (x==34)?1:y
      : "+r"(y) : "r" (x));
  return y;
}

1.1.2. 约束

每个 PTX 寄存器类型都有一个单独的约束字母

"h" = .u16 reg
"r" = .u32 reg
"l" = .u64 reg
"q" = .u128 reg
"f" = .f32 reg
"d" = .f64 reg

例子

asm("cvt.f32.s64 %0, %1;" : "=f"(x) : "l"(y));

生成

ld.s64 rd1, [y];
cvt.f32.s64 f1, rd1;
st.f32 [x], f1;

请注意，约束 "q" 仅在支持 __int128 的平台上受支持。

约束 "n" 可用于具有已知值的立即整数操作数。例子

asm("add.u32 %0, %0, %1;" : "=r"(x) : "n"(42));

生成

add.u32 r1, r1, 42;

约束 "C" 可用于 ‘const char 数组’类型的操作数，其中数组内容在编译时已知。它旨在允许基于编译时计算自定义 PTX 指令模式（请参阅示例）。以下是 "C" 约束的规范

'C'(constant-expression)

constant-expression 在编译期间求值，并应生成变量 V 的地址，其中

V 具有静态存储持续时间。
V 具有 ‘const char 数组’ 类型。
V 是常量初始化的。
如果 V 是静态类成员，则 V 的初始化声明是类中的声明。

在转换期间，编译器将用 V 初始化器的内容替换对汇编器模板中操作数的引用，除了最后一个尾随零。此约束不允许使用约束修饰符。此约束只能在设备代码中使用。

（斜体术语是 C++ 标准术语和/或 GNU 内联汇编规范中的术语）。

以下是使用 C 约束来基于编译时计算生成不同 PTX 指令模式的示例

constexpr int mode_rz = 0;
constexpr int mode_rn = 1;

template <int mode>
struct helper;

template<> struct helper<mode_rz> {
    static constexpr const char mode[] = ".rz";
};

template<> struct helper<mode_rn> {
    static constexpr const char mode[] = ".rn";
};

template <int rounding_mode>
__device__ float compute_add(float a, float b) {
    float result;
    asm ("add.f32%1 %0,%2,%3;" : "=f"(result)
                            : "C"(helper<rounding_mode>::mode),
                              "f"(a), "f"(b));
    return result;
}

__global__ void kern(float *result, float a, float b) {
    *result++ = compute_add<mode_rn>(a,b); // generates add.f32.rn
    *result   = compute_add<mode_rz>(a,b); // generates add.f32.rz
}

其他示例（在 C++17 或更高版本的方言中编译）

struct S1 {
static constexpr char buf1[] = "Jumped";
static constexpr char buf2[] = {'O', 'v', 'e', 'r', 0};
};

template <const char *p1, const char *p2, const char *p3>
__device__ void doit() {
asm volatile ("%0 %1 %2" : : "C"(p1), "C"(p2), "C"(p3));
}

struct S2 {
static const char buf[];
};
const char S2::buf[] = "this";

const char buf3[] = "Jumped";
extern const char buf4[];

__global__ void foo() {
    static const char v1[] = "The";
    static constexpr char v2[] = "Quick";
    static const char v3[] = { 'B' , 'r' , 'o', 'w', 'n', 0 };
    static constexpr char v4[] = { 'F', 'o', 'x', 0 };

    //OK: generates 'The Quick Brown Fox Jumped Over' in PTX
    asm volatile ("%0 %1 %2 %3 %4 %5" : :  "C"(v1) , "C"(v2), "C"(v3),  "C"(v4), "C"(S1::buf1), "C"(S1::buf2) );

    //OK: generates 'Brown Fox Jumped' in PTX
    doit<v3, v4, buf3>();


    //error cases

    const char n1[] = "hi";

    //error: argument to "C" constraint is not a constant expression
    asm volatile ("%0" :: "C"(n1));

    //error: S2::buf was not initialized at point of declaration
    asm volatile ("%0" :: "C"(S2::buf));

    //error: buf4 was not initialized
    asm volatile ("%0" :: "C"(buf4));
}

没有用于 8 位宽 PTX 寄存器的约束字母。接受 8 位宽类型的 PTX 指令类型允许操作数比指令类型大小更宽。例子

__device__ void copy_u8(char* in, char* out) {
    int d;
    asm("ld.u8 %0, [%1];" : "=r"(d) : "l"(in) : "memory");
    *out = d;
}

生成

ld.u8 r1, [rd1];
st.u8 [rd2], r1;

使用上面未指定的约束字符串的行为是未定义的。

1.2. 陷阱

虽然 asm() 语句非常灵活和强大，但您可能会遇到一些陷阱——这些陷阱在本节中列出。

1.2.1. 命名空间冲突

如果多次在代码中调用和内联立方体函数（如前所述），则会生成关于临时寄存器 t1 重复定义的错误。要避免此错误，您需要

不内联立方体函数，或者，

将 t1 用法嵌套在 {} 内，以便它对每次调用都有单独的作用域，例如

__device__ int cube (int x)
{
  int y;
  asm("{\n\t"                        // use braces for local scope
      " reg .u32 t1;\n\t"            // temp reg t1,
      " mul.lo.u32 t1, %1, %1;\n\t"  // t1 = x * x
      " mul.lo.u32 %0, t1, %1;\n\t"  // y = t1 * x
      "}"
      : "=r"(y) : "r" (x));
  return y;
}

请注意，您也可以类似地对 asm() 语句内的本地标签使用花括号。

1.2.2. 内存空间冲突

由于 asm() 语句无法知道寄存器位于哪个内存空间中，因此用户必须确保使用适当的 PTX 指令。对于 sm_20 及更高版本，asm() 语句的任何指针参数都作为通用地址传递。

1.2.3. 不正确的优化

编译器假定 asm() 语句除了更改输出操作数外，没有副作用。为了确保在生成 PTX 期间不删除或移动 asm，您应该使用 volatile 关键字，例如

asm volatile ("mov.u32 %0, %%clock;" : "=r"(x));

通常，任何写入的内存都将指定为输出操作数，但是如果用户内存上有隐藏的读取或写入（例如，通过操作数间接访问内存位置），或者如果您想停止围绕在 PTX 生成期间执行的 asm() 语句的任何内存优化，您可以在第三个冒号后添加 “memory” 损坏规范。例如

asm volatile ("mov.u32 %0, %%clock;" : "=r"(x) :: "memory");
asm ("st.u32 [%0], %1;" :: "l"(p), "r"(x) : "memory");

1.2.4. 不正确的 PTX

编译器前端不解析 asm() 语句模板字符串，也不知道它的含义，甚至不知道它是否是有效的 PTX 输入。因此，如果字符串中存在任何错误，则在 ptxas 之前不会显示。例如，如果您使用 “r” 约束传递一个值，但在 add.f64 中使用它，您将从 ptxas 获得解析错误。同样，不支持操作数修饰符。例如，在

asm("mov.u32 %0, %n1;" : "=r"(n) : "r"(1));

“%n1” 中的 ‘n’ 修饰符不受支持，并将传递给 ptxas，在那里它可能导致未定义的行为。有关更多编译器相关详细信息，请参阅文档 nvcc.pdf。

1.3. 错误检查

以下是编译器将对 inlinePTXasm 执行的一些错误检查

不允许单个 asm 操作数有多个约束字母，例如
```
asm("add.s32 %0, %1, %2;" : "=r"(i) : "rf"(j), "r"(k));
```
错误：一个 asm 操作数在一个 __device__/__global__ 函数中只能指定一个约束字母
只允许标量变量作为 asm 操作数。具体来说，不允许像 ‘struct’ 类型变量这样的聚合，例如
```
int4 i4;
asm("add.s32 %0, %1, %2;" : "=r"(i4) : "r"(j), "r"(k));
```
错误：asm 操作数必须具有标量类型
PTX asm 约束隐含的类型和大小必须与关联的操作数匹配。大小不匹配的示例
对于 ‘char’ 类型变量 “ci”，
asm("add.s32 %0,%1,%2;":"=r"(ci):"r"(j),"r"(k));
错误：asm 操作数类型大小 (1) 与约束 ‘r’ 隐含的类型/大小不匹配

为了在上面的 asm 语句中使用 ‘char’ 类型变量 “ci”、“cj” 和 “ck”，可以使用类似于以下代码段的代码段，
int temp = ci; asm("add.s32 %0,%1,%2;":"=r"(temp):"r"((int)cj),"r"((int)ck)); ci = temp;
另一个类型不匹配的示例

对于 ‘float’ 类型变量 “fi”，
asm("add.s32 %0,%1,%2;":"=r"(fi):"r"(j),"r"(k));
错误：asm 操作数类型大小 (4) 与约束 ‘r’ 隐含的类型/大小不匹配

2. 通知

2.1. 通知

本文档仅供参考，不得视为对产品的特定功能、条件或质量的保证。NVIDIA Corporation (“NVIDIA”) 对本文档中包含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对使用此类信息或因使用此类信息而可能导致的侵犯第三方专利或其他权利的后果或用途不承担任何责任。本文档不构成对开发、发布或交付任何材料（如下定义）、代码或功能的承诺。

NVIDIA 保留随时修改、增强、改进和对本文档进行任何其他更改的权利，恕不另行通知。

客户应在下订单前获取最新的相关信息，并应验证此类信息是否为最新且完整。

NVIDIA 产品根据订单确认时提供的 NVIDIA 标准销售条款和条件进行销售，除非 NVIDIA 和客户的授权代表签署的个别销售协议（“销售条款”）另有约定。NVIDIA 特此明确反对将任何客户通用条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档未直接或间接形成任何合同义务。

NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或故障可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对在此类设备或应用中包含和/或使用 NVIDIA 产品不承担任何责任，因此此类包含和/或使用由客户自行承担风险。

NVIDIA 不作任何陈述或保证，保证基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每个产品的全部参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合客户计划的应用和用途，并为应用执行必要的测试，以避免应用或产品的默认设置。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同条件和/或要求。NVIDIA 对可能基于或归因于以下原因的任何默认设置、损坏、成本或问题不承担任何责任：(i) 以任何与本文档相悖的方式使用 NVIDIA 产品；或 (ii) 客户产品设计。

在本文档下，未授予任何 NVIDIA 专利权、版权或其他 NVIDIA 知识产权的明示或暗示许可。NVIDIA 发布的关于第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成 NVIDIA 对其的保证或认可。使用此类信息可能需要获得第三方专利或其他知识产权的许可，或获得 NVIDIA 专利或其他 NVIDIA 知识产权的许可。

只有在事先获得 NVIDIA 书面批准的情况下，才允许复制本文档中的信息，并且复制时不得进行修改，必须完全遵守所有适用的出口法律和法规，并且必须附带所有相关的条件、限制和通知。

本文档和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”，单独称为“材料”）均按“原样”提供。NVIDIA 对材料不作任何明示、暗示、法定或其他形式的保证，并且明确否认对非侵权性、适销性和特定用途适用性的所有暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、偶然、惩罚性或后果性损害，无论何种原因造成，也无论责任理论如何）承担责任，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，但 NVIDIA 对本文所述产品的客户承担的累计和累积责任应根据产品的销售条款进行限制。

2.2. OpenCL

OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。

2.3. 商标

NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。