0

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

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

  2. 根据 W[idx] 更新 V[idx]。但是，只有在步骤 1 中使用 V[idx] 处的值来计算 W 之后，才应该对 V[idx] 进行更新。

假设我在内核中有 3 次 for 循环迭代。如果一个线程在迭代 1 中并尝试使用迭代 1 的 V[2] 在迭代 1 中计算 W[idx = 18]。另一个线程假设在迭代 2 中并尝试在迭代 2 中计算 W[2] a,b,c,d 并在迭代 2 中计算 V[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 个计算之间创建一个同步点。并且在第二次计算结束时。我想找到一种方法，只创建一个内核，并在存在同步的情况下在该内核内执行迭代。

4

1 回答 1

0

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

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

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

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

你可以尝试的技巧

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

  • group_barrier：如果您设法以某种方式制定问题，使得生产者-消费者依赖关系不会跨越两个工作组之间的边界，则可以group_barrier()用于同步
  • atomic_ref：如果您以某种方式知道您的 SYCL 实现/驱动程序/硬件都保证您的生产者工作组在消费者工作组之前或期间执行，那么您可以在全局内存中拥有一个原子标志来存储该值是否已经更新。您可以使用atomic_ref存储/加载来实现全局内存中的自旋锁之类的东西。
  • vec多个缓冲区：如果在第二个内核结束时将更新的内容存储在临时缓冲区而不是原始缓冲区中，则可以合并两个内核。两个内核完成后，翻转原始缓冲区和临时缓冲区以进行下一次迭代。
于 2021-08-30T00:36:47.660 回答