Сокращение в CUDA

Question

Я только начинаю изучать программирование на CUDA, и у меня возникло недоразумение по поводу сокращения.

Я знаю, что глобальная память имеет большую задержку при посещении по сравнению с разделяемой памятью, но могу ли я использовать глобальную память для (по крайней мере) имитации поведения, подобного разделяемой памяти?

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

    __global__ void parallelSum(unsigned int* array) {

    unsigned int totalThreadsNum = gridDim.x * blockDim.x;
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;

    int i = totalThreadsNum / 2;
    while (i != 0) {
            if (idx < i) {
                array[idx] += array[idx + i];
        }
        __syncthreads();
        i /= 2;
    }
}

Я сравнил результат этого кода и результат, сгенерированный последовательно на хосте, и странная вещь: иногда результаты совпадают, но иногда они, очевидно, отличаются. Есть ли какая-либо причина, связанная с использованием глобальной памяти здесь?

Answer 1

Лучше всего начать с примера сокращения в примерах CUDA. Пример сканирования также хорош для изучения принципов параллельных вычислений на архитектуре пропускной способности.

Тем не менее, если вы на самом деле просто хотите использовать оператор сокращения в своем коде, вам следует взглянуть на Thrust (вызовы с хоста, кроссплатформенность) и CUB (специфично для CUDA GPU).

Чтобы посмотреть на ваши конкретные вопросы:

• Нет причин, по которым вы не можете использовать глобальную память для сокращения, пример кода в наборе инструментов проходит различные уровни оптимизации, но в каждом случае данные запускаются в глобальной памяти.
• Ваш код неэффективен (см. Пример в наборе инструментов для более подробной информации об эффективности работы!).
• Ваш код пытается связаться между потоками в разных блоках без надлежащей синхронизации; __syncthreads() синхронизирует только потоки в определенном блоке, а не в разных блоках (это было бы невозможно, по крайней мере, в общем, поскольку вы склонны переподписывать GPU, что означает, что не все блоки будут работать в данный момент времени).

Последний пункт самый важный. Если поток в блоке X хочет прочитать данные, записанные в блоке Y, то вам нужно разбить это на два запуска ядра, поэтому типичное параллельное сокращение использует многофазный подход: уменьшить пакеты внутри блоков, а затем уменьшить между пакетами,

Answer 2

Том уже ответил на этот вопрос. В своем ответе он рекомендует использовать Thrust или CUB для выполнения сокращений в CUDA.

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

#define CUB_STDERR

#include <stdio.h>

#include <thrust/device_ptr.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>

#include <cub/device/device_reduce.cuh>

#include "TimingGPU.cuh"
#include "Utilities.cuh"

using namespace cub;

/********/
/* MAIN */
/********/
int main() {

    const int N = 8388608;

    gpuErrchk(cudaFree(0));

    float *h_data       = (float *)malloc(N * sizeof(float));
    float h_result = 0.f;

    for (int i=0; i<N; i++) {
        h_data[i] = 3.f;
        h_result = h_result + h_data[i];
    }

    TimingGPU timerGPU;

    float *d_data;          gpuErrchk(cudaMalloc((void**)&d_data, N * sizeof(float)));
    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    /**********/
    /* THRUST */
    /**********/
    timerGPU.StartCounter();
    thrust::device_ptr<float> wrapped_ptr = thrust::device_pointer_cast(d_data);
    float h_result1 = thrust::reduce(wrapped_ptr, wrapped_ptr + N);
    printf("Timing for Thrust = %f\n", timerGPU.GetCounter());

    /*******/
    /* CUB */
    /*******/
    timerGPU.StartCounter();
    float           *h_result2 = (float *)malloc(sizeof(float));
    float           *d_result2; gpuErrchk(cudaMalloc((void**)&d_result2, sizeof(float)));
    void            *d_temp_storage = NULL;
    size_t          temp_storage_bytes = 0;

    DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N);
    gpuErrchk(cudaMalloc((void**)&d_temp_storage, temp_storage_bytes));
    DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N);

    gpuErrchk(cudaMemcpy(h_result2, d_result2, sizeof(float), cudaMemcpyDeviceToHost));

    printf("Timing for CUB = %f\n", timerGPU.GetCounter());

    printf("Results:\n");
    printf("Exact: %f\n", h_result);
    printf("Thrust: %f\n", h_result1);
    printf("CUB: %f\n", h_result2[0]);

}

Пожалуйста, обратите внимание, что CUB может быть несколько быстрее, чем Thrust, из-за другой базовой философии, поскольку CUB оставляет детали, критичные для производительности, такие как точный выбор алгоритма и степень параллелизма, в руках пользователя. Таким образом, эти параметры могут быть настроены для максимизации производительности для конкретной архитектуры и приложения.

Сравнение для вычисления евклидовой нормы массива сообщается на CUB в действии - несколько простых примеров с использованием библиотеки шаблонов CUB.