英特尔oneAPI-英特尔异构统一编程与接口初识及其深度学习应用

---

1 高性能计算的今天

高性能计算(High performance computing，HPC) 通常使用不同的设备、处理器进行计算。HPC的市场增长逾加迅速，在云计算、大数据、AI等领域大放异彩。如今，全球芯片短缺似乎无法避免，而高性能计算需求愈加旺盛，随着深度学习等需要高性能算力的模型快速发展，硬件价格的上涨、计算能力的限制、模型设计的日益发展，让有限的设备发挥高性能的计算能力，成为无法避免的问题与挑战。

HPC 可以提供给使用者比常规的计算机设备更快的处理大量数据的能力，从而更快地获得所需要的计算结果。 通常情况下，HPC 提供的算力比当前较为先进的笔记本电脑强大一百万倍。这种强大的计算能力使使用者能够运行复杂运算、大型分析计算等。总之，HPC 赋予了卓越的性能，使使用者，无论是公司还是个人，都能够以更少的支出做更多的事情。

---

2 oneAPI的生正逢时

one 代表统一

oneAPI 是一种开放的、跨行业的、基于标准的统一编程模型，可提供跨加速器架构的通用开发人员体验，以实现更快的应用程序性能、更高的生产力和更大的创新。 oneAPI 倡议鼓励在 oneAPI 规范和整个生态系统中兼容的 oneAPI 实现上进行协作。

对于oneAPI，其提供了高性能异构计算的统一编程，打造异构计算的新纪元，英特尔对于oneAPI的特征如此表述：

• 告别专用锁定
• 充分实现硬件价值
• 快速正确开发高性能代码

以上三点是英特尔对于其oneAPI的特征定义，可以带来几点信息，同时也是oneAPI面向解决方案。oneAPI是接口，也是一种标准。

• 对于编程与接口的规范，可以使代码应用在不同的架构中，避免编程架构带来的应用瓶颈；
• 对于异构编程的支持，CPU、GPU、FPGA等不同硬件平台，实现高性能计算及其算法加速的设计；
• 对于库、工具和框架的完善，通过简洁直白、高效易懂的方式实现异构开发与计算。

对不同的开发者，oneAPI也提供了不同类型的工具包：

• 对于大多数开发人员提供英特尔® oneAPI 基础工具包
  用例：使用此基础工具集开发跨英特尔® CPU、GPU 和 FPGA的高性能、以数据为中心的应用程序。
• 对于 HPC 开发人员提供英特尔® oneAPI HPC 工具包
  用例：跨共享和分布式内存计算系统构建、分析和扩展应用程序。
• 面向数据科学家和 AI 开发人员提供英特尔® oneAPI 人工智能分析工具包
  用例：使用 Python* 工具和框架加速端到端数据科学和机器学习管道。
• 面向深度学习推理开发人员提供英特尔® OpenVINO 工具套件分发版（由 oneAPI 提供支持）
  用例：将高性能推理应用程序从边缘部署到云端。
• 面向视觉创作者、科学家和工程师提供英特尔® oneAPI 渲染工具包
  用例：创建突破可视化界限的高保真、逼真的体验。
• 面向边缘设备和物联网开发人员提供英特尔® oneAPI 物联网工具包
  用例：在网络边缘运行的应用程序和解决方案的快速开发。
• 对于系统工程师提供英特尔® 系统启动工具包
  用例：利用硬件和软件洞察力增强系统可靠性，并优化功耗和性能。

目前，oneAPI最新版本为2022.1

英特尔® oneAPI 工具包 发布 2022.1
使用一流的编译器、性能库、框架以及分析和调试工具在 CPU 和 XPU 上构建、分析和优化高性能、跨架构应用程序。

同时，我们也可以看到在其官网的最新更新中的内容，以获得最新支持：

• 使用 OpenMP 让 CPU 和 GPU 同时运行
• 幕后故事：用 SciVis 专业知识武装研究人员 [19:01]
• 使用 HPO方法加快模型开发
• 点播视频：简化并加快 3D 体积计算 [58:58] 英特尔® VTune™ 指南（新配置）
• 使用数据并行C++/SYCL 求解二维热传导方程
• LLVM 和 GCC 中针对英特尔® CPU 和 GPU 的 SIMD 矢量化

