1. 前言

在人工智能興起的今天，大規(guī)模、高性能計(jì)算已成為社會(huì)發(fā)展的剛需。動(dòng)輒千萬節(jié)點(diǎn)規(guī)模的社交網(wǎng)絡(luò)、交通網(wǎng)絡(luò)，語言聊天模型中的大規(guī)模神經(jīng)網(wǎng)絡(luò)，以及航空航天等涉及大規(guī)模計(jì)算的場(chǎng)景，都少不了并行計(jì)算的支持。并行計(jì)算是一種一次可執(zhí)行多個(gè)指令的算法，目的是提高計(jì)算速度，及通過擴(kuò)大問題求解規(guī)模，解決大型而復(fù)雜的計(jì)算問題。理論上，使用并行計(jì)算可以將性能提升至單線程計(jì)算的任意倍數(shù)，這大大提高了計(jì)算機(jī)的性能，極大地便利了人們的日常生活。

2. Intel oneAPI概述

Intel oneAPI 是一個(gè)跨行業(yè)、開放、基于標(biāo)準(zhǔn)的統(tǒng)一的編程模型，它為跨 CPU、GPU、FPGA、專用加速器的開發(fā)者提供統(tǒng)一的體驗(yàn)。oneAPI的開放規(guī)范基于行業(yè)標(biāo)準(zhǔn)和現(xiàn)有開發(fā)者編程模型，廣泛適用于不同架構(gòu)和來自不同供應(yīng)商的硬件。oneAPI 行業(yè)計(jì)劃鼓勵(lì)生態(tài)系統(tǒng)內(nèi)基于oneAPI規(guī)范的合作以及兼容 oneAPI的實(shí)踐。
Intel oneAPI 產(chǎn)品是英特爾基于oneAPI 的實(shí)現(xiàn)，它包括了 oneAPI 標(biāo)準(zhǔn)組件如直接編程工具（Data Parallel C++）、含有一系列性能庫的基于 API 的編程工具，以及先進(jìn)的分析、調(diào)試工具等組件。開發(fā)人員從現(xiàn)在開始就可以在英特爾 DevCloud for oneAPI 上對(duì)基于多種英特爾架構(gòu)（包括英特爾至強(qiáng)可擴(kuò)展處理器、帶集成顯卡的英特爾酷睿處理器、英特爾 FPGA 如英特爾 Arria、Stratix 等）的代碼和應(yīng)用進(jìn)行測(cè)試。
Intel oneAPI 功能豐富，涵蓋了計(jì)算機(jī)編程各個(gè)方面，是世界范圍內(nèi)公認(rèn)的優(yōu)秀編程模型。Intel oneAPI包括以下套件：

接下來從其精細(xì)的內(nèi)存和任務(wù)管理、并行化、安全保證、高性能操作、程序分析等方面，結(jié)合具體詳實(shí)登錄例子，介紹 Intel oneAPI的豐富功能。

3. 內(nèi)存和任務(wù)管理

3.1 Queue

oneapi使用DPC++語言，DPC++基于傳統(tǒng)C++語言，支持CPU、GPU、FPGA等異構(gòu)計(jì)算設(shè)備和分布式計(jì)算，提供了一種統(tǒng)一的編程模型以實(shí)現(xiàn)異構(gòu)計(jì)算，而無需學(xué)習(xí)多種不同的編程語言和API。

Intel OneAPI中的queue是一種基于DPC++編程語言的命令隊(duì)列，允許開發(fā)人員將命令（比如內(nèi)存拷貝、計(jì)算任務(wù)等）添加到隊(duì)列中，并在異構(gòu)設(shè)備上執(zhí)行這些命令。queue支持多種不同的執(zhí)行選項(xiàng)，例如同步執(zhí)行和異步執(zhí)行，可以幫助開發(fā)人員實(shí)現(xiàn)更加高效的計(jì)算。

• 創(chuàng)建queue并選擇設(shè)備

queue q(gpu_selector_v); //使用了gpu_selector_v選擇器，指定計(jì)算設(shè)備
std::cout << "Device: " << myQueue.get_device().get_info<info::device::name>() << "\n"; //打印設(shè)備

