在 OpenCL 中减少矩阵行答案

【问题标题】：Reduction of matrix rows in OpenCL在 OpenCL 中减少矩阵行
【发布时间】：2012-11-14 18:01:35
【问题描述】：

我有一个在 GPU 中存储为一维数组的矩阵，我正在尝试制作一个 OpenCL 内核，它将在该矩阵的每一行中使用归约，例如：

假设我的矩阵是 2x3，元素为 [1, 2, 3, 4, 5, 6]，我想要做的是：

[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]

执行此操作的实际缩减内核是那个（实际上是我在 stackoverflow 中从这里得到的）：

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

【问题讨论】：

如果您向我们提供问题，我们会更容易回答。但是，我猜您想知道如何将此示例扩展到 2D/a 矩阵。您可以先并行地对每一行进行归约，然后对最后一列进行归约。

标签： matrix sum opencl gpgpu reduction

【解决方案1】：

我想一种解决方案是修改你的归约内核，这样它就可以对数组的一部分进行归约。

__kernel void
sum2(__global float *inVector,
     __global float *outVector,
     unsigned int   inVectorOffset,
     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[inVectorOffset + 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);
}

然后您可以减少矩阵的一行，将行的开头作为 inVectorOffset 提供，并将行中元素的数量作为 inVectorSize 提供。

【讨论】：