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

将 Jacobi 迭代方法从 CUDA 迁移到 SYCL

Intel
starIconstarIconstarIconstarIconstarIcon

5.00/5 (2投票s)

2022年7月27日

CPOL

24分钟阅读

viewsIcon

5945

本文档演示了如何将用 CUDA* 编写的线性代数雅可比迭代法迁移到 SYCL* 异构编程语言。

雅可比迭代法

雅可比迭代法用于在数值线性代数中找到对角占优形式 Ax = b 的线性方程组的近似数值解。该算法从 x 的初始估计值开始，并迭代更新它直到收敛。如果矩阵 A 是对角占优的，则雅可比方法保证收敛。

CUDA 到 SYCL 的迁移方法

本文档介绍了 CUDA 到 SYCL 迁移的两种方法

  • 第一种方法是使用 Intel® DPC++ 兼容性工具自动将 CUDA 源迁移到 SYCL 源。该工具迁移 80% 到 90% 的代码，并为其余部分生成警告，其余部分必须手动迁移到 SYCL。我们将查看 Intel DPC++ 兼容性工具生成的警告，并学习如何迁移 Intel DPC++ 兼容性工具未迁移的代码。这种方法有助于加速 CUDA 源到 SYCL 的迁移，并且已被证明对于大型代码库特别有用。
  • 第二种方法是通过分析 CUDA 源并将所有 CUDA 特定的调用替换为等效的 SYCL 调用来进行手动迁移。这种方法有助于 CUDA 开发人员理解 SYCL 编程。迁移完成后，我们使用 VTuneTM Profiler 和 Intel® Advisor Roofline 进行性能分析，以了解性能瓶颈。然后我们着手优化代码以提高性能。有关更多详细信息，请查阅 SYCL 2020 规范。

以下流程图显示了用于 CUDA 到 SYCL 迁移的方法

使用 Intel® DPC++ 兼容性工具进行迁移

Intel DPC++ 兼容性工具及其使用方法

Intel DPC++ 兼容性工具是 Intel® oneAPI Base Toolkit 的一个组件，它帮助开发人员将用 CUDA 编写的程序迁移到用 DPC++ 编写的程序。

尽管 Intel DPC++ 兼容性工具会自动迁移大部分代码，但要完全迁移仍需要一些手动工作。该工具会输出警告以指示何时何地需要手动干预。这些警告具有“DPCT10XX”形式的指定 ID，可以在 开发人员指南和参考 中查阅。此指南包含所有警告的列表、其描述和修复建议。

使用 Intel DPC++ 兼容性工具将 CUDA 迁移到 SYCL

Intel DPC++ 兼容性工具帮助将基于 CUDA 的代码传输到 SYCL 并生成可读代码，同时保留原始代码中的原始标识符。该工具还检测并将标准 CUDA 索引计算转换为 SYCL。此示例的目标是使用 Intel DPC++ 兼容性工具执行从 CUDA 到 SYCL 的迁移过程，并演示迁移后的 SYCL 代码在不同 GPU 和 CPU 设备中实现的便携性。该工具通过拦截构建过程并将 CUDA 代码替换为 SYCL 对应代码来工作。

此 CUDA 源主要通过将 CUDA 表达式替换为等效的 SYCL 表达式，并将内核调用转换为向 SYCL 队列提交带有 lambda 表达式的 parallel_for 来迁移到 SYCL。

雅可比迭代的 Intel DPC++ 兼容性工具迁移代码可在 sycl_dpct_output 中找到。

要确保您拥有 CUDA 版本和所需的工具，请参阅 Intel DPC++ 兼容性工具系统要求。