• 命令執(zhí)行：

使用submit函數(shù)，并向其中傳遞一個(gè)lambda表達(dá)式。這個(gè)lambda表達(dá)式接受一個(gè)handler對(duì)象作為參數(shù)，用于向隊(duì)列中添加要執(zhí)行的命令：

myQueue.submit([&](handler& h) {
  // 添加要執(zhí)行的命令
  h.single_task([]() { std::cout << "Hello, world!" << std::endl; });
});

并行執(zhí)行任務(wù)：

myQueue.submit([&](handler& h) {
  h.parallel_for(range<1>(N), [=](id<1> i) {
    // CODE THAT RUNS ON DEVICE
  });
});

myQueue.submit([&](handler& h) {
h.parallel_for(nd_range<1>(range<1>(1024),range<1>(64)), [=](nd_item<1> item){
    // CODE THAT RUNS ON DEVICE
});
});

以下是利用oneapi queue的bfs運(yùn)算的部分代碼：

q.submit([&](handler& h) {
    h.parallel_for(nd_range<1>(N, B), [=](nd_item<1> item) {
        auto tid = item.get_global_id()[0];
        for (uint i = tid; i < nodeNum; i += N) {
            uint id = activeNodesD[i];
            uint edgeIndex = nodePointersD[id];
            uint sourceValue = valueD[id];
            uint finalValue;
            for (uint i = 0; i < degreeD[id]; i++) {
                finalValue = sourceValue + 1;
                uint vertexId;
                vertexId = edgeListD[edgeIndex + i];
                if (1!=valueD[vertexId]) {
                    valueD[vertexId] = 1;
                    labelD[vertexId] = 1;
                }
    }
        }
    });
}).wait();

3.2 Device Memory

Intel OneAPI提供了一些用于在設(shè)備內(nèi)存上進(jìn)行操作的函數(shù)和方法，可以幫助開發(fā)人員更加輕松地管理和執(zhí)行異構(gòu)計(jì)算，如分配和釋放設(shè)備內(nèi)存、復(fù)制數(shù)據(jù)和執(zhí)行計(jì)算任務(wù)等。

• 在device端分配和釋放內(nèi)存

在Intel OneAPI中，可以使用 malloc_device 函數(shù)在設(shè)備上動(dòng)態(tài)分配內(nèi)存， free 函數(shù)來釋放內(nèi)存。

void* malloc_device(size_t bytes, const queue& q);
// bytes：分配內(nèi)存大?。ㄗ止?jié)），q：使用的設(shè)備隊(duì)列
// 例如：
int* deviceData = malloc_device<int>(vertexArrSize,myqueue);
free(deviceData, myQueue);

• 從主機(jī)內(nèi)存復(fù)制到設(shè)備內(nèi)存

myQueue.memcpy(deviceData, hostData, 1000 * sizeof(int));

• 從設(shè)備內(nèi)存復(fù)制到主機(jī)內(nèi)存

myQueue.memcpy(hostData, deviceData, 1000 * sizeof(int));

• 在主機(jī)和設(shè)備之間動(dòng)態(tài)共享內(nèi)存

void* malloc_shared(size_t bytes, const queue& q);
// bytes：分配內(nèi)存大?。ㄗ止?jié)），q：使用的設(shè)備隊(duì)列
// 例如：
int* sharedData  = malloc_shared<int>(vertexArrSize,myqueue);
free(sharedData , myQueue);

• 初始化

memset 函數(shù)對(duì)設(shè)備內(nèi)存進(jìn)行初始化操作，該函數(shù)用于將一段內(nèi)存空間中的每一個(gè)字節(jié)都設(shè)置成指定的值。

void* memset(void* ptr, int value, size_t num);
// ptr：指向要被設(shè)置的內(nèi)存空間的指針，value：要設(shè)置的值，以整數(shù)形式傳遞，num：要設(shè)置的字節(jié)數(shù)。
// 例如：
memset(deviceData, 0, bytes);

以下是利用device memory的內(nèi)存初始化部分：

