1

我正在研究一个大型 cuda 内核，我注意到内核每个线程使用 43 个寄存器。为了了解发生了什么，我编写了一个较小的程序来计算寄存器的使用情况。我注意到每当我使用 时if，寄存器的使用率都会上升。小代码如下：

#include <limits.h>
#include <stdio.h>
#include <fstream>
#include <iostream>
#include <cstdlib>
#include <stdint.h>

using namespace std;

__global__ void test_ifs(unsigned int* result){
  unsigned int k = 0;
  for(int j=0;j<MAX_COMP;j++){
    //if(j <= threadIdx.x%MAX_COMP){                                                                                                                                                                                                          
      k += j;
      //}                                                                                                                                                                                                                                     
  }
  result[threadIdx.x] = k;
}

int main(){
  unsigned int* result;
  cudaError_t e1 = cudaMalloc((void**) &result, THREADSPERBLOCK*sizeof(unsigned int));
  if(e1 == cudaSuccess){
    test_ifs<<<1, THREADSPERBLOCK>>>(result);
    cudaError_t e2 = cudaGetLastError();
    if(e2 == cudaSuccess){
    }
    else{
      cout << "kernel failed to launch" << endl;
    }
  }
  else{
    cout << "Failed to allocate results memory" << endl;
  }
}

当我编译这段代码时，每个线程使用 5 个寄存器

ptxas info    : Compiling entry function '_Z8test_ifsPj' for 'sm_20'
ptxas info    : Function properties for _Z8test_ifsPj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 5 registers, 40 bytes cmem[0]

但是，如果我取消注释if，每个线程使用 8 个寄存器。谁能向我解释发生了什么事？

ptxas info    : Compiling entry function '_Z8test_ifsPj' for 'sm_20'
ptxas info    : Function properties for _Z8test_ifsPj
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 8 registers, 40 bytes cmem[0]
4

1 回答 1

1

您在此示例中看到的行为是由于编译器优化造成的。在简单循环的情况下，可以在编译时计算循环的结果并将整个循环代码替换为一个常量，而在包含 if 语句的情况下，循环的结果取决于一个变量，其值不是编译器知道，并且循环必须保留。

为了证明确实如此，让我们看一下内核的一个稍微修改过的版本：

#define MAX_COMP (32)

template<unsigned int s>
__global__ void test_ifs(unsigned int * result){
    unsigned int k = 0;
    for(int j=0;j<MAX_COMP;j++){
        switch (s) {
            case 1:
                if (j <= threadIdx.x%MAX_COMP){ k += j; }
                break;            

            case 0:
                { k += j; }
        }
    }
    result[threadIdx.x] = k;
}

template __global__ void test_ifs<0>(unsigned int *);
template __global__ void test_ifs<1>(unsigned int *);

以及它发出的 PTX。对于第一种情况：

    .entry _Z8test_ifsILj0EEvPj (
        .param .u32 __cudaparm__Z8test_ifsILj0EEvPj_result)
    {
    .reg .u16 %rh<3>;
    .reg .u32 %r<6>;
    .loc    14  4   0
$LDWbegin__Z8test_ifsILj0EEvPj:
    .loc    14  16  0
    mov.u32     %r1, 496;  <--- here the loop has been replaced with 496
    ld.param.u32    %r2, [__cudaparm__Z8test_ifsILj0EEvPj_result];
    mov.u16     %rh1, %tid.x;
    mul.wide.u16    %r3, %rh1, 4;
    add.u32     %r4, %r2, %r3;
    st.global.u32   [%r4+0], %r1;
    .loc    14  17  0
    exit;
$LDWend__Z8test_ifsILj0EEvPj:
    } // _Z8test_ifsILj0EEvPj

第二种情况循环保持不变：

    .entry _Z8test_ifsILj1EEvPj (
        .param .u32 __cudaparm__Z8test_ifsILj1EEvPj_result)
    {
    .reg .u32 %r<11>;
    .reg .pred %p<4>;
    .loc    14  4   0
$LDWbegin__Z8test_ifsILj1EEvPj:
    cvt.u32.u16     %r1, %tid.x;
    and.b32     %r2, %r1, 31;
    mov.s32     %r3, 0;
    mov.u32     %r4, 0;
$Lt_1_3842:
 //<loop> Loop body line 4, nesting depth: 1, iterations: 32
    .loc    14  7   0
    add.u32     %r5, %r3, %r4;
    setp.le.u32     %p1, %r3, %r2;
    selp.u32    %r4, %r5, %r4, %p1;
    add.s32     %r3, %r3, 1;
    mov.u32     %r6, 32;
    setp.ne.s32     %p2, %r3, %r6;
    @%p2 bra    $Lt_1_3842;
    .loc    14  16  0
    ld.param.u32    %r7, [__cudaparm__Z8test_ifsILj1EEvPj_result];
    mul24.lo.u32    %r8, %r1, 4;
    add.u32     %r9, %r7, %r8;
    st.global.u32   [%r9+0], %r4;
    .loc    14  17  0
    exit;
$LDWend__Z8test_ifsILj1EEvPj:
    } // _Z8test_ifsILj1EEvPj

您不应断定差异始终是由于编译器优化造成的，因为这在很大程度上取决于代码和编译器。但在这种情况下，这就是区别。

于 2012-07-20T05:56:21.370 回答