CUDA中的原子Saxpy

Question

CUDA中的原子Saxpy

3

我在CUDA中遇到了以下问题。

假设我们有一个索引列表，其中一些或全部索引可能出现多次：

inds = [1, 1, 1, 2, 2, 3, 4]

使用这些索引，我想在一个浮点数数组x上执行原子saxpy操作（并行）。我不担心操作的顺序。也就是说，对于浮点数a和k，我想要做到这一点：

x[i] = x[i]*a + k;

如果 inds 中没有重复的索引，那么这将变得微不足道。

我的当前解决方案（不起作用）如下：

// assume all values in adr are greater than or equal to 0.
// also assume a and k are strictly positive.

__device__ inline void atomicSaxpy(float *adr,
                                   const float a,
                                   const float k){

  float old = atomicExch(adr, -1.0f); // first exchange
  float new_;
  if (old <= -1.0f){
    new_ = -1.0f;
  } else {
    new_ = old*a + k;
  }

  while (true) {
    old = atomicExch(adr, new_); // second exchange
    if (old <= -1.0f){
      break;
    }
    new_ = old*a + k;
  }
}

这在许多情况下似乎返回了正确的答案。

以下是我认为当您没有得到正确答案时会发生的情况：

1. 在第一次交换中，`old` 获得了 `-1.0f` 的值。=> `new_ = -1.0f` 2. 在第二次交换中，`old` 再次获得 `-1.0f` 的值。 3. 函数退出时根本没有任何外部影响。

另一种稍微不同的方法是：

__device__ inline void atomicSaxpy(float *adr,
                                   const float ia,
                                   const float k){

  float val;

  while (true) {
    val = atomicExch(adr, -1.0f);
    if (val > 1.0f){
      break;
    }
    atomicExch(adr, val*ia + k);
  }
}

这段文字的意思是：“这个函数在我的机器上经常死锁，即使是像上面的例子数据这么简单的输入也会死锁。有没有可能重新编写这个函数以正确地运行？假设k=0.1和a=0.95，并且初始值args的所有索引都为0.5，那么结果应该是：”

[0.5, 0.7139374999999998, 
 0.6462499999999999, 0.575, 0.575, ...]

我使用Python计算了这些值，它们在CUDA中可能会有所不同。这是算法应该表现的样例，而不是运行到竞争条件问题时的好样本集。

参考文献

这里有一个线程，他们使用atomicExch实现了atomicAdd（此时已经存在于浮点数中）：

https://devtalk.nvidia.com/default/topic/458062/atomicadd-float-float-atomicmul-float-float-/

一个例子看起来像这样：

__device__ inline void atomicAdd(float* address, float value) {
  float old = value;  
  float new_old;

  do {
    new_old = atomicExch(address, 0.0f);
    new_old += old;
  }
  while ((old = atomicExch(address, new_old)) != 0.0f);
};

这似乎更容易一些，但我无法完全看到如何适应它。

其他解决方案

能够以这种方式解决此问题对于我未来与内存IO相关的问题有几个优点。出于这个原因，我想知道是否有可能。

一个可能的不同方法是计算每个索引在CPU上出现的次数，然后在GPU上执行“常规” saxpy。我假设还有其他可能性，但我仍然对这个特定问题的答案感兴趣。

- inconvergent

1

你可能会更好地将其视为前缀和。 - talonmies

你能否更新问题，并提供你展示的示例输入的期望答案？ - talonmies

@talonmies，请查看更新后的问题。 - inconvergent

@talonmies 我稍微考虑了一下前缀和公式。我看不出如何在事先不知道数组中每个索引出现次数的情况下完成这项任务。你能详细说明一下你的想法吗？ - inconvergent

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- rompetroll · Accepted Answer

如果这是一个非并行问题，你只需这样做：

*adr = *adr * a + k;

因为有多个线程在操作 adr，所以我们应该使用原子操作进行读写。

float adrValue = atomicExch(adr, -1.0f)
float newValue = adrValue * a + k
atomicExch(adr, newValue)

然而，我们必须意识到另一个线程在我们读取步骤（ln1）和写入步骤（ln3）之间更新了 adr 的可能性。

因此，我们目前的三步操作是非原子的。

为了使其原子化，我们应该使用比较并交换（atomicCAS）来确保仅在自从我们读取后其值未更改时才更新内存。我们可以简单地重复我们的步骤，在每次迭代中以当前 adr 值作为计算输入，直到 step3 返回预期的锁定值 -1.0f。

do {
    float adrValue = atomicExch(adr, -1.0f)
    float newValue = adrValue * a + k
    adrValue = __int_to_float(atomicCAS(adr, 
                                        __float_as_int(-1.0f),
                                        __float_as_int(newValue)))
} while (adrValue != -1.0f)

请考虑上面的伪代码。