CUDA循环在三角形区域展开

Question

是否可以在三角形区域展开循环，例如：

for (int i = 0; i < ROW_LENGTH; i++)
{
    for (int j = 0; j < i; j++)
    {
        // Some array operation here
    }
}

其中ROW_LENGTH是编译时定义的常量？就目前而言，我认为这是不可能的，因为 i 随着程序的执行而改变（更重要的是，它在编译时不是常量）。我想您可以将 2D 数组视为 1D 数组，从 0 迭代到 (ROW_LENGTH^2)/2，然后尝试一些数学技巧来获取索引，但是额外的操作破坏了首先循环展开。

Answer 1

CUDA 7.0 编译器将在我的测试中展开它。循环索引在编译时都是已知的，所以没有理由不能这样做。

考虑以下代码，它将 a 的三角形部分设置为 1。

#define ROW_LENGTH 4
__global__ void triUnrollTest1(float* a) {
   #pragma unroll
   for (int i = 0; i < ROW_LENGTH; i++)
   {
      #pragma unroll
      for (int j = 0; j < i; j++)
      {
         a[i * ROW_LENGTH + j] = 1.f;
      }
   }
}

因为 ROW_LENGTH 只有 4 个，我们可以自己展开这个：

__global__ void triUnrollTest2(float* a) {
   a[1 * ROW_LENGTH + 0] = 1.f;
   a[2 * ROW_LENGTH + 0] = 1.f;
   a[2 * ROW_LENGTH + 1] = 1.f;
   a[3 * ROW_LENGTH + 0] = 1.f;
   a[3 * ROW_LENGTH + 1] = 1.f;
   a[3 * ROW_LENGTH + 2] = 1.f;
}

使用 CUDA 7.0 为 SM 35 编译： nvcc -arch=sm_35 -c triUnroll.cu

然后转储 SASS 汇编程序： cuobjdump --dump-sass triUnroll.o

我们得到：

code for sm_35
        Function : _Z14triUnrollTest1Pf
.headerflags    @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
                                                          /* 0x08b8b8a0b010a000 */
/*0008*/                   MOV R1, c[0x0][0x44];          /* 0x64c03c00089c0006 */
/*0010*/                   MOV R0, c[0x0][0x140];         /* 0x64c03c00281c0002 */
/*0018*/                   IADD R2.CC, R0, 0x10;          /* 0xc0840000081c0009 */
/*0020*/                   MOV32I R0, 0x3f800000;         /* 0x741fc000001fc002 */
/*0028*/                   IADD.X R3, RZ, c[0x0][0x144];  /* 0x60804000289ffc0e */
/*0030*/                   ST.E [R2], R0;                 /* 0xe4800000001c0800 */
/*0038*/                   ST.E [R2+0x10], R0;            /* 0xe4800000081c0800 */
                                                          /* 0x080000b810b8b8b8 */
/*0048*/                   ST.E [R2+0x14], R0;            /* 0xe48000000a1c0800 */
/*0050*/                   ST.E [R2+0x20], R0;            /* 0xe4800000101c0800 */
/*0058*/                   ST.E [R2+0x24], R0;            /* 0xe4800000121c0800 */
/*0060*/                   ST.E [R2+0x28], R0;            /* 0xe4800000141c0800 */
/*0068*/                   EXIT;                          /* 0x18000000001c003c */
/*0070*/                   BRA 0x70;                      /* 0x12007ffffc1c003c */
/*0078*/                   NOP;                           /* 0x85800000001c3c02 */
        .....................................


        Function : _Z14triUnrollTest2Pf
.headerflags    @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
                                                          /* 0x08b8b8a0b010a000 */
/*0008*/                   MOV R1, c[0x0][0x44];          /* 0x64c03c00089c0006 */
/*0010*/                   MOV R0, c[0x0][0x140];         /* 0x64c03c00281c0002 */
/*0018*/                   IADD R2.CC, R0, 0x10;          /* 0xc0840000081c0009 */
/*0020*/                   MOV32I R0, 0x3f800000;         /* 0x741fc000001fc002 */
/*0028*/                   IADD.X R3, RZ, c[0x0][0x144];  /* 0x60804000289ffc0e */
/*0030*/                   ST.E [R2], R0;                 /* 0xe4800000001c0800 */
/*0038*/                   ST.E [R2+0x10], R0;            /* 0xe4800000081c0800 */
                                                          /* 0x080000b810b8b8b8 */
/*0048*/                   ST.E [R2+0x14], R0;            /* 0xe48000000a1c0800 */
/*0050*/                   ST.E [R2+0x20], R0;            /* 0xe4800000101c0800 */
/*0058*/                   ST.E [R2+0x24], R0;            /* 0xe4800000121c0800 */
/*0060*/                   ST.E [R2+0x28], R0;            /* 0xe4800000141c0800 */
/*0068*/                   EXIT;                          /* 0x18000000001c003c */
/*0070*/                   BRA 0x70;                      /* 0x12007ffffc1c003c */
/*0078*/                   NOP;                           /* 0x85800000001c3c02 */
        .....................................

显然两者是一样的，而且展开得很好。有趣的是，当我不小心用 6.5 编译我的第一个答案时，编译器没有展开，所以我想在这种情况下更新是值得的！

CUDA循环在三角形区域展开

CUDA loop unrolling on triangular region

cuda

loop-unrolling