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

【问题标题】：Can anybody help me with atomicmin function syntax for cuda?任何人都可以帮助我了解 cuda 的 atomicmin 函数语法吗？
【发布时间】：2019-03-13 11:32:09
【问题描述】：

我无法获得使用 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 ) ;     
            }
        }
    }

【问题讨论】：

没有内置的，但是您可以使用一些漂亮的技巧来创建自己的：stereopsis.com/radix.html
这个 Devtalk 论坛帖子似乎给出了一个解决方案，我没有测试它，试试看：devtalk.nvidia.com/default/topic/492068/atomicmin-with-float
docs.nvidia.com/cuda/cuda-c-programming-guide/… -- 文档清楚地表明对于任何浮点类型（单或双）都没有 atomicMin
@talonmies 我认为这个问题没有你回答的那么字面意思。 OP 可能知道它不存在，并正在寻求有关如何实现它的帮助。
要求“正确的语法”听起来不像.....

标签： cuda gpu atomic minimum double-compare-and-swap

【解决方案1】：

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

至少有 2 个选项。

您可以重构代码以将原子替换为 classical parallel reduction。
如the programming guide 中所述，可以使用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);
}

This 相关的问题和答案也可能很有趣，尽管它主要关注的是float。

其他几个cmets：

通过切换到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);
```
我认为应该更快

【讨论】：