degreeD = malloc_device<SIZE_TYPE>(vertexArrSize,streamStatic);
isActiveD = malloc_device<uint>(vertexArrSize,streamStatic);
isStaticActive = malloc_device<uint>(vertexArrSize,streamStatic);
isOverloadActive = malloc_device<uint>(vertexArrSize,streamStatic);

streamStatic.memcpy(degreeD, degree, vertexArrSize * sizeof(SIZE_TYPE));
streamStatic.memcpy(isActiveD, label, vertexArrSize * sizeof(uint));

streamStatic.memset(isStaticActive, 0, vertexArrSize * sizeof(uint));
streamStatic.memset(isOverloadActive, 0, vertexArrSize * sizeof(uint));

4. 多線程及正確性保障

4.1 oneTBB

oneTBB（oneAPI Threading Building Blocks）是一個(gè)面向任務(wù)的C++庫，旨在使底層復(fù)雜的線程和鎖機(jī)制對(duì)用戶透明化?？缙脚_(tái)、跨體系結(jié)構(gòu)特性使其更加靈活，一份代碼即可以兼容多種不同的軟硬件環(huán)境。這使得用戶能夠更專注于任務(wù)本身，只需使用少量抽象層的接口代碼即可實(shí)現(xiàn)大規(guī)模的圖計(jì)算任務(wù)。

// 該代碼用于大規(guī)模圖計(jì)算任務(wù)中填充邊數(shù)組的多線程執(zhí)行
oneapi::tbb::parallel_for( oneapi::tbb::blocked_range<size_t>(0,overloadNodeSize),
    [=](const oneapi::tbb::blocked_range<size_t>& r) {
        for(size_t i=r.begin(); i!=r.end(); ++i){
			// 按照CSR格式計(jì)算對(duì)應(yīng)的邊與結(jié)點(diǎn)
            unsigned int thisNode = overloadNodeList[i];
            unsigned int thisDegree = degree[thisNode];
            EDGE_POINTER_TYPE fromHere = activeOverloadNodePointers[i];
            EDGE_POINTER_TYPE fromThere = nodePointers[thisNode];

            // 計(jì)算on-demand結(jié)點(diǎn)的邊數(shù)據(jù)
            for (unsigned int j = 0; j < thisDegree; j++) {
                overloadEdgeList[fromHere + j] = edgeArray[fromThere + j];
            }
        }
    }
);

4.2 高性能處理——reduction

oneAPI reduce為用戶提供了包括group在內(nèi)多粒度的reduction接口，使得用戶可以便捷地實(shí)現(xiàn)多層次的任務(wù)管理，利用多線程的異步執(zhí)行機(jī)制，輔以恰當(dāng)?shù)耐剑梢詷O大地提升程序的效率。

q.submit([&](handler &h) {                                                                                 // 提交任務(wù)
   h.parallel_for(nd_range<1>(N, B), [=](nd_item<1> item) [[intel::reqd_sub_group_size(sub_group_size)]] { // 并行處理
     // 執(zhí)行pagerank算法
     auto sg = item.get_sub_group();
     auto tid = item.get_global_id()[0];
     int gid = sg.get_group_id()[0] + B / 32 * floor(tid / B);
     SIZE_TYPE cnt_disenableD = 0;
     // 沒有依賴關(guān)系，自動(dòng)并行執(zhí)行
     for (SIZE_TYPE i = tid; i < vertexArrSize; i += N){
       if (inactiveNodeD[i])
         continue;
       uint edgeIndex = nodePointers[i];
       T tempSum = 0;
       for (uint j = edgeIndex; j < edgeIndex + degree[i]; j++){
         uint srcNodeIndex = edgeArray[j];
         if (outDegree[srcNodeIndex]){
           T tempValue = output[srcNodeIndex] / outDegree[srcNodeIndex];
           tempSum += tempValue;
         }
       }
       valueD[i] = (1.0 - beta) + beta * tempSum;
     }
     for (SIZE_TYPE i = tid; i < vertexArrSize; i += N){
       if (inactiveNodeD[i])
         continue;
       T diff = abs(valueD[i] - output[i]);
       output[i] = valueD[i];
       if (diff < 0.001){
         inactiveNodeD[i] = true;
         cnt_disenableD++;
       }
     }
     item.barrier(access::fence_space::local_space);                        // 同步
     SIZE_TYPE sum = reduce_over_group(sg, cnt_disenableD, sycl::plus<>()); // 調(diào)用reduction接口，統(tǒng)計(jì)活躍節(jié)點(diǎn)數(shù)
     if (sg.get_local_id()[0] == 0)
       disenableD[gid] = sum;
   });
 }).wait(); // 等待任務(wù)完成，完成前阻塞

