Produit Cuda Dot

Question

J'essaie d'implémenter le noyau de produit de point classique pour les réseaux de double précision avec un calcul atomique de la somme finale sur les différents blocs. J'ai utilisé l'atomicadd pour la double précision comme indiqué à la page 116 du guide de programmation. Étant donné que chaque fois que je exécute mon noyau avec les mêmes données, je reçois des résultats différents. Je serai reconnaissant si quelqu'un pouvait repérer l'erreur ou fournir une solution alternative! Voici mon noyau:

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

Et voici ma fonction d'appareil 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

Vous utilisez le cuda_atomicAdd fonction incorrectement. Cette section de votre noyau:

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

est le coupable. Ici, vous ajoutez atomiquement à dot_res. alors non atomiquement Positionner dot_res Avec le résultat, il revient. Le résultat de retour de cette fonction est le valeur antérieure de l'emplacement étant mis à jour atomiquement, et il a fourni des «informations» ou une utilisation locale de l'appelant uniquement. Vous ne l'a pas affecté à ce que vous êtes atomiquement mis à jour, ce qui va complètement à l'objectif d'utiliser l'accès à la mémoire atomique en premier lieu. Faites quelque chose comme ça à la place:

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

Answer 2

Obtenir une bonne réduction à l'aide du code CUDA ad hoc peut être délicat, alors voici une solution alternative à l'aide d'un algorithme de poussée, qui est incluse avec la boîte à outils 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

N'a pas vérifié votre code cette profondeur, mais voici quelques conseils.
Je ne conseillerais que l'utilisation de Thrust si vous n'utilisez votre GPU que pour de telles tâches génériques, car si un problème complexe se pose, les gens n'ont aucune idée de programmer efficacement parallèle sur le GPU.

1. Commencez un nouveau noyau de réduction parallèle pour résumer le produit DOT.
   Étant donné que les données sont déjà sur l'appareil, vous ne verrez pas de diminution des performances qui démarrent un nouveau noyau.

2. Votre noyau ne semble pas évoluer sur le nombre maximum de blocs possibles sur le GPU le plus récent. Si c'était le cas et que votre noyau serait en mesure de calculer le produit DOT de millions de valeurs, les performances diminueraient considérablement en raison de l'opération atomique sérialisée.

3. Erreur des débutants: vos données d'entrée sont-elles et l'accès à la mémoire partagée gamme vérifiée? Ou êtes-vous sûr que les données d'entrée sont toujours multiples de votre taille de bloc? Sinon vous lirez les ordures. La plupart de mes mauvais résultats étaient dus à cette faute.

4. Optimisez votre réduction parallèle. Ma thèse ou Optimisations Mark Harris

Non testé, je viens de l'écrire dans le bloc-notes:

/*
 * @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: supprimé des points inutiles.