按照以下步骤将 CUDA 雅可比迭代示例迁移到 SYCL

  1. 确保系统已安装 Nvidia CUDA SDK（在默认路径中），并且您已从 Intel® oneAPI Base Toolkit 安装了 Intel DPC++ 兼容性工具。
  2. 设置环境变量，setvars.sh 脚本位于 oneAPI 安装的根文件夹中，通常是 /opt/intel/oneapi/
    . /opt/intel/oneapi/setvars.sh
  3. 从以下位置获取雅可比迭代方法的 CUDA 实现：JacobiCUDA_Sample。
  4. 转到 CUDA 源文件夹并使用工具 intercept-build 生成编译数据库。这将创建一个 JSON 文件，其中包含所有编译器调用，并存储输入文件的名称和编译器选项。
    intercept-build make
  5. 使用 Intel DPC++ 兼容性工具迁移代码；它会将结果存储在迁移文件夹 dpct_output 中。
    Intel DPC++ 兼容性工具选项可简化迁移和调试
    --keep-original-code 在生成的 SYCL 文件的注释中保留原始 CUDA 代码。允许轻松比较原始 CUDA 代码与生成的 SYCL 代码。
    --comments 插入解释生成代码的注释
    ---always-use-async-handler 始终使用异步异常处理程序创建 cl::sycl::queue 
    dpct -p compile_commands.json
  6. 验证迁移并处理通过查阅 诊断参考 生成的任何 Intel DPC++ 兼容性工具警告，以获取有关 Intel DPC++ 兼容性工具警告的详细信息。
  7. 修改 makefile 以在适当的时候使用 DPCPP 编译器，并删除 CUDA 特定的编译标志。

有关更多信息，请参阅 Intel® DPC++ 兼容性工具最佳实践。

实现未迁移的 SYCL 代码

一旦 Intel DPC++ 兼容性工具迁移了代码，未迁移的代码可以通过警告来识别。这些警告有一个指定的 ID，可以通过查阅开发人员指南和参考来手动解决。

雅可比迭代的 Intel DPC++ 兼容性工具完整迁移代码可在 sycl_dpct_migrated 中找到。

迁移和手动变通方法生成的警告

  • DPCT1025：SYCL 队列在创建时忽略了标志和优先级选项。
cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);

DPCPP 中 CUDA 流的等效项将是队列，因此 Intel DPC++ 兼容性工具创建 SYCL 队列时，忽略了标志和优先级选项。

sycl::queue *stream1;
  • DPCT1065：如果没有访问全局内存，请考虑将 sycl::nd_item::barrier() 替换为 sycl::nd_item::barrier(sycl::access::fence_space::local_space) 以获得更好的性能。
    cg::sync(cta);

在内核内部，如果没有访问全局内存，Intel DPC++ 兼容性工具建议替换 barrier() 以获得更好的性能。在这种情况下，用户应该检查内存访问并进行修改。

  • DPCT1007：DPC++ 兼容性工具不支持此 CUDA API 的迁移。
cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

atomicAdd(sum, temp_sum);

许多 CUDA 设备属性没有 SYCL 等效项，略有不同，或者目前不受支持。在许多情况下，这将导致检索到不正确的值。因此，用户必须手动审查和纠正对设备的信息查询。

以下是上述代码的变通方法

sub_group tile_sg = item_ct1.get_sub_group();
int tile_sg_size = tile_sg.get_local_range().get(0);
atomic_ref<double, memory_order::relaxed, memory_scope::device,
      access::address_space::global_space> at_sum { *sum };

at_sum.fetch_add(temp_sum);
  • DPCT1039：生成的代码假定“sum”指向全局内存地址空间。如果它指向局部内存地址空间，请将“dpct::atomic_fetch_add”替换为
"dpct::atomic_fetch_add<double, sycl::access::address_space::local_space>"
     atomicAdd(sum, temp_sum);

为了获得更好的内存访问和性能，我们需要将访问地址空间指定为全局或局部空间。

sycl::atomic<int> at_sum(sycl::make_ptr<int,sycl::access::address_space::global_space>((int*)   sum));

sycl::atomic_fetch_add<int>(at_sum, temp_sum);
  • DPCT1049：传递给 SYCL 内核的工作组大小可能超出限制。要获取设备限制，请查询 info::device::max_work_group_size。如果需要，调整工作组大小。

Intel DPC++ 兼容性工具建议查询 info::device::max_work_group_size 以获取设备限制并相应地调整工作组大小。

cgh.parallel_for(nd_range<3>(nblocks * nthreads, nthreads),

     [=](nd_item<3>item_ct1) [[intel::reqd_sub_group_size(ROWS_PER_CTA)]] {

         JacobiMethod(A, b, conv_threshold, x_new, x, d_sum, item_ct1, x_shared_acc_ct1.get_pointer(), b_shared_acc_ct1.get_pointer());

        });
  • DPCT1083：迁移代码中的局部内存大小可能与原始代码不同。检查迁移代码中分配的内存大小是否正确。