详细内容可以参考 oneAPI主页
具体 Intel® oneAPI Toolkits 工具包下载页面参考这个链接

---

3 初探oneAPI

All the features of DPC++, and SYCL, are supported by both the open-source and commercial versions of the DPC++ compilers.

这意味着DPC++是一个完全开源的架构，是一个跨平台编译器。

oneAPI 由主机和设备组成，这和大多数异构统一编程模型类似。主机通常是CPU，设备则是一个或多个 GPU、FPGA 和其他设备，可以作为主机的处理器，也可以作为软件的目标。运行在主机上的应用程序和运行在设备上的功能通过内存进行通信，并提供了缓冲区对象、统一寻址、统一共享内存三种内存共享机制。

oneAPI Data Parallel C++ (DPC++) 是 oneAPI 的直接编程语言和相关的直接编程 API，DPC++支持C++编程语言，兼容g++，它提供了定义数据并行函数并在设备上启动它们所需的特性。DPC++ 建立在 The Khronos Group 的 SYCL 规范之上。SYCL是OpenCL的高级编程模型，作为基于纯C++的单源特定于域的嵌入式语言(DSEL), 用于提高编程效率，是一种基于C++异构平行编程框架，用来加速高性能计算。SYCL 语言支持定义可卸载到设备的数据并行函数，并定义用于编排卸载函数的运行时 API 和类，开放标准SYCL、OpenCL与NVIDIA的CUDA类似。

DPC++ 是 C++ 模型。
DPC++ 程序员应该对 C++11 感到熟悉。
DPC++ 基于 C++ 特性的异构平台和并行性。
DPC++ 将为未来的 ISO C++ 标准提供有价值的输入

同时，英特尔对如 PyTorch，Numpy，SciPy 等进行了优化，作为库和框架，下图即为oneAPI对AI的优化架构。

详细的介绍文档可以参考官方文档

---

4 异构计算的小试牛刀

4.1 oneAPI与深度学习

oneAPI为深度学习的开发者提供了英特尔® 发行版 OpenVINO™ 工具套件，让开发人员可以轻松编程，提供更多深度学习模型、设备可移动性和更高的推理性能，同时减少代码更改。有着针对推理优化的 Open Model Zoo 预训练模型。该工具套件与大部分热门框架都已经兼容，并最大限度地提高了性能。

可以通过官方介绍可以安装OpenVINO在Linux系统中，提供多种安装方式：

Install OpenVINO Runtime using an Installer
Install OpenVINO from PyPI
Install OpenVINO Runtime from APT
Install OpenVINO Runtime from YUM
Install OpenVINO Runtime from Anaconda Cloud
Install OpenVINO with Docker

需要注意的是， OpenVINO 在2022.1 版本以后，仅支持Ubuntu 18.04 长期支持 (LTS)，64 位和Ubuntu 20.04 长期支持 (LTS)，64 位等操作系统，不再支持 CentOS 7.6，64 位。

在安装好 OpenVINO 开发套件后，可以安装OpenVINO Development Tools，例如使用caffe架构下的开发模型，执行如下命令：

pip install openvino-dev[caffe]

可以下载一些提供的预训练模型，如googlenet-v1等

omz_downloader --name <model_name> --output_dir <models_dir>