4.3 原子操作

提高計(jì)算機(jī)處理速度的常見方法是使用多線程，即多個(gè)線程同時(shí)訪問共享資源，各自完成自己的計(jì)算。理想情況下，不同線程訪問和使用的共享資源是保持一致的。然而，每個(gè)線程訪問共享資源時(shí)，都需要確保該時(shí)刻沒有其他任何線程修改了這一資源，否則不同線程將訪問或使用到不同的資源，導(dǎo)致程序結(jié)果不正確。因此需要特殊的操作避免這種情況發(fā)生。

oneAPI提供了原子操作(atomic)的接口，所謂原子操作是指不會(huì)被線程調(diào)度機(jī)制打斷的操作；這種操作一旦被一個(gè)線程開始，就一直由這個(gè)線程運(yùn)行到結(jié)束，中間不會(huì)有其他進(jìn)程運(yùn)行。

以下面代碼展示了oneAPI支持的原子操作：

#include <algorithm>
#include <stdio.h>                  /* for EOF */
#include <string.h>                 /* for strchr() */
#include<CL/sycl.hpp>

using namespace sycl;
using namespace std;

int main() {
    sycl::queue q;
    int size = 1000;
    //int sum = 0;
    int* pre_sum = new int[1];
    pre_sum [0]= 0;
    int* value = new int[size];
    int* valueD = malloc_device<int>(size, q);
    int* sum = malloc_device<int>(1, q);
    q.memcpy(sum, pre_sum, sizeof(int)).wait();
    //sum[0] = 0;
    int* sum_noatomic = malloc_device<int>(1, q);
    q.memcpy(sum_noatomic, pre_sum, sizeof(int)).wait();
    //sum_noatomic[0] = 0;
    //int* valueD =;
    for (int i = 0; i < size; i++) {
        value[i] = i + 1;
    }
    q.memcpy(valueD, value, size * sizeof(int)).wait();
    cout << "before buffer \n";
    //buffer<int> buf(sum, 1);
	int NUM_THREADS = 50;
	
	buffer<int> buf(pre_sum, 1);

    q.submit([&](handler& h) {
        ext::oneapi::atomic_accessor acc(buf, h, ext::oneapi::relaxed_order, ext::oneapi::system_scope);
        h.parallel_for(sycl::range(NUM_THREADS), [=](sycl::id<1> ind) {
            //int ind = item.get_global_id(0);
            for (int index
                = ind; index < size; index += NUM_THREADS) {
                acc[0] += valueD[index];    
            }
            }

        );
        }).wait();
        host_accessor host_acc(buf,read_only);

        cout << "atomic sum = " << host_acc[0] << "\n";

        //no atomic test begin
			int* sum_noatomic_tmp = malloc_shared<int>(NUM_THREADS*sizeof(int), q);
			q.memset(sum_noatomic_tmp,0,NUM_THREADS*sizeof(int));
			q.submit([&](handler& h) {
            h.parallel_for(sycl::range(NUM_THREADS), [=](sycl::id<1> index) {
				int global_id=index[0];
				int sum=0;
				for(int i=global_id;i<size;i+=NUM_THREADS)
					sum+=valueD[i];
				sum_noatomic_tmp[global_id]=sum;
                });
            }).wait();
			
			pre_sum[0]=0;
			for(int i=0;i<NUM_THREADS;i++)
				pre_sum[0]+=sum_noatomic_tmp[i];
            cout << "noatomic sum = " << pre_sum[0] << "\n";
        //no atomic test end

        free(sum, q);
        free(valueD, q);
		free(sum_noatomic_tmp,q);
        delete []value;      
        return 0;
}