在迁移代码中，某些类型的大小与原始代码中不同；例如，sycl::float3 与 float3。因此，应在迁移代码中验证局部内存的分配大小。

  • 当使用宏指定块大小并用于创建 sycl::range 时，扩展值应改回宏。在这种情况下，工具会保留原始宏的注释。

分析 CUDA 源

雅可比迭代方法的 CUDA 实现可在以下位置获取：JacobiCUDA_Sample。

雅可比迭代方法实现的 CUDA 源位于以下文件中。

  • main.cpp——主机代码，用于
    • 设置 CUDA 流
    • 在 GPU 上分配内存
    • 在 CPU 上初始化数据
    • 将数据复制到 GPU 内存进行计算
    • 启动 GPU 上的计算
    • 验证并打印结果
  • jacobi.cu——用于在 GPU 上运行的雅可比迭代方法计算的内核代码
    • 定义内核
    • 分配共享局部内存
    • 用于线程块分区的协作组
    • 使用 warp 原语
    • 使用原子操作累加瓦片总和值

CUDA 代码只有两个文件：main.cpp 和 jacobi.cu。main.cpp 包含雅可比方法的 CPU 实现内存分配、初始化、内核启动、内存复制和使用 SDK 计时器计算执行时间。

jacobi.cu 包含内核、雅可比方法和最终误差。雅可比方法涉及将向量加载到共享内存中以加快内存访问速度，并将线程块划分为瓦片。然后，使用 warp 级原语在每个分区瓦片中执行输入数据的缩减。然后，这些中间结果通过原子添加累加到最终求和变量中。这也导致更快的实现，避免了不必要的块级同步。

此示例应用程序演示了使用流捕获、原子操作、共享内存和协作组等关键概念的 CUDA 雅可比迭代方法。

从 CUDA 迁移到 SYCL

在本节中，我们将通过分析 CUDA 代码并识别相关的 SYCL 特性，将 CUDA 代码迁移到 SYCL。CUDA 和 SYCL 的底层概念相似，但理解每种语言的术语对于将 CUDA 代码迁移到 SYCL 代码至关重要。有关更多详细信息，请查阅 SYCL 2020 规范。

main.cpp 和 jacobi.cu 中的 CUDA 代码将迁移到 main.cpp 和 jacobi.cpp 中的 SYCL 版本。

CUDA 头文件和 SYCL 头文件

在 CUDA 实现中，使用 cuda_runtime.h 头文件，它定义了 CUDA 运行时 API 的公共主机函数、内置类型定义，以及 CUDA 语言扩展和设备内在函数的函数叠加。

  #include <cuda_runtime.h>

SYCL 实现使用此单个头文件来包含 API 接口和 API 中定义的所有依赖类型。

  #include <CL/sycl.hpp>

CUDA 流和 SYCL 队列

CUDA 流是一系列操作，它们按照主机代码发出的顺序在设备上执行。主机将 CUDA 操作（例如，内核启动、内存复制）放入流中并立即继续。然后设备在资源空闲时从流中调度工作。同一流中的操作是先进先出 (FIFO) 有序的。另一方面，不同的流可能会相对于彼此或并发地无序执行其命令。

SYCL 具有将主机程序连接到单个设备的队列。程序通过队列向设备提交任务，并可以监视队列以等待完成。与 CUDA 流类似，SYCL 队列异步提交命令组以执行。但是，SYCL 是一个更高级别的编程模型，数据传输操作是从提交到任何队列的内核的依赖关系中隐式推断出来的。

在 CUDA 实现中，第一步是创建一个新的异步流。flag 参数决定了流的行为。cudaStreamNonBlocking 指定在创建的流中运行的工作可以与流 0（NULL 流）中的工作并发运行，并且创建的流不应与流 0 执行任何隐式同步。CUDA 流用于执行异步 memset 和 memcpy 以实现并发模型，然后使用指定的 stream 启动内核，以便它们在调用后立即返回到主机线程。CUDA 流的设置如下（main.cpp）：

cudaStream_t stream1;

checkCudaErrors(cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking));

在 SYCL 中，我们以与 CUDA 流类似的方式使用队列；队列异步提交命令组以执行。SYCL 运行时自动处理跨多个设备中多个队列的不同命令组（内核 + 依赖项）的执行顺序。我们可以使用 in_order 队列属性设置 SYCL 队列，并使用 default_selector()；这将确定要使用的 SYCL 设备。默认选择器选择第一个可用的 SYCL 设备。队列所需的系统资源在其超出范围后会自动释放。in_order 队列将确保内核计算仅在 memcpy 操作完成后才开始，并且不会发生内核执行重叠。

