Cub设备扫描使用自定义扫描OP失败

Question

我正在使用

template <typename T>
struct MultAddFunctor
{
    const T factor;
    MultAddFunctor(T factor) : factor(factor) {}
    
    __device__ __forceinline__
    T operator()(const T &a, const T &b) const {
        return factor*a + b;
    }
};

否则时，我的代码几乎与文档中的示例代码相同（除了我释放了分配的内存并添加了其他同步以排除该问题为问题）。当factor是1.0时，这会产生正确的结果（这只是前缀sum）。当factor是其他东西（例如0.8）时，结果对于前12个值是正确的，但之后大大差异。例如，如果被扫描的数组只是一堆1.0 s，我会得到以下结果：

        CUDA     Serial
   0     1.000     1.000 ✅
   1     1.800     1.800 ✅
   2     2.440     2.440 ✅
   3     2.952     2.952 ✅
   4     3.362     3.362 ✅
   5     3.689     3.689 ✅
   6     3.951     3.951 ✅
   7     4.161     4.161 ✅
   8     4.329     4.329 ✅
   9     4.463     4.463 ✅
  10     4.571     4.571 ✅
  11     4.656     4.656 ✅
  12     6.313     4.725 ❌
  13     6.050     4.780 ❌
  14     5.840     4.824 ❌
  15     5.672     4.859 ❌
...

在元素12中，值突然跳跃，然后减少，即使这应该只是只有继续以一致的速度越来越大。

起初，我认为这是由于操作的非交换性，但文档明确指出这很好。我还认为factor字段本身可能无法正确访问设备，但是即使我在等式中硬编码a 0.8，它仍然不正确（尽管，尽管因子可能总是在全局内存中，因此将来移动factor陷入共享/本地会更好）。

扫描还可以计算错误结果的其他原因是什么？

Answer 1

发生故障的原因是Cub并行扫描如其他并行扫描实现我是我意识到，需要一个关联的二进制扫描操作员。

这直接源于关联的定义，以添加为例。附加的关联属性说，

(a+b)+c = a+(b+c)

当此处应用时，这意味着在并行设置中，我可以将二进制OP应用于（a+b）首先或（b+c）首先，然后将二进制OP应用于其余步骤，我应该得到相同的结果。这是关联性，而不是交流性。这里的操作不是关联的。为了证明这一点，我们实际上执行了操作员在每种情况下指示的数学（想象/假装factor为1000），并且我们得到：

step 1: (1000*a+b)+c = a+(1000*b+c)
step 2: 1000(1000*a+b)+c = 1000*a+(1000*b+c)

平等不保持。