5. Vtune性能分析

下面介紹oneapi常用的分析工具：Vtune和Advisor.
Vtune是用于檢測(cè)和優(yōu)化性能瓶頸的一款軟件工具，可用于CPU、GPU、FPGA系統(tǒng)的代碼測(cè)試。測(cè)試代碼是一段運(yùn)行單源最短路徑算法（簡稱SSSP）的代碼，使用的圖數(shù)據(jù)集規(guī)模為995節(jié)點(diǎn)、24087條邊。可執(zhí)行文件名稱為MAIN_TEST1，數(shù)據(jù)集名稱為email_Eu_core.bwcsr.啟動(dòng)vtune分析的命令如下：

$ vtune -collect io ../MAIN_TEST1 --input ../email_Eu_core.bwcsr --source 0 --type sssp

-collect選項(xiàng)表示收集的是CPU與外設(shè)IO過程的數(shù)據(jù)。程序正確執(zhí)行后，vtune會(huì)產(chǎn)生下面幾個(gè)方面的數(shù)據(jù)。

1. CPU 使用情況

如下，包括程序運(yùn)行使用的CPU時(shí)間、CPI、總線程數(shù)等信息。

CPU Time: 0.308s
Effective Time: 0.308s
Spin Time: 0s
Overhead Time: 0s
Instructions Retired: 921,800,000
CPI Rate: 0.919
Total Thread Count: 4
Paused Time: 0s

2. PCIe傳輸情況匯總

在觀察分析數(shù)據(jù)前，簡要介紹一下PCIe的情況：
PCI-Express(peripheral component interconnect express)是一種高速串行計(jì)算機(jī)擴(kuò)展總線標(biāo)準(zhǔn)，由英特爾在2001年提出。它基于點(diǎn)到點(diǎn)拓?fù)?，單?dú)的串行鏈路將每個(gè)設(shè)備連接到根系統(tǒng)（主機(jī)），在數(shù)據(jù)傳輸速度上做出了重大升級(jí)。

PCIe Traffic Summary
    Inbound PCIe Read, MB/sec: 18.511
        L3 Hit, %: 0.000
        L3 Miss, %: 100.000
         | A significant portion of inbound I/O read requests misses the L3
         | cache. To reduce inbound read latency and to avoid induced DRAM and
         | UPI traffic, make sure both the device and the memory it accesses
         | reside on the same socket and/or consider optimizations that localize
         | I/O data in the L3.
         |

上面這一段描述了L3緩存讀取數(shù)據(jù)的命中率和失誤率。上述情況是因?yàn)槭褂玫臄?shù)據(jù)集遠(yuǎn)小于平常測(cè)試的數(shù)據(jù)集。

Inbound PCIe Write, MB/sec: 7.341
    L3 Hit, %: 20.662
    L3 Miss, %: 79.338
     | A significant portion of inbound I/O write requests misses the L3
     | cache. To reduce inbound write latency and to avoid induced DRAM and
     | UPI traffic, make sure both the device and the memory it accesses
     | reside on the same socket and/or consider optimizations that localize
     | I/O data in the L3.
     |
    CPU/IO Conflicts, %: 0.000
    Average Latency, ns: 195.302

上面這一段描述了L3緩存寫入數(shù)據(jù)的命中率和失誤率。失誤率有所降低可能是因?yàn)樯厦孀x入時(shí)將部分?jǐn)?shù)據(jù)提前導(dǎo)入了三級(jí)緩存。

    Outbound PCIe Read, MB/sec: 0.301
     | Non-zero outbound read traffic caused by loads from memory-mapped I/O
     | devices may significantly limit system throughput. Explore MMIO accesses
     | to locate the code reading the memory of I/O devices through MMIO space.
     |
Outbound PCIe Write, MB/sec: 3.711

上面兩段描述了PCIe出棧讀取和寫入的速率。

3. 帶寬利用情況

