Завершение неактивных потоков во время сокращения

Question

Большинство сокращений, которые я когда-либо видел, выглядят так:

for( i = N; i > 0; i /=2 ) {
    if( tid < i )
        assign-shared;
    __syncthreads();
}
if( tid == 0 )
    copy-value-to-global;

Я только что изменил это на:

for( i = N; i > 0; i /= 2 ) {
    if( tid >= i )
        return;
    assign-shared;
    __syncthreads();
}
copy-value-to-global;

и заметил существенное повышение производительности.Есть ли недостаток в том, что потоки, которые больше не участвуют в восстановлении, возвращаются рано?

Answer 1

Второй сегмент кода обеспечивает лучшую производительность, поскольку неиспользуемые деформации не должны возвращаться и выполнять проверку ветвления.

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

Answer 2

Долан, в своем комментарии выше, поднимает вопрос о том, что схема, предложенная Уильямом Перселлом, заходит в тупик, согласно Могу ли я использовать __syncthreads () после удаления потоков? . Что касается этой проблемы, я бы сказал, что в соответствии с условными syncthreads & deadlock (или нет) код не будет тупиковым на большинстве графических процессоров, поскольку они поддерживают ранний выход, поскольку в этих графических процессорах аппаратное обеспечение поддерживает активный поток рассчитывать для каждого блока: этот счет затем используется для синхронизации барьера, а не начальный счет потока для блока.

Я рассмотрел пример reduce4 CUDA SDK и изменил его в соответствии с вопросом ОП. А именно, я сравниваю две функции __global__:

ORIGINAL

template <class T>
__global__ void reduce4(T *g_idata, T *g_odata, unsigned int N)
{
    extern __shared__ T sdata[];

    unsigned int tid    = threadIdx.x;                              // Local thread index
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;       // Global thread index - Fictitiously double the block dimension

    // --- Performs the first level of reduction in registers when reading from global memory. 
    T mySum = (i < N) ? g_idata[i] : 0;
    if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x];
    sdata[tid] = mySum;

    // --- Before going further, we have to make sure that all the shared memory loads have been completed
    __syncthreads();

    // --- Reduction in shared memory. Only half of the threads contribute to reduction.
    for (unsigned int s=blockDim.x/2; s>32; s>>=1)
    {
        if (tid < s) { sdata[tid] = mySum = mySum + sdata[tid + s]; }
        // --- At the end of each iteration loop, we have to make sure that all memory operations have been completed
        __syncthreads();
    }

    // --- Single warp reduction by loop unrolling. Assuming blockDim.x >64
    if (tid < 32) {
        sdata[tid] = mySum = mySum + sdata[tid + 32]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid + 16]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  8]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  4]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  2]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  1]; __syncthreads();
    }

    // --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of
    //     individual blocks
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

ОБНОВЛЕНО

template <class T>
__global__ void reduce4_deadlock_test(T *g_idata, T *g_odata, unsigned int N)
{
    extern __shared__ T sdata[];

    unsigned int tid    = threadIdx.x;                              // Local thread index
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;       // Global thread index - Fictitiously double the block dimension

    // --- Performs the first level of reduction in registers when reading from global memory. 
    T mySum = (i < N) ? g_idata[i] : 0;
    if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x];
    sdata[tid] = mySum;

    // --- Before going further, we have to make sure that all the shared memory loads have been completed
    __syncthreads();

    // --- Reduction in shared memory. Only half of the threads contribute to reduction.
    for (unsigned int s=blockDim.x/2; s>32; s>>=1)
    {
        if (tid >= s) return;
        sdata[tid] = mySum = mySum + sdata[tid + s];
        // --- At the end of each iteration loop, we have to make sure that all memory operations have been completed
        __syncthreads();
    }

    // --- Single warp reduction by loop unrolling. Assuming blockDim.x >64
    if (tid < 32) {
        sdata[tid] = mySum = mySum + sdata[tid + 32]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid + 16]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  8]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  4]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  2]; __syncthreads();
        sdata[tid] = mySum = mySum + sdata[tid +  1]; __syncthreads();
    }

    // --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of
    //     individual blocks
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
    }

Я проверил, что модифицированный код не блокируется на GT210, GT540M и Kepler K20c. Тем не менее, на карте Kepler ускорение модифицированной версии не так важно (времена в ms):

N          Original          Modified
131072     0.021             0.019
262144     0.030             0.032
524288     0.052             0.052
1048576    0.091             0.080
2097152    0.165             0.146
4194304    0.323             0.286
8388608    0.637             0.555
16777216   1.264             1.122
33554432   2.514             2.189

Я не проверял время для других архитектур, но, вероятно, риск застрять в тупике для некоторых графических процессоров не стоит достижимого ускорения (при условии, что достижимое ускорение остается того же порядка).

Answer 3

Поскольку вы уже выполняете оператор if со своим исходным кодом, я не вижу никаких недостатков.

Если результаты вашего оператора if не имели пространственной локализации (как правило, тот же результат по всему блоку), вы можете не увидеть никакого ускорения.Кроме того, ускорение может зависеть от возможностей вашего устройства: более ранние устройства CUDA могут не дать вам повышения производительности.