1

我有一个在 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);
}
4

1 回答 1

0

我想一种解决方案是修改你的缩减内核，这样它就可以缩减数组的一部分。

__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 提供。

于 2012-11-16T09:16:19.897 回答