14

私は多くのモジュラー計算を行うことになっているGPUアルゴリズムに取り組んでいます。特に、有限体の行列に対するさまざまな演算は、長期的には次のようなプリミティブ演算になります。（a * b --c * d）mod mまたは（a * b + c）mod mここで、a、b、c、およびdはmを法とする剰余であり、mは32ビットの素数です。

実験を通じて、ハードウェアのGPUでは整数モジュロ（％）および除算演算がサポートされていないため、アルゴリズムのパフォーマンスは低速のモジュラー演算によってほとんど制限されることがわかりました。

誰かがCUDAで効率的なモジュラー計算を実現する方法を教えてくれたらありがたいですか？

これがCUDAにどのように実装されているかを確認するために、次のコードスニペットを使用します。

__global__ void mod_kernel(unsigned *gout, const unsigned *gin) {

unsigned tid = threadIdx.x;
unsigned a = gin[tid], b = gin[tid * 2], m = gin[tid * 3];

typedef unsigned long long u64;

__syncthreads();
unsigned r = (unsigned)(((u64)a * (u64)b) % m);
__syncthreads();
gout[tid] = r;
}

このコードは機能しないはずです。モジュール式の削減がCUDAにどのように実装されているかを確認したかっただけです。

これをcuobjdump--dump-sassで分解すると（アドバイスをくれたnjuffaに感謝します！）、次のように表示されます。

