Дополнительное использование регистра с помощью if

Question

Я работал над большим ядром 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]

Answer 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

Не следует делать вывод, что различия всегда будут связаны с оптимизацией компилятора, потому что это сильно зависит от кода и компилятора. Но в этом и заключается разница.