接下来对下载的预训练模型进行优化为 IR (Intermediate Representation）格式，

mkdir ~/ir
mo --input_model ~/models/public/googlenet-v1/googlenet-v1.caffemodel --data_type FP16 --output_dir ~/ir

通过下载测试图像或视频，OpenVINO工具套件包括几个示例图像和视频，可以使用它们来运行代码示例和演示应用程序，详细见这里

最后，就可以运行测试了，首先设置环境变量，之后运行测试，这里以CPU为例：

source  <INSTALL_DIR>/setupvars.sh
cd ~/inference_engine_cpp_samples_build/intel64/Release
./classification_sample_async -i ~/Downloads/dog.bmp -m ~/ir/googlenet-v1.xml -d CPU

运行程序完成后，可以看到前 10 个类别的标签和置信度，就像这个样子：

官方提供了更多实例，详细见这里

4.3 认识一下oneDNN

作为一位深度学习的开发者，这里继续选择深度学习模型的高性能计算，并使用oneAPI的oneDNN，作为一个入门教程，带大家使用一下英特尔的异构编程平台。

简而言之，就是oneAPI 深度神经网络库 (oneDNN) 可以用于深度学习并支持卷积神经网络、循环神经网络、归一化、激活函数计算、多精度类型（8 位整数、16 位、32 位和 bfloat16 浮点数据类型）等，基本上常规的神经网络计算操作都有支持，应该是有着比较完善的算子库。

官方给出了oneDNN的系统构成：
oneDNN由4部分组成：

• dnnl::primitive
• dnnl::engine
• dnnl::stream
• dnnl::memory

其中，dnnl::primitive对相关特定计算进行封装，即算子，例如CONV等； dnnl::engine则是不同的计算设备如CPU、GPU 等，大多数算子为指定设备进行创建和优化； dnnl::stream对执行上下文进行封装；dnnl::memory则是对存储相关进行封装，如张量维度、数据类型和内存格式等。

对于深度学习模型的执行，则需要将相关操作和参数等传递dnnl::primitive::execute()；和TVM、TensorRT类似，oneAPI也使用引擎定义在不同设备运行的执行单元，

关于数据类型的定义，oneAPI支持非常丰富的数据类型：

值的注意的是，oneAPI支持Bfloat16 ( bf16)，这是基于 IEEE 32 位单精度浮点数据类型 ( f32) 的16 位浮点数据类型，与f32不同，f32尾数是 23 位，但bf16只有7位，这种定义显然会造成一定程度的精度下降。

对深度学习一些计算操作进行融合以加快运行速度，一个例子如下图。

当然，还有诸多算子操作，这里不在赘述，详细可以参考官方的解释文档，点击这里啦

4.4 动手跑一下oneDNN的一个例子

先看一下支持的设备，优化比较好的肯定是自家设备啦～
注意，oneDNN只提供 C 和 C++ API

oneDNN 对以下架构提供实验性支持：Arm* 64 位架构 (AArch64)、NVIDIA* GPU、OpenPOWER* Power ISA (PPC64)、IBMz* (s390x) 和 RISC-V。

首先安装oneAPI，我的设备是linux，所以在官网下载工具包，点击oneAPI Toolkits，我的选择是Intel® oneAPI Base Toolkit，按照提示下载安装。

wget https://registrationcenter-download.intel.com/akdlm/irc_nas/18487/l_BaseKit_p_2022.1.2.146.sh 
sudo sh ./l_BaseKit_p_2022.1.2.146.sh -a --silent --eula accept

有一个小细节，由于我是远程登录的服务器，这个安装是有gui界面的，需要使用ssh -X连接，其实英特尔可以直接搞一个命令行也许也蛮友好的。

Base toolkit离线安装包大概有3G，需要等待～

clone一下代码库，可能会存在网络较慢的问题，等待一下，下载好代码后，进行编译。
注意，这里使用的是linux下的GCC编译器，这需要你提前下载并安装, 当然，也可以使用Clang、Intel C/C++ Compiler

关于cmake的编译选项，可以参考这里，点击它！

git  clone  https : //github.com/oneapi-src/oneDNN.git
cd oneDNN
mkdir -p build
cd build
cmake .. 
make -j # 这里可以选择多核心编译，如 -j4

以下就是编译完成啦

编译速度很快，安装完后，可以测试一下，测试需要一段时间，等待一下～

ctest

测试完成！
使用conda建立虚拟环境，并安装相关依赖库

sudo cmake --build . --target install

好啦，到这里，环境的安装就大功告成！

接下来，对官方文档代码进行复现，这可以让我们了解其基础知识。

回到oneDNN的目录下，在examples下有着一些代码，这里主要看 getting_started.cpp 这个代码文件

首先是头文件：

#include <cmath>
#include <numeric>
#include <stdexcept>
#include <vector>
#include "oneapi/dnnl/dnnl.hpp"
#include "oneapi/dnnl/dnnl_debug.h"
#include "example_utils.hpp"
using namespace dnnl;

在主函数中，我们可以看出他的运行逻辑：

int main(int argc, char **argv) {
    int exit_code = 0;

    engine::kind engine_kind = parse_engine_kind(argc, argv);
    try {
        getting_started_tutorial(engine_kind);
    } catch (dnnl::error &e) {
        std::cout << "oneDNN error caught: " << std::endl
                  << "\tStatus: " << dnnl_status2str(e.status) << std::endl
                  << "\tMessage: " << e.what() << std::endl;
        exit_code = 1;
    } catch (std::string &e) {
        std::cout << "Error in the example: " << e << "." << std::endl;
        exit_code = 2;
    }

    std::cout << "Example " << (exit_code ? "failed" : "passed") << " on "
              << engine_kind2str_upper(engine_kind) << "." << std::endl;
    return exit_code;
}

这段代码主要是函数getting_started_tutorial()，其中，engine_kind是运行设备定义，这里我们在CPU上运行。

void getting_started_tutorial(engine::kind engine_kind) {

对引擎，即设备的种类和索引进行定义，此外，还需要定义流来执行。

	engine eng(engine_kind, 0);
    stream engine_stream(eng);

接下来定义数据，做过图像的图像一定对创建一个 NHWC 格式的 4D 张量十分熟悉啦，我们这里也创建一个类似图像张量的数据。

	const int N = 1, H = 13, W = 13, C = 3;	
    const int stride_N = H * W * C;
    const int stride_H = W * C;
    const int stride_W = C;
    const int stride_C = 1;

还有一些数据的辅助操作：

// 将逻辑索引映射到物理偏移量的辅助函数
auto  offset  =  [ = ]( int  n ,  int  h ,  int  w ,  int  c )  { 
    return  n  *  stride_N  +  h  *  stride_H  +  w  *  stride_W  +  c  *  stride_C ; 
}; 

// 图片大小
const  int  image_size  =  N  *  H  *  W  *  C; 

// 为图像分配一个缓冲区
std :: vector < float >  image ( image_size ); 

// 用一些值初始化图像
for  ( int  n  =  0 ;  n  <  N ;  ++ n ) 
    for  ( int  h  =  0 ;  h  <  H ;  ++ h ) 
        for  ( int  w  =  0 ;  w  <  W ;  + +w ) 
            for  ( int  c  =  0 ;  c  <  C ;  ++ c )  { 
                int  off  =  offset ( n ,  h ,  w ,  c );  // 获取像素
                image[ off ]  =  - std :: cos ( off  /  10.f ); 
            }

在定义好数据后，需要将其定义在dnnl::memory对象中，首先初始化内存描述：

auto src_md = memory::desc(
            {N, C, H, W}, // logical dims, the order is defined by a primitive
            memory::data_type::f32, // tensor's data type
            memory::format_tag::nhwc // memory format, NHWC in this case
    );

准备好内存描述符和引擎后，让我们为 ReLU 算子创建输入和输出内存对象，并定义算子。

    auto src_mem = memory(src_md, eng);
    write_to_dnnl_memory(image.data(), src_mem);
    auto dst_mem = memory(src_md, eng);
    // ReLU 操作描述符（没有特定于引擎或实现的信息）
	auto  relu_d  =  eltwise_forward :: desc ( 
       	prop_kind :: forward_inference ,  algorithm :: eltwise_relu , 
        src_md ,  // 操作的内存描述符
        0.f ,  // alpha 参数在 ReLU 的情况下表示负斜率
        0.f  // 在 ReLU 的情况下 beta 参数被忽略
	); 

在依赖库中有他的算子库，挂上去就可以：

auto relu_pd = eltwise_forward::primitive_desc(relu_d, eng);

这一句可以创建一个可以在内存对象上执行以计算操作的 primitive（这个应该被翻译为「原语」吗，感觉怪怪的hh）

auto relu = eltwise_forward ( relu_pd );

在最后一步，就是执行primitive，后得出结果，源代码也有验证～

	relu.execute(engine_stream,  {{DNNL_ARG_SRC, src_mem}, {DNNL_ARG_DST, dst_mem},  });
    engine_stream.wait();
    std::vector<float> relu_image(image_size);
    read_from_dnnl_memory(relu_image.data(), dst_mem);

完成！

Example passed.

5 使用DPCT迁移CUDA代码到DPC++

由于现有的相当一部分深度学习相关的代码为基于NVIDIA的显卡，其计算和加速使用CUDA框架，oneAPI也提供了一个兼容性转换工具dpct，通过它可以实现讲CUDA代码转换为DPC++代码。

测试环境可以在本地电脑中安装 oneAPI 环境，也可以使用devCLoud进行测试与调优，作为一个具有官方环境的云服务，注册即可使用。在本地使用DPCT时，需要要设置英特尔® DPC++ 兼容性工具环境，请运行以下命令：

在 Linux (sudo) 上：source /opt/intel/oneapi/setvars.sh
在 Linux（用户）上：source ~/intel/oneapi/setvars.sh
在 Windows 上：\Program Files (x86)\Intel\oneAPI\setvars.bat

devCLoud给出的连接步骤及使用非常清晰，点赞👍

这里使用英特尔® DPC++ 兼容性工具迁移使用 CUDA 版本 8.0、9.x、10.x 和 11.0–11.6 实现的程序。使用下面的命令查看命令选项：

dpct --help

英特尔® DPC++ 兼容性工具可识别代码中在文件迁移期间可能需要注意的位置，以使代码 DPC++ 兼容或正确。同时，注释会插入到生成的源文件中，并在输出中显示为警告，例如；

我们在这里对矩阵相乘这个经典代码为例，编写基于GPU的多线程矩阵相乘代码，原代码为test_matrix.cu，是基于CUDA的代码，利用DPC++ 兼容性工具，命令 dpct 进行转换成为适用于英特尔的代码。

完整的 test_matrix.cu 如下，在设备端实现两个矩阵的相乘，相信这段代码伴随注释并不会很难懂，这里就不过多介绍了。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>

using namespace std;

// Print GPU information
void GetCudaInf() {
	int deviceCount;
	cudaGetDeviceCount(&deviceCount);
	int dev;
	for (device = 0; dev < deviceCount; device++) {
		int driver_version(0), runtime_version(0);
		cudaDeviceProp deviceProp;
		cudaGetDeviceProperties(&deviceProp, device);
		if (device == 0)
			if (deviceProp.minor = 9999 && deviceProp.major == 9999)
				cout << endl;
		cudaDriverGetVersion(&driver_version);
		cudaRuntimeGetVersion(&runtime_version);
		cout << "=> GPU device " << device << ": " << deviceProp.name << endl;
		cout << "=> CUDA driver version:" << driver_version / 1000 << "." << (driver_version % 1000) / 10 << endl;
		cout << "=> CUDA runtime version:" << runtime_version / 1000 << "." << (runtime_version % 1000) / 10 << endl;
	}
}

// M(row, col) = *(M.elements + row * M.width + col)
struct Matrix {
	int width;
	int height;
	float *elements;
};

// Get the position of matrix A as (row, col) element
__device__ float getElement(Matrix *A, int row, int col) {
	return A->elements[row * A->width + col];
}

// Set a value to the element at position (row, col) of matrix A
__device__ void setElement(Matrix *A, int row, int col, float value) {
	A->elements[row * A->width + col] = value;
}

// Matrix multiplication kernel function, 2-D, computes one element Cij per thread
__global__ void matMulKernel(Matrix *A, Matrix *B, Matrix *C) {
	float Cvalue = 0.0;
	int row = threadIdx.y + blockIdx.y * blockDim.y;
	int col = threadIdx.x + blockIdx.x * blockDim.x;
	for (int i = 0; i < A->width; ++i) {
		Cvalue += getElement(A, row, i) * getElement(B, i, col);
	}
	setElement(C, row, col, Cvalue);
}

int main() {
	// Initialization
	GetCudaInf();
	int width = 1 << 10;
	int height = 1 << 10;
	Matrix *A, *B, *C;
	cudaMallocManaged((void**)&A, sizeof(Matrix));
	cudaMallocManaged((void**)&B, sizeof(Matrix));
	cudaMallocManaged((void**)&C, sizeof(Matrix));
	int nBytes = width * height * sizeof(float);
	cudaMallocManaged((void**)&A->elements, nBytes);
	cudaMallocManaged((void**)&B->elements, nBytes);
	cudaMallocManaged((void**)&C->elements, nBytes);
	A->height = height;
	A->width = width;
	B->height = height;
	B->width = width;
	C->height = height;
	C->width = width;
	for (int i = 0; i < width * height; ++i) {
		A->elements[i] = 1.0;
		B->elements[i] = 2.0;
	}
	// Define the block size of the kernel is (32, 32) and the grid size is (32, 32)
	dim3 blockSize(32, 32);
	dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
		(height + blockSize.y - 1) / blockSize.y);
	// Run kernel
	matMulKernel << < gridSize, blockSize >> > (A, B, C);
	cudaDeviceSynchronize();
	// Check the result
	float maxError = 0.0;
	for (int i = 0; i < width * height; ++i)
		maxError = fmax(maxError, fabs(C->elements[i] - 2 * width));
	cout << endl;
	cout << "=> The error : " << maxError << endl;
	return 0;
}

通过如下命令进行转换：

dpct --in-root=. test_matrix.cu

转换完成后， test_matrix.dp.cpp 文件应该出现在当前路径下的 dpct_output 目录，该文件现在是 DPC++ 源文件。由于代码转换并不能都能成功，所以需要看代码中的插入注释，按照注释内容尝试进行修改。

这段代码中的注释主要为CUDA信息打印时的函数定义不同导致，剩余代码均转换成功，如：

if (device == 0)
            /*
			DPCT1005:2: The SYCL device version is different from CUDA Compute Compatibility. You may need to rewrite this code.
			*/
            if (deviceProp.set_minor_version(
                        9999 && deviceProp.get_major_version() == 9999))
                cout << endl;
        /*
		DPCT1043:0: The version-related API is different in SYCL. An initial code was generated, but you need to adjust it.
		*/
        driver_version = dpct::get_current_device()
                                 .get_info<sycl::info::device::version>();
        /*
		DPCT1043:1: The version-related API is different in SYCL. An initial code was generated, but you need to adjust it.
		*/

完整的 test_matrix.dp.cpp 代码如下：

#include <CL/sycl.hpp>
#include <dpct/dpct.hpp>
#include <iostream>
#include <stdio.h>
#include <cmath>

using namespace std;

// Print GPU information
void GetCudaInf() {
	int deviceCount;
    deviceCount = dpct::dev_mgr::instance().device_count();
    int dev;
	for (device = 0; dev < deviceCount; device++) {
		int driver_version(0), runtime_version(0);
        dpct::device_info deviceProp;
        cudaGetDeviceProperties(&deviceProp, device);
		if (device == 0)
            /*
			DPCT1005:2: The SYCL device version is different from CUDA Compute Compatibility. You may need to rewrite this code.
			*/
            if (deviceProp.set_minor_version(
                        9999 && deviceProp.get_major_version() == 9999))
                cout << endl;
        /*
		DPCT1043:0: The version-related API is different in SYCL. An initial code was generated, but you need to adjust it.
		*/
        driver_version = dpct::get_current_device()
                                 .get_info<sycl::info::device::version>();
        /*
		DPCT1043:1: The version-related API is different in SYCL. An initial code was generated, but you need to adjust it.
		*/
        runtime_version = dpct::get_current_device()
                                  .get_info<sycl::info::device::version>();
        cout << "=> GPU device " << device << ": " << deviceProp.name << endl;
		cout << "=> CUDA driver version:" << driver_version / 1000 << "." << (driver_version % 1000) / 10 << endl;
		cout << "=> CUDA runtime version:" << runtime_version / 1000 << "." << (runtime_version % 1000) / 10 << endl;
	}
}

// M(row, col) = *(M.elements + row * M.width + col)
struct Matrix {
	int width;
	int height;
	float *elements;
};

// Get the position of matrix A as (row, col) element
float getElement(Matrix *A, int row, int col) {
	return A->elements[row * A->width + col];
}

// Set a value to the element at position (row, col) of matrix A
void setElement(Matrix *A, int row, int col, float value) {
	A->elements[row * A->width + col] = value;
}

// Matrix multiplication kernel function, 2-D, computes one element Cij per thread
void matMulKernel(Matrix *A, Matrix *B, Matrix *C, sycl::nd_item<3> item_ct1) {
	float Cvalue = 0.0;
    int row = item_ct1.get_local_id(1)
            + item_ct1.get_group(1) * item_ct1.get_local_range().get(1);
    int col = item_ct1.get_local_id(2)
            + item_ct1.get_group(2) * item_ct1.get_local_range().get(2);
    for (int i = 0; i < A->width; ++i) {
		Cvalue += getElement(A, row, i) * getElement(B, i, col);
	}
	setElement(C, row, col, Cvalue);
}

int main() {
 dpct::device_ext &dev_ct1 = dpct::get_current_device();
 sycl::queue &q_ct1 = dev_ct1.default_queue();
    // Initialization
	// GetCudaInf();
	int width = 1 << 10;
	int height = 1 << 10;
	Matrix *A, *B, *C;
    A = sycl::malloc_shared<Matrix>(1, q_ct1);
    B = sycl::malloc_shared<Matrix>(1, q_ct1);
    C = sycl::malloc_shared<Matrix>(1, q_ct1);
    int nBytes = width * height * sizeof(float);
    A->elements = (float *)sycl::malloc_shared(nBytes, q_ct1);
    B->elements = (float *)sycl::malloc_shared(nBytes, q_ct1);
    C->elements = (float *)sycl::malloc_shared(nBytes, q_ct1);
    A->height = height;
	A->width = width;
	B->height = height;
	B->width = width;
	C->height = height;
	C->width = width;
	for (int i = 0; i < width * height; ++i) {
		A->elements[i] = 1.0;
		B->elements[i] = 2.0;
	}
	// Define the block size of the kernel is (32, 32) and the grid size is (32, 32)
    sycl::range<3> blockSize(1, 32, 32);
    sycl::range<3> gridSize(1, (height + blockSize[1] - 1) / blockSize[1],
            (width + blockSize[2] - 1) / blockSize[2]);
    // Run kernel
	matMulKernel << < gridSize, blockSize >> > (A, B, C);
    dev_ct1.queues_wait_and_throw();
    // Check the result
	float maxError = 0.0;
	for (int i = 0; i < width * height; ++i)
		maxError = fmax(maxError, fabs(C->elements[i] - 2 * width));
	cout << endl;
	cout << "=> The error : " << maxError << endl;
	return 0;
}

有关更复杂的示例说明，请参阅英特尔® DPC++ 兼容性工具开发人员指南和参考的 迁移项目部分。

6 总结

本文对高性能计算的背景进行引入，介绍了英特尔异构统一编程接口「oneAPI」，主要介绍了oneAPI的架构及其在深度学习中的应用，涵盖内容从相关概念的解读与梳理，到开发环境的配置安装，DPC++兼容工具的使用，以及深度学习中的实例运行教程等。总之，oneAPI是一个相对完善的异构编程接口，为开发人员带来了不错的便利工具，相对完善的文档、程序样例和开源的获取方式会推动高性能计算领域的进一步发展。

笔者认知有限，如有错误欢迎指出。
参考内容：
https://www.intel.cn/content/www/cn/zh/developer/tools/oneapi/toolkits.html#openvino-kit
https://baike.baidu.com/item/%E9%AB%98%E6%80%A7%E8%83%BD%E8%AE%A1%E7%AE%97/329249
https://zhuanlan.zhihu.com/p/103299140
https://www.intel.cn/content/www/cn/zh/developer/tools/oneapi/overview.html#gs.z02tlf
https://blog.csdn.net/weixin_64235560/article/details/124517039
https://space.bilibili.com/1840724569/channel/collectiondetail?sid=282053