时间戳

首页 > 解决方案 > 我如何避免这种原子操作的竞争条件？

c - 我如何避免这种原子操作的竞争条件？

问题描述

以下面的代码片段为例：

    __global__ void my_kernel(float *d_min, uint32_t *d_argmin, float *d_input, uint32_t N)
    {
        uint32_t ii = blockDim.x * blockIdx.x + threadIdx.x;

        if (ii>=N)
            return;

        float cost_ii = someCostFunction( d_input[ii] );
        float old_val = atomicMin( d_min, cost_ii );
        if (old_val != cost_ii)
        {
            *d_argmin = ii;
        } 
    }

让我们假设d_min由调用者初始化说，9999.9。线程 0 计算较小的成本100.0，并安全地自动调整最小值。同时，在另一个线程块中，cost_ii计算出另一个线程块，结果为cost_ii = 10.0。它阻塞atomicMin指令，然后设置它。

所以分配给的顺序d_min如下：

d_min = 9999.9  // from external initialization
d_min = 100.0   // from thread 0 in block 0
d_min = 10.0    // from thread 0 in block 1

但后面的if语句以不同的顺序执行：

d_argmin = 1024 // from thread 0 in block 1
d_argmin = 0    // from thread 0 in block 0

如何避免atomicMin()执行时间和d_argmin设置时间之间的竞争条件。

标签： ccudarace-condition

解决方案

推荐阅读