扭曲改组以减少任何长度的数组答案

【问题标题】：warp shuffling to reduction of arrays with any length扭曲改组以减少任何长度的数组
【发布时间】：2016-11-23 01:14:06
【问题描述】：

我正在研究执行矢量点积 (A x B) 的 Cuda 内核。我假设每个向量的长度是 32 (32,64, ...) 的倍数，并将块大小定义为等于数组的长度。块中的每个线程将 A 的一个元素乘以 B 的相应元素（线程 i ==>psum = A[i]xB[i]）。在乘法之后，我使用了以下函数，这些函数使用了 warp shuffle 技术来执行归约并计算所有乘法的总和。

__inline__ __device__
float warpReduceSum(float val) {
    int warpSize =32;
    for (int offset = warpSize/2; offset > 0; offset /= 2)
        val += __shfl_down(val, offset);
    return val;
}

__inline__ __device__
float blockReduceSum(float val) {
    static __shared__ int shared[32]; // Shared mem for 32 partial sums
    int lane = threadIdx.x % warpSize;
    int wid = threadIdx.x / warpSize;
    val = warpReduceSum(val);         // Each warp performs partial reduction
    if (lane==0) 
        shared[wid]=val;              // Write reduced value to shared memory
    __syncthreads();                  // Wait for all partial reductions
    //read from shared memory only if that warp existed
    val = (threadIdx.x < blockDim.x / warpSize) ? shared[lane] : 0;
    if (wid==0) 
        val = warpReduceSum(val);     // Final reduce within first warp
    return val;
}

我只是调用 blockReduceSum(psum) ，其中 psum 是两个元素乘以一个线程。

当数组的长度不是 32 的倍数时，这种方法不起作用，所以我的问题是，我们可以更改此代码，使其也适用于任何长度吗？还是不可能，因为如果数组的长度不是 32 的倍数，则某些扭曲的元素属于多个数组？

【问题讨论】：

标签： parallel-processing cuda

【解决方案1】：

首先，根据您使用的 GPU，仅使用 1 个块执行点积可能效率不高（只要您不在 1 个内核中批处理多个点积，每个点积由单个块完成） .

回答您的问题：您可以通过调用内核来重用您编写的代码，线程数是高于N（数组长度）的最接近的 32 倍数，并在之前引入 if 语句打电话给blockReduceSum 会这样：

__global__ void kernel(float * A, float * B, int N) {
    float psum = 0;
    if(threadIdx.x < N) //threadIDx.x because your are using single block, you will need to change it to more general id once you move to multiple blocks
        psum = A[threadIdx.x] * B[threadIdx.x];
    blockReduceSum(psum);
    //The rest of computation
}

这样，没有关联数组元素但由于使用__shfl而需要存在的线程将贡献0。

【讨论】：