5 分钟创建你的第一个 DevCloud oneAPI 应用
5.00/5 (3投票s)
使用 DPC++ 和 oneAPI 编写应用程序并在 Intel DevCloud 上部署的快速演示
引言
英特尔的 oneAPI 是一个开放的、基于标准的编程模型，使 C++ 开发人员能够以统一的方式针对不同的加速器架构。这与 Vulkan、CUDA 或 D3D12 等专门针对 GPU 的 API，或者 Verilog 和 VHDL 等针对 FPGA 的 API 不同。使用 oneAPI 编程时，相同的代码可用于 CPU、GPU 和 FPGA。您只需要一个正常工作的 oneAPI 工具包的安装，其中包含 DPC++，这是英特尔对 Clang 编译器的扩展，用于将 C++ 代码编译为目标各种受支持的加速器类型。
使用 oneAPI 的另一个好处是能够将应用程序部署到英特尔的 DevCloud，这是一个提供对配备强大 CPU 和 FPGA 的计算集群访问权限的沙箱环境。熟悉管理 FPGA 工具链复杂性的从业者可能会乐于在一个近乎即插即用的环境中部署软件，针对 Arria 10 等 FPGA，而无需立即投资硬件。在本文中，我们将从头开始使用 oneAPI 和 DPC++ 编写一个小程序，然后将其部署并在 DevCloud 的硬件上运行。
Sobel 卷积
Sobel 算子是一个简单的边缘检测滤波器，当与图像卷积时，会产生一个边缘被强化的相应图像。它主要用于计算机视觉或渲染应用程序，作为执行推理或渲染各种视觉效果的下游管道阶段的先导。要执行 Sobel 滤波器，我们需要用以下滤波器对图像进行两次卷积：
|1 0 -1|
|2 0 -2|
|1 0 -1|
和
| 1 2 1|
| 0 0 0|
|-1 -2 -1|
直观地说，第一个核会“检测” x 方向上的边缘，第二个核会检测 y 方向上的边缘。实际上，在单个像素上应用核是通过有限差分计算的定向梯度。这样产生的梯度可以合并以返回每个像素处定向导数的范数。范数越大，边缘就越明显，如下面的图像所示。
对于 3x3 核，应用于单个像素需要 9 次乘法和 8 次加法来累积结果，总共 17 次操作（不考虑 SIMD 或 MADD 或任何其他此类操作融合，这只是一个粗略的成本估算）。如果图像的尺寸为 w x h，那么我们预计两次卷积需要大约 34wh 次操作。但是，有一个技巧，它利用了卷积算子是可分离的这一事实，只要卷积矩阵的秩为 1。在 x 导数有限差分核的情况下，我们可以将其分解为以下外积：
|1 0 -1| |1|
|2 0 -2| = |2| * [1 0 -1]
|1 0 -1| |1|
随后，图像可以分两步进行卷积，首先是 3x1 水平核，然后是 1x3 垂直核。这两个核需要 10 次操作，这意味着整个 Sobel 卷积可以在大约 20wh 次操作中完成，而不是 34wh 次操作（但需要一些额外的中间内存）。
我们将编写的 DPC++ 代码将执行以下步骤：
-
加载图像
-
初始化设备队列
-
将图像转换为灰度
-
并行地
-
在两次顺序卷积中执行水平梯度计算
-
在两次顺序卷积中执行垂直梯度计算
-
-
将第二步的两个结果合并到输出图像中
-
将图像写入磁盘
加载图像
为了读写图像，我们将分别使用单文件库 stb_image 和 stbimagewrite。将这些文件放入您的源目录后，创建一个 main.cpp 文件，并在顶部包含以下头文件。
#include <CL/sycl.hpp>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"
#define STB_IMAGE_IMPLEMENTATION
#include "stb_image.h"
然后，我们可以像这样加载图像：
int channels;
int width;
int height;
uint8_t* image = stbi_load(path, &width, &height, &channels, 3);
cl::sycl::buffer<uint8_t, 1> image_buffer{image, width * height * channels};
stbi_load 函数执行后，内存已在主机上可用。但是，我们还构建了一个缓冲区类型来管理稍后对图像的读写屏障。
初始化设备队列
为了提交硬件加速命令，我们需要构建一个 cl::sycl::queue。队列用于将命令和内存编排到设备。
cl::sycl::queue queue{cl::sycl::default_selector};
除了默认选择器（它会尝试根据一组启发式方法选择“最佳”设备）之外，您还可以指定 gpu_selector、fpga_selector、host_selector 甚至您自己定义的自定义选择器。如果没有 cl::sycl::queue 抽象，我们将需要实现自定义逻辑来与每种加速器类型的不同设备驱动程序进行交互。编写 SYCL 代码的主要好处是能够使用单一统一的接口来定位所有这些设备。
将图像转换为灰度
为了评估边缘梯度，我们首先需要将 3 通道彩色图像转换为灰度。对于单个像素，我们使用以下函数执行此映射：
float luminance(uint8_t r, uint8_t g, uint8_t b)
{
float r_lin = static_cast<float>(r) / 255;
float g_lin = static_cast<float>(g) / 255;
float b_lin = static_cast<float>(b) / 255;
// Perceptual luminance (CIE 1931)
return 0.2126f * r_lin + 0.7152 * g_lin + 0.0722 * b_lin;
}
如果您不熟悉色彩理论，只需记住绿色比红色和蓝色在感知上更亮，而蓝色是三种颜色成分中感知上最不亮的。
然后，我们可以像这样并行计算每个像素的亮度：
// This constructor indicates that the memory should be allocated by the runtime
cl::sycl::buffer<float, 1> greyscale_buffer{width * height};
queue.submit([&greyscale_buffer, &image_buffer, image, width, height](
cl::sycl::handler& h) {
// A discard_write is a write access that doesn't need to preserve existing
// memory contents
auto data = greyscale_buffer.get_access<cl::sycl::access::mode::discard_write>(h);
auto image_data = image_buffer.get_access<cl::sycl::access::mode::read>(h);
h.parallel_for(cl::sycl::range<1>(width * height),
[image_data, data](cl::sycl::id<1> idx) {
int offset = 3 * idx[0];
data[idx[0]] = luminance(image_data[offset],
image_data[offset + 1],
image_data[offset + 2]);
});
});
SYCL 缓冲区的方法 get_access<M> 允许我们声明提交给队列的代码将以特定方式访问内存。然后，SYCL 运行时会依次对队列提交进行排序，并进行任何必要的内存同步。此外，请注意，尽管我们使用纯 C++ 编写了亮度函数，但编译器能够将其编译为适合我们目标设备的执行代码。
水平和垂直卷积
接下来，我们需要执行计算边缘梯度所需的卷积，我们将把它们存储在两个缓冲区中：
cl::sycl::buffer<float, 1> dx{width * height};
cl::sycl::buffer<float, 1> dy{width * height};
现在，进行水平卷积：
// Open a new scope so that dx_tmp is deallocated once no longer needed
{
cl::sycl::buffer<float, 1> dx_tmp{width * height};
// Extract a 3x1 window around (x, y) and compute the dot product
// between the window and the kernel [1, 0, -1]
queue.submit([&greyscale_buffer, &dx_tmp, width, height](cl::sycl::handler& h) {
auto data = greyscale_buffer.get_access<cl::sycl::access::mode::read>(h);
auto out = dx_tmp.get_access<cl::sycl::access::mode::discard_write>(h);
h.parallel_for(cl::sycl::range<2>(width, height),
[data, width, out](cl::sycl::id<2> idx) {
int offset = idx[1] * width + idx[0];
float left = idx[0] == 0 ? 0 : data[offset - 1];
float right = idx[0] == width - 1 ? 0 : data[offset + 1];
out[offset] = left - right;
});
});
// Extract a 1x3 window around (x, y) and compute the dot product
// between the window and the kernel [1, 2, 1]
queue.submit([&dx, &dx_tmp, width, height](cl::sycl::handler& h) {
auto data = dx_tmp.get_access<cl::sycl::access::mode::read>(h);
auto out = dx.get_access<cl::sycl::access::mode::discard_write>(h);
h.parallel_for(
cl::sycl::range<2>(width, height),
[data, width, height, out](cl::sycl::id<2> idx) {
// Convolve vertically
int offset = idx[1] * width + idx[0];
float up = idx[1] == 0 ? 0 : data[offset - width];
float down = idx[1] == height - 1 ? 0 : data[offset + width];
float center = data[offset];
out[offset] = up + 2 * center + down;
});
});
}
需要注意的主要重要一点是，操作的依赖关系图由我们在每个队列提交中包含的内存访问屏障隐式定义。例如，虽然我们在第一次 3x1 卷积和灰度转换之间没有显式同步，但 SYCL 保证了这两个步骤之间存在一个“发生在此之后”的关系，因为灰度内存产生的数据会被卷积读取。
然后，垂直卷积以相同的方式执行，但使用不同的核：
{
cl::sycl::buffer<float, 1> dy_tmp{width * height};
queue.submit([&greyscale_buffer, &dy_tmp, width, height](
cl::sycl::handler& h) {
auto data = greyscale_buffer.get_access<cl::sycl::access::mode::read>(h);
auto out = dy_tmp.get_access<cl::sycl::access::mode::discard_write>(h);
// Create a scratch buffer for the intermediate computation
h.parallel_for(cl::sycl::range<2>(width, height),
[data, width, out](cl::sycl::id<2> idx) {
// Convolve horizontally
int offset = idx[1] * width + idx[0];
float left = idx[0] == 0 ? 0 : data[offset - 1];
float right = idx[0] == width - 1 ? 0 : data[offset + 1];
float center = data[offset];
out[offset] = left + 2 * center + right;
});
});
queue.submit([&dy, &dy_tmp, width, height](cl::sycl::handler& h) {
auto data = dy_tmp.get_access<cl::sycl::access::mode::read>(h);
auto out = dy.get_access<cl::sycl::access::mode::discard_write>(h);
h.parallel_for(
cl::sycl::range<2>(width, height),
[data, width, height, out](cl::sycl::id<2> idx) {
// Convolve vertically
int offset = idx[1] * width + idx[0];
float up = idx[1] == 0 ? 0 : data[offset - width];
float down = idx[1] == height - 1 ? 0 : data[offset + width];
out[offset] = up - down;
});
});
}
请注意，垂直梯度和水平梯度之间没有依赖关系，因此 SYCL 可以并行执行它们。
合并梯度
对于每个像素，我们可以在 x 和 y 轴上投影梯度，因此计算梯度的大小是一个简单的问题：
// Allocate a memory region shared between the host and device queue
uint8_t* out = reinterpret_cast<uint8_t*>(
cl::sycl::malloc_shared(width * height, queue));
queue.submit([&dx, &dy, width, height, out](cl::sycl::handler& h) {
auto dx_data = dx.get_access<cl::sycl::access::mode::read>(h);
auto dy_data = dy.get_access<cl::sycl::access::mode::read>(h);
h.parallel_for(cl::sycl::range<1>(width * height),
[dx_data, dy_data, out](cl::sycl::id<1> idx) {
float dx_val = dx_data[idx[0]];
float dy_val = dy_data[idx[0]];
// NOTE: if deploying to an accelerated device, math
// functions MUST be used from the sycl namespace
out[idx[0]] = cl::sycl::sqrt(dx_val * dx_val + dy_val * dy_val) * 255;
});
});
在这里，我们将数据写回统一内存分配，在重新缩放到 8 位灰度格式之后。请注意，因为我们从 dx 和 dy 缓冲区进行读取，所以这项工作将与水平和垂直边缘卷积完成后进行排序。
输出结果
最后，我们准备好读取结果并将其写入磁盘。
queue.wait();
stbi_write_png("edges.png", weidth, height, 1, out, width);
// Reclaim now unused memory
stbi_image_free(image);
cl::sycl::free(out, queue);
这里需要 `wait`，因为与之前不同，我们没有 get_access 请求来创建隐式屏障，并且我们直接从主机内存读取。
部署到 DevCloud
现在我们有了一个可工作的程序，我们可以将源代码和任何构建脚本部署到 Intel oneAPI DevCloud。注册该程序后，您应该会收到一封电子邮件，其中包含有关如何获取 SSH 凭据以登录 DevCloud 平台的说明。按照说明操作后，您应该能够使用以下命令上传源文件和测试图像：
scp main.cpp devcloud:~
scp stb_image.h devcloud:~
scp stb_image_write.h devcloud:~
scp peppers.png devcloud:~
这会将您的源文件上传到您分配的 DevCloud 用户帐户的主目录。然后，您可以登录以部署您的程序：
ssh devcloud
创建一个脚本，如下所示编译和运行您的程序：
#!/usr/bin/env bash
# run.sh
# Ensure the correct environment variables are set
source /opt/intel/inteloneapi/setvars.sh
# Compile our program
dpcpp main.cpp -o sobel -std=c++17 -fsycl -lOpenCL
# Invoke the program with our test image
./sobel peppers.png
然后，我们可以使用 qsub（队列提交）命令在 DevCloud 上可用的各种计算节点上调用我们的脚本。要查看可用主机列表，请运行 pbsnodes 命令。这将按 ID 列出节点，并附带有关处理器类型和可用加速器等其他信息。
例如，要将作业提交到具有 GPU 的主机，我们可以运行以下命令：
qsub -l nodes=1:gpu:ppn=2 -d . run.sh
简而言之，这些选项表示我们想在具有 GPU 的单个节点上运行我们的脚本，我们想完全占用该节点（ppn=2 选项），我们希望工作目录是当前目录，我们希望节点调用 run.sh 脚本。
要查看作业状态，您可以运行 qstat，它将产生类似于以下内容的输出：
Job ID Name User Time Use S Queue
------------------------- ---------------- --------------- -------- - -----
681510.v-qsvr-1 run.sh u47956 0 R batch
Job ID 可以作为参数传递给 qdel 命令，以取消挂起的作业。
作业完成后，您将在当前目录中找到 run.sh.o681510 和 run.sh.e068150 等文件，分别对应脚本的标准输出和标准错误输出。如果我们的程序运行成功，您还应该会有一个 edges.png 图像，您可以检查其正确性。使用 exit 命令注销，并使用 scp 将图像传输回您的主机：
scp devcloud:~/edges.png .
结论
在本文中，我们使用 SYCL 运行时和英特尔 DPC++ 编译器提供的附加扩展开发了一个 C++ 应用程序。该应用程序演示了统一的编程模型如何能够针对不同的体系结构，以及 SYCL 运行时为协调内存访问和使用隐式依赖关系图编写并行代码而提供的抽象。
最后，我们展示了如何将我们的代码部署并测试到英特尔 DevCloud 提供的不同硬件配置。要了解更多关于 SYCL 运行时、英特尔 DevCloud 或英特尔 DPC++ 编译器，鼓励您在英特尔 DevZone 此处阅读更多内容。
历史
- 2020 年 11 月 9 日：初始版本