/*0098*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
/*00a0*/     /*0x1c315c4350000000*/     IMUL.U32.U32.HI R5, R3, R7;
/*00a8*/     /*0x1c311c0350000000*/     IMUL.U32.U32 R4, R3, R7;
/*00b0*/     /*0xfc01dde428000000*/     MOV R7, RZ;
/*00b8*/     /*0xe001000750000000*/     CAL 0xf8;
/*00c0*/     /*0x00000007d0000000*/     BPT.DRAIN 0x0;
/*00c8*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;

bar.red.popcへの2つの呼び出しの間に、いくつかの高度なアルゴリズム（約50命令以上）を実装する0xf8プロシージャへの呼び出しがあることに注意してください。mod（％）演算が遅いことを驚くことではありません

4

3 に答える 3

13

少し前に、GPUでモジュラー演算を試してみました。Fermi GPUでは、倍精度演算を使用して、コストのかかるdivおよびmod演算を回避できます。たとえば、モジュラー乗算は次のように実行できます。

// fast truncation of double-precision to integers
#define CUMP_D2I_TRUNC (double)(3ll << 51)
// computes r = a + b subop c unsigned using extended precision
#define VADDx(r, a, b, c, subop) \
    asm volatile("vadd.u32.u32.u32." subop " %0, %1, %2, %3;" :  \
            "=r"(r) : "r"(a) , "r"(b), "r"(c));

// computes a * b mod m; invk = (double)(1<<30) / m
__device__ __forceinline__ 
unsigned mul_m(unsigned a, unsigned b, volatile unsigned m,
    volatile double invk) { 

   unsigned hi = __umulhi(a*2, b*2); // 3 flops
   // 2 double instructions
   double rf = __uint2double_rn(hi) * invk + CUMP_D2I_TRUNC;
   unsigned r = (unsigned)__double2loint(rf);
   r = a * b - r * m; // 2 flops

   // can also be replaced by: VADDx(r, r, m, r, "min") // == umin(r, r + m);
   if((int)r < 0) 
      r += m;
   return r;
}

ただし、これは31ビット整数モジュロでのみ機能し（1ビットが重要でない場合）、事前に「invk」を事前に計算する必要もあります。これにより、私が達成できる最小限の指示が得られます。

SHL.W R2, R4, 0x1;
SHL.W R8, R6, 0x1;
IMUL.U32.U32 R4, R4, R6;
IMUL.U32.U32.HI R8, R2, R8;
I2F.F64.U32 R8, R8;
DFMA R2, R2, R8, R10;
IMAD.U32.U32 R4, -R12, R2, R4;
ISETP.GE.AND P0, pt, R4, RZ, pt;
@!P0 IADD R4, R12, R4;

アルゴリズムの説明については、私の論文 gpu_resultantsをご覧ください。（x y --z w）modmのような他の操作もここで説明されています。

好奇心から、モジュラー乗算を使用して結果のアルゴリズムのパフォーマンスを比較しました。

unsigned r = (unsigned)(((u64)a * (u64)b) % m);

mul_mで最適化されたバージョンに対して。

デフォルトの％演算を使用したモジュラー演算：

low_deg: 11; high_deg: 2481; bits: 10227
nmods: 330; n_real_pts: 2482; npts: 2495

res time: 5755.357910 ms; mod_inv time: 0.907008 ms; interp time: 856.015015 ms; CRA time: 44.065857 ms
GPU time elapsed: 6659.405273 ms;

mul_mを使用したモジュラー演算：

low_deg: 11; high_deg: 2481; bits: 10227
nmods: 330; n_real_pts: 2482; npts: 2495

res time: 1100.124756 ms; mod_inv time: 0.192608 ms; interp time: 220.615143 ms; CRA time: 10.376352 ms
GPU time elapsed: 1334.742310 ms;

したがって、平均して約5倍高速です。また、多数のmul_mod操作（ saxpyの例など）を備えたカーネルを使用して生の算術パフォーマンスを評価するだけでは、スピードアップが見られない場合があることにも注意してください。しかし、制御ロジック、同期バリアなどを備えた実際のアプリケーションでは、スピードアップは非常に顕著です。

于 2012-09-04T12:40:59.717 に答える
9

ハイエンドのFermiGPU（GTX 580など）は、このための出荷カードの中で最高のパフォーマンスを提供する可能性があります。符号付き除算とモジュロの処理には追加のオーバーヘッドがあるため、最高のパフォーマンスを得るには、すべての32ビットオペランドを「unsignedint」型にする必要があります。

コンパイラは、除数が固定された除算とモジュロ用の非常に効率的なコードを生成します。私が覚えているように、通常、FermiとKeplerでのマシン命令命令は約3〜5です。生成されたSASS（マシンコード）は、cuobjdump--dump-sassで確認できます。いくつかの異なる除数のみを使用する場合は、定数除数でテンプレート関数を使用できる可能性があります。

FermiとKepler全体で、可変除数を使用した符号なし32ビット演算に対して生成される16個のインラインSASS命令の順序を確認する必要があります。コードは整数乗算のスループットによって制限され、FermiクラスのGPUはハードウェアソリューションと競合します。整数乗算スループットが低下しているため、現在出荷されているKeplerクラスのGPUではパフォーマンスがいくらか低下しています。

[質問を明確にした後、後で追加：]

一方、符号なし64ビット除算と可変除数を使用したモジュロは、FermiとKeplerでは約65命令のサブルーチンと呼ばれます。彼らは最適に近いように見えます。Fermiでは、これはハードウェア実装とまだかなり競争力があります（64ビット整数除算は、これを組み込み命令として提供するCPUでは正確に超高速ではないことに注意してください）。以下は、説明で説明されている種類のタスクのために、しばらく前にNVIDIAフォーラムに投稿したコードです。コストのかかる除算を回避しますが、かなり大きなバッチのオペランドが同じ除算器を共有していることを前提としています。倍精度演算を使用します。これは、（コンシューマーカードとは対照的に）TeslaクラスのGPUで特に高速です。コードの大まかなテストのみを行いました。デプロイする前に、これをより注意深くテストすることをお勧めします。

// Let b, p, and A[i] be integers < 2^51
// Let N be a integer on the order of 10000
// for i from 1 to N
// A[i] <-- A[i] * b mod p

/*---- kernel arguments ----*/
unsigned long long *A;
double b, p; /* convert from unsigned long long to double before passing to kernel */
double oop;  /* pass precomputed 1.0/p to kernel */

/*---- code inside kernel -----*/
double a, q, h, l, rem;
const double int_cvt_magic = 6755399441055744.0; /* 2^52+2^51 */

a = (double)A[i];

/* approximate quotient and round it to the nearest integer */
q = __fma_rn (a * b, oop, int_cvt_magic);
q = q - int_cvt_magic;

/* back-multiply, representing p*q as a double-double h:l exactly */
h = p * q;
l = __fma_rn (p, q, -h);

/* remainder is double-width product a*b minus double-double h:l */
rem = __fma_rn (a, b, -h);
rem = rem - l;

/* remainder may be negative as quotient rounded; fix if necessary */
if (rem < 0.0) rem += p;

A[i] = (unsigned long long)rem;
于 2012-09-04T01:13:00.747 に答える
1

mod演算を効率的に実行するためのトリックがありますが、mだけが基数2の場合。

たとえば、x mod y == x＆（y-1）、ここでyは2^nです。ビット演算の実行が最速です。

そうでなければ、おそらくルックアップテーブル？以下は、効率的なモジュロ実装の説明に関するリンクです。それを最大限に活用するには、自分で実装する必要があるかもしれません。

modの効率的な計算

于 2012-09-04T03:00:05.293 に答える