| Bandwidth Domain                                                                                                             | Platform Maximum | Observed Maximum | Average | % of Elapsed Time with High BW Utilization(%)  |
|------------------------------------------------------------------------------------------------------------------------------|------------------|------------------|---------|------------------------------------------------|
| DRAM, GB/sec                                                                                                                 | 70               | 3.600            | 3.066   | 0.0%                                           |
| DRAM Single-Package, GB/sec                                                                                                  | 35               | 3.200            | 1.802   | 0.0%                                           |
| UPI Utilization Single-link, (%)  100                          3.200    2.475                                           0.0% | 100              | 3.200            | 2.475   | 0.0%                                           |
| PCIe Bandwidth, MB/sec                                                                                                       | 40               | 45.100           | 29.533  | 80.7%                                          |

如上表，四條總線上的帶寬利用情況，分別給出了它們的平臺(tái)理論最大值、觀測(cè)最大值、平均值和實(shí)耗時(shí)間占總帶寬利用率的比例。如圖，PCIe總線上的實(shí)耗時(shí)間占比最高，數(shù)據(jù)主要在這條總線上進(jìn)行傳輸。

4. 熱點(diǎn)函數(shù)

上圖是根據(jù)CPU占用時(shí)間列出的熱點(diǎn)函數(shù)排名，可見占用時(shí)間最高的是一個(gè)未知模塊，其次是一個(gè)函數(shù)，函數(shù)名是一串?dāng)?shù)字，如果想要看到自己寫的函數(shù)的名稱可以嘗試在編譯時(shí)加上-g選項(xiàng)。

除了上面所列，還有兩條核利用率相關(guān)的數(shù)據(jù)如下：
Effective Physical Core Utilization: 3.9% (0.781 out of 20)
較低的物理核利用率可能是由于負(fù)載不均、線程切換開銷、線程通信與同步開銷等。

Effective Logical Core Utilization: 2.1% (0.828 out of 40)
想要優(yōu)化邏輯核利用率，首先應(yīng)從優(yōu)化物理核利用率上考慮，提高上述利用率可以提升處理器吞吐量和多線程應(yīng)用的性能。

5. 其他

主要是一些平臺(tái)和硬件信息，如下：

Collection and Platform Info

如下是應(yīng)用在控制臺(tái)輸入的指令：

Application Command Line: ../MAIN_TEST1 "--input" "../email_Eu_core.bwcsr" "--source" "0" "--type" "sssp" 

如下是運(yùn)行時(shí)使用的操作系統(tǒng)名稱和版本。

Operating System: DISTRIB_DESCRIPTION="Ubuntu 22.04.1 LTS"
Result Size: 7.3 MB 

如下是數(shù)據(jù)收集開始和結(jié)束時(shí)間：

Collection start time: 14:11:23 12/11/2022 UTC
Collection stop time: 14:11:24 12/11/2022 UTC
Collector Type: Driverless Perf system-wide sampling
CPU
    Name: Intel(R) Xeon(R) Processor code named Cascadelake
    Frequency: 2.195 GHz
    Logical CPU Count: 40
    Max DRAM Single-Package Bandwidth: 35.000 GB/s
	如下是硬件的緩存情況：
    Cache Allocation Technology
        Level 2 capability: not detected
        Level 3 capability: available

以上只是Vtune的命令行測(cè)試輸出，一般用于遠(yuǎn)程測(cè)試程序。如果在本地使用Vtune測(cè)試，或者將遠(yuǎn)程的分析結(jié)果文件下載到本地用Vtune打開，就可以看到分析結(jié)果的圖形界面如下：（下圖是PCIe帶寬利用分布圖）

文章來源地址http://www.zghlxwxcb.cn/news/detail-789683.html

到了這里，關(guān)于Intel oneAPI——讓高性能計(jì)算觸手可及的文章就介紹完了。如果您還想了解更多內(nèi)容，請(qǐng)?jiān)谟疑辖撬阉鱐OY模板網(wǎng)以前的文章或繼續(xù)瀏覽下面的相關(guān)文章，希望大家以后多多支持TOY模板網(wǎng)！