如何实现适配SYCL Kernel的可定制Grid类apply_all_occupied方法？

要解决这两个问题，核心是让用户能明确指定每个buffer的访问模式，同时让functor直接接收展开的、类型明确的accessor而非元组。可以通过让用户传递绑定了访问模式的buffer对象，结合C++模板折叠表达式实现参数展开来完成。

步骤1：定义带访问模式的buffer包装方式 #

我们可以用自定义结构体绑定buffer引用和访问模式，搭配辅助函数简化创建：

template<typename T>
struct BufferWithAccess {
    sycl::buffer<T>& buf;
    sycl::access_mode mode;
};

// 辅助函数，快速生成带访问模式的buffer包装对象
template<typename T>
BufferWithAccess<T> make_buffer_access(sycl::buffer<T>& buf, sycl::access_mode mode) {
    return {buf, mode};
}

步骤2：实现Grid类的apply_all_occupied方法 #

利用C++17折叠表达式逐个处理参数，生成对应accessor后展开传递给用户functor：

class Grid {
public:
    template <typename TFunctor, typename... TBufferArgs>
    sycl::event apply_all_occupied(sycl::queue& q, TFunctor&& function, TBufferArgs&&... buffer_args) {
        return q.submit([this, 
                         function = std::forward<TFunctor>(function), 
                         ... args = std::forward<TBufferArgs>(buffer_args)]
                        (sycl::handler& cgh) mutable {
            // 获取Grid内部cells的read访问器
            auto cells_accessor = d_cells.get_access<sycl::access_mode::read>(cgh);

            // 为每个带访问模式的buffer生成accessor
            auto get_accessor = [&cgh](auto&& arg) {
                using ArgType = std::decay_t<decltype(arg)>;
                if constexpr (std::is_same_v<ArgType, BufferWithAccess<typename ArgType::value_type>>) {
                    return arg.buf.get_access(arg.mode, cgh);
                } else {
                    // 兼容直接传buffer的场景，默认read模式
                    return arg.get_access<sycl::access_mode::read>(cgh);
                }
            };

            auto accessors = std::make_tuple(get_accessor(std::forward<TBufferArgs>(args))...);

            // 遍历所有cell，仅对已占用cell执行用户逻辑
            cgh.parallel_for(d_cells.range(), 
                [cells_accessor, accessors, function = std::move(function)]
                (sycl::id<3> id, sycl::kernel_handler kh) {
                    auto cell = cells_accessor.at(kh, id);
                    if (cell.is_occupied()) {
                        // 展开accessor元组，传递给用户functor
                        std::apply([&](auto&&... accs) {
                            function(kh, cell, std::forward<decltype(accs)>(accs)...);
                        }, accessors);
                    }
                }
            );
        });
    }

private: 
    sycl::buffer<Cell> d_cells;
};

步骤3：用户调用示例 #

用户可以明确指定每个buffer的访问模式，functor直接接收类型明确的accessor：

// 初始化统计用buffer
std::vector<unsigned> count_occupied(1, 0);
sycl::buffer<unsigned> d_count_occupied{count_occupied.data(), count_occupied.size()};

// 用户自定义统计逻辑
auto count_fn = [](sycl::kernel_handler kh, const Cell& cell, sycl::accessor<unsigned, 1, sycl::access_mode::write> count_acc) {
    sycl::atomic_ref<
        unsigned,
        sycl::memory_order::relaxed,
        sycl::memory_scope::device>
        count{count_acc[0]};
    count++;
};

// 调用Grid方法，指定访问模式
sycl::queue q;
Grid grid;
grid.apply_all_occupied(q, count_fn, make_buffer_access(d_count_occupied, sycl::access_mode::write))
    .wait_and_throw();

方案优势 #

• 支持自定义访问模式：用户可为每个buffer单独指定read/write等模式，满足不同操作需求。
• 避免元组传递：用户functor直接接收类型明确的accessor，代码可读性、易用性大幅提升。
• 兼容性强：可选兼容直接传递buffer的场景，降低用户学习成本。

内容的提问来源于stack exchange，提问作者Fantastic Mr Fox