sycl::queue q{sycl::default_selector(),sycl::property::queue::in_order()};

更多信息可在 SYCL 队列 中找到。

GPU 设备上的内存分配—cudaMalloc 和 sycl::malloc_device

我们必须首先在 GPU 设备上分配内存，以便将其用于将数据复制到 GPU 内存，从而使其可用于 GPU 上的计算。cudaMalloc 函数可以从主机或设备调用，以在设备上分配内存，非常类似于主机的 malloc。使用 cudaMalloc 分配的内存必须使用 cudaFree 释放。

在 CUDA 中，使用 cudaMalloc 函数在 GPU 上进行内存分配如下：

checkCudaErrors(cudaMalloc(&d_b, sizeof(double) * N_ROWS));
checkCudaErrors(cudaMalloc(&d_A, sizeof(float) * N_ROWS * N_ROWS));

在 SYCL 中，使用 sycl::malloc_device 函数在加速器设备上进行内存分配如下：

d_b     = sycl::malloc_device<double>(N_ROWS, q); 
d_A     = sycl::malloc_device<float>(N_ROWS * N_ROWS, q);

sycl::malloc_device 成功时返回指向指定设备上新分配内存的指针。此内存无法在主机上访问。由 sycl::malloc_device 分配的内存必须使用 sycl::free 释放，以避免内存泄漏。

有关统一共享内存 (USM) 概念和内存分配的更多信息可在 SYCL USM 中找到。

将内存从主机复制到 GPU 内存

一旦在 GPU 上分配了内存，我们必须将内存从主机复制到设备，以便数据可在设备上进行计算。

在 CUDA 中，使用 cudaMemsetAsync 将内存从主机复制到 GPU，如下所示。内存相对于主机异步复制，因此主机将传输放入流中，调用可能立即返回。操作可以通过传递非零 stream 参数来选择性地与流关联。如果 stream 非零，则操作可能与其他流中的操作重叠。

checkCudaErrors(cudaMemsetAsync(d_x, 0, sizeof(double) * N_ROWS, stream1));

checkCudaErrors(cudaMemsetAsync(d_x_new, 0, sizeof(double) * N_ROWS, stream1));

checkCudaErrors(cudaMemcpyAsync(d_A, A, sizeof(float) * N_ROWS * N_ROWS, cudaMemcpyHostToDevice, stream1));

checkCudaErrors(cudaMemcpyAsync(d_b, b, sizeof(double) * N_ROWS, cudaMemcpyHostToDevice, stream1));

CUDA 流通过 cudaStreamSynchronize 进行同步，它会阻塞主机，直到流中所有发出的 CUDA 调用完成。

checkCudaErrors(cudaStreamSynchronize(stream));

在 SYCL 中，我们使用 memcpy 将内存从主机复制到设备内存。要初始化内存，可以使用 memset 初始化向量数据，如下所示：

q.memset(d_x,     0, sizeof(double) * N_ROWS);

q.memset(d_x_new, 0, sizeof(double) * N_ROWS);

q.memcpy(d_A, A, sizeof(float)  * N_ROWS * N_ROWS);

q.memcpy(d_b, b, sizeof(double) * N_ROWS);

第一个参数是带有值的内存地址指针；这必须是 USM 分配。SYCL memcpy 将数据从指针源复制到目标。源和目标可以是主机或 USM 指针。

内存是异步复制的，但在可以使用任何内存之前，我们需要通过使用以下方式进行同步，以确保复制完成：

q.wait();

wait() 将阻塞调用线程的执行，直到提交到队列的所有命令组都已完成执行。

有关 SYCL memcpy 和异步复制以及数据同步的更多信息可在 SYCL 队列 和 memcpy 以及 wait 中找到。

这就完成了主机端 CUDA 代码 (main.cpp) 到 SYCL 的迁移

main.cpp 的 CUDA 主机代码可在 main.cpp 中找到。

主机代码 main.cpp 的 SYCL 代码可在 main.cpp 中找到。

以下部分解释了 CUDA 内核代码 (jacobi.cu) 到 SYCL 的迁移。

将计算卸载到 GPU

