Prodotto Cuda Dot

Question

Sto cercando di implementare il classico kernel dot-prodotti per array a doppia precisione con calcolo atomico della somma finale attraverso i vari blocchi. Ho usato l'atomicadd per la doppia precisione come indicato nella pagina 116 della Guida alla programmazione. Probabilmente sto facendo qualcosa di sbagliato. Le somme parziali attraverso i thread in ogni blocco sono calcolate correttamente ma post -parole L'operazione atomica non sembra funzionare correttamente Poiché ogni volta che eseguo il mio kernel con gli stessi dati, ricevo risultati diversi. Sarò grato se qualcuno potesse individuare l'errore o fornire una soluzione alternativa! Ecco il mio kernel:

__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]);
    }
}

Ed ecco la funzione del mio dispositivo 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

Stai usando il cuda_atomicAdd funzione in modo errato. Questa sezione del tuo kernel:

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

è il colpevole. Qui, aggiungi atomicamente a dot_res. poi non atomicamente impostare dot_res Con il risultato restituisce. Il risultato di ritorno da questa funzione è il valore precedente della posizione che viene aggiornata atomicamente e fornita solo per "informazioni" o uso locale del chiamante. Non lo assegni a ciò che sei atomicamente aggiornato, che sconfigge completamente lo scopo di utilizzare l'accesso alla memoria atomica in primo luogo. Fare qualcosa del genere invece:

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

Answer 2

Ottenere una riduzione giusta usando il codice CUDA ad hoc può essere complicato, quindi ecco una soluzione alternativa usando un algoritmo di spinta, che è incluso con il toolkit CUDA:

#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 3

Non ho controllato il tuo codice quella profondità, ma ecco alcuni consigli.
Consiglio di consigli solo di utilizzare la spinta se usi la tua GPU solo per tali attività generiche, poiché se sorgerà un problema complesso, le persone non hanno idea di programmare in modo efficiente parallelo alla GPU.

1. Avvia un nuovo kernel di riduzione parallela per riassumere il prodotto DOT.
   Poiché i dati sono già sul dispositivo non vedrai una diminuzione delle prestazioni che avvia un nuovo kernel.

2. Il tuo kernel sembra non ridimensionare il numero massimo di possibili blocchi sulla GPU più recente. Se lo facesse e il tuo kernel sarebbe in grado di calcolare il prodotto punto di milioni di valori, le prestazioni diminuirebbero drasticamente a causa dell'operazione atomica serializzata.

3. Errore per principianti: i tuoi dati di input e l'accesso alla memoria condivisa gamma controllata? O sei sicuro che i dati di input sono sempre più della dimensione del blocco? Altrimenti leggerai la spazzatura. La maggior parte dei miei risultati errati erano dovuti a questo difetto.

4. Ottimizza la tua riduzione parallela. La mia tesi o Optimizzazioni Mark Harris

Non testato, l'ho appena scritto in blocco note:

/*
 * @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!
}

EDIT: punti rimossi inutili.