基于openAPI的C++/SYCL实现并行排序算法）
一、 任务需求
使用基于oneAPI的C++/SYCL实现⼀个高效的并行归并排序。需要考虑数据的分割和合并以及线程之间的协作。
- 项目简介: 该项目旨在使用 SYCL（基于 C++ 的高性能并行编程模型）来加速归并排序算法过程。通过利用 GPU 和其他硬件加速器的并行计算能力，大幅提高归并排序的效率。
- 在实际实现中，归并排序可使用共享内存来加速排序过程。
- 具体来说，可以利用共享内存来存储临时数据，减少对全局内存的访问次数，从而提高排序的效率。另外，在合并操作中，需要考虑同步机制来保证多个线程之间的数据⼀致性。
- 技术栈及实现方案:
- SYCL：用于并行编程和加速计算
- Intel® oneAPI：提供工具和库来优化性能和简化开发过程
- Intel® AI Analytics Toolkit：包含加速AI和机器学习工作流的工具
- C++：项目的主要编程语言
二、问题分析
归并排序是⼀种分治算法，其基本原理是将待排序的数组分成两部分，分别对这两部分进行排序，然后将已排序的子数组合并为⼀个有序数组。可考虑利用了异构并行计算的特点，将排序和合并操作分配给多个线程同时执行，以提高排序效率。具体实现过程如下：
- 将待排序的数组分割成多个较小的子数组，并将这些⼦数组分配给不同的线程块进行处理。
- 每个线程块内部的线程协作完成子数组的局部排序。
- 通过多次迭代，不断合并相邻的有序⼦数组，直到整个数组有序。
在实际实现中，归并排序可使用共享内存来加速排序过程。具体来说，可以利用共享内存来存储临时数据，减少对全局内存的访问次数，从而提高排序的效率。另外，在合并操作中，需要考虑同步机制来保证多个线程之
间的数据⼀致性。
需要注意的是，在实际应用中，要考虑到数组大小、线程块大小、数据访问模式等因素，来设计合适的算法和参数设置，以充分利用目标计算硬件GPU的并行计算能力，提高排序的效率和性能。
三、oneAPI 介绍
随着科学技术的迅猛发展，高性能计算在人工智能、药物研制、智慧医疗、计算化学等领域发挥着日益重要的作用。然而，随着后摩尔时代的来临，计算机系统结构进入了百花齐放、百家争鸣的繁荣时期，CPU、GPU、FPGA和AI芯片等相互补充。硬件的多样性带来了软件设计与开发的复杂性，高性能计算并行程序的计算效率以及在不同计算平台之间的可移植性日益变得重要。为了解决这一问题，Intel推出了oneAPI。
3.1 oneAPI诞生背景
随着科技的迅猛发展，高性能计算在人工智能、药物研制、智慧医疗、计算化学等领域的作用日益显著。然而，随着后摩尔时代的到来，计算机系统结构迎来了繁荣时期，CPU、GPU、FPGA和AI芯片等各异的硬件相互补充。这种硬件多样性导致了软件设计与开发的复杂性的增加，高性能计算并行程序的计算效率和在不同计算平台之间的可移植性变得日益关键。为了解决这一问题，Intel推出了oneAPI。
3.2 oneAPI 简介
- Intel oneAPI 是一个跨行业、开放、基于标准的统一的编程模型，旨在提供一个适用于各类计算架构的统一编程模型和应用程序接口。其核心思想是使开发者只需编写一次代码，便可在跨平台的异构系统上运行，支持的底层硬件架构包括 CPU、GPU、FPGA、神经网络处理器以及其他专为不同应用设计的硬件加速器等。这意味着，oneAPI不仅提高了开发效率，同时具备一定的性能可移植性。通过采用这一编程模型，开发者能够更灵活地利用不同类型的硬件，充分发挥各种计算资源的潜力，从而更好地适应不同应用场景的需求。
- 使用Intel oneAPI编程具有以下好处：
- 跨硬件平台：允许开发者针对不同类型的硬件（包括非Intel硬件）编写统一的代码，增加代码的可移植性和灵活性。
- 性能优化：提供针对Intel硬件优化的库和工具，可以帮助开发者提高应用程序的性能。
- 简化开发流程：通过统一的API和工具链，简化了对多种硬件的编程和优化过程，降低了学习曲线。
- 支持现代编程语言：支持包括C++、Python在内的多种现代编程语言，适用于广泛的开发需求。
- 促进创新：通过提供高级的编程模型和工具，使开发者能够更容易地探索和实现新的算法和应用。
简而言之，oneAPI 是一种通用的可以适配Intel的多种异构硬件的并行编程库，个人理解可以用于HPC场景以及深度学习的卷积优化、向量传播等场景，并降低适配不同硬件的编译适配成本。
OneAPI提供了云平台DevCloud可以免费注册后使用，即使笔记本不是Intel的也可以使用Intel OneAPI配套硬件的体验，还无需自行安装toolkit。
jupyter.oneapi.devcloud.intel.com/ 这是我在本次实验中使用的云平台链接，本次实验主要基于Jupyter Lab进行，使用简单，用Jupyter Notebook构建了可视化界面，即使不会使用命令行进行编译也可以直接复制官方教程中的脚本进行编译，免去了CPP的编译痛苦。
3.3 oneAPI整体架构
3.3.1 oneAPI 开放式规范
oneAPI 这一开放式规范包括一种跨架构的编程语言 Data Parallel C++（DPC++）、一套用于API编程的函数库以及底层硬件接口（oneAPI Level Zero），如下图1所示。有了这些组件，英特尔和其它企业就能创建他们自己的 oneAPI 实现来支持他们自己的产品，或基于 oneAPI 进行新产品开发。
3.3.2 SYCL规范
SYCL第一次是在2014年引入，它是一种基于C++异构平行编程框架，用来加速高性能计算，机器学习，内嵌计算，以及在相当宽泛的处理器构架之上的计算量超大的桌面应用。这些处理器包括了CPU, GPU, FPGA, 和张量加速器。
2021年2月9号 , 科纳斯组织（Khronos® Group），作为一个由工业界主流公司组成的创建先进的互联标准的开放协会，宣布了SYCL 2020最终版规范的批准和发布。这个规范是单源C++并行编程的开放标准。作为多年来规范开发的一个主要的里程碑，SYCL 2020是在SYCL 1.2.1的功能的基础之上建立的，用以进一步改善可编程性，更小的代码尺寸，和高效的性能。基于C++17之上的SYCL 2020, 使得标准C++应用的加速更为容易， 而且推动使之与ISO C++的路线图变得更为一致。SYCL 2020 将会进一步加速在多平台上的采用和部署，包括使用除了OpenCLTM之外的多样的加速API 后端。
SYCL 2020集成了超过40项新的特征，包括了为简化代码所做的更新，和更小的代码尺寸。一些主要增加的内容包括：
- 统一的共享存储（USM）使得带有指向器的代码，可以在不需要缓冲与存取器的情况下自然地工作
- 工作组和子工作组算法，在工作项目中增加了有效率的并行操作
- 类模板参数推导(CTAD)与模版减量指南，简化了类模板实例化
- 扩展的互操作性通过各种后端加速API实现高效加速
- 统一的共享存储（USM）使得带有指向器的代码，可以在不需要缓冲与存取器的情况下自然地工作;
- 并行的减量增加了一种内置的减量操作，来减少样板代码以达到具有内置的减量操作加速硬件上的最大性能。
3.3.3 DPC++介绍
oneAPI包含一种全新的跨架构编程语言 DPC++，DPC++基于 C++编写，由一组C++类、模板与库组成，同时兼容 Kronos 的 SYCL 规范，图2给出了DPC++与SYCL、C++关系。同时，intel DPC++兼容性工具可以实现将CUDA代码迁移到DPC++上，其中大约会有80%-90%的代码实现了自动迁移并提供内联注释，很大程度上帮助开发人员减轻代码移植的负担。
DPC++是一种单一源代码语言，其中主机代码和异构加速器内核可以混合在同一源文件中。在主机上调用 DPC++程序，并将计算加载到加速器。程序员使用熟悉的C++和库结构，并添加诸如工作目标队列、数据管理缓冲区和并行性并行的函数，以指导计算和数据的哪些部分应该被加载。
3.3.4 oneAPI工具包
oneAPI 编程模式兼容性堪称达到了历史最强。目前在各个领域应用比较广泛的高性能计算开发工具如 Fortran，在 AI 领域的 Python，以及像 OpenMP 这样不同领域使用的语言都可以做到无缝对接，同时，oneAPI 也支持一些主流的 AI 工具包，包括 Hadoop、Spark、TensorFlow、PyTorch、PaddlePaddle、OpenVINO 等等，形成更适合人工智能时代的软件栈。oneAPI有六个工具包，几乎涵盖了高性能计算、物联网、渲染、人工智能、大数据分析这些领域。
- Intel® oneAPI Base Toolkit：这个工具包是 oneAPI 其他产品的基础，包含了几个在 Parallel Studio中常用的软件以及 icc 编译器、MPI、DPCPP 等。这个工具包使开发人员都可以跨CPU、GPU和FPGA构建、测试和部署以性能为中心、以数据为中心的应用程序。
- Intel® oneAPI HPC Toolkit ：这个工具包提供可扩展的快速C ++、Fortran、OpenMP和MPI应用程序。从某种程度上来说 Intel® oneAPI Base Toolkit 加 Intel® oneAPI HPC Toolkit 基本就包含Intel Parallel Studio XE的功能了。
- Intel® oneAPI IoT Toolkit ：这个工具包主要用于建立可在网络边缘运行的高性能、高效、可靠的解决方案，属于物联网领域。
- Intel® AI Analytics Toolkit ：这个工具包提供优化的深度学习框架和高性能Python库，加速端到端机器学习和数据科学库。这些组件是使用 oneAPI 库构建的，用于低级计算优化。这可以最大化从预处理到机器学习的性能。
- Intel® oneAPI Rendering Toolkit：它主要用于创建高性能、高保真的可视化应用程序，适用于各种渲染领域。
- Intel® Distribution of OpenVINO™ Toolkit：这个工具包用于从设备到云部署高性能推理应用程序。该工具包基于卷积神经网络（CNN），可将工作负载扩展到整个英特尔®硬件（包括加速器），并最大限度地提高性能。该工具包可以使深度学习推理从边缘到云，加速人工智能工作负载,包括计算机视觉、音频、演讲,语言,和推荐系统。支持异构执行在英特尔架构和AI加速器CPU、iGPU,英特尔Movidius视觉处理单元(VPU)、FPGA,和英特尔高斯 & 神经加速器(Intel® GNA)。
四、oneAPI并行编程介绍
4.1 编程框架
**平台模型：**oneAPI的平台模型基于SYCL*平台模型。它指定控制一个或多个设备的主机。主机是计算机，通常是基于CPU的系统，执行程序的主要部分，特别是应用范围和命令组范围。主机协调并控制在设备上执行的计算工作。设备是加速器，是包含计算资源的专门组件，可以快速执行操作的子集，通常比系统中的CPU效率更高。每个设备包含一个或多个计算单元，可以并行执行多个操作。每个计算单元包含一个或多个处理元素，充当单独的计算引擎。图3所示为平台模型的可视化描述。一个主机与一个或多个设备通信。每个设备可以包含一个或多个计算单元。每个计算单元可以包含一个或多个处理元素。
**执行模型：**执行模型基于SYCL*执行模型。它定义并指定代码（称为内核kernel）如何在设备上执行并与控制主机交互。主机执行模型通过命令组协调主机和设备之间的执行和数据管理。命令组（由内核调用、访问器accessor等命令组成）被提交到执行队列。访问器(accessor)形式上是内存模型的一部分，它还传达执行的顺序要求。使用执行模型的程序声明并实例化队列。可以使用程序可控制的有序或无序策略执行队列。有序执行是一项英特尔扩展。设备执行模型指定如何在加速器上完成计算。从小型一维数据到大型多维数据集的计算通过ND-range、工作组、子组（英特尔扩展）和工作项的层次结构中进行分配，这些都在工作提交到命令队列时指定。需注意的是，实际内核代码表示为一个工作项执行的工作。内核外的代码控制执行的并行度多大；工作的数量和分配由 ND-range和工作组的规格控制。下图4描述了ND-range、工作组、子组和工作项之间的关系。总工作量由ND-range的大小指定。工作的分组由工作组大小指定。本例显示了ND-range的大小X*Y*Z，工作组的大小 X’* Y’*Z’，以及子组的大小X’。因此，有X*Y*Z工作项。有(X*Y*Z)/ (X’*Y’*Z’)工作组和(X*Y*Z)/X’子组。
**内存模型：*oneAPI 的内存模型基于 SYCL 内存模型。它定义主机和设备如何与内存交互。它协调主机和设备之间内存的分配和管理。内存模型是一种抽象化，旨在泛化和适应不同主机和设备配置。在此模型中，内存驻留在主机或设备上，并由其所有，通过声明内存对象来指定如图5所示。内存对象有两种：缓冲器和图像。这些内存对象通过访问器在主机和设备之间进行交互，访问器传达期望的访问位置（如主机或设备）和特定的访问模式（如读或写)。
在oneAPI内存模型中的Buffer Model：缓冲区(Buffer)将数据封装在跨设备和主机的 SYCL 应用中。访问器(Accessor)是访问缓冲区数据的机制。设备(device)和主机(host)可以共享物理内存或具有不同的内存。当内存不同时，卸载计算需要在主机和设备之间复制数据。DPC++不需要您管理数据复制。通过创建缓冲区(buffer)和访问器(accessor)，DPC++能够确保数据可供主机和设备使用，而无需您介入。DPC++ 还允许您明确地显式控制数据移动，以实现最佳性能。需要注意的是，在这种内存模式下，若有多个内核程序使用相同的缓冲区，访问器需要根据依赖关系，以对内核执行进行排序以避免争用缓冲区而出现错误（通过主机访问器或缓冲区破坏实现。
**编程模型：*面向 oneAPI 的内核编程模式基于 SYCL 内核编程模型。它支持主机和设备之间的显式并行性。并行性是显式的，因为程序员决定在主机和设备上执行什么代码；它不是自动的。内核代码在加速器上执行。采用 oneAPI 编程模型的程序支持单源，这意味着主机代码和设备代码可以在同一个源文件中。但主机代码中所接受的源代码与设备代码在语言一致性和语言特性方面存在差异。SYCL 规范详细定义了主机代码和设备代码所需的语言特性。
4.2 编程流程
DPC++程序设计大致可分为以下5个步骤：
(1) 申请Host内存
(2) 创建SYCL缓冲区并为其定义访问缓冲区内存的方法。
设备(device)和主机(host)可以共享物理内存或具有不同的内存。当内存不同时，卸载计算需要在主机和设备之间复制数据。而通过创建缓冲区(buffer)和访问器(accessor)的方式，DPC++就不需要您管理数据复制，其能够确保数据可供主机和设备使用，而无需您介入。DPC++ 还允许您明确地显式控制数据移动，以实现最佳性能。
//创建vector1向量的SYCL缓冲区；
buffer vector1_buffer(vector1,R);
定义了访问缓冲区内存的accessor；
accessor v1_accessor (vector1_buffer,h,read_only);
（3）创建队列以向Device（包括Host）提交工作（包括选择设备和排队）
q.submit([&](handler& h) {
//COMMAND GROUP CODE
});
可以通过选择器(selector)选择 CPU、GPU、FPGA和其他设备。使用默认的 q，这意味着 DPC++ 运行时会使用默认选择器（default selector）选择功能最强大的设备。
（4）调用oneAPI的核函数在Device上完成指定的运算。
该内核将应用于索引空间中的每个点，内核封装在C++ lambda函数中。DPC++中内核的形式如下：
h.parallel_for(range<1>(1024), [=](id<1> i){
A[i] = B[i] + C[i];
});
在该循环中，每个迭代都是完全独立的，并且不分顺序。使用 parallel_for 函数表示并行内核。
（5）将SYCL缓冲区的数据读到Host端。
五、并⾏排序算法实现
并行化归并排序算法的核心思想是将数据集划分为较小的子集，然后在这些子集上并行地执行归并排序操作。为了实现这一目标，我们可以使用 oneAPI 的任务并行模型。
首先，我们需要定义一个并行归并排序类，该类将负责实现并行归并算法。以下是该类的基本结构：
template <typename T>
class ParallelMergeSort {
public:
ParallelMergeSort() {}
void operator()(sycl::handler &cgh) {
sycl::nd_range<1> ndRange = cgh.get_nd_range();
cgh.parallel_for<class ParallelMergeSortKernel>(
ndRange, [=](sycl::nd_item<1> item) {
int globalId = item.get_global_id(0);
mergeSort(data_, 0, data_.size() - 1);
});
}
void setData(std::vector<T> &data) {
data_ = data;
}
std::vector<T> getData() {
return data_;
}
private:
std::vector<T> data_;
// 在此处添加代码：归并排序算法的实现
};
在上述代码中，我们定义了一个模板类 ParallelMergeSort，它包含一个调用运算符 operator()，其中我们将实现并行排序算法的逻辑。
接下来，我们需要在 operator() 函数中实现并行排序的逻辑。在 operator() 函数中，我们首先获取全局范围 ndRange，然后使用 parallel_for 函数来实现并行计算。以下是在 operator() 函数中添加的代码段：
cgh.parallel_for<class ParallelMergeSortKernel>(
ndRange, [=](sycl::nd_item<1> item) {
// 获取当前工作项的全局索引
int globalId = item.get_global_id(0);
// 在此处添加代码：调用归并排序算法进行排序
// e.g.mergeSort(data_, 0, data_.size() - 1);
});
在上述代码中，我们使用 parallel_for 函数来进行并行计算。在 lambda 表达式中，我们获取当前工作项的全局索引 get_global_id，然后在此处添加代码来调用归并排序算法进行排序。
private:
std::vector<T> data_;
void merge(std::vector<T>& arr, int l, int m, int r)
{
int i, j, k;
int n1 = m - l + 1;
int n2 = r - m;
std::vector<T> L(n1), R(n2);
for (i = 0; i < n1; i++)
L[i] = arr[l + i];
for (j = 0; j < n2; j++)
R[j] = arr[m + 1+ j];
i = 0;
j = 0;
k = l;
while (i < n1 && j < n2)
{
if (L[i] <= R[j])
{
arr[k] = L[i];
i++;
}
else
{
arr[k] = R[j];
j++;
}
k++;
}
while (i < n1)
{
arr[k] = L[i];
i++;
k++;
}
while (j < n2)
{
arr[k] = R[j];
j++;
k++;
}
}
void mergeSort(std::vector<T>& arr, int l, int r)
{
if (l < r)
{
int m = l+(r-l)/2;
mergeSort(arr, l, m);
mergeSort(arr, m+1, r);
merge(arr, l, m, r);
}
}
};
在上述代码中，我们定义了两个私有函数：mergeSort 和 merge函数实现了归并排序算法的递归调用逻辑。
void operator()(sycl::handler &cgh) {
// 获取全局范围
sycl::nd_range<1> ndRange = cgh.get_nd_range();
// 并行排序
cgh.parallel_for<class ParallelMergeSortKernel>(
ndRange, [=](sycl::nd_item<1> item) {
// 获取当前工作项的全局索引
int globalId = item.get_global_id(0);
mergeSort(data_, 0, data_.size() - 1);
});
}
在上述代码中，我们在 parallel_for 函数中调用 mergeSort 函数来实现归并排序算法的并行计算。每个工作项将负责对数据的一个子集进行排序操作。
最后，我们需要使用 ParallelMergeSort 类来实现完整的排序过程。以下是使用 oneAPI 编程模型和工具链实现并行归并排序算法的示例代码：
#include <CL/sycl.hpp>
#include <iostream>
#include <vector>
namespace sycl = cl::sycl;
template <typename T>
class ParallelMergeSort {
public:
ParallelMergeSort() {}
void operator()(sycl::handler &cgh) {
sycl::nd_range<1> ndRange = cgh.get_nd_range();
cgh.parallel_for<class ParallelMergeSortKernel>(
ndRange, [=](sycl::nd_item<1> item) {
int globalId = item.get_global_id(0);
mergeSort(data_, 0, data_.size() - 1);
});
}
void setData(std::vector<T> &data) {
data_ = data;
}
std::vector<T> getData() {
return data_;
}
private:
std::vector<T> data_;
void merge(std::vector<T>& arr, int l, int m, int r)
{
int i, j, k;
int n1 = m - l + 1;
int n2 = r - m;
std::vector<T> L(n1), R(n2);
for (i = 0; i < n1; i++)
L[i] = arr[l + i];
for (j = 0; j < n2; j++)
R[j] = arr[m + 1+ j];
i = 0;
j = 0;
k = l;
while (i < n1 && j < n2)
{
if (L[i] <= R[j])
{
arr[k] = L[i];
i++;
}
else
{
arr[k] = R[j];
j++;
}
k++;
}
while (i < n1)
{
arr[k] = L[i];
i++;
k++;
}
while (j < n2)
{
arr[k] = R[j];
j++;
k++;
}
}
void mergeSort(std::vector<T>& arr, int l, int r)
{
if (l < r)
{
int m = l+(r-l)/2;
mergeSort(arr, l, m);
mergeSort(arr, m+1, r);
merge(arr, l, m, r);
}
}
};
int main() {
std::vector<int> data = {9, 4, 2, 7, 5, 1, 8, 3, 6};
sycl::default_selector selector;
sycl::queue queue(selector);
sycl::buffer<int> buffer(data.data(), sycl::range<1>(data.size()));
queue.submit([&](sycl::handler &cgh) {
auto acc = buffer.get_access<sycl::access::mode::read_write>(cgh);
cgh.single_task<class SortTask>([=]() {
ParallelMergeSort<int> parallelMergeSort;
parallelMergeSort.setData(acc.get_pointer());
parallelMergeSort(cgh);
acc.get_pointer() = parallelMergeSort.getData();
});
});
auto result = buffer.get_access<sycl::access::mode::read>();
for (int i = 0; i < result.get_range()[0]; i++) {
std::cout << result[i] << " ";
}
std::cout << std::endl;
return 0;
}
上述代码中，我们首先初始化一个整数向量 data，然后创建一个 sycl::buffer 对象，用于存储数据。接下来，我们使用 oneAPI 的队列和设备来提交并行计算任务。在任务中，我们将数据传递给 ParallelMergeSort 类，并调用它来执行并行归并排序算法。最后，我们从缓冲区中获取排序后的结果，并将其打印出来。
这样，我们就成功地利用英特尔 oneAPI 技术实现了高性能的并行归并排序算法。通过利用异构计算资源，我们可以加速排序过程，提高算法的执行效率。
六、结论
在这篇文章中，我们采用了SYCL运行时，利用Intel DPC++编译器提供的额外扩展来开发一个C++应用程序。该应用程序展示了统一编程模型在不同体系结构上的运用，以及SYCL运行时为协调对内存的访问和使用隐式依赖图编写并行代码提供的抽象。
七、收获与感悟
- 学会了如何运用SYCL进行有效的并行编程；
- 学会了如何利用oneAPI中的工具和库来优化图像处理算；
- 理解并应用并行计算的原理以加速图像卷积运算；
- 在不同的硬件架构上实施并测试性能优化；
- 深化了对高性能计算和并行处理的理解。
致谢
最后，感谢英特尔提供的DevCloud平台支持，感谢英特尔亚太研发有限公司技术团队提供的技术支持。
扫一扫
3174
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
