Comment puis-je définir des valeurs à plusieurs emplacements de mémoire, de manière atomique?

Je pense que vous manquez le point de l'opération implémentée ici. En a + = b , l'opération logique est a = a + b , mais avec CAS vous évitez les modifications fallacieuses de a entre sa lecture et sa écrivez. b est utilisé une fois et pas de problème.

Dans a = b + c , aucune des valeurs n'apparaît deux fois, il n'est donc pas nécessaire de se protéger contre les modifications intermédiaires.

merci à tous les gars répondez-moi! J'ai maintenant la solution. nous pouvons combiner les deux variables dans une structure. ainsi nous pouvons transférer "deux variables avec deux adresses" dans "une structure avec une adresse". voici le code:

x=15000.000000
y=20000

et la solution est

#include <stdio.h>
struct pair_t
{
    float x;
    int y;
};

__device__ float single(double *address,double val)
{   

    unsigned long long int *address_as_ull =(unsigned long long int*)address;
    unsigned long long int assumed;
    unsigned long long int old = *address_as_ull;

    do
    {
        assumed = old;
        old = atomicCAS(address_as_ull,assumed,__double_as_longlong(val + __longlong_as_double(assumed)));
    }while(assumed !=old);
    return __longlong_as_double(old);
}



__device__ void myadd(pair_t *address, double val_1 ,int val_2)
{   
    union myunion
    {  
        pair_t p;
        unsigned long long int ull;
    };

    unsigned long long int *address_as_ull;
    address_as_ull = (unsigned long long int *)address;

    union myunion assumed;
    union myunion old_value;
    union myunion new_value;

    old_value.p = *(pair_t *)address_as_ull;

    do
    {
        assumed = old_value;
        // cirtical area begin--------------------
        new_value.p.x = assumed.p.x+val_1;
        new_value.p.y = assumed.p.y+val_2;
        // cirtical area end----------------------

        old_value.ull = atomicCAS(address_as_ull,assumed.ull,new_value.ull);
    }while(assumed.ull !=old_value.ull);
}


__global__ void kernel (pair_t *p)
{
    myadd(p,1.5,2);
}

int main()
{
    pair_t p;
    p.x=0;
    p.y=0;
    pair_t *d_p = NULL;
    cudaMalloc((pair_t **)&d_p, sizeof(pair_t));
    cudaMemcpy(d_p, &p, sizeof(pair_t), cudaMemcpyHostToDevice);

    kernel<<<100, 100>>>(d_p);

    cudaMemcpy(&p, d_p, sizeof(pair_t), cudaMemcpyDeviceToHost);

    cudaDeviceSynchronize();
    printf("x=%lf\n", p.x);
    printf("y=%d\n", p.y);
    cudaDeviceReset();
    return 0;
}

maintenant tout ira bien ~

En général, vous ne pouvez pas faire cela. Le matériel ne prend pas en charge les modifications atomiques à plusieurs emplacements de la mémoire. Bien que vous puissiez contourner cela si les deux variables sont suffisamment petites pour tenir dans la taille d'une seule opération atomique, cette approche échouera si vous avez plus de 8 octets au total. Vous rencontrerez le "trop ​​de lait" problème.

Une chose que vous pourriez faire est d'avoir une sorte de protocole de synchronisation pour accéder à ces deux valeurs. Par exemple, vous pouvez utiliser un mutex, que seul un thread peut obtenir, pour savoir en toute sécurité que personne d'autre ne change les valeurs pendant que ce thread travaille dessus. Voir: Évitez de prendre trop de temps pour terminer le scénario "trop ​​de lait" .

Bien sûr, cela coûte assez cher dans un réglage GPU. Vous feriez probablement mieux de faire l'une des choses suivantes (en augmentant l'ordre de favorabilité):

• Utilisez un pointeur ou un index dans un tableau plus grand, et au lieu de changer la structure de manière atomique, changez le pointeur de manière atomique. Cela résout le problème de concurrence, mais ralentit les accès.
• Modifiez votre algorithme afin que les accès puissent être séparés et ne soient pas obligés de se produire de manière atomique.
• Modifiez davantage votre algorithme pour éviter que plusieurs threads n'écrivent dans une seule structure de données complexe.