sycl 内核调用很慢

Posted 2023-03-10

【中文标题】sycl 内核调用很慢

【英文标题】：Sycl kernel-call very slow

【发布时间】：2020-05-09 19:40:00

【问题描述】：

我是 ***、sycl 和 gpu 编程的新手。我有一个使用基本 sycl 内核的项目。逻辑是有效的，所以我在问题中跳过它。同样在编译和执行过程中也没有报错。

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

我测量的时间：（发布 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 gpus 上实现了 sycl/computecpp 并且预计会这么慢？

【参考方案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 代码，值得一试。

【参考方案2】：

当您想将 3 或 4 个数字相加或相乘时，GPU 非常糟糕。为此，您最好使用经过优化的 CPU，并且您可能有一个 AVX 扩展，该扩展已针对矢量数学进行了优化。因此，您应该将cl::sycl::gpu_selector 替换为cl::sycl::cpu_selector。我不确定sycl是否使用AVX，当你有一个时，它肯定会使用多线程。

但是当您尝试添加 500'000 个数字时，GPU 会比 CPU 快得多。

This video解释得很好。