c++ - Sycl 内核调用非常慢

Question

我是 stackoverflow、sycl 和 gpu 编程的新手。我有一个带有基本 sycl 内核的项目。逻辑是有效的，所以我在问题中跳过它。另外在编译和执行过程中也没有错误。

现在最大的问题是 sycl 代码的调用非常慢。首先，我认为这是一些内存复制或类似操作，所以除了您在下面看到的内容之外，我遗漏了任何内容(最低限度，注释是代码在不是最小内核时所在的位置)。

我测量的时间:(发布 x64)

显示 Visual Studio 调试器，空内核调用的函数总时间:~100 毫秒

使用 Cuda Nsight，OpenCl 内核执行时间:~5 us

5 us 的内核 gpu 时间与空内核的预期一样快。
但是我的代码中 c++ 函数的总时间慢了 100 毫秒。

这里可能有什么问题？或者 sycl 开销预计会这么慢？(我真的很怀疑)

我的努力:

我将我的 compute++.exe 标志从 -O2 更改为 -O3，这将总时间缩短了大约 5 到 10 毫秒。

我将内核设置为最低限度

dll 函数内的代码:

 { //scope

    sycl::gpu_selector gpuSel;
    sycl::queue myQueue(gpuSel);

    //....buffers

    auto ra = range<1>(size);

    myQueue.submit([&](sycl::handler& hd)
    {
            //....get_access<access::mode::read>

            auto kernel = ([=](cl::sycl::id<1> id)
            {
                    //...some vector math
            });

            hd.parallel_for<someName>(ra, kernel);             
    });

    myQueue.wait();
}

我在用:

Visual Studio 2019

ComputeCpp 社区 2.0.0

最新的 Cuda 驱动程序

NVIDIA Gtx 980 ptx64(实验性 ComputeCpp 支持)

计算++调用:

"..\compute++.exe" -sycl -D_ALLOW_COMPILER_AND_STL_VERSION_MISMATCH -O3 -mllvm -inline-threshold=1000 -intelspirmetadata -sycl-target ptx64 -std=c++14 -I"../Codeplay/ComputeCpp/include" -I"../NVIDIA GPU Computing Toolkit/CUDA/v10.2/include"  -sycl-ih something.cpp.sycl -c something.cpp

总结:
sycl 内核的总执行时间很慢。
我可以在这里做些什么来改进它还是因为在 Nvidia gpu 上实现了 sycl/computecpp 并且预计会这么慢？

Answer 1

首先，我要指出这是一组非常简单的 SYCL 代码，因此如果您想衡量性能，它可能不是一个非常相关的示例。这是一篇研究论文，展示了 ComputeCpp 与 CUDA 执行缩减算法基准的可比性能，see slide 40 for the chart .您还将在演示文稿中看到，性能提升会根据正在处理的数据集的大小呈指数增长。这对于 HPC 编程来说通常是一样的，因为 GPU 的好处通常只有在处理更大的数据集时才能看到。

您看到的不同之处在于 ComputeCpp 使用 OpenCL 回调，而 NVIDIA OpenCL 驱动程序在使用这些回调时似乎确实引入了开销。这是一个 relevant post不久前关于这个

如果您要编写一个使用回调的简单 OpenCL 内核，它将表现出相同的行为。

我还要补充一点，我们有 implemented NVIDIA support for the DPC++ compiler直接使用 CUDA 并且没有看到相同级别的开销。您可以在我们的博客文章中找到更多相关信息，如果您想在 NVIDIA 硬件上运行 SYCL 代码，那么值得一试。