MaxwellとNVIDIAアーキテクチャのためのLOP3ベースの命令を書くには？

Question

マクスウェルアーキテクチャNVIDIA blogに応じLOP3呼ばPTXアセンブリ内の新しい命令を導入しました：

「とは、複数の入力に対して複雑な論理演算 を実行するときの手順を保存することができます。」 GTC 2016

で

、いくつかのCUDA開発者は、このような命令でTegraはX1のプロセッサ用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.b32PTX instructionは、3つの変数に多かれ少なかれ任意のブール（論理）操作を行うことができますA、B、およびCを含む。

実際の操作を実行するには、「ルックアップテーブル」即値引数（immLut - 8ビット数）を指定する必要があります。 the documentationに示すように、所定の操作F(A,B,C)に必要immLut引数を計算する方法は、実際の所望式CためAため0xF0、Bため0xCC、及び0xAAの値を置換することです。例えば、我々が計算したいとします

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

その後、我々はによってimmLut引数を計算します：Fに指定された式は、引数A、Bの治療、ブール式であることを

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

を注意し、ブール値としてCを返し、真偽結果（F）を生成する。ただし、immLutを計算する式はビット単位で論理演算です。

通常のCUDA C/C++コードにPTX命令を使用することが望ましいかどう上記例えば、immLutは、おそらく最も一般的な（そしておそらく最も簡単な）方法をが0x54

の計算値を有するであろうがあろうinline PTXを使用してください。インラインPTX is documentedとそれを使用する方法について他の質問があります（this oneなど）ので、ここではこれを繰り返さないでください。

上記の例の作業例です。この特定のPTX命令は、cc5でのみ使用可能であることに注意してください。0以上のアーキテクチャをサポートしているので、少なくともそのレベルのターゲットでコンパイルしてください。 immLutので

$ 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 
$ 

はPTXコードにおいて即値定数であり、Iは、関数のパラメータとしてこれを渡すために、インラインPTXを使用しない方法を知っている - テンプレートを使用しても。 provided linkに基づいて、そのプレゼンテーションの著者は、特定の望ましい即時値（おそらく0xE2と0x2E）のために別々に定義された関数を使用したようです。また、関数の戻り値として操作の結果を返すように関数を記述することにしました。リンクしたプレゼンテーションの作者は、戻り値を関数パラメータを介して返すように見えます。いずれの方法も有効であるはずです。実際には、__LOP3...というコードを通常の関数ではなくマクロと書いているようです。