SYCL Command Graph与FFT节点组合时的死锁问题求解

我们使用SYCL command_graph进行数据处理，处理流程支持运行时配置。其中一个可选处理步骤是采用Intel oneAPI MKL DFT实现的实部转复数FFT。当FFT节点被加入处理图且并非图中唯一节点时，会出现死锁问题。

可复现的最小示例代码（MRE） #

#include <iostream>
#include <vector>
#include <complex>
#include <sycl/sycl.hpp>
#include <oneapi/mkl/dft.hpp>

int main(int, char const**)
{
  sycl::queue q(sycl::gpu_selector_v);

  const int batch_size = 4;
  const int signal_size = 16;
  const int output_size = signal_size / 2 + 1;

  // Allocate input and output buffers
  auto* input = sycl::malloc_device<double>(signal_size * batch_size, q);
  auto* output = sycl::malloc_device<std::complex<double>>(output_size * batch_size, q);

  // descriptor for batched 1D FFTs
  oneapi::mkl::dft::descriptor<oneapi::mkl::dft::precision::DOUBLE, oneapi::mkl::dft::domain::REAL> desc(signal_size);
  desc.set_value(oneapi::mkl::dft::config_param::NUMBER_OF_TRANSFORMS, batch_size);
  desc.set_value(oneapi::mkl::dft::config_param::PLACEMENT, oneapi::mkl::dft::config_value::NOT_INPLACE);
  desc.set_value(oneapi::mkl::dft::config_param::FWD_DISTANCE, signal_size);
  desc.set_value(oneapi::mkl::dft::config_param::BWD_DISTANCE, output_size);
  desc.commit(q);
  auto desc_ptr = &desc;

  // Create a graph
  auto modifiable = sycl::ext::oneapi::experimental::command_graph<sycl::ext::oneapi::experimental::graph_state::modifiable>(q.get_context(), q.get_device());

  auto nodeFFT = modifiable.add([=](sycl::handler& cgh) {
    cgh.host_task([=]() {
      auto compute_event = oneapi::mkl::dft::compute_forward(*desc_ptr, input, output);
      compute_event.wait();
    });
  });

  // if any device node is added a deadlock will occur
  bool create_deadlock = true;
  if (create_deadlock) {
    auto nodeDummy = modifiable.add([=](sycl::handler& cgh) {
      cgh.single_task([=]() {
        input[0] = 42.0;
      });
    });
    modifiable.make_edge(nodeFFT, nodeDummy);
  }

  // Finalize the graph
  auto executable = modifiable.finalize();

  // Execute the graph
  q.ext_oneapi_graph(executable).wait_and_throw();
  std::cout << "Graph executed successfully.\n" << std::flush;

  // Clean up
  sycl::free(input, q);
  sycl::free(output, q);
  return 0;
}

死锁时的线程调用栈信息 #

1   ntdll!ZwWaitForAlertByThreadId 0x7ffca3be5844
2   ntdll!RtlSleepConditionVariableSRW 0x7ffca3b14d4e
3   SleepConditionVariableSRW 0x7ffca1319618 
4   MSVCP140D!?_Winerror_message *std * *YAKKPEADK *Z 0x7ffc56602107 
5   _Cnd_timedwait 0x7ffc566022de 
6   sycl8d!?verifyReductionProps *detail *_V1 *sycl * *YAXAEBVproperty_list *23 * *Z 0x7ffbab7baa0d
7   sycl8d!?verifyReductionProps *detail *_V1 *sycl * *YAXAEBVproperty_list *23 * *Z 0x7ffbab7ad999 
8   sycl8d!?finalize *handler *_V1 *sycl * *AEAA?AVevent *23 *XZ 0x7ffbab814dff 
9   sycl8d!.sycl_unregister_lib 0x7ffbab73867f 
10  sycl8d!.sycl_unregister_lib 0x7ffbab7370a5 
11  sycl8d!.sycl_unregister_lib 0x7ffbab742dfe 
12  sycl8d!?submit_with_event_impl *queue *_V1 *sycl * *AEAA?AVevent *23 *V?$function *$$A6AXAEAVhandler *_V1 *sycl * * *Z *std * *AEBVSubmissionInfo *detail *23 *AEBUcode_location *823 *_N *Z 0x7ffbab86b933 
13  mkl_sycl_dftd.5!??1?$error_handler *V?$complex *N *std * * *detail *vm *mkl *oneapi * *UEAA *XZ 0x7ffba1b86380 
14  mkl_sycl_dftd.5!??1?$error_handler *V?$complex *N *std * * *detail *vm *mkl *oneapi * *UEAA *XZ 0x7ffba1b85651 
15  mkl_sycl_dftd.5!??1error_handler_base *detail *vm *mkl *oneapi * *UEAA *XZ 0x7ffba1aa8ff7 
16  mkl_sycl_dftd.5!??1error_handler_base *detail *vm *mkl *oneapi * *UEAA *XZ 0x7ffba1aa7a4e 
17  mkl_sycl_dftd.5!??$compute_forward *V?$descriptor *$0CE *$0CB * *dft *mkl *oneapi * *V?$complex *N *std * *V56 * *dft *mkl *oneapi * *YA?AVevent *_V1 *sycl * *AEAV?$descriptor *$0CE *$0CB * *012 *PEAV?$complex *N *std * *1AEBV?$vector *Vevent *_V1 *sycl * *V?$allocator *Vevent *_V1 *sycl * * *std * * *8 * *Z 0x7ffba1a508cd 
... <More>

解决方案 #

方法1：直接将MKL FFT操作关联到command_graph handler #

原代码在host_task内调用compute_forward并手动wait，会阻塞SYCL命令执行线程，导致后续device节点无法处理。正确做法是将MKL操作直接绑定到command_graph的handler，让runtime自动管理依赖：

auto nodeFFT = modifiable.add([=](sycl::handler& cgh) {
  // 无需host_task，直接把MKL计算关联到当前handler
  oneapi::mkl::dft::compute_forward(*desc_ptr, input, output, cgh);
});

这种方式下，MKL的FFT操作会成为command_graph的一部分，依赖关系由SYCL runtime自动处理，不会出现线程阻塞。

方法2：使用独立队列执行host_task内的MKL操作 #

如果必须用host_task包装MKL调用，可以创建独立的host队列，避免占用主队列的执行线程：

// 创建独立的host队列，与主队列共享上下文和设备
sycl::queue host_q(q.get_context(), q.get_device(), sycl::host_selector_v);

auto nodeFFT = modifiable.add([=](sycl::handler& cgh) {
  cgh.host_task([=]() {
    // 指定使用独立的host队列执行MKL操作
    auto compute_event = oneapi::mkl::dft::compute_forward(*desc_ptr, input, output, host_q);
    compute_event.wait();
  });
});

这样host_task内的等待不会阻塞主队列处理device节点的任务。

方法3：临时调整节点依赖方向（仅作为测试用 workaround） #

如果上述方法暂时无法实施，可以尝试反转节点依赖顺序（需符合业务逻辑）：

if (create_deadlock) {
  auto nodeDummy = modifiable.add([=](sycl::handler& cgh) {
    cgh.single_task([=]() {
      input[0] = 42.0;
    });
  });
  // 修改依赖：先执行dummy节点，再执行FFT节点
  modifiable.make_edge(nodeDummy, nodeFFT);
}

这种方法不保证通用，仅能在特定场景下规避死锁。

内容的提问来源于stack exchange，提问作者schoeffmann