如何为 Maxwell 和 NVIDIA 架构编写基于 LOP3 的指令？

Question

Maxwell Architecture 在 PTX 汇编中引入了一条名为 LOP3 的新指令，它根据 NVIDIA blog:

"Can save instructions when performing complex logic operations on multiple inputs."

Tegra X1 处理器 (Maxwell) 的 GTC 2016, some CUDA developers managed to accelerated the 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 instruction 可以对 3 个变量 A、B 和 C 执行或多或少的任意布尔（逻辑）运算。

为了设置要执行的实际操作，我们必须提供一个"lookup-table"立即数参数（immLut——一个8位的量）。如 the documentation 中所示，为给定操作 F(A,B,C) 计算必要的 immLut 参数的方法是将 0xF0 的值替换为 A，0xCC 用于 B，0xAA 用于实际所需的方程式中的 C。例如假设我们要计算：

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

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

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

请注意，F 的指定方程是布尔方程，将参数 A、B 和 C 视为布尔值，并生成 true/false 结果 (F)。但是，计算 immLut 的方程式是按位逻辑运算。

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

如果希望在普通 CUDA C/C++ 代码中使用 PTX 指令，可能最常见（也可以说是最简单）的方法是使用 inline PTX. Inline PTX is documented, and there are other questions discussing how to use it (such as this one)，所以我将此处不再赘述。

这是上述示例案例的一个有效示例。请注意，此特定 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 将其作为函数参数传递 - 即使使用模板。根据您的 provided link，该演示文稿的作者似乎还为特定的所需立即值使用了单独定义的函数——在他们的情况下大概是 0xE2 和 0x2E。另外，请注意，我选择编写我的函数，以便它 return 作为函数 return 值的操作结果。您链接的演示文稿的作者似乎正在通过函数参数传回 return 值。任何一种方法都应该可行。（事实上，他们似乎将 __LOP3... 代码编写为功能性宏而不是普通函数。）

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

如何为 Maxwell 和 NVIDIA 架构编写基于 LOP3 的指令？

How to write LOP3 based instructions for Maxwell and up NVIDIA Architecture?

cuda

nvidia