两个一维向量点积的约简算法

Question

我一直在尝试制定一种算法，通过约简获得 CUDA 程序中两个向量的点积，但似乎卡住了：/

本质上，我正在尝试在 CUDA 中编写这段代码：

for (int i = 0; i < n; i++)
    h_h += h_a[i] * h_b[i];

其中h_a 和h_b 是浮点数组，h_h 对点积求和。

我正在尝试在这里使用归约 - 到目前为止，我已经掌握了这个...

__global__ void dot_product(int n, float * d_a, float * d_b){

     int i = threadIdx.x;

     for (int stride = 1; i + stride < n; stride <<= 1) {
         if (i % (2 * stride) == 0){
             d_a[i] += d_a[i + stride] * d_b[i + stride];
         }
         __syncthreads();
     }
}

如果我将主线更改为d_a[i] += d_a[i + stride];，它可以很好地总结数组。从我收集的信息来看，我似乎在这里遇到了一个平行的问题。有人能指出我的问题吗？

我的内核调用是：

dot_product<<<1, n>>>(n, d_a, d_b);，其中n 是每个数组的大小。

Answer 1

这里有两个问题：

1. 正如 cmets 中所指出的，您永远不会计算第一个元素的乘积（这是一个小问题）
2. 您的点积计算不正确。并行减少应该执行相应元素的单个乘积的总和。您的代码在并行归约的每个阶段执行产品，因此产品在相加时再次相乘。这是不正确的。

你想做这样的事情：

__global__ void dot_product(int n, float * d_a, float * d_b){

     int i = threadIdx.x;

     d_a[i] = d_a[i] * d_b[i]; // d_a now contains products
     __syncthreads();

     for (int stride = 1; i + stride < n; stride <<= 1) {
         if (i % (2 * stride) == 0){
             d_a[i] += d_a[i + stride]; // which are summed by reduction
         }
         __syncthreads();
     }
}

[免责声明：在浏览器中编写，从未编译或测试，使用风险自负]