Soit la structure définie comme suit :

typedef struct S { 
    float x;
    float y;
} T;

Et l'opération struct_add est définie comme suit :

__device__ T struct_add(T a1, T a2) {
    T result;
    result.x = a1.x + a2.x;
    result.y = a1.y + a2.y;
}

Si je veux appliquer struct_add de manière atomique, comment puis-je l'implémenter dans CUDA ? Par exemple, a, b et c doivent être additionnés à l'aide de struct_add, et le résultat doit être stocké dans d. (où le type de a, b, c et d est T)

J'ai entendu dire que le "contrôle de verrouillage et d'accès" via une boucle while n'est pas recommandé. Existe-t-il un moyen approprié de mettre en œuvre cela?

0
Square 17 oct. 2020 à 01:13

1 réponse

Meilleure réponse

Il n'y a pas de méthode atomique générale fournie par CUDA qui couvre les mises à jour atomiques de structure arbitraire. Quelques possibilités :

  1. Parce que vous souhaitez spécifiquement mettre à jour deux éléments 32 bits adjacents, vous pouvez utiliser une opération atomique 64 bits généralisée qui serait une variante de ce qui est décrit ici.

  2. Une autre alternative est celle que vous avez déjà mentionnée, implémentant essentiellement une section critique.

  3. Enfin, une autre approche possible peut être la réduction parallèle, bien que cette n'est pas exactement analogue à l'utilisation atomique

Dans le sens de la suggestion 1 ci-dessus, voici une modification du code à partir de cette réponse qui peut indiquer comment vous pouvez utiliser un atomic 64 bits :

$ cat t56.cu
#include <stdio.h>
#define DSIZE 512
#define nTPB 256

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

typedef union {
  float floats[2];
  unsigned long long int ulong;    // for atomic update
} my_atomics;

__device__ my_atomics test;

__device__ unsigned long long int my_atomicAdd_2floats(unsigned long long int* address, float val0, float val1)
{
    my_atomics loctest;
    unsigned long long old = *address;
    do {
      loctest.ulong = old;
      my_atomics loc;
      loc.floats[0] = val0 + loctest.floats[0];
      loc.floats[1] = val1 + loctest.floats[1];
      old = atomicCAS(address, loctest.ulong,  loc.ulong);}
    while (old != loctest.ulong);
    return old;
}


__global__ void min_test(const float* data)
{

    int idx = (blockDim.x * blockIdx.x) + threadIdx.x;
    if (idx < DSIZE)
      my_atomicAdd_2floats(&(test.ulong), data[idx], (float)idx);
}

int main() {

  float *d_data, *h_data;
  my_atomics my_init;
  my_init.floats[0] = 0.0f;
  my_init.floats[1] = 0.0f;

  h_data = (float *)malloc(DSIZE * sizeof(float));
  if (h_data == 0) {printf("malloc fail\n"); return 1;}
  cudaMalloc((void **)&d_data, DSIZE * sizeof(float));
  cudaCheckErrors("cm1 fail");
  for (int i = 0; i < DSIZE; i++) h_data[i] = 1.0f;
  cudaMemcpy(d_data, h_data, DSIZE*sizeof(float), cudaMemcpyHostToDevice);
  cudaCheckErrors("cmcp1 fail");
  cudaMemcpyToSymbol(test, &(my_init.ulong), sizeof(unsigned long long int));
  cudaCheckErrors("cmcp2 fail");
  min_test<<<(DSIZE+nTPB-1)/nTPB, nTPB>>>(d_data);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");

  cudaMemcpyFromSymbol(&(my_init.ulong), test, sizeof(unsigned long long int));
  cudaCheckErrors("cmcp3 fail");

  printf("device float0 result = %f\n", my_init.floats[0]);
  printf("device float1 result = %f\n", my_init.floats[1]);

  float host_val0 = 0.0f;
  float host_val1 = 0.0f;
  for (int i=0; i<DSIZE; i++) {
          host_val0 += h_data[i];
          host_val1 += (float)(i);}
  printf("host float0 result = %f\n", host_val0);
  printf("host float1 result = %f\n", host_val1);
  return 0;
}
$ nvcc -arch=sm_35 -o t56 t56.cu -Wno-deprecated-gpu-targets
$ cuda-memcheck ./t56
========= CUDA-MEMCHECK
device float0 result = 512.000000
device float1 result = 130816.000000
host float0 result = 512.000000
host float1 result = 130816.000000
========= ERROR SUMMARY: 0 errors
$

Je ne garantis pas que le code ci-dessus est sans défaut. Je suggère de le tester soigneusement avant de l'utiliser.

1
Robert Crovella 17 oct. 2020 à 01:14