基于oneAPI SYCL与Nvidia插件的单GPU核函数并发问题

可行性结论 #

完全可行，但你遇到的串行/并发差异是由NVIDIA oneAPI插件的编译与调度策略导致的，并非SYCL标准或硬件限制。

问题原因分析 #

你观察到的现象核心在于：

• 当复用同一个lambda时，编译器将其编译为单一kernel实例，NVIDIA插件可以识别为可复用的任务，进而调度到GPU的并发流中执行。
• 当使用两个独立lambda时，编译器会生成两个独立的kernel编译单元，插件默认的调度策略（尤其是在有序队列下）可能会将这两个kernel的任务串行化执行，导致你的死锁场景——第一个kernel等待第二个修改内存，而第二个被阻塞在串行队列中无法启动。

你的补充测试也验证了这一点：独立lambda的delay()任务耗时翻倍，说明它们被串行调度，而同一lambda的任务则并发执行。

解决方案 #

以下几种方法可以解决独立lambda的串行问题，避免死锁：

1. 使用无序队列（Out-of-Order Queues） #

默认SYCL队列是有序队列（in-order），虽然标准不禁止并发，但NVIDIA插件可能对有序队列中的独立kernel任务采用串行调度。显式创建无序队列，允许运行时自由调度任务并发：

std::vector<sycl::queue> queues;
// 创建无序队列，允许任务并发执行
queues.emplace_back(sycl::gpu_selector_v, sycl::property::queue::in_order(false));
sycl::device d = queues[0].get_device();
queues.emplace_back(d, sycl::property::queue::in_order(false));

修改队列属性后，两个独立lambda的核函数会被调度到并发的CUDA流中，从而避免死锁。

2. 统一核函数的编译单元 #

将两个独立lambda的逻辑封装到同一个模板函数或函数对象中，让编译器生成同一个kernel实例，复用并发调度逻辑：

// 封装统一的核函数逻辑
template<int Flag>
void kernel_logic(volatile int* ptr) {
    if constexpr (Flag == 0) {
        *ptr = 0xaaaaaaaa;
        while (*ptr != 0xbbbbbbbb);
    } else {
        while (*ptr != 0xaaaaaaaa);
        *ptr = 0xbbbbbbbb;
    }
}

// 提交任务时复用同一模板
queues[0].single_task([=](){ kernel_logic<0>(ptr); });
queues[1].single_task([=](){ kernel_logic<1>(ptr); });

这种方式让编译器将两个任务识别为同一模板的不同实例，插件更可能调度它们并发执行。

3. 调整NVIDIA插件的环境变量 #

部分环境变量可以修改插件的调度行为：

• SYCL_PI_CUDA_ENABLE_GRAPHS=0：禁用CUDA图优化，避免插件将任务串行化到图中。
• SYCL_PI_CUDA_STREAM_PER_QUEUE=1：确保每个队列对应独立的CUDA流，强制任务并发。

在运行程序前设置这些环境变量，例如：

export SYCL_PI_CUDA_ENABLE_GRAPHS=0
./your_program

验证建议 #

修改后可以再次用nsys profile分析任务调度，确认两个核函数是否在不同CUDA流中并发执行；同时观察死锁是否消失，耗时是否回归到并发水平。

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