Продукт CUDA Dot

Question

Я пытаюсь реализовать классическое ядро ​​точечного произведения для массивов двойной точности с атомарным вычислением окончательной суммы по различным блокам.Я использовал atomicAdd для двойной точности, как указано на странице 116 руководства по программированию. Возможно, я делаю что-то не так. Частичные суммы по потокам в каждом блоке вычисляются правильно, но после слов атомная операция, кажется, не работает должным образомтак как каждый раз, когда я запускаю свое ядро ​​с одними и теми же данными, я получаю разные результаты.Буду признателен, если кто-то обнаружит ошибку или предложит альтернативное решение!Вот мое ядро:

__global__ void cuda_dot_kernel(int *n,double *a, double *b, double *dot_res)
{
    __shared__ double cache[threadsPerBlock]; //thread shared memory
    int global_tid=threadIdx.x + blockIdx.x * blockDim.x;
    int i=0,cacheIndex=0;
    double temp = 0;
    cacheIndex = threadIdx.x;
    while (global_tid < (*n)) {
        temp += a[global_tid] * b[global_tid];
        global_tid += blockDim.x * gridDim.x;
    }
    cache[cacheIndex] = temp;
    __syncthreads();
    for (i=blockDim.x/2; i>0; i>>=1) {
        if (threadIdx.x < i) {
            cache[threadIdx.x] += cache[threadIdx.x + i];
        }
        __syncthreads();
    }
    __syncthreads();
    if (cacheIndex==0) {
        *dot_res=cuda_atomicAdd(dot_res,cache[0]);
    }
}

А вот функция моего устройства atomicAdd:

__device__ double cuda_atomicAdd(double *address, double val)
{
    double assumed,old=*address;
    do {
        assumed=old;
        old= __longlong_as_double(atomicCAS((unsigned long long int*)address,
                    __double_as_longlong(assumed),
                    __double_as_longlong(val+assumed)));
    }while (assumed!=old);

    return old;
}

Answer 1

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

#include <thrust/inner_product.h>
#include <thrust/device_ptr.h>

double do_dot_product(int n, double *a, double *b)
{
  // wrap raw pointers to device memory with device_ptr
  thrust::device_ptr<double> d_a(a), d_b(b);

  // inner_product implements a mathematical dot product
  return thrust::inner_product(d_a, d_a + n, d_b, 0.0);
}

Answer 2

Вы используете функцию cuda_atomicAdd неправильно.Этот раздел вашего ядра:

if (cacheIndex==0) {
    *dot_res=cuda_atomicAdd(dot_res,cache[0]);
}

является виновником.Здесь вы атомарно добавляете к dot_res.затем не атомарно установить dot_res с результатом, который он возвращает.Возвращаемым результатом этой функции является предыдущее значение местоположения, которое атомарно обновляется, и оно предоставляется только для «информации» или локального использования вызывающей стороны.Вы не назначаете это на то, что вы обновляете атомарно, что полностью противоречит цели использования атомарного доступа к памяти.Сделайте что-то вроде этого:

if (cacheIndex==0) {
    double result=cuda_atomicAdd(dot_res,cache[0]);
}

Answer 3

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

1. Запустите новое ядро ​​параллельного сокращения для суммирования точечного продукта.
   Поскольку данные уже находятся на устройстве, при запуске вы не увидите снижения производительностиновое ядро.

2. Кажется, ваше ядро ​​не масштабируется по максимально возможному количеству блоков в новейшем графическом процессоре.Если это произойдет, и ваше ядро ​​сможет рассчитать скалярное произведение миллионов значений, производительность резко снизится из-за сериализованной атомарной операции.

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

4. оптимизируйте параллельное сокращение. Мой тезис или Оптимизации Марк Харрис

Не проверено, я просто записал это в блокноте:

/*
 * @param inCount_s unsigned long long int Length of both input arrays
 * @param inValues1_g double* First value array
 * @param inValues2_g double* Second value array
 * @param outDots_g double* Output dots of each block, length equals the number of blocks
 */
__global__ void dotProduct(const unsigned long long int inCount_s,
    const double* inValuesA_g,
    const double* inValuesB_g,
    double* outDots_g)
{
    //get unique block index in a possible 3D Grid
    const unsigned long long int blockId = blockIdx.x //1D
            + blockIdx.y * gridDim.x //2D
            + gridDim.x * gridDim.y * blockIdx.z; //3D


    //block dimension uses only x-coordinate
    const unsigned long long int tId = blockId * blockDim.x + threadIdx.x;

    /*
     * shared value pair products array, where BLOCK_SIZE power of 2
     *
     * To improve performance increase its size by multiple of BLOCK_SIZE, so that each threads loads more then 1 element!
     * (outDots_g length decreases by same factor, and you need to range check and initialize memory)
     * -> see harris gpu optimisations / parallel reduction slides for more informations.
     */
    __shared__ double dots_s[BLOCK_SIZE];


    /*
     * initialize shared memory array and calculate dot product of two values, 
     * shared memory always needs to be initialized, its never 0 by default, else garbage is read later!
     */
    if(tId < inCount_s)
        dots_s[threadIdx.x] = inValuesA_g[tId] * inValuesB_g[tId];
    else
        dots_s[threadIdx.x] = 0;
    __syncthreads();

    //do parallel reduction on shared memory array to sum up values
    reductionAdd(dots_s, dots_s[0]) //see my thesis link

    //output value
    if(threadIdx.x == 0)
        outDots_g[0] = dots_s[0];

    //start new parallel reduction kernel to sum up outDots_g!
}

Редактировать: убраны лишние точки.