CUDA: как сложить все элементы массива в одно число в GPU?

Прежде всего, позвольте мне заявить, что я полностью осознаю, что мой вопрос уже задан: Сокращение блока в CUDA Однако, как я надеюсь прояснить, мой вопрос является продолжением этого, и у меня есть особые потребности, которые делают Решение, найденное этим ОП, не подходит.

Итак, позвольте мне объяснить. В моем текущем коде я запускаю ядро ​​Cuda на каждой итерации цикла while, чтобы выполнить некоторые вычисления над значениями массива. В качестве примера, подумайте об этом следующим образом:

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    iteration++;
}

Однако затем я должен выполнить, казалось бы, трудную задачу для GPU. На каждой итерации цикла while, который вызывает ядро, я должен суммировать все значения, сгенерированные в odata и сохранить результат в intмассив называется resultв позиции в таком массиве, которая соответствует текущей итерации. Это должно быть выполнено внутри ядра или, по крайней мере, все еще в графическом процессоре, потому что из-за ограничений производительности я могу получить только result Массив в самом конце после всех итераций.

Неправильная наивная попытка выглядит примерно так:

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata, int* result)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    result[iteration] = 0;
    for(int j=0; j < max_iterations; j++)
    {
        result[iteration] += odata[j];            
    }

    iteration++;
}

Конечно, приведенный выше код не работает из-за того, что графический процессор распределяет код по потокам. Чтобы научиться правильно это делать, я читал здесь на сайте другие вопросы о сокращении массивов с использованием CUDA. В частности, я нашел упоминание в очень хорошем pdf от NVIDIA на эту тему, которое также обсуждается в предыдущем вопросе SO, который я упоминал в начале: http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

Однако, хотя я полностью понимаю шаги кода, описанного на таких слайдах, а также общие оптимизации, я не понимаю, как этот подход может суммировать массив до одного числа, если код фактически выдает полный массив (и одно из неясных измерений). Может ли кто-нибудь пролить свет на это и показать мне пример того, как это будет работать (например, как получить одно число из выходного массива)?

Теперь, возвращаясь к этому вопросу, который я упоминал в начале ( сокращение блока в CUDA). Обратите внимание, что принятый ответ просто предлагает прочитать PDF-файл, который я связал выше, - который не говорит о том, что делать с выходным массивом, сгенерированным кодом. В комментариях ОП упоминается, что он / она смог завершить работу, суммируя выходной массив в ЦП - что я не могу сделать, так как это означало бы загрузку выходного массива на каждой итерации моего цикла while. Наконец, третий ответ в этой ссылке предлагает использовать библиотеку для достижения этой цели, но я заинтересован в изучении собственного способа сделать это.

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

Источник

2 ответа

Решение

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

На реальный вопрос здесь, вы можете делать то, что вы хотите, если вы делаете что-то вроде этого:

__global__ void kernel(....., int* iter_result, int iter_num) {

    // Your calculations first so that each thread holds its result

    // Block wise reduction so that one thread in each block holds sum of thread results

    // The one thread holding the adds the block result to the global iteration result
    if (threadIdx.x == 0)
        atomicAdd(iter_result + iter_num, block_ressult);
}

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

Источник

Если вы добавите 2 смежных числа и сохраните результат в любом из слотов, где вы сохраняете эти числа, вам нужно будет только запустить, несколько раз одно и то же ядро, чтобы продолжать уменьшать в 2 раза суммы массива, как в этом примере:

Массив для суммирования значений:

[·1,·2,·3,·4,·5,·6,·7,·8,·9,·10]

Сначала запустите n/2 потоков, суммируйте смежные элементы массива и сохраните их слева от каждого, теперь массив будет выглядеть так:

[·3,2,·7,4,·11,6,·15,8,·19,10]

Запустите то же ядро, запустите н /4 нити, теперь добавьте каждые 2 элемента и сохраните его в крайнем левом элементе, массив теперь будет выглядеть так:

[·10,2,7,4,·26,6,15,8,·19,10]

Запустите то же ядро, запустите н /8 потоков, теперь добавьте каждые 4 элемента и сохраните в крайнем левом элементе массива, чтобы получить:

[·36,2,7,4,26,6,15,8,·19,10]

Запустите последний раз, один поток, чтобы добавить каждые 8 ​​элементов и сохранить в крайнем левом элементе массива, чтобы получить:

[55,2,7,4,26,6,15,8,19,10]

Таким образом, вам нужно всего лишь запустить ваше ядро ​​с некоторыми потоками в качестве параметров, чтобы в конце получить избыточность, в первом элементе (55) посмотрите на "точки" (·), чтобы увидеть, какие элементы в массиве "активны". суммировать их, каждый пробег.

Источник
Другие вопросы по тегам