任意の長さの配列を減らすためのワープ・シャッフル

Question

私はベクトル・ドット積（A×B）を実行するCudaカーネルに取り組んでいます。私は各ベクトルの長さが32の倍数（32,64、...）であると仮定し、ブロックサイズを配列の長さと等しくなるように定義しました。ブロック内の各スレッドは、Aの1つの要素にBの対応する要素（スレッドi ==> psum = A [i] xB [i]）を乗算します。乗算後、ワープシャッフル技術を使って減算を行い、すべての乗算の合計を計算する次の関数を使用しました。

__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; 
} 

psumは2つの要素をスレッドで乗算したものです。単にblockReduceSum（psum）を呼び出します。

配列の長さが32の倍数でない場合、このアプローチは機能しません。だから私の質問は、このコードを変更して任意の長さにすることができますか？配列の長さが32の倍数でない場合、複数のワープに複数の配列に属する要素があるため、不可能ですか？

Answer 1

まず、使用しているGPUによっては、1ブロックだけでドットプロダクトを実行すると、効率的ではありません（1つのカーネルで複数のドットプロダクトをバッチしない限り）。あなたはblockReduceSumに呼び出す前に、ちょうどスレッドの数がN（配列の長さ）よりも32以上の最も近い倍数であると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を貢献していきます。