在 SYCL 中使用 std::atomic_ref / cl::sycl::atomic_ref 进行读后写和读后写依赖

Question

假设我们有 2 个长度为 n 的向量 V 和 W。我在SYCL中启动一个内核，对V的每个实体执行3次for循环迭代。for循环的描述如下：

首先，循环根据当前迭代中 V 的 4 个随机值计算索引 (W[idx]) 处的 W 值。即，W[idx] = 和 (V[a] + V[b] + V[c]+ V[d])。其中 a、b、c 和 d 不是连续的，而是为每个 idx 定义的。
基于W[idx]更新V[idx]。 但是，只有在 V[idx] 的值已在步骤 1 中用于计算 W 后，才能更新 V[idx]。

假设我在内核中有 3 次 for 循环迭代。如果一个线程在迭代 1 中并尝试使用迭代 1 的 V[2] 来计算迭代 1 中的 W[idx = 18]。另一个线程假设在迭代 2 中并尝试在迭代 2 中计算 W[2] a,b,c,d 并计算 V[2] 已经在迭代 2.

如果第二个线程领先于第一个线程，第二个线程将在迭代2更新V[2]的值。在这种情况下，当第一个线程想要使用第一次迭代的V[2]时，我如何确保这是 Syncd。在 SYCL 中。在这种情况下使用 atomic_ref 会有所帮助，考虑到第二个线程的目标是仅在线程 [1] 使用 V[2] 之后才写入它。还要注意的是，第一次迭代的 V[2] 也需要在某些其他线程中的第一次迭代运行中计算其他一些 W。仅当在所有必需实例中使用了第一次迭代的 V[2] 时，如何确保第二次迭代中 V[2] 的值在第二次迭代中得到更新？这是源代码：

void jacobi_relaxation(cl::sycl::queue& q, ProblemVar& obj, int current_level) {
  for (int iterations = 1; iterations <= mu1; iterations++) {
    // TODO   =>        v(k+1) = [(1 - omega) x I + omega x D^-1 x(-L-U)] x v(k) + omega x
    // D^-1
    // x
    // f
    //
    // step 1 =>        v* = (-L-U) x v
    // step 2 =>        v* = D^-1 x (v* + f)
    // step 3 =>        v = (1-omega) x v + omega x v*

    q.submit([&](cl::sycl::handler& h) {
      // Accessor for current_level matrix CSR values
      auto row = obj.A_sp_dict[current_level].row.get_access<cl::sycl::access::mode::read>(h);
      auto col = obj.A_sp_dict[current_level].col.get_access<cl::sycl::access::mode::read>(h);
      auto val = obj.A_sp_dict[current_level].values.get_access<cl::sycl::access::mode::read>(h);
      auto diag_indices
          = obj.A_sp_dict[current_level].diag_index.get_access<cl::sycl::access::mode::read>(h);

      auto vec = obj.vecs_dict[current_level].get_access<cl::sycl::access::mode::read>(h);
      auto f = obj.b_dict[current_level].get_access<cl::sycl::access::mode::read>(h);
      cl::sycl::accessor<double, 1, cl::sycl::access::mode::write> vec_star{
          obj.temp_dict[current_level], h, cl::sycl::noinit};

      // Require 2 kernels as we perform Jacobi Relaxations
      h.parallel_for(
          cl::sycl::range<1>{obj.num_dofs_per_level[current_level]}, [=](cl::sycl::id<1> idx) {
            // double diag_multiplier = 0.0;
            vec_star[idx[0]] = 0.0;
            for (std::int32_t i = row[idx[0]]; i < row[idx[0] + 1]; i++) {

              vec_star[idx[0]] += -1.0 * val[i] * vec[col[i]];

            }
            
            vec_star[idx[0]] = (1.0 / val[diag_indices[idx[0]]]) * (vec_star[idx[0]] + f[idx[0]])
                               + vec[idx[0]]; // step 2
          });
    });
    q.wait();
    q.submit([&](cl::sycl::handler& h) {
      // Accessor for current_level vector
      auto vec = obj.vecs_dict[current_level].get_access<cl::sycl::access::mode::read_write>(h);
      auto vec_star
          = obj.temp_dict[current_level].get_access<cl::sycl::access::mode::read_write>(h);

      h.parallel_for(cl::sycl::range<1>{obj.num_dofs_per_level[current_level]},
                     [=](cl::sycl::id<1> idx) {
                       vec[idx[0]] = (1.0 - omega) * vec[idx[0]] + omega * vec_star[idx[0]]; // step
                                                                                             // 3
                       vec_star[idx[0]] = 0.0;
                     });
    });
    q.wait();
  }
}

如果你看到，对于每次迭代，我都被迫启动 2 个内核，以便我可以在 2 个计算之间创建一个同步点。并且在第二次计算结束时。我想找到一种方法，只创建一个内核，并在存在同步的情况下在该内核内执行迭代。

Answer 1

首先，重要的是要了解SYCL 所做的同步保证。与许多其他异构模型（例如 OpenCL）一样，SYCL 只允许在一个工作组内进行同步，而不允许与来自其他工作组的工作项进行同步。这里的背景是不需要硬件、驱动程序或 SYCL 实现来并行执行工作组，以便它们独立向前推进。相反，堆栈可以自由地以任何顺序执行工作组——在极端情况下，它可以按顺序一个接一个地执行工作组。一个简单的例子是，如果你是在单核 CPU 上。在这种情况下，SYCL 实现的后端线程池的大小可能仅为 1，因此 SYCL 实现可能只是按顺序迭代所有工作组。

这意味着很难制定跨越多个工作组的生产者-消费者算法[其中一个工作项产生另一个工作项等待的值]，因为生产者工作组总是被调度在消费者工作组之后运行，如果可用的硬件资源阻止两者同时运行，可能会导致死锁。

因此，在内核的所有工作项之间实现同步的规范方法是将内核分成两个内核，正如您所做的那样。

我不确定您是否只是为了代码示例而这样做，或者它是否也在您的生产代码中，但我想指出内核之间和内核之后的 q.wait() 调用似乎不必要。 queue::wait() 导致主机线程等待提交的操作完成，但对于这个用例，如果您知道内核运行是有序的就足够了。 SYCL 缓冲区访问器模型会自动保证这一点，因为 SYCL 实现会检测到两个内核都是读写 vec_star，因此在 SYCL 任务图中插入了一个依赖边。通常，为了性能，除非绝对必要，否则您希望避免主机同步，并让设备异步完成所有排队的工作。

你可以尝试的技巧

原则上，在某些特殊情况下，您可以尝试其他方法。但是，对于大多数用例，我不认为它们是比仅使用两个内核更好的选择。

group_barrier：如果您以某种方式设法制定问题，使生产者-消费者依赖关系不会跨越两个工作组之间的边界，则可以使用 group_barrier() 进行同步
atomic_ref：如果你不知何故知道你的 SYCL implementation/driver/hardware 都保证你的生产者工作组在消费者工作组之前或期间执行，你可以在全局内存中有一个原子标志来存储该值是否已经更新。您可以使用 atomic_ref store/load 在全局内存中实现类似自旋锁的功能。
多个缓冲区：如果在第二个内核的末尾将更新的 vec 存储在临时缓冲区而不是原始缓冲区中，则可能会合并两个内核。两个内核完成后，翻转原始缓冲区和临时缓冲区以进行下一次迭代。

在 SYCL 中使用 std::atomic_ref / cl::sycl::atomic_ref 进行读后写和读后写依赖

Use of std::atomic_ref / cl::sycl::atomic_ref in SYCL for read after write and write after read dependencies

c++

multithreading

sycl