red.shared指令

Question

我正在学习CUDA，在阅读了PTX手册后，我发现有一个名为red.shared的指令，它可以在经度上执行归约。

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

Answer 1

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