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

2 番目のケースでは、ループはそのまま残ります。

    .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 に答える