cuda - red.shared 命令

Question

私は CUDA を学習しています。PTX マニュアルを調べたところ、ワープ全体でリダクションを実行する red.shared という命令があることがわかりました。

ハードウェアがリダクションをネイティブでサポートしているかどうかに興味があります。もしそうなら、CUDAコードでどのように使用できますか? おそらく誰かがそれを実験しましたか？

Answer 1

実際、好奇心から「赤」の命令も試してみました。これが Kepler でどのようになっているのかはわかりませんが、Fermi アーキテクチャでは「赤い」命令が別の命令のシーケンスにマップされるだけです。たぶん彼らは将来のGPUのためにそれを残しました. これが私が遊んだコードです：

#define WS 32
#define HF 16

__global__ void test_red_kernel(unsigned *g_R, const unsigned *g_U) {

  extern __shared__ unsigned shared[];

  unsigned thid = threadIdx.x, bidx_x = blockIdx.x;
  unsigned *r = shared;
  unsigned ofs = bidx_x << 7, thid_in_warp = thid & WS-1;

  unsigned a = (g_U + ofs)[thid];

  volatile unsigned *t = (volatile unsigned *)r + HF + UMUL(thid >> 5,
        WS + HF + 1) + thid_in_warp;

  t[-HF] = 0;
t[0] = a;
// warp reduction
a = a + t[-HF], t[0] = a;
a = a + t[-8], t[0] = a;
a = a + t[-4], t[0] = a;
a = a + t[-2], t[0] = a;
a = a + t[-1], t[0] = a;

CU_SYNC

volatile unsigned *t2 = r + HF + UMUL(WS*4 >> 5, WS + HF + 1);

if(thid < 4) {

    unsigned loc_ofs = HF + WS-1 + UMUL(thid, WS + HF + 1);
    unsigned a2;

    volatile unsigned *ps = t2 + thid;
    ps[-2] = 0;

    a2 = r[loc_ofs]; ps[0] = a2;
    a2 = a2 + ps[-2], ps[0] = a2;
    a2 = a2 + ps[-1], ps[0] = a2;
}

CU_SYNC

a = a + t2[(thid >> 5) - 1];

unsigned b;      
asm volatile("mov.u32 %r11, shared;" : );
asm volatile("red.shared.add.u32 [%r11], %0;" :
            "+r"(b) : );

b = r[0]; // results of 'red.shared', compare it with a

(g_R + ofs)[thid] = a - b; 
}

「red」命令がハードウェアにどのように実装されているかを確認するには、生成された「cubin」ファイルで cuobjdump ツールを使用できます (nvcc でオプション -keep を使用)。