CUDA 内核代码位于 jacobi.cu 中。计算发生在两个内核中：Jacobi 方法和最终误差。这些计算被卸载到设备。在 Jacobi 方法和最终误差计算中，我们都使用了共享内存、协作组和归约。向量被加载到共享内存中，以实现更快、更频繁的块内存访问。协作组用于进一步将工作组划分为子组。由于上述计算发生在子组内部，这消除了对块屏障的需求，并且适用于具有较低粒度的归约算法，使每个线程运行效率更高或有效地分配工作。归约使用 sync() 对不同的线程块进行同步，而不是整个网格，因此实现避免了同步块，速度更快。shift group left 是一种 SYCL 原语，用于在子组内进行计算，以累加所有线程值并传递给第一个线程。所有子组和通过原子加法累加。

最终误差用于计算 CPU 和 GPU 计算之间的误差总和，以验证输出。Warpsum 加上 x 减去 1 的绝对值（每个线程值都加上），然后所有 Warpsum 值都累加到 blocksum。最终误差存储在 g_sum 中。

在 CUDA 中，一组线程称为一个 thread block 或简称一个 block。这等同于 SYCL 中的 work-group 概念。block 和 work-group 都可以访问同一级别的层次结构，并公开类似的同步操作。

CUDA 流是 CUDA 操作的序列，从主机代码提交。这些操作按提交顺序异步执行。如果没有给出 CUDA 流，则会创建一个默认的 CUDA 流，并且所有操作都提交到默认流。

与 CUDA 流类似，SYCL 队列异步提交命令组以执行。但是，SYCL 数据传输操作是从提交到任何队列的内核的依赖关系中隐式推断出来的。

在 CUDA 中，内核使用以下参数启动：nblocks 指定网格的维度和大小，nthreads 指定每个块的维度和大小，第三个参数指定每个块动态分配的共享内存的字节数，stream 指定关联的流，这是一个可选参数，默认为 0。

JacobiMethod<<<nblocks, nthreads, 0, stream>>> (A, b, conv_threshold, x, x_new, d_sum);

在 SYCL 中，single_task、parallel_for 和 parallel_for_work_group 等内核构造都将函数对象或 lambda 函数作为其参数之一。函数对象或 lambda 函数中的代码在设备上执行。

q1.submit([&](handler &cgh) {

cgh.parallel_for(nd_range<3>(nblocks * nthreads, nthreads), [=](nd_item<3> item_ct1) {

          JacobiMethod(A, b, conv_threshold, x, x_new, d_sum, item_ct1, x_shared_acc_ct1.get_pointer(), b_shared_acc_ct1.get_pointer(), stream_ct1);

          });

});

队列设置完成后，在我们的命令组中，我们使用 parallel_for 提交一个内核。此函数将在多个工作项上并行执行内核。nd_range 指定一个 1、2 或 3 维的工作项网格，每个工作项执行内核函数，这些工作项在工作组中一起执行。nd_range 由两个 1、2 或 3 维范围组成：全局工作大小（指定工作项的完整范围）和局部工作大小（指定每个工作组的范围）。

nd_item 描述了 sycl::nd_range 中一个点的位置。nd_item 通常作为参数传递给 parallel_for 中的内核函数。除了包含工作组和全局空间中工作项的 ID 外，nd_item 还包含定义索引空间的 sycl::nd_range。

CUDA 线程块和 SYCL 工作组

在 CUDA 中，协作组提供设备代码 API，用于定义、划分和同步线程组。我们通常需要定义和同步小于线程块的线程组，以便实现更高的性能和设计灵活性。

thread_block 的实例是 CUDA 线程块中线程组的句柄，您按如下方式初始化它：

cg::thread_block cta = cg::this_thread_block();

执行该行的每个线程都有自己的 block 变量实例。具有相同 CUDA 内置变量 blockIdx 值的线程属于同一线程块组。

在 SYCL 中，给定内核的单次执行被组织成工作组和工作项。每个工作组包含相同数量的工作项，并通过工作组 ID 唯一标识。此外，在工作组内，工作项可以通过其局部 ID 标识，局部 ID 与工作组 ID 的组合等同于全局 ID。

auto cta = item_ct1.get_group();

SYCL 的 get_group 返回组 ID 的组成元素，表示工作组在给定维度中在整个 nd_range 中的位置。

共享局部内存访问

