在 5 分钟内将 CUDA 应用移植到 oneAPI 和 DPC++

引言

CUDA是Nvidia拥有的并行计算平台和编程模型，用于在GPU上运行软件。它被广泛应用于研究人员和行业从业者，以加速计算密集型工作负载，而无需采用与传统软件开发相比完全陌生的工作流程和编程模型。采用CUDA的额外好处是能够立即访问广泛的现有库，以及使用多种工具来调试和可视化CUDA代码。

在本文中，我们将解释如何将CUDA代码移植到Intel的oneAPI工具包，特别是将CUDA内核移植到Intel的DPC++编译器。 "oneAPI"工具包指的是数据并行C++（简称DPC++）编程模型，以及一系列旨在支持高性能计算应用程序的API。 DPC++是基于LLVM的Clang编译器构建的，通过SYCL（一个旨在允许C++应用程序针对异构系统的开放标准）扩展了现代C++的功能。

为什么要移植CUDA到oneAPI？

鉴于CUDA在图像分析、机器学习等领域的广泛使用，你可能会想知道我们为什么要做这样的移植。简而言之，Intel的平台有一些引人注目的优势值得考虑。

首先，DPC++可以轻松地针对FPGA加速器，就像针对GPU一样。

其次，DPC++建立在Clang和Khronos发布的开源标准之上。Intel非常热衷于将DPC++的工作贡献给LLVM项目，这将对各种并行STL算法的价值产生直接影响。

第三，值得将代码移植到DPC++，至少是为了理解通用编程模型的工作方式，这可能会转化为对未来需要加速的代码如何更好地进行架构设计的新见解。

也许最大的潜在好处是能够将oneAPI软件部署到Intel DevCloud，这是一个提供CPU、GPU和FPGA供你使用的云环境。特别是，许多可用的硬件都是尖端的，可能不适合在家中或办公室进行实验。例如，只需几条命令，你就可以轻松地将你的应用程序与Arria 10 FPGA和Xeon Platinum进行基准测试。人们也有主观原因倾向于编写DPC++代码，即DPC++程序读起来就像语法正确的C++，而无需来自CUDA的你可能熟悉的外国语法或属性。

CUDA应用程序

首先，我们需要选择一个CUDA应用程序进行演示移植。在这里，我们将移植经典的曼德勃罗集分形生成器，因为我们更感兴趣的是学习DPC++编程模型本身。简而言之，让我们快速浏览一下CUDA代码。首先，我们需要例程来乘法两个复数，加法两个复数，以及计算复数的平方模。

struct complex
{
    float r;
    float i;
};

// __device__ := Invoke this function from device and execute it on device
__device__ complex operator*(complex a, complex b)
{
    return {a.r * b.r - a.i * b.i, a.r * b.i + a.i * b.r};
}

__device__ complex operator+(complex a, complex b)
{
   return {a.r + b.r, a.i + b.i};
}

__device__ float sqr_magnitude(complex c)
{
    return c.r * c.r + c.i * c.i;
}

在CUDA中，我们打算在加速器设备上调用的函数需要__device__属性。接下来，我们将编写计算与每个像素关联的曼德勃罗集"值"的函数。

constexpr static uint32_t max_iterations = 12000u;

__device__ uint32_t mandelbrot_pixel(complex c)
{
    // Evaluate iteratively z_{n + 1} = z_n^2 + c
    // Terminate if the max iterations are reached or if the norm exceeds 2
    complex z = {};

    uint32_t i = 0;
    for (; i != max_iterations; ++i) {
        complex z_next = z * z + c;
        if (sqr_magnitude(z_next) > 4.0) {
            return i;
        } else {
            z = z_next;
        }
    }

    return i;
}

简而言之，这个函数接受一个常数c，将变量z初始化为0，然后不断评估z_next = z^2 + c; z = z_next，直到z_next的模数超过2。函数返回发生此事件所需的迭代次数。接下来，我们需要内核函数，它将计算并写出对应于每次调用的像素的颜色。

