OpenCL 基准测试 - 关于参数变化的建议

Question

我想在 radeon HD 7970 Tahiti XT 上使用 OpenCL（来自 AMD link）执行关于两阶段 Sum 减少的运行时基准测试。

最初，我采用了第一个版本的代码，其中我没有使用第一个循环，该循环执行从大小为N 的输入数组到大小为NworkItems 的输出数组的缩减。这是进入内核代码的第一个循环：

  int global_index = get_global_id(0);
  float accumulator = 0;
  // Loop sequentially over chunks of input vector
  while (global_index < length) {
    float element = buffer[global_index];
    accumulator += element;
    global_index += get_global_size(0);
  }

因此，在第一个版本中，我将运行时间测量为输入数组大小（等于线程总数）和不同大小的工作组的函数。结果如下：

现在，我想做一个基准测试，在上面使用这个初始循环。但我不知道我必须改变哪些参数。

来自this link，有人说 AMD 建议工作组大小为 64 的倍数（NVIDIA 为 32）。

此外，从this other link 的最后评论中，建议将工作组大小设置为：WorkGroup size = (Number of total threads) / (Compute Units)。 在我的 GPU 卡上，我有 32 个计算单元。

因此，为了比较第二个版本（使用第一个缩减循环）中的运行时间，我想获得一些建议，以了解哪些参数会有趣地变化。例如，我可以为比率(N size of input array) / (total NworkItems) 取不同的值，为WorkGroup size 取一个固定值（见上面的表达式），

或者相反，即我应该改变WorkGroup size的值并固定比率(N size of input array) / (total NworkItems)？

Answer 1

您应该汇总本地数据而不是分散数据，以帮助内存传输（合并数据访问）。所以改用这个：

  int chunk_size = length/get_global_size(0)+(length%get_global_size(0) > 0); //Will give how many items each work item needs to process
  int global_index = get_group_id(0)*get_local_size(0)*chunk_size + get_local_id(0); //Start at this address for this work item
  float accumulator = 0;

  for(int i=0; i<chunk_size; i++)
    // Loop sequentially over chunks of input vector
    if (global_index < length) {
      float element = buffer[global_index];
      accumulator += element;
      global_index += get_local_size(0);
    }
  }

您还应该使用 2 的幂的大小来帮助缓存。