cuda - red.shared 指令

Question

我正在学习 CUDA，在查看了 PTX 手册后，我发现有一条名为 red.shared 的指令可以在扭曲中执行缩减。

我很好奇硬件是否具有对减少的本机支持。如果是这样，如何在 CUDA 代码中使用它？也许有人已经尝试过了？

Answer 1

实际上恰好是这样，我出于好奇也尝试了“红色”指令。我不知道这在开普勒上是怎么回事，但在费米架构上，“红色”指令只是映射到另一个指令序列。也许他们把它留给了未来的 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; 
}

要查看“红色”指令是如何在硬件中实现的，您可以在生成的“cubin”文件上使用 cuobjdump 工具（使用选项 -keep 和 nvcc）