Кто-нибудь может мне помочь с синтаксисом функции atomicmin для cuda?

Question

Я не могу получить правильный синтаксис для использования atomicMin. Я хотел бы использовать эту функцию для работы с двойным, а не целым числом.

__global__ void npd(int *a, int *g)         
    {   
        int index = threadIdx.x;

        __shared__ int d[N];

        d[threadIdx.x]=a[index];        

        __syncthreads();        

        int dd;
        int inn;
        int u;

        if( 0==threadIdx.x )
        { 
            for( int u = 0; u<16; u++ )
            {
                atomicMin( g, d ) ;     
            }
        }
    }

Answer 1

Функция atomicMin , определенная CUDA , не поддерживает использование с величинами с плавающей запятой.Ссылаясь на документацию, мы видим, что единственные доступные прототипы предназначены для int, unsigned int и unsigned long long int (последний требует компиляции и запуска на графическом процессоре с вычислительной способностью 3,5 или выше).

Есть как минимум 2 варианта.

1. Вы можете реорганизовать свой код, чтобы заменить атомики на классическое параллельное сокращение .

2. Как указано в Руководство по программированию , "произвольная" атомика может быть создана с использованием atomicCAS (сравнение и замена) плюс какой-то цикл.

Вот одна из возможных реализаций для double:

__device__ double atomicMin_double(double* address, double val)
{
    unsigned long long int* address_as_ull = (unsigned long long int*) address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
            __double_as_longlong(fmin(val, __longlong_as_double(assumed))));
    } while (assumed != old);
    return __longlong_as_double(old);
}

Этот связанный вопрос и ответы также могут представлять интерес, хотя ипрежде всего имеет значение float.

Несколько других комментариев:

• путем переключения на float вместо double Я считаю, что можно упроститьОперация atomicMin (или atomicMax), как указано в ответе, на который я ссылался выше, возможно, с несколькими оговорками (например, без данных NaN, INF, например).Я полагаю, что iee754 float следует правилу заказа для двух величин A и B, таких, что если A > B, то *reinterpret_cast<int*>(&A) > *reinterpret_cast<int*>(&B).Я не уверен, что double следует аналогичному правилу с long long (возможно, кто-то еще сможет сказать).

• в вашем коде, этот цикл может работать на локальномсначала количество, затем сделайте одну атомарную операцию в конце, например:

      double v = *g;
      for( int u = 0; u<16; u++ )
      {
          v = min(v,d);     
      }
      atomicMin_double(g, v);
  

  , что, я думаю, должно быть быстрее