Параллельное сокращение и поиск индекса по CUDA

Question

У меня есть массив значений 20 КБ, и я уменьшаю его на 50 блоков по 400 потоков в каждом. num_blocks = 50 и block_size = 400.

Мой код выглядит так:

getmax <<< num_blocks,block_size >>> (d_in, d_out1, d_indices);

__global__ void getmax(float *in1, float *out1, int *index)
{
    // Declare arrays to be in shared memory.
    __shared__ float max[threads];

    int nTotalThreads = blockDim.x;    // Total number of active threads
    float temp;
    float max_val;
    int max_index;
    int arrayIndex;

    // Calculate which element this thread reads from memory
    arrayIndex = gridDim.x*blockDim.x*blockIdx.y + blockDim.x*blockIdx.x + threadIdx.x;
    max[threadIdx.x] = in1[arrayIndex];
    max_val = max[threadIdx.x];
    max_index = blockDim.x*blockIdx.x + threadIdx.x;
    __syncthreads();

    while(nTotalThreads > 1)
    {
        int halfPoint = (nTotalThreads >> 1);
        if (threadIdx.x < halfPoint) 
        {
            temp = max[threadIdx.x + halfPoint];
            if (temp > max[threadIdx.x]) 
            {
                max[threadIdx.x] = temp;
                max_val = max[threadIdx.x];            
            }
        }
        __syncthreads();

        nTotalThreads = (nTotalThreads >> 1);    // divide by two.
    }

    if (threadIdx.x == 0)
    {
        out1[num_blocks*blockIdx.y + blockIdx.x] = max[threadIdx.x];
    }

    if(max[blockIdx.x] == max_val )
    {
        index[blockIdx.x] = max_index;    
    }
}

Проблема / проблема здесь в том, что в какой-то момент значение nTotalThreads не равно степени 2, что приводит к значению мусора для индекса. Массив out1 дает мне максимальное значение в каждом блоке, которое является правильным и проверенным. Но значение индекса неверно. Например: максимальное значение в первом блоке происходит при index = 40, но ядро ​​дает значения индекса как 15. Аналогично, значение max во втором блоке равно 440, но ядро ​​дает 416.

Есть предложения ??

Answer 1

Должно быть легко убедиться, что nTotalThreads всегда имеет степень 2.

Сделайте первое сокращение в особом случае, когда nTotalThreads получит степень 2. Например, поскольку вы начинаете с 400 потоков вблок, сделать первое сокращение с 256 потоков.Потоки 0-199 уменьшатся с двух значений, а потокам 200-255 просто не нужно будет делать сокращение на этом начальном этапе.С этого момента все будет в порядке.

Answer 2

Второе предложение Джеффа.

Взгляните на функцию сокращения библиотеки CUDA Thrust. Показано, что она обладает эффективностью 95 +% по сравнению с сильно настроенными ядрами и довольно гибкая.и простой в использовании.

Answer 3

Вы уверены, что вам действительно нужно, чтобы «проблема» «nTotalThreads» не была степенью 2? Это делает код менее читабельным, и я думаю, что это может также повлиять на производительность. В любом случае, если вы подставите

nTotalThreads = (nTotalThreads >> 1);

с

nTotalThreads = (nTotalThreads +1) >> 1;

это должно решить одну ошибку, связанную с этой «проблемой».

Francesco

Answer 4

проверьте мое ядро ​​.Вы можете поместить свои blockresults в массив (который может быть в глобальной памяти) и получить результат в глобальной памяти

И посмотреть, как я вызываю его в коде хоста:

sumSeries<<<dim3(blockCount),dim3(threadsPerBlock)>>>(deviceSum,threadsPerBlock*blockCount);