matrix - OpenCL での行列行の削減

Question

GPU に 1D 配列として格納されている行列があります。たとえば、この行列のすべての行でリダクションを使用する OpenCL カーネルを作成しようとしています。

私の行列が要素[1、2、3、4、5、6]を持つ2x3であると考えてみましょう。私がやりたいことは次のとおりです。

[1, 2, 3] = [ 6]
[4, 5, 6]   [15]

明らかに削減について話しているので、実際の戻り値は行ごとに複数の要素になる可能性があります。

[1, 2, 3] = [3, 3]
[4, 5, 6]   [9, 6]

次に、別のカーネルまたは CPU で実行できる最終的な計算です。

さて、これまでのところ、削減を行うカーネルですが、次のように配列のすべての要素を使用しています。

[1, 2, 3] = [21]
[4, 5, 6]

これを行うための実際のリダクション カーネルはその 1 つです (これは、実際にはスタック オーバーフローでここから取得しました)。

__kernel void
sum2(__global float *inVector, __global float *outVector,
     const unsigned int inVectorSize, __local float *resultScratch)
{
  const unsigned int localId = get_local_id(0);
  const unsigned int workGroupSize = get_local_size(0);

  if (get_global_id(0) < inVectorSize)
    resultScratch[localId] = inVector[get_global_id(0)];
  else
    resultScratch[localId] = 0;

  for (unsigned int a = workGroupSize >> 1; a > 0; a >>= 1)
  {
    barrier(CLK_LOCAL_MEM_FENCE);
    if (a > localId)
      resultScratch[localId] += resultScratch[localId + a];
  }

  if (localId == 0)
    outVector[get_group_id(0)] = resultScratch[0];
  barrier(CLK_LOCAL_MEM_FENCE);
}