在 CUDA 中，共享内存是片上内存，比局部内存和全局内存快得多。共享内存延迟大约比未缓存的全局内存延迟低 100 倍。线程可以访问由同一线程块内其他线程从全局内存加载到共享内存中的数据。内存访问可以通过线程同步来控制，以避免竞争条件 (__syncthreads)。

__shared__ double x_shared[N_ROWS];

__shared__ double b_shared[ROWS_PER_CTA + 1];

在 SYCL 中，共享局部内存 (SLM) 在每个工作组中都是片上内存；SLM 具有比全局内存高得多的带宽和低得多的延迟。由于它可供工作组中的所有工作项访问，因此 SLM 可以根据工作组大小容纳数百个工作项之间的数据共享和通信。工作组中的工作项可以显式地将数据从全局内存加载到 SLM 中。数据在工作组的生命周期内保留在 SLM 中以加快访问速度。在工作组完成之前，SLM 中的数据可以由工作项显式地写回到全局内存。工作组完成执行后，SLM 中的数据将失效。

accessor<double, 1, access_mode::read_write, access::target::local> x_shared_acc_ct1(range<1>(N_ROWS), cgh);

accessor<double, 1, access_mode::read_write, access::target::local> b_shared_acc_ct1(range<1>(ROWS_PER_CTA + 1), cgh);

CUDA 线程块同步和 SYCL 屏障同步

同步用于同步共享相同资源的线程的状态。

在 CUDA 中，所有线程组都支持同步。我们可以通过调用其集体 sync() 方法或调用 cooperative_groups::sync() 函数来同步一个组。这些函数在组中的所有线程之间执行屏障同步。

cg::sync(cta);

在 SYCL 中，为了同步内存状态，我们使用 item::barrier(access::fence_space) 操作。它确保工作组中的每个工作项都到达屏障调用。换句话说，它保证工作组在代码中的某个点上同步。

item_ct1.barrier();

item::barrier 在特定空间发出内存屏障——它可以是 access::fence_space::local_space、::global_space 或 ::global_and_local。屏障确保指定空间的状态在工作组中的所有工作项之间保持一致。

CUDA 协作组和 SYCL 子组

CUDA 协作组和 SYCL 子组旨在扩展编程模型，允许内核动态组织线程组，以便线程协作并共享数据以执行集体计算。

在 CUDA 中，协作组为您提供了通过划分现有组来创建新组的灵活性。这使得在更精细的粒度上进行协作和同步成为可能。cg::tiled_partition() 函数将线程块划分为多个瓦片。

cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

执行分区的每个线程都将获得一个 32 线程组的句柄（在 tile32 中）。

在 SYCL 中，子组允许将工作组划分为映射到低级硬件并提供额外调度保证。子组是 SYCL 执行模型的扩展，它在 work_group 和 work_item 之间处于分层位置。SYCL 实现通常将子组映射到低级硬件功能：例如，子组中的工作项通常在支持向量指令的硬件上以 SIMD 方式执行。

sub_group tile_sg = item_ct1.get_sub_group();

设备支持的子组大小集是特定于设备的，单个内核可以在编译时请求特定的子组大小。此子组大小是一个编译时常量。为了获得最佳性能，我们需要尝试将最佳子组大小与硬件上的计算单元大小匹配。如果未设置，编译器将尝试选择子组的最佳大小。

[[intel::reqd_sub_group_size(SIZE)]]

CUDA Warp 原语和 SYCL 组算法

引入原语是为了使 warp 级编程安全有效。CUDA 以单指令、多线程 (SIMT) 方式执行线程组。我们可以通过利用 warp 执行来实现高性能。

在 CUDA 中，使用 thread_block_tile::shfl_down() 来简化我们的 warp 级归约，并消除了对共享内存的需求。

for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {

    rowThreadSum += tile32.shfl_down(rowThreadSum, offset);

    }

每次迭代都会使活动线程数减半，并且每个线程将其部分和添加到块的第一个线程。

在 SYCL 中，与 CUDA shfl_down 等效的是 shift_group_left，它通过将值固定数量的工作项向左移动，将组中工作项持有的值直接移动到组中的另一个工作项。

for (int offset = tile_sg.get_local_range().get(0) / 2; offset > 0; offset /= 2) {

    rowThreadSum += shift_group_left(tile_sg,rowThreadSum, offset);

    }

