О компактной работе в cuddpp

Question

Следующая функция ядра - это компактная операция в cudpp, cudad library (http://gpgpu.org/developer/cudpp).

Мой вопрос: почему разработчик повторяет часть письма 8 раз? И почему это может улучшить производительность?

И почему один поток обрабатывает 8 элементов, почему не каждая нить обрабатывает один элемент?

template <class T, bool isBackward> 
__global__ void compactData(T     *d_out, 
         size_t    *d_numValidElements, 
         const unsigned int *d_indices, // Exclusive Sum-Scan Result 
         const unsigned int *d_isValid, 
         const T   *d_in, 
         unsigned int  numElements) 
{ 
    if (threadIdx.x == 0) 
    { 
     if (isBackward) 
      d_numValidElements[0] = d_isValid[0] + d_indices[0]; 
    else 
     d_numValidElements[0] = d_isValid[numElements-1] + d_indices[numElements-1]; 
    } 

    // The index of the first element (in a set of eight) that this 
    // thread is going to set the flag for. We left shift 
    // blockDim.x by 3 since (multiply by 8) since each block of 
    // threads processes eight times the number of threads in that 
    // block 
    unsigned int iGlobal = blockIdx.x * (blockDim.x << 3) + threadIdx.x; 

    // Repeat the following 8 (SCAN_ELTS_PER_THREAD) times 
    // 1. Check if data in input array d_in is null 
    // 2. If yes do nothing 
    // 3. If not write data to output data array d_out in 
    // the position specified by d_isValid 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal];  
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
    iGlobal += blockDim.x; 
    if (iGlobal < numElements && d_isValid[iGlobal] > 0) { 
     d_out[d_indices[iGlobal]] = d_in[iGlobal]; 
    } 
} 

Answer 1

Мой вопрос, почему разработчик повторяет написание часть 8 раз? И почему это может улучшить производительность?

Как указано в @torrential_coding, разворот цикла может помочь в производительности. В частности, в таком случае, когда петля очень плотная (в ней мало логики). Однако кодер должен был использовать поддержку CUDA для автоматического разворачивания цикла, а не вручную.

И почему один поток обрабатывает 8 элементов, почему не каждая нить обрабатывает один элемент?

Там могут быть некоторые небольшой прирост производительности в расчете только полный индекс iGlobal и проверка threadIdx.x нуля на каждые 8 ​​операций, а не для каждой операции, которые должны быть сделано, если каждое ядро ​​сделал только один элемент.