OpenCl 全局工作项操作优先级

Question

我想知道下面代码sn-ps（简单的二维矩阵乘法例程）的索引计数的优先级。

kernel void mmul(
   const int N,
   global float* A,
   global float* B,
   global float* C)
 {
     int k;
     int i = get_global_id(0);
     int j = get_global_id(1);
     float tmp;
     if ((i < N) && (j < N))
     {
         tmp = 0.0f;
         for (k = 0; k < N; k++)
             tmp += A[i*N+k] * B[k*N+j];
         C[i*N+j] = tmp;
     }
}

如果您使用“k”计数器查看 for 循环内部，您可以看到全局工作项“i”和“j”放置在同一行中。我想知道它们中的哪些在计算“i”和“j”的索引（例如 1、2、3、4、...、n）方面具有优先权。我不明白这将如何工作，因为我是 OpenCl 的新手，如果我只是使用普通的 C 或 Python，我会使用嵌套的 for 循环来进行这种类型的操作。

有人能解释一下全局工作项是如何工作的吗？

谢谢。

Answer 1

您应该更多地关注内存读/写优先级，而不是工作项发布顺序。要对内存操作强制执行优先级/顺序，请使用 mem_fence(in-workitem) 、 barrier(in-workgroup) 甚至内核（所有工作项同步点）。使用故意的空 for 循环或原子函数不能保证内存写入/读取优先级。只有内存栅栏/屏障/内核可以。

任何工作项（开始/结束运行）都没有优先级，但它们被分组并在有许多线程来运行它们的计算单元上执行。不能保证工作项 i,j 将在 i+1,j+1 之前执行，但可以保证如果它们在同一个工作组中（大小为 16），它们将在同一个计算单元（内核共享 L1 缓存）中执行,16 例如）在使用 Nvidia 和 Amd gpus 时。

在同一计算单元中执行会增加同时发布的机会，这当然不是优先事项，但共享 L1 缓存等资源意味着高性能。

即使在同一个工作组中，也不能保证本地工作项是否在其他工作项之前发布，但如果它们在同一个 SIMD 单元上（例如 Amd gpu 中的 16 宽部分），它们更有可能同时发生.