CUDA 原子操作和 SYCL 原子操作

原子操作是指在不受到任何其他线程干扰的情况下执行的操作。原子操作通常用于防止竞争条件，这是多线程应用程序中的常见问题。

在 CUDA 中，atomicAdd() 读取全局或共享内存中某个地址的一个字，向其添加一个数字，然后将结果写回同一地址。在操作完成之前，没有其他线程可以访问此地址。原子函数不充当内存屏障，不暗示内存操作的同步或排序约束。

if (tile32.thread_rank() == 0) {

    atomicAdd(&b_shared[i % (ROWS_PER_CTA + 1)], -rowThreadSum);

    }

在 SYCL 中，等效的是 atomic_ref，它支持 float 和 double 数据类型。支持的排序集是特定于设备的，但每个设备都保证至少支持 memory_order::relaxed。

if (tile32_sg.get_local_id()[0] == 0) {    

    atomic_ref<double, memory_order::relaxed, memory_scope::device, access::address_space::local_space> at_h_sum { b_shared[i % (ROWS_PER_CTA + 1)]};

      at_h_sum.fetch_add(-rowThreadSum);

    }

模板参数 space 允许是 access::address_space::generic_space、access::address_space::global_space 或 access::address_space::local_space。

这就完成了内核端 CUDA 代码 (jacobi.cu) 到 SYCL 的迁移

jacobi.cu 的 CUDA 内核代码可在 jacobi.cu 中找到。

jacobi.cpp 的 SYCL 内核代码可在 jacobi.cpp 中找到。

这就结束了所有 CUDA 到 SYCL 的迁移。我们现在有 main.cpp 和 jacobi.cpp 源文件，可以使用相应的 SYCL 编译器为任何 GPU 进行编译，而不是只有只能在 Nvidia GPU 上运行的 CUDA 源。

使用此 SYCL 源，我们可以使用 oneAPI DPC++ 编译器 编译以在 Intel® GPU 或 CPU 上运行雅可比迭代。或者我们可以使用 开源 LLVM 编译器 或 hipSYCL 编译器 编译以在 Nvidia GPU/AMD GPU 上运行。

SYCL 允许我们的源代码在不同供应商的 CPU 和 GPU 之间可移植，而不是被锁定在特定供应商的硬件上。

性能分析工具

Intel® VTune™ Profiler

Intel® VTune Profiler 是一款用于串行和多线程应用程序的性能分析工具。它有助于分析算法选择并识别应用程序如何以及在何处从可用硬件资源中受益。数据收集器使用操作系统计时器分析您的应用程序，中断进程，以 10 毫秒的采样间隔收集所有活动指令地址的样本，并为每个样本捕获调用序列（堆栈）。默认情况下，收集器不收集系统范围的性能数据，而只关注您的应用程序。有关更多详细信息，请查阅 Intel® VTune™ Profiler 入门。

要收集分析数据，可以在命令行中运行以下脚本

#!/bin/bash

source /opt/intel/oneapi/setvars.sh

#Vtune GPU Hotspot script

bin="jacobiSYCL"

prj_dir="vtune_data"

echo $bin

rm -r ${prj_dir}

echo "Vtune Collect hotspots"

vtune -collect gpu-hotspots -result-dir ${prj_dir} $(pwd)/${bin}

echo "Vtune Summary Report"

vtune -report summary -result-dir ${prj_dir} -format html -report-output $(pwd)/vtune_${bin}.html

确保上述脚本 "vtune_report.sh" 文件与应用程序二进制文件位于同一位置，如果您的二进制文件名称不同，请对脚本中的二进制文件名称进行任何必要的更改，运行脚本以收集 VTune Profiling 数据并生成 html 报告，HTML 报告将如下所示

图 1：VTuneTM Profiler 指标。

图 1 是 VTune Profiler 的快照，它表示雅可比迭代 SYCL 迁移代码的总运行时间。总运行时间为 13.764 秒，其中 GPU 时间为 5.041 秒。报告还描述了 GPU 的空闲时间；在整个执行期间，GPU 核心带宽仅利用了 50%，这有很大的改进空间。

Intel® Advisor Roofline

屋顶线图是应用程序性能与硬件限制（包括内存带宽和计算峰值）之间关系的视觉表示。屋顶线需要来自调查和行程计数以及 flops 分析类型的数据。您可以选择单独运行这些分析，也可以使用快捷命令逐个运行它们。有关更多详细信息，请查阅 Intel® Advisor 入门。

