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

Question

我想在 radeon HD 7970 Tahiti XT 上使用 OpenCL（从此AMD link）执行关于两阶段总和缩减的运行时基准测试。

最初，我采用了第一个版本的代码，其中我没有使用第一个循环，该循环执行从大小为 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 建议 WorkGroup 的大小为 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 的幂的大小来帮助缓存。

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

OpenCL benchmark - advice about parameters to vary

benchmarking

reduction

opencl