在SYCL中使用std::atomic_ref / cl::sycl::atomic_ref实现读写依赖关系

Question

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

1. 首先，循环根据当前迭代中V的4个随机值计算索引处的W值(Widx)。即，Widx = sum (Va + Vb + Vc+ Vd)。其中，a、b、c和d不是连续的，而是为每个idx.

定义的。

1. 基于Widx更新Vidx。但是，只有在步骤1中使用了Vidx的值来计算W.

之后，才能对Vidx进行更新。

假设我在内核中有3个for循环迭代。如果一个线程在迭代1中，并试图在迭代1时使用迭代1的V2来计算Widx = 18。另一个线程可以说是在迭代2中，并尝试从a、b、c、d计算迭代2中的W2，并在iteration2上计算V2。

如果第二个线程位于第一个线程的前面，那么第二个线程将在迭代2中更新V2的值。在这种情况下，当第一个线程想要使用第一次迭代的V2时，如何确保这是Syncd。在SYCL。在这种情况下将使用atomic_ref help，考虑到第二个线程的目标是只在线程1使用它之后才编写V2。还需要注意的是，第一次迭代的V2也需要计算其他一些W以及在其他线程中运行的第一次迭代。如何确保第二次迭代中的V2值在第二次迭代中得到更新，只有在第一次迭代的V2已在所有需要的实例中使用时才能得到更新？以下是源代码：

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个内核，以便在两个计算之间创建一个同步点。在第二次计算结束时。我想找到一种方法，我只创建一个内核，并在该内核内与同步一起执行迭代。

Answer 1

首先，了解SYCL提供的同步保证是很重要的。与许多其他异构模型(如OpenCL)一样，SYCL只允许工作组内的同步，而不允许与其他工作组的工作项同步。这里的背景是，不需要硬件、驱动程序或SYCL实现并行执行工作组，使它们能够独立地向前推进。相反，堆栈可以任意顺序地执行工作组--在极端情况下，它可以一个一个地依次执行工作组。一个简单的例子是，如果你是在一个单一的核心CPU。在这种情况下，SYCL实现的后端线程池可能只有1大小，因此SYCL实现可能只是顺序地遍历所有工作组。

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

因此，实现内核所有工作项之间的同步的标准方法是将内核拆分到两个内核中，正如您所做的那样，是。

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

技巧你可以试试

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

• group_barrier：如果您设法解决了这样的问题，使生产者-消费者依赖关系不会跨越两个工作组之间的界限，那么如果您知道您的SYCL实现/驱动程序/硬件都能保证您的生产者工作组在使用者工作组之前或期间执行，您就可以在全局内存中使用一个原子标志来存储该值是否已经更新。您可以使用memory.
• Multiple存储/加载来实现全局atomic_ref缓冲区中的自旋锁:如果在第二个内核的末尾将更新的vec存储在临时缓冲区中而不是原始缓冲区中，则可以组合这两个内核。在两个内核完成后，您将为下一次迭代翻转原始和临时缓冲区。