要收集分析数据，可以在命令行中运行以下脚本

#!/bin/bash

source /opt/intel/oneapi/setvars.sh

#Advisor Roofline script

bin="jacobiSYCL"

prj_dir="./roofline_data"

echo $bin

rm -r ${prj_dir}

advisor --collect=survey --project-dir=${prj_dir} --profile-gpu -- ./${bin} -q

advisor --collect=tripcounts --project-dir=${prj_dir} --flop --profile-gpu -- ./${bin} -q

advisor --report=roofline --gpu --project-dir=${prj_dir} --report-output=./roofline_gpu_${bin}.html -q

确保上述脚本“roofline_report.sh”文件与应用程序二进制文件位于同一位置，如果您的二进制文件名称不同，请对脚本中的二进制文件名称进行任何必要的更改，运行脚本以收集 Intel Advisor Roofline 数据并生成 html 报告，HTML 报告将如下所示

图 2：GPU 屋脊图。

GPU Roofline 图表显示了应用程序在内存和计算方面的性能。x 轴表示算术强度，y 轴表示计算性能。该报告显示，使用 DRAM 可达到的最大带宽为 34.58 GB/秒，GTI 带宽为 76 GB/秒，L3 带宽为 202.4 GB/秒，SLM 带宽为 202.43 GB/秒——矢量 FMA 和矢量加法的单精度和双精度可达到的最大计算性能。

在图表中，每个点代表应用程序中的一个循环或函数。点的位置表示循环或函数的性能，这受其优化和算术强度的影响。点的大小和颜色表示循环或函数占总应用程序时间的多少。大的红点占用时间最多，是最佳优化候选对象。小的绿点占用时间相对较少，因此可能不值得优化。

优化 SYCL 代码以提高性能

归约操作优化

shift_group_left 将组中工作项持有的值直接移动到组中的另一个工作项，通过将值向左移动固定数量的工作项。

for (int offset = tile_sg.get_local_range().get(0) / 2; offset > 0; offset /= 2) {

    rowThreadSum += shift_group_left(tile_sg, rowThreadSum, offset);

    }

以上代码片段展示了优化前的 Jacobi 迭代代码。这里，shift_group_left 已被 reduce_over_group 替换，以获得更好的性能。

reduce_over_group 通过组合组中工作项直接持有的值，在内部实现数组元素的广义求和。工作组归约的值数量等于组的大小，每个工作项提供一个值。

rowThreadSum = reduce_over_group(tile_sg, rowThreadSum, sycl::plus<double>());

以上代码片段展示了 Jacobi SYCL 优化代码。通过使用 reduce_over_group API，在 Intel GPU 上应用程序的运行时观察到大约 35% 的缩减。

原子操作优化

fetch_add 将操作数原子地添加到此 atomic_ref 引用的对象的值，并将结果分配给引用的对象的值。在这里，原子 fetch_add 用于将所有子组值求和到 temp_sum 变量中。

if (tile_sg.get_local_id()[0] == 0) {

     atomic_ref<double, memory_order::relaxed, memory_scope::device, access::address_space::global_space> at_sum{*sum};
      at_sum.fetch_add(temp_sum);

    }

以上代码片段展示了优化前的 Jacobi 迭代代码。优化后的实现从代码中删除了 fetch_add，并将 temp_sum 的值直接添加到全局变量 at_sum。

if (tile_sg.get_local_id()[0] == 0) {
     atomic_ref<double, memory_order::relaxed, memory_scope::device,
                  access::address_space::global_space>
     at_sum{*sum};
     at_sum += temp_sum;
    }

通过删除 fetch_add，在 Intel GPU 上应用程序的运行时观察到大约 5% 的缩减。

前面的代码描述了 Jacobi SYCL 优化代码。

雅可比迭代的 SYCL 迁移优化代码可在 sycl_migrated_optimized 中找到。

源代码链接

CUDA 源代码 Github 链接
SYCL 源代码—手动迁移 1-1 映射 Github 链接
SYCL 源代码—手动迁移并应用优化 Github 链接
SYCL 源代码—DPCT 输出，包含未迁移的代码 Github 链接
SYCL 源代码—DPCT 输出，包含已实现未迁移的代码 Github 链接

参考文献

© . All rights reserved.