经纱洗牌，以减少任何长度的阵列

我正在执行矢量点产品（A×B）的Cuda内核。我假定每个向量的长度是32（32,64，...）的倍数，并且将块大小定义为等于数组的长度。块中的每个线程都将A的一个元素与B的相应元素相乘（线程i ==> psum = A [i] xB [i]）。在乘法之后，我使用了使用warp shuffling技术执行还原的以下函数，并计算所有乘法的和。经纱洗牌，以减少任何长度的阵列

__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的倍数，那么一些warp的元素属于多个数组？

来源

2016-11-21 starrr

首先，根据您使用的GPU，仅使用1个块执行点积可能不会非常有效（只要您不在一个内核中批量生成多个点积，每个点都由一个块完成）。

要回答你的问题：你可以打电话给blockReduceSum之前重用你已经通过只是线程是32高于N（数组的长度），并引入if声明最接近的倍数叫号内核编写的代码这将是这样的：

__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的总和。

来源

2016-11-22 17:14:06 ptrendx

经纱洗牌，以减少任何长度的阵列

回答

相关问题