__global__ void mandelbrot(uint8_t* output, int width, int height)
{
    // Remap workgroup and thread ID to an x-y coordinate on a 2D raster
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= width || y >= height)
    {
        return;
    }

    // Remap pixel values to a range from [-2, 1] for the real part and
    // [-1.5, 1.5] for the imaginary part
    complex c = {static_cast<float>(x) / width * 3.f - 2.f,
    static_cast<float>(y) / height * 3.f - 1.5f};

    // Evaluate the mandelbrot iterations for a single thread and write out the
    // result after first normalizing to the range [0, 256)
    uint32_t iterations = mandelbrot_pixel(c);

    // Tonemap color
    uint32_t color = iterations * 6;

    // For stylistic reasons, draw saturated values as black
    output[y * width + x] = color >= 256 ? 0 : color;
}

曼德勃罗集函数使用__global__属性来指示它打算由主机调用。此内核中采用的模式相当常见，即使用块大小、块索引和线程索引来将内核的特定调用与输出光栅中的像素相关联。此像素的坐标用于计算颜色，然后将其写入输出缓冲区。由于分派的每个调用都针对输出光栅中的唯一像素，因此每个调用都可以独立于所有其他调用进行操作，而无需锁、原子操作或任何其他同步原语。

最后，我们需要一个main函数来分配用于输出的设备内存，分派我们的内核，分配主机内存，读回输出，最后将输出写入图像。为了生成图像，我们将使用经典的单头文件/源文件stb_image_write.h，来自经典的stb库集合，以保持简单。

int main(int argc, char* argv[])
{
    constexpr static int width          = 512;
    constexpr static int height         = 512;
    constexpr static size_t buffer_size = width * height;

    // Allocate a 512x512 256-bit greyscale image on device
    uint8_t* buffer;
    cudaMalloc(&buffer, buffer_size);

    // Operate with 8x8 workgroup sizes (1 AMD wavefront, 2 NVIDIA warps)
    dim3 workgroup_dim{8, 8};
    dim3 workgroup_count{width / workgroup_dim.x, height / workgroup_dim.y};

    mandelbrot<<<workgroup_count, workgroup_dim>>>(buffer, width, height);

    // Flush all work queued to device
    cudaDeviceSynchronize();

    // Write back device memory to host memory and deallocate device memory
    uint8_t* host_buffer = reinterpret_cast<uint8_t*>(std::malloc(buffer_size));
    cudaMemcpy(host_buffer, buffer, width * height, cudaMemcpyDeviceToHost);
    cudaFree(buffer);

    // Write out results to an image
    int result = stbi_write_png("mandelbrot.png", width, height, 1, host_buffer, width);

    std::free(host_buffer);

    return 0;
}

最后，如果您正在跟随操作，请确保在文件顶部包含以下所需头文件

#include <cmath>
#include <cstdint>
#include <cstdlib>
#include <iostream>

#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"

编译并评估上述代码后，您应该能够生成以下PNG图像

移植到DPC++

为了从CUDA移植到DPC++，我们可以费力地将CUDA代码"翻译"成DPC++。幸运的是，Intel提供了DPC++兼容性工具来简化移植过程。虽然截至本文撰写之时，该工具仍处于"beta"状态，但我移植CUDA曼德勃罗集代码没有任何问题。

首先，您需要确保您的机器上同时安装了DPC++和兼容性工具。最简单的方法是安装oneAPI工具包。编译器和兼容性工具都包含在基础工具包中。如果您需要针对CUDA后端，您可能需要自己构建支持CUDA的工具链，因为支持CUDA的工具链仍处于实验阶段。要了解如何做到这一点，请参考这里的文档。此外，如果您目前使用的是CUDA 11，则必须同时安装CUDA 10或9，兼容性工具才能运行。

接下来，打开您选择的操作系统的shell后，您需要调用一个shell脚本来本地修改各种环境变量，以确保Intel oneAPI库和可执行文件可以被定位。在类Unix平台上，该脚本名为setvars.sh，位于工具包的安装根目录（通常是/opt/intel/oneapi或~/intel/oneapi）。在Windows上，该脚本提供为setvars.bat，并且同样位于安装根目录。

验证PATH设置正确后，dpct兼容性工具应该可用。对于我们只有一个main.cu文件的简单示例，以下命令足以执行转换并将输出发送到同一目录中的dpct_output文件夹。

dpct --extra-arg="-std=c++17" main.cu --out-root dpct_output

dpct_output的目录内容通常是cpp源文件，扩展名为.dp.cpp。此外，您可能会看到各种yaml文件，其中列出了对项目中各种文件的代码替换。虽然它们不参与编译，但它们有助于理解所做的操作并解决出现的任何问题。

要编译代码并进行测试，请调用以下命令

mkdir build
cd build
dpcpp ../main.dp.cpp -o mandelbrot -lsycl -lOpenCL

在Windows上，您需要输出一个扩展名为.exe的可执行文件。在同一个终端中，执行mandelbrot程序应该会生成一个与我们之前使用CUDA生成的图像完全相同的图像。

您可能会发现的一个小麻烦是，在新终端中或从文件浏览器调用上面生成的可执行文件可能会导致运行时错误，抱怨找不到各种共享库。这是因为默认情况下，dpcpp对sycl库使用动态链接，这有助于程序在oneAPI安装将来升级时接收被动更新。要解决此问题，您可以选择将库与可执行文件一起打包到同一目录中，或者修改库加载路径。

部署到Intel DevCloud平台

为了结束我们的移植，让我们将应用程序部署到Intel的DevCloud。这将使我们能够试验Intel提供的硬件。首先，请通过以下DevCloud注册页面创建一个帐户。之后，按照随后通过电子邮件发送给您的唯一登录链接，并SSH到预配的DevCloud实例。登录后的重定向页面应包含有关如何在您的操作系统上执行此连接的说明。大部分情况下，这相当于在您的SSH配置中有一个Host条目，将devcloud重新映射到您的凭证的代理连接。

之后，我们可以使用scp将我们的源文件传输到DevCloud实例

scp -r dpct_output devcloud:~/mandelbrot

此外，您还需要一个Makefile和一个脚本来运行您的应用程序。以下Makefile可用于编译我们的示例

CXX = dpcpp
CXXFLAGS = -o
LDFLAGS = -lOpenCL -lsycl
EXE_NAME = mandelbrot
SOURCES = main.dp.cpp
BINDIR = bin

all: main

main:
[ -d $(BINDIR) ] || mkdir $(BINDIR)
$(CXX) $(CXXFLAGS) $(BINDIR)/$(EXE_NAME) $(SOURCES) $(LDFLAGS)

run:
$(BINDIR)/$(EXE_NAME)

clean:
rm -rf $(BINDIR)/$(EXE_NAME)

还需要一个调用make并执行编译程序的脚本（这里我们将其命名为run.sh，但您可以选择自己的名称并相应地调整以下命令）

#!/bin/bash
source /opt/intel/inteloneapi/setvars.sh
make run

有了这些，我们现在就可以将作业提交到DevCloud中的各种硬件队列了。有关与作业队列交互的完整文档 在此 提供。作为演示，以下命令将分派我们的请求，运行它，并读回结果。

# ON DEVCLOUD
# Queue submission with a job label, working directory, and script to run
qsub -N mandelbrot -d . run.sh

# Show job submission status
qstat

# ON HOST
scp devcloud:~/mandelbrot/mandelbrot.png .
# Verify the image looks correct

结论

在本文中，我们演示了如何将现有的CUDA应用程序移植到DPC++，对其进行编译，并在DevCloud上运行。假设您熟悉这些命令和原始CUDA程序的简单性，这样的移植和部署到边缘硬件可以在几分钟内完成。更复杂的项目可能需要此处未涵盖的额外步骤：例如，在Visual Studio项目或由CMake等工具生成的编译器命令数据库上调用兼容性工具。此外，DevCloud的各种功能也未涵盖，例如针对特定类别的硬件或计算节点的能力，以及执行计时执行的脚本。要利用这些功能并了解更多关于Intel的DPC++编译器，请参考Intel Developer Zone上的文档。

历史

• 2020 年 11 月 9 日：初始版本