CUDA中的内联PTX汇编

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

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

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

本应用说明描述了如何将PTX汇编语言语句内联到CUDA代码中。

1.1. 汇编器(ASM)语句

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

asm("membar.gl;");

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

1.1.1. 参数

当我们在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;
}

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

__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)

常量表达式在编译时被求值，并应生成变量V的地址，其中：

• V 具有静态存储周期。

• V 的类型是 'const char 数组'。

• V 是常量初始化的。

• 如果 V 是一个静态类成员，那么 V 的初始化声明就是类内部的声明。

During translation, the compiler will replace a reference to the operand within the Assembler Template with the contents of V’s initializer, except for the last trailing zero. No constraint modifiers are allowed for this constraint. This constraint can only be used in device code.

（斜体中的术语是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. 命名空间冲突

如果在代码中多次调用并内联之前描述的cube函数，会产生关于临时寄存器t1重复定义的错误。为避免此错误，您需要：

• 不内联cube函数，或者，

• 将t1的使用嵌套在{}中，以便每次调用都有独立的作用域，例如：

  __device__ int cube (int x)
  {
    int y;
    asm("{\n\t"                        // 使用大括号创建局部作用域
        " reg .u32 t1;\n\t"            // 临时寄存器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));
  

  错误：在__device__/__global__函数中，一个asm操作数只能指定一个约束字母

• 仅允许标量变量作为asm操作数。特别是不允许使用如'struct'类型的聚合变量，例如：

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

  错误：asm操作数必须为标量类型

• PTX汇编约束所隐含的类型和大小必须与相关操作数匹配。以下是一个大小不匹配的示例：

  对于'char'类型变量"ci"：

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

  错误：asm操作数类型大小(1)与约束'r'隐含的类型/大小不匹配

  为了在上述汇编语句中使用'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公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

2.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

2.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。