65.9K
CodeProject 正在变化。 阅读更多。
Home

使用 oneAPI 将 CUDA 转换为 SYCL

Sergiu Ovidiu Oprea
starIconstarIconstarIconstarIconstarIcon

5.00/5 (1投票)

2022年4月13日

CPOL

9分钟阅读

viewsIcon

8491

本文将着手介绍将 CUDA 转换为 SYCL 的过程。

作为开发者，我们一直在致力于利用专用架构来加速我们的应用程序。每个设备可能都需要特定的优化程序来实现最佳性能。这种情况通常会让我们接触到各种编程语言和厂商特定的库。因此，跨架构开发应用程序极具挑战性。

但如果我们关注性能和效率，我们就需要随着新硬件的出现，定期在这些新硬件上重用我们的代码。我们必须确保不浪费任何晶体管、电阻或半导体。

为了实现高性能和高效率，我们需要一个统一且简化的编程模型，使我们能够为手头的任务选择最佳硬件。我们需要一种高级、开放标准、异构的编程语言，它既基于标准的演进，又可扩展。它必须能够提高开发人员的生产力，同时在各种架构上提供一致的性能。oneAPI 规范解决了这些挑战。

oneAPI 规范包括 Data Parallel C++ (DPC++), oneAPI 对 Khronos SYCL 标准的实现。它还包括特定的库和硬件抽象层。oneAPI 技术咨询委员会一直在根据行业标准迭代完善 oneAPI 规范。此外，Intel oneAPI 工具包通过提供编译器、优化库、Intel® DPC++ 兼容工具 (DPCT) 以及高级分析和调试工具来提供规范的实现。

本文将演示如何使用 DPCT 将现有的 Compute Unified Device Architecture (CUDA) 应用程序迁移到 SYCL。我们将从 SYCL 规范的高层概述开始，并描述兼容工具的工作原理。然后，我们将展示如何迁移简单的 CUDA 代码到 SYCL。

通过 Jupyter Notebook 进行的实际演示将展示串行步骤。Jupyter Notebook 补充了本文，使我们能够运行下面描述的代码并将其用作沙盒。Notebook 还提供

  • 完整的 CUDA 实现
  • 迁移生成的 SYCL 代码
  • 我们在此教程中手动优化的版本

首先，让我们来了解一下 SYCL 和 Intel DPC++ 兼容工具。

什么是 SYCL？

SYCL（发音为 sickle）是一个免费、开放的单一源 C++ 标准。它规定了一个抽象层，允许在异构架构上进行编程。

通用的异构编程模型遵循国际标准化组织 (ISO) C++ 规范。这种标准化使得我们的代码能够无缝地在多个设备上运行。

图 1 – 目前可用的 SYCL 实现

想象一下使用 Nvidia 图形处理单元 (GPU) 来加速我们单一源 C++ 应用程序的部分功能。Nvidia 提供了 CUDA，这是一个通用并行编程模型，用于加速 Nvidia GPU 上的代码。

但是，如果我们想使用另一个厂商的 GPU 或现场可编程门阵列 (FPGA) 而不是 Nvidia GPU 呢？我们必须将我们的 CUDA 代码迁移到新的架构。这个过程可能既繁琐又耗时。但是，借助兼容工具，我们可以将代码迁移到 SYCL。它平均能自动转换 90-95% 的代码，显著提高生产力。

Intel DPC++ 兼容工具

Intel DPC++ 兼容工具 (DPCT) 帮助开发者将现有的 CUDA 迁移到 SYCL。它缩短了迁移时间，生成了人类可读的代码，并指出了需要手动干预的代码部分。

兼容工具提供了一套丰富的选项来控制迁移过程。例如，我们可以选择在生成的代码中使用统一共享内存 (USM) 还是缓冲区和访问器。

接下来我们将观察 DPCT 的实际应用。

示例：将向量加法 CUDA 迁移到 SYCL

为了提供迁移过程的实际概览，本文使用了一个简单的 CUDA 向量加法实现。我们将仔细查看兼容工具生成的代码。主要关注 CUDA 和 SYCL 差异最大的代码部分。

我们将使用 Intel® oneAPI Base Toolkit 中的兼容工具和 Intel® oneAPI DPC++/C++ 编译器来完成此任务。要安装该工具包，请遵循 oneAPI 安装指南。

使用以下工作流程成功将现有 CUDA 应用程序迁移到 SYCL

  1. 使用 intercept-build 工具拦截 Makefile 发出的命令，并将它们保存在 JSON 格式的编译数据库文件中。此步骤对于单一源项目是可选的。
  2. 使用 DPCT 将 CUDA 代码迁移到 SYCL。
  3. 验证生成代码的正确性，并在警告消息明确指出时手动完成迁移。请参阅 Intel DPC++ 兼容工具开发者指南和参考以修复警告。
  4. 使用 Intel oneAPI DPC++/C++ 编译器编译代码，运行程序，然后检查输出。

然后，您可以使用 Intel 的 oneAPI 分析和调试工具，包括 Intel® VTune Profiler，进一步优化您的代码。

以向量加法为例。向量加法涉及将向量 A 和 B 的元素相加到向量 C 中。CUDA 内核计算方式如下：

__global__ void vector_sum(const float *A, const float *B, float *C, const int num_elements){
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < num_elements) C[idx] = A[idx] + B[idx];}

在 CUDA 中，一组线程是线程块，相当于 SYCL 中的工作组。但是，我们以不同的方式计算线程索引。在 CUDA 中，我们使用内置变量来标识线程（请参阅我们上面如何计算 idx 变量）。迁移到 SYCL 后，相同的内核如下所示：

void vector_sum(const float *A, const float *B, float *C, const int num_elements, sycl::nd_item<3> item_ct1){
    int idx = item_ct1.get_local_range().get(2) * item_ct1.get_group(2) + item_ct1.get_local_id(2);
    if (idx < num_elements) C[idx] = A[idx] + B[idx];}

像 CUDA 线程一样，SYCL 中的工作项在全局空间中有一个全局标识符，或在工作组内有一个局部标识符。我们可以从 nd_item 变量获取这些标识符。因此，我们不再需要显式计算全局标识符。

但是，本次演示展示了我们如何在 SYCL 中执行此操作，因此我们看到了与 CUDA 内置变量的相似之处。请注意，由于 CUDA 中的 dim3 类型，nd_items 是三维的。在这种情况下，我们可以将 nd_items 全部设为一维。此操作将工作项映射到向量中的每个元素。

要运行 CUDA 内核，我们必须设置块大小和所需的块数。在 SYCL 中，我们必须定义执行范围。如下面的代码所示，我们使用 nd_range 变量来完成此操作，该变量结合了全局范围和局部范围。全局范围表示工作项的总数，而局部范围是工作组的大小。

我们还必须确保不要超过设备工作组的最大大小，因为这样做可能会触发 DPCT1049 警告。我们在以下代码中通过设置设备可以处理的最大块大小来解决此挑战。请查看 Jupyter Notebook 以了解更多信息。

const int num_elements = 512;

dpct::device_info prop;
dpct::dev_mgr::instance().get_device(0).get_device_info(prop);
const size_t max_block_size = prop.get_max_work_group_size();
const size_t block_size = std::min<size_t>(max_block_size, num_elements);

range<1> global_rng(num_elements);
range<1> local_rng(block_size);
nd_range<1> kernel_rng(global_rng, local_rng);

要调用我们的 SYCL 内核，我们使用 parallel_for 和执行范围将内核提交到队列。每个工作项都会调用一次内核。在这种情况下，我们有相同数量的工作项来处理每个向量元素。让我们看看它是如何工作的：

dpct::get_default_queue().parallel_for(kernel_rng, [=](nd_item<1> item_ct1) {
    vector_sum(d_A, d_B, d_C, num_elements, item_ct1);});

到目前为止，我们已经探讨了如何实现和运行内核。但是，在运行内核之前，我们需要考虑内存分配并将数据复制到设备。

  1. 首先，我们在主机上为操作数向量分配内存并对其进行初始化。
  2. 然后，我们在设备上做同样的事情。CUDA 使用 cudaMalloc 例程。默认情况下，DPCT 将此例程迁移到 malloc_device，它使用统一共享内存 (USM)。
  3. 现在，我们使用 memcpy 命令将向量从主机内存复制到设备。

完成这些步骤后，我们运行内核。执行完成后，我们将结果复制回主机。然后我们检查结果的正确性。最后，通过分别调用 free 和 sycl::free 来释放主机和设备上的内存。

现实世界用例：将 CUDA 代码迁移到 SYCL

柏林祖斯研究所 (ZIB) 使用 DPC++ 兼容工具成功地将海啸模拟应用程序 easyWave 从 CUDA 移植到了 SYCL。1 其结果在 Intel CPU、GPU 和 FPGA 架构上取得了强劲的性能，并且在 Nvidia P100 上的性能接近 CUDA 性能（误差在 5% 以内）。2 更多详情请参阅：ZIB：使用海啸模拟 easyWave 的 oneAPI 案例研究。

斯德哥尔摩大学 使用 oneAPI 和 SYCL 改进了 GROMACS（一个用于模拟蛋白质、脂质和核酸的分子动力学软件包，用于设计新的抗癌药物、COVID-19 等）在异构架构上的并行化。oneAPI 的开放且基于标准的编程、对 OpenMP 的支持以及一流的 OpenCL 实现对此项工作有所帮助。该团队使用 DPC++ 兼容工具将 GROMACS 的 Nvidia CUDA 代码移植到 SYCL，创建了一个新的、单一的、可移植的、跨架构就绪的代码库。这大大简化了开发，并为在多架构（CPU 和 GPU）、多厂商环境中的部署提供了灵活性。了解更多：GROMACS 添加 SYCL 支持的经验。

Bittware 使用 oneAPI 创建了一个可在其 FPGA 上运行的单一代码库。这简化了设计、测试和实施，使得应用程序可以在几天内而不是几周内投入使用。

高阶穷举连锁分析 是一种计算复杂度非常高的生物信息学应用程序，它搜索基因标记之间的相关性，例如单核苷酸多态性 (SNP)（在至少 1% 的人群中出现的 DNA 序列中的单个碱基变化）和表型（例如疾病状态）。寻找基因型和表型之间的新关联有助于改善预防性护理、个性化治疗以及为更多疾病开发更好的药物。来自 INESC-ID Lisboa (Instituto de Engenharia de Sistemas e Computadores: Investigação e Desenvolvimento em Lisboa) 的专家使用 DPC++ 兼容工具和 Intel® DevCloud 将 OpenMP 和 CUDA 代码迁移到 SYCL，现在正在 Intel Iris Xᵉ Max (独立 GPU) 上运行代码。

结论

使用 Intel 的兼容工具，我们轻松地将现有的 CUDA 迁移到了 SYCL。现在，我们可以无缝地在多个设备上运行代码，跨越架构和厂商界限。

这种方法提高了我们的生产力，使我们能够专注于应用程序的性能。想象一下，如果没有 DPCT，迁移我们的代码将需要多少时间和精力。

立即尝试工具包中的高级分析和调试工具。

资源

1Nvidia CUDA 代码被移植到 Data Parallel C++ (DPC++)，它是 oneAPI 的 SYCL 实现，以创建新的跨架构就绪代码。

2详情请参阅 XPUG 演示文稿：ZIB：使用海啸模拟 EasyWave 的 oneAPI 案例研究，从 CUDA 到 DPC++ 再回到 Nvidia GPU 和 FPGA – 配置：计算域：约 2000 x 1400 个单元；10 小时模拟时间。同一代码可在 CPU、Intel GPU 和 FPGA 上生成有效数据。使用 Intel® DPC++ 兼容工具将 CUDA 代码迁移到 Data Parallel C++ 后，在 DevCloud Coffee Lake Gen9.5 GT2 iGPU 上进行 oneAPI 性能演进，并使用开源 Intel LLVM w/ CUDA 支持（Codeplay 贡献）进行构建。在 Nvidia P100-SXM2-16GB 上进行的典型应用程序运行显示，迁移后的 DPC++ 代码仅比 CUDA 代码慢 4%。结果：相同的 DPC++ 代码几乎无需修改即可针对不同的平台。• 性能与特定架构的 CUDA 代码相当。有关工作负载和配置，请访问www.Intel.com/PerformanceIndex。结果可能有所不同。Intel 不控制或审计第三方数据。您应该查阅其他来源来评估准确性。

© . All rights reserved.