如何为 Maxwell 及更高版本的 NVIDIA 架构编写基于 LOP3 的指令？

Question

Maxwell Architecture 在 PTX 汇编中引入了一条新指令，称为 LOP3，根据NVIDIA 博客：

“在对多个输入执行复杂逻辑运算时可以保存指令。”

在GTC 2016 上，一些 CUDA 开发人员设法使用此类指令加速了Tegra X1 处理器（Maxwell）的atan2f功能。

但是，下面的函数中定义的.cu文件导致未定义的定义__SET_LT和__LOP3_0xe2。

我是否必须在.ptx文件中定义它们？如果是这样，如何？

float atan2f(const float dy, const float dx) 
{
 float flag, z = 0.0f;
 __SET_LT(flag, fabsf(dy), fabsf(dx));

 uint32_t m, t1 = 0x80000000; 
 float t2 = float(M_PI) / 2.0f;

 __LOP3_0x2e(m, __float_as_int(dx), t1, __float_as_int(t2));
 float w = flag * __int_as_float(m) + float(M_PI)/2.0f; 

 float Offset = copysignf(w, dy);
 float t = fminf(fabsf(dx), fabsf(dy)) / fmaxf(fabsf(dx), fabsf(dy));

 uint32_t r, b = __float_as_int(flag) << 2;
 uint32_t mask = __float_as_int(dx) ^ __float_as_int(dy) ^ (~b);
 __LOP3_0xe2(r, mask, t1, __floast_as_int(t));

 const float p = fabsf(__int_as_float(r)) - 1.0f;
 return ((-0.0663f*(-p) + 0.311f) * (-p) + float(float(M_PI)/4.0)) * (*(float *)&r) + Offset;
}

编辑：

宏定义最终是：

#define __SET_LT(D, A, B) asm("set.lt.f32.f32 %0, %1, %2;" : "=f"(D) : "f"(A), "f"(B))
#define __SET_GT(D, A, B) asm("set.gt.f32.f32 %0, %1, %2;" : "=f"(D) : "f"(A), "f"(B))
#define __LOP3_0x2e(D, A, B, C) asm("lop3.b32 %0, %1, %2, %3, 0x2e;" : "=r"(D) : "r"(A), "r"(B), "r"(C))
#define __LOP3_0xe2(D, A, B, C) asm("lop3.b32 %0, %1, %2, %3, 0xe2;" : "=r"(D) : "r"(A), "r"(B), "r"(C))

Answer 1

的lop3.b32 PTX指令可以在3个变量A，B，和C.执行更多或更少的任意的布尔（逻辑）操作

为了设置要执行的实际操作，我们必须提供一个“查找表”立即参数（immLut-- 一个 8 位数量）。如文档中所述，计算immLut给定操作的必要参数的一种方法F(A,B,C)是在实际所需的方程中替换0xF0for A、0xCCforB和0xAAfor的值C。例如假设我们要计算：

F = (A || B) && (!C)   ((A or B) and (not-C))

然后我们将immLut通过以下方式计算参数：

immLut = (0xF0 | 0xCC) & (~0xAA)

注意，对于指定的方程F是一个布尔方程，处理所述参数A，B和C为布尔值，并产生一个真/假结果（F）。但是，要计算的方程immLut是按位逻辑运算。

对于上面的示例，immLut计算值为0x54

如果希望在普通 CUDA C/C++ 代码中使用 PTX 指令，可能最常见（也可以说是最简单）的方法是使用内联 PTX。内联 PTX有文档记录，还有其他问题讨论如何使用它（例如这个），所以我不会在这里重复。

这是上述示例案例的一个有效示例。请注意，此特定 PTX 指令仅在 cc5.0 和更高架构上可用，因此请确保至少针对该级别的目标进行编译。

$ cat t1149.cu
#include <stdio.h>

const unsigned char A_or_B_and_notC=((0xF0|0xCC)&(~0xAA));

__device__ int my_LOP_0x54(int A, int B, int C){
  int temp;
  asm("lop3.b32 %0, %1, %2, %3, 0x54;" : "=r"(temp) : "r"(A), "r"(B), "r"(C));
  return temp;
}

__global__ void testkernel(){

  printf("A=true, B=false, C=true,   F=%d\n", my_LOP_0x54(true, false, true));
  printf("A=true, B=false, C=false,  F=%d\n", my_LOP_0x54(true, false, false));
  printf("A=false, B=false, C=false, F=%d\n", my_LOP_0x54(false, false, false));
}


int main(){

  printf("0x%x\n", A_or_B_and_notC);
  testkernel<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_50 -o t1149 t1149.cu
$ ./t1149
0x54
A=true, B=false, C=true,   F=0
A=true, B=false, C=false,  F=1
A=false, B=false, C=false, F=0
$

由于immLut是 PTX 代码中的直接常量，我知道无法使用内联 PTX 将其作为函数参数传递 - 即使使用模板。根据您提供的链接，该演示文稿的作者似乎还为特定的所需立即数使用了单独定义的函数——在他们的情况下大概是 0xE2 和 0x2E。另外，请注意，我选择编写我的函数，以便它返回操作的结果作为函数的返回值。您链接的演示文稿的作者似乎通过函数参数将返回值传回。任何一种方法都应该是可行的。（实际上，他们似乎将__LOP3...代码编写为函数宏而不是普通函数。）

另请参阅此处了解 8 位真值表 ( immLut) 如何在源代码级别为 LOP3 工作的方法。