Продукт CUDA dot дает неправильные результаты - PullRequest
Новая
       7

Продукт CUDA dot дает неправильные результаты

0 голосов
/ 28 мая 2018

Я написал код точечного продукта с CUDA для вычисления точечного произведения двух двойных векторов.Ядро было вызвано N потоков (N <1024) 1 блок.Но это не может дать правильные результаты.Я не могу понять это. </p> 

__global__ void dotthread(double* a, double *b,double *sum, int N)
  {
    int tid = threadIdx.x;  
    sum[0] = sum[0]+a[tid] * b[tid];  //every thread write to the sum[0] 
   __syncthreads();  
  }

Ответы [ 2 ]

0 голосов
/ 29 мая 2018

Я написал две версии процедур точечного продукта.Первый использует функцию atomiAdd, второй выделяет одну общую переменную для каждого блока.Время вычислений составляет 3,33 мс и 0,19 мс соответственно, по сравнению с 0,17 мс и 411,43 мс для произведения точечной редукции и произведения из одной нити. 

GPU Device 0: "GeForce GTX 1070 Ti" with compute capability 6.1
2000000flow array allocated 2147483647
 naive elapsedTimeIn Ms 411.437042 Ms
sum is 6.2831853071e+06

thread atomic add elapsedTimeIn Ms 3.3371520042 Ms
sum is 6.2831853071e+06

 cache reduction elapsedTimeIn Ms 0.1764480025 Ms
sum is 6.2831853072e+06

 cache atomicadd elapsedTimeIn Ms 0.1914239973 Ms
sum is 6.2831853072e+06

__global__ void dotatomicAdd(double* a, double *b,double *sum, int N)
{
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    while (tid < N){
        double t=a[tid] * b[tid];
        atomicAdd(sum,t);
        tid+=blockDim.x * gridDim.x;
    }
}


__global__ void dotcache(double* a, double *b,double *c, int N)
{

    __shared__ double cache;
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;
    double  temp = 0.0;
    cache=0.0;
    __syncthreads();

    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    atomicAdd(&cache,temp);
    __syncthreads();

    if (cacheIndex == 0) c[blockIdx.x] = cache;
}
0 голосов
/ 28 мая 2018

Давайте посмотрим на две из трех ваших строк кода: 

    sum[0] = sum[0]+a[tid] * b[tid];  //every thread write to the sum[0] 
    __syncthreads();

Первая строка содержит гонку памяти.Каждый поток в блоке будет одновременно пытаться записать в сумму [0].В модели исполнения cuda нет ничего, что могло бы помешать этому.Не существует автоматической сериализации или защиты памяти, которая может остановить это поведение.

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

Код, который вы написали, необратимо нарушен.Канонический способ выполнения такого рода операций - параллельное сокращение.Это можно сделать разными способами.Это, вероятно, наиболее описанный и документированный параллельный алгоритм для графических процессоров.Если вы установили инструментарий CUDA, у вас уже есть и полный рабочий пример, и исчерпывающий документ, описывающий алгоритм, как он будет реализован с использованием общей памяти.Я предлагаю вам изучить его.

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

...