cuda - 任何人都可以帮助我了解 cuda 的 atomicmin 函数语法吗？

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

标签： cudagpuatomicminimumdouble-compare-and-swap

Answer 1

CUDA 定义的atomicMin函数不支持使用浮点数。参考文档，我们看到唯一可用的原型是、和（最后一个需要在计算能力为 3.5 或更高的 GPU 上编译和运行）。intunsigned intunsigned long long int

至少有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 数据）。我相信 iee754float遵循两个数量的排序规则A，B如果A > B, then *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);
  

  我认为应该更快