0

在第一个代码示例 (kernel_conv) 中，我编写了一个简单的卷积，它与预期的结果 [1,1,2,1,1] 一起工作。

然后我使用元素内核来对向量的所有条目求和。但是，如果我运行第二个示例 (kernel_sum)，我会得到结果 [3,0,0]，但会期望 [6,0,0]。

这两个例子有什么区别？为什么第一个示例中的变量 y 被更新，而在第二个示例中它似乎被覆盖了？

import numpy as np 
import cupy as cp 

kernel_conv = cp.ElementwiseKernel(
    'raw float32 x', 'raw float32 y',
    ''' int idx = i*2 + 1;
        for(size_t j=0;j<3;j++){
          y[idx - 1 + j] += x[j];
        }
    ''', 'conv')

x = cp.asarray(np.array([1,1,1]),dtype=np.float32)
y = cp.zeros((5,),dtype=np.float32)
z = kernel_conv(x,y,size=2)
print(z)

kernel_sum = cp.ElementwiseKernel(
  'raw float32 x', 'raw float32 y',
  ''' 
      y[0] += x[i]
  ''', 'summe')

x = cp.asarray(np.array([1, 2, 3]), dtype=np.float32)
y = cp.zeros((3,),dtype=np.float32)
z = kernel_sum(x,y,size=3)
print(z)
4

1 回答 1

1

不正确的结果kernel_sum是由于数据竞争造成的。在这种情况下，3 个线程尝试同时写入全局内存 ( y[0]) 的同一地址。为避免数据竞争，您应该 1）使用atomicAdd或 2）使用cupy.ReductionKernel减少。

实际上，kernel_conv也有数据竞赛。运行的第一个线程y[2] += x[2]可能与运行的第二个线程发生冲突y[2] += x[0]。由于第一个在实际执行中略有滞后，所以结果不受影响，但这是一个时序问题，一般不能保证*。要更正它，您可以atomicAdd再次使用此处，或者您也可以更改通过多个线程运行计算的方式（例如，启动 5 个线程，每个线程计算 的不同元素y）。

* 确实，在上面的例子中kernel_conv，我猜想当所有线程都在同一个warp中运行时，正确性是有保证的，即线程数不大于32。这是因为同一个warp中的所有线程都在运行同步直到被控制流发散。如果线程数设置为更大的值，则数据竞争可能发生在扭曲边界处。

于 2019-09-28T00:16:13.003 回答