DPC++是Data Parallel C++（數(shù)據(jù)并行C++）的首字母縮寫(xiě)，它是Intel為了將SYCL引入LLVM和oneAPI所開(kāi)發(fā)的開(kāi)源項(xiàng)目。SYCL是為了提高各種加速設(shè)備上的編程效率而開(kāi)發(fā)的一種高級(jí)別的編程模型，簡(jiǎn)單來(lái)說(shuō)它是一種跨平臺(tái)的抽象層，用戶(hù)不需要關(guān)心底層的加速器具體是什么，按照標(biāo)準(zhǔn)編寫(xiě)統(tǒng)一的代碼就可以在各種平臺(tái)上運(yùn)行?？梢哉f(shuō)SYCL大大提高了編寫(xiě)異構(gòu)計(jì)算代碼的可移植性和編程效率，已經(jīng)成為了異構(gòu)計(jì)算的行業(yè)標(biāo)準(zhǔn)。值得一提的是SYCL并不是由多個(gè)單詞的首字母的縮寫(xiě)。DPC++正是建立在SYCL和現(xiàn)代C++語(yǔ)言之上，具體來(lái)說(shuō)是建立在C++17標(biāo)準(zhǔn)之上的。
寫(xiě)本篇文章的目是為了討論現(xiàn)代C++語(yǔ)言在DPC++中的應(yīng)用，算是對(duì)《現(xiàn)代C++語(yǔ)言核心特性解析》一書(shū)的補(bǔ)充，而不是要探究異構(gòu)計(jì)算的原理，因?yàn)檫@是一個(gè)龐大的話(huà)題，需要資深專(zhuān)家才好駕馭。
關(guān)于實(shí)驗(yàn)環(huán)境，我選擇的是本地安裝Intel oneApi Toolkit，因?yàn)楸镜毓ぞ哂闷饋?lái)還是更加方便一些。不過(guò)，如果讀者朋友們的硬件條件不允許，那么我們可以注冊(cè)使用DevCloud。DevCloud是Intel公司提供的遠(yuǎn)程開(kāi)發(fā)環(huán)境，包含了最新的Intel 硬件和軟件集群。
我們提供的服務(wù)有：成都網(wǎng)站建設(shè)、做網(wǎng)站、微信公眾號(hào)開(kāi)發(fā)、網(wǎng)站優(yōu)化、網(wǎng)站認(rèn)證、唐河ssl等。為近千家企事業(yè)單位解決了網(wǎng)站和推廣的問(wèn)題。提供周到的售前咨詢(xún)和貼心的售后服務(wù)，是有科學(xué)管理、有技術(shù)的唐河網(wǎng)站制作公司
數(shù)據(jù)并行編程既可以被描述為一種思維方式，也可以被描述為一種編程方式。 數(shù)據(jù)由一組并行的處理單元進(jìn)行操作。 每個(gè)處理單元都是能夠?qū)?shù)據(jù)進(jìn)行計(jì)算的硬件設(shè)備。這些處理單元可能存在于單個(gè)設(shè)備上，也可能存在于我們計(jì)算機(jī)系統(tǒng)中的多個(gè)設(shè)備上。 我們可以指定代碼以?xún)?nèi)核的形式處理我們的數(shù)據(jù)。
內(nèi)核是數(shù)據(jù)并行編程中一個(gè)重要的概念，它的功能是讓設(shè)備上的處理單元執(zhí)行計(jì)算。這個(gè)術(shù)語(yǔ)在SYCL、OpenCL、CUDA 和 DPC++都有使用到。
異構(gòu)系統(tǒng)是包含多種類(lèi)型的計(jì)算設(shè)備的任何系統(tǒng)。 例如，同時(shí)具有CPU和GPU的系統(tǒng)就是異構(gòu)系統(tǒng)。現(xiàn)在已經(jīng)有很多中這樣的計(jì)算設(shè)備了，包括 CPU、GPU、FPGA、DSP、ASIC和AI 芯片。異構(gòu)系統(tǒng)的出現(xiàn)帶來(lái)了一個(gè)很大的挑戰(zhàn)，就是剛剛提到的這些設(shè)備，每一種都具有不同的架構(gòu)，也具有不同的特性，這就導(dǎo)致對(duì)每個(gè)設(shè)備有不同編程和優(yōu)化需求，而DPC++開(kāi)發(fā)一個(gè)動(dòng)機(jī)就是幫助解決這樣的挑戰(zhàn)。
因?yàn)楫悩?gòu)計(jì)算很重要，一直以來(lái)計(jì)算機(jī)架構(gòu)師致力于限制功耗、減少延遲和提高吞吐量的工作。從1990年到2006年，由于處理器性能每?jī)傻饺攴环ㄖ饕且驗(yàn)闀r(shí)鐘頻率每?jī)赡攴环?，?dǎo)致那個(gè)時(shí)候應(yīng)用程序的性能都跟著有所提升。這種情況在2006年左右結(jié)束，一個(gè)多核和多核處理器的新時(shí)代出現(xiàn)了。由于架構(gòu)向并行處理的轉(zhuǎn)變?yōu)槎嗳蝿?wù)系統(tǒng)帶來(lái)了性能提升，但是在不改變編程代碼的情況下，并沒(méi)有為大多數(shù)現(xiàn)有的單個(gè)應(yīng)用程序帶來(lái)性能提升。在這個(gè)新時(shí)代，GPU等加速器因?yàn)槟軌蚋咝У募铀賾?yīng)用程序變得比以往任何時(shí)候都流行。這催生了一個(gè)異構(gòu)計(jì)算時(shí)代，誕生了大量的具有自己的專(zhuān)業(yè)處理能力的加速器以及許多不同的編程模型。它們通過(guò)更加專(zhuān)業(yè)化的加速器設(shè)計(jì)可以在特定問(wèn)題上提供更高性能的計(jì)算，因?yàn)樗鼈儾槐厝ヌ幚硭袉?wèn)題。這是一個(gè)經(jīng)典的計(jì)算機(jī)架構(gòu)權(quán)衡。它通常意味著加速器只能支持為處理器設(shè)計(jì)的編程語(yǔ)言的子集。事實(shí)上，在DPC++中，只有在內(nèi)核中編寫(xiě)的代碼才能在加速器中運(yùn)行。
加速器架構(gòu)可以分為幾大類(lèi)，這些類(lèi)別會(huì)影響我們對(duì)編程模型、算法以及如何高效使用加速器的決策。例如，CPU是通用代碼的最佳選擇，包括標(biāo)量和決策代碼，并且通常內(nèi)置向量加速器。GPU則是尋求加速向量和密切相關(guān)的張量。DSP尋求是以低延遲加速特定數(shù)學(xué)運(yùn)算，通常用于處理手機(jī)的模擬信號(hào)等。AI加速器通常用于加速矩陣運(yùn)算，盡管有些加速器也可能加速圖。FPGA和ASIC特別適用于加速計(jì)算空間問(wèn)題。
一方面因?yàn)镈PC++具有可移植性、高級(jí)性和非專(zhuān)有性，同時(shí)滿(mǎn)足現(xiàn)代異構(gòu)計(jì)算機(jī)體系結(jié)構(gòu)的要求。另一方面，它可以讓跨主機(jī)和計(jì)算設(shè)備的代碼使用相同的編程環(huán)境，即現(xiàn)代C++的編程環(huán)境。最后，計(jì)算機(jī)體系結(jié)構(gòu)的未來(lái)包括跨越標(biāo)量、向量、矩陣和空間 (SVMS) 操作的加速器，需要對(duì)包括 SVMS 功能在內(nèi)的異構(gòu)機(jī)器的支持。并且這種支持應(yīng)該涵蓋高度復(fù)雜的可編程設(shè)備，以及可編程性較低的固定功能或?qū)Ｓ玫脑O(shè)備。
在開(kāi)始討論現(xiàn)代C++語(yǔ)言在DPC++中的應(yīng)用之前，讓我們先看一遍完整的代碼，順便測(cè)試我們的實(shí)驗(yàn)環(huán)境：
重慶分公司
重慶分公司
#include
constexpr int N = 16;
using namespace sycl;
class IntelGPUSelector : public device_selector {
public:
int operator()(const device& Device) const override {
const std::string DeviceName = Device.get_info();
const std::string DeviceVendor = Device.get_info();
return Device.is_gpu() && (DeviceName.find("Intel") != std::string::npos) ? 100 : 0;
}
};
int main() {
IntelGPUSelector d;
queue q(d);
int* data = malloc_shared(N, q);
q.parallel_for(N, [=](auto i) {
data[i] = i;
}).wait();
for (int i = 0; i < N; i++) std::cout << data[i] << " ";
free(data, q);
}
編譯運(yùn)行上面的代碼，如果沒(méi)有問(wèn)題應(yīng)該輸出：
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
簡(jiǎn)單解釋一下這段代碼，sycl是DPC++的實(shí)體的命名空間，用using namespace sycl;打開(kāi)命名空間可以簡(jiǎn)化后續(xù)代碼。IntelGPUSelector是一個(gè)繼承了device_selector的設(shè)備選擇器，其中device_selector是純虛類(lèi)，它有個(gè)純虛函數(shù)int operator()(const device& Device) const需要派生類(lèi)來(lái)實(shí)現(xiàn)，該函數(shù)會(huì)遍歷計(jì)算機(jī)上的計(jì)算設(shè)備，并且返回使用設(shè)備的優(yōu)先級(jí)，返回?cái)?shù)字越高優(yōu)先級(jí)越高，這里選擇Intel的GPU作為首選的計(jì)算設(shè)備，注意這個(gè)函數(shù)使用了override來(lái)說(shuō)明其目的是覆蓋虛函數(shù)。queue的目的是指定工作的目標(biāo)位置，這里設(shè)置的是Intel的GPU。函數(shù)模板malloc_shared分配了可在設(shè)備上使用的工作內(nèi)存。成員函數(shù)parallel_for執(zhí)行并行計(jì)算。值得注意的是free調(diào)用的是sycl::free而不是C運(yùn)行時(shí)庫(kù)的free。在這段代碼中，比較明顯使用了現(xiàn)在C++語(yǔ)法的地方是函數(shù)parallel_for的實(shí)參，
[=](auto i) { data[i] = i; }
這是一個(gè)lambda表達(dá)式。
Ⅳ DPC++和lambda表達(dá)式
如果要選出一個(gè)對(duì)DPC++最重要的現(xiàn)代C++語(yǔ)言特性，我覺(jué)得lambda表達(dá)式應(yīng)該可以被選上。因?yàn)樵贒PC++的代碼中，內(nèi)核代碼一般都是以lambda表達(dá)式的形式出現(xiàn)。比如上面的例子就是將lambda表達(dá)式作為對(duì)象傳入到Intel的GPU設(shè)備上然后進(jìn)行計(jì)算的。在這個(gè)lambda表達(dá)式中，[=]是捕獲列表，它可以捕獲當(dāng)前定義作用域內(nèi)的變量的值，這也是它可以在函數(shù)體內(nèi)使用data[i]的原因。捕獲列表[=]之后的是形參列表(auto i)，注意這里的形參類(lèi)型使用的是auto占位符，也就是說(shuō)，我們將形參類(lèi)型的確認(rèn)工作交給了編譯器。我們一般稱(chēng)這種lambda表達(dá)式為泛型lambda表達(dá)式。當(dāng)然，如果在編譯時(shí)選擇C++20標(biāo)準(zhǔn)，我們還可以將其改為模板語(yǔ)法的泛型lambda表達(dá)式：
[=](T i) { data[i] = i; }
lambda表達(dá)式的捕獲列表功能非常強(qiáng)大，除了捕獲值以外，還可以捕獲引用，例如：
[&](auto i) { data[i] = i; }
以上代碼會(huì)捕獲當(dāng)前定義作用域內(nèi)的變量的引用，不過(guò)值得注意的是，由于這里的代碼會(huì)交給加速核心運(yùn)行，捕獲引用并不是一個(gè)正確的做法，會(huì)導(dǎo)致編譯出錯(cuò)。另外一般來(lái)說(shuō)，我們并不推薦直接捕獲所有可捕獲的對(duì)象，而是有選擇的捕獲，例如：
[data](auto i) { data[i] = i; }
當(dāng)然，除了使用lambda表達(dá)式，我們也可以選擇其他形式的代碼來(lái)運(yùn)行設(shè)備，比如使用仿函數(shù)：
struct AssginTest {
void operator()(auto i) const { data_[i] = i; }
int* data_;
};
AssginTest functor{data};
q.parallel_for(N, functor).wait();
但是很明顯，這種方法沒(méi)有使用lambda表達(dá)式來(lái)的簡(jiǎn)單直接。
之所以能夠讓parallel_for這么靈活的接受各種形式的實(shí)參，是因?yàn)?code>parallel_for本身是一個(gè)成員函數(shù)模板：
template
event parallel_for(range<1> NumWorkItems,
_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
_CODELOCARG(&CodeLoc);
return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc);
}
其中KernelFunc就是傳入的lambda表達(dá)式或者仿函數(shù)，KernelType是KernelFunc的類(lèi)型。
如果從這里的代碼一路運(yùn)行跟蹤下去，會(huì)發(fā)現(xiàn)它們都是用模板傳遞實(shí)參類(lèi)型，直到submit_impl：
sycld.dll!cl::sycl::queue::submit_impl
dpcpp.exe!cl::sycl::queue::submit
dpcpp.exe!cl::sycl::queue::parallel_for_impl
dpcpp.exe!cl::sycl::queue::parallel_for
這是因?yàn)閟ycld.dll是一個(gè)二進(jìn)制模塊，它無(wú)法以模板的形式提供代碼，所有的類(lèi)型必須確定下來(lái)，為了解決這個(gè)問(wèn)題，cl::sycl::queue::submit_impl使用了std::function：
event submit_impl(function_class CGH,
const detail::code_location &CodeLoc);
函數(shù)模板cl::sycl::queue::parallel_for_impl將KernelFunc封裝到另外一個(gè)lambda表達(dá)式對(duì)象中，并且通過(guò)function_class來(lái)傳遞整個(gè)lambda表達(dá)式：
template
event parallel_for_impl(
range NumWorkItems, KernelType KernelFunc,
const detail::code_location &CodeLoc = detail::code_location::current()) {
return submit(
[&](handler &CGH) {
CGH.template parallel_for(NumWorkItems,
KernelFunc);
},
CodeLoc);
}
其中function_class就是std::function。注意這里CGH.template parallel_for需要說(shuō)明符template否則尖括號(hào)會(huì)解析出錯(cuò)。DPC++通過(guò)這樣一系列的操作，最大限度的保留了用戶(hù)編程的靈活性。
DPC++代碼中大量的運(yùn)用了C++17標(biāo)準(zhǔn)才引入的模板推導(dǎo)特性，關(guān)于這些特性我們還是從一個(gè)DPC++的小例子開(kāi)始：
int main() {
IntelGPUSelector d;
queue q(d);
std::vector v1(N);
std::array v2;
{
buffer buf1(v1);
buffer buf2(v2);
q.submit([&](handler& h) {
accessor a1(buf1, h, write_only);
accessor a2(buf2, h, write_only);
h.parallel_for(N, [=](auto i) {
a1[i] = i;
a2[i] = i;
});
});
}
for (int i = 0; i < N; i++) std::cout << v1[i] << v2[i] << " ";
}
這段代碼沒(méi)有使用malloc_shared分配內(nèi)存，取而代之的是使用buffer和accessor，其中buffer用于封裝數(shù)據(jù)，accessor用于訪(fǎng)問(wèn)數(shù)據(jù)。這里以buffer為例解析DPC++對(duì)模板推導(dǎo)的使用。
首先觀察buffer的兩個(gè)實(shí)例，它們的構(gòu)造函數(shù)的實(shí)參分別是std::vector和std::array類(lèi)型。之所以能夠這樣調(diào)用構(gòu)造函數(shù)，并不是因?yàn)?code>buffer為這兩個(gè)類(lèi)型重載了它的構(gòu)造函數(shù)，而是因?yàn)槠錁?gòu)造函數(shù)使用了模板。這里涉及到一個(gè)C++17標(biāo)準(zhǔn)新特性——類(lèi)模板的模板實(shí)參推導(dǎo)。在以往，類(lèi)模板的實(shí)例化必須是顯式傳入模板實(shí)參，否則會(huì)造成編譯出錯(cuò)。在新的標(biāo)準(zhǔn)中，類(lèi)模板的模板實(shí)參已經(jīng)可以根據(jù)構(gòu)造函數(shù)來(lái)推導(dǎo)了。來(lái)看一下buffer的構(gòu)造函數(shù)：
template 0) &&
(dimensions <= 3)>>
class buffer {
public:
...
template ,
typename = EnableIfContiguous>
buffer(Container &container, AllocatorT allocator,
const property_list &propList = {})
: Range(range<1>(container.size())) {
impl = std::make_shared(
container.data(), get_count() * sizeof(T),
detail::getNextPowerOfTwo(sizeof(T)), propList,
make_unique_ptr>(
allocator));
}
template ,
typename = EnableIfContiguous>
buffer(Container &container, const property_list &propList = {})
: buffer(container, {}, propList) {}
...
};
代碼buffer buf1(v1);會(huì)執(zhí)行
buffer(Container &container, const property_list &propList = {})
這條構(gòu)造函數(shù)，值得注意的是該構(gòu)造函數(shù)并沒(méi)有實(shí)際的實(shí)現(xiàn)代碼，而是通過(guò)委托構(gòu)造函數(shù)的方法調(diào)用了
buffer(Container &container, AllocatorT allocator, const property_list &propList = {})
委托構(gòu)造函數(shù)是C++11引入的特性，它可以讓某個(gè)構(gòu)造函數(shù)將構(gòu)造的執(zhí)行權(quán)交給另外的構(gòu)造函數(shù)?；氐侥０逋茖?dǎo)，這里通過(guò)構(gòu)造函數(shù)會(huì)推導(dǎo)出Container是std::vector的推導(dǎo)結(jié)果是1，而后面兩個(gè)模板參數(shù)是用來(lái)檢查前兩個(gè)模板參數(shù)是否正確的，這里大量的使用了模板元編程的技巧：
template
using EnableIfOneDimension = typename detail::enable_if_t<1 == dims>;
template
using EnableIfContiguous =
detail::void_t().data())> (*)[],
const T (*)[]>::value>,
decltype(std::declval().size())>;
首先它們都是使用using定義的別名模板，它們的目的分別是檢查dims是否為1和Container是否為連續(xù)的。第一個(gè)別名模板很簡(jiǎn)單，直接檢查dims是否為1，detail::enable_if_t就是std::enable_if_t。第二個(gè)檢查連續(xù)性的方法稍微麻煩一些，簡(jiǎn)單來(lái)說(shuō)就是檢查容器對(duì)象的成員函數(shù)data()返回值的類(lèi)型的數(shù)組指針是否能和const T (*)[]轉(zhuǎn)換，這里主要檢查兩點(diǎn)，第一容器具有data()成員函數(shù)，第二返回類(lèi)型的指針和T const T (*)[]轉(zhuǎn)換。事實(shí)上，在標(biāo)準(zhǔn)容器中，只有連續(xù)容器有data()成員函數(shù)，其他的都會(huì)因?yàn)闆](méi)有data()而報(bào)錯(cuò)，例如：
no member named 'data' in 'std::list'
仔細(xì)閱讀上面代碼的朋友應(yīng)該會(huì)發(fā)現(xiàn)另外一個(gè)問(wèn)題，那就是沒(méi)有任何地方可以幫助編譯器推導(dǎo)出buffer的類(lèi)模板形參T。這就不得不說(shuō)DPC++將C++17關(guān)于模板推導(dǎo)的新特性用的淋漓盡致了。實(shí)際上在代碼中，有這樣一句用戶(hù)自定義推導(dǎo)指引的代碼：
template
buffer(Container &, const property_list & = {})
->buffer;
用戶(hù)自定義推導(dǎo)指引是指程序員可以指導(dǎo)編譯器如何通過(guò)函數(shù)實(shí)參推導(dǎo)模板形參的類(lèi)型。最后在這個(gè)例子中，需要注意一下，buffer在析構(gòu)的時(shí)候才會(huì)將緩存的數(shù)據(jù)寫(xiě)到v1和v2，所以這里用了單獨(dú)的作用域。
~buffer_impl() {
try {
BaseT::updateHostMemory();
} catch (...) {
}
}
本篇文章從幾個(gè)簡(jiǎn)單的DPC++的例子展開(kāi)，逐步探究了DPC++對(duì)于現(xiàn)代C++語(yǔ)言特性的運(yùn)用，其中比較重要的包括lambda表達(dá)式、泛型和模板推導(dǎo)，當(dāng)然DPC++運(yùn)用的新特性遠(yuǎn)不止這些。從另一方面來(lái)看，這些新特性的加入確實(shí)的幫助DPC++完成了過(guò)去無(wú)法完成的工作，這也是近幾年C++的發(fā)展趨勢(shì)，越來(lái)越多的代碼庫(kù)開(kāi)始引入新的特性，并且有一些非常”神奇“的代碼也孕育而生。DPC++就是其中之一，光是閱讀DPC++中使用新特性的代碼就已經(jīng)足夠讓人拍案叫絕了，更何況還有代碼的組織架構(gòu)、底層的抽象等等。我知道，單單一篇文章并不能討論清楚DPC++中現(xiàn)代C++語(yǔ)言的特性，所以王婆賣(mài)瓜的推薦自己寫(xiě)的書(shū)《現(xiàn)代C++語(yǔ)言核心特性解析》和盛格塾課程《現(xiàn)代C++42講》，相信看完這本書(shū)或者經(jīng)過(guò)課程訓(xùn)練后朋友們會(huì)對(duì)現(xiàn)代C++語(yǔ)言的特性有一個(gè)比較深入的理解。
1.DPC++ Part 1: An Introduction to the New Programming Model [https://simplecore-ger.intel.com/techdecoded/wp-content/uploads/sites/11/Webinar-Slides-DPC-Part-1-An-Introduction-to-the-New-Programming-Model-.pdf]
2.Data Parallel C++: Mastering DPC++ for Programming of Heterogeneous Systems Using C++ and SYCL preview [https://resource-cms.springernature.com/springer-cms/rest/v1/content//data/v1]
3.Intel? DevCloud [https://software.intel.com/en-us/devcloud/oneapi]
4.New, Open DPC++ Extensions Complement SYCL and C++ [https://insidehpc.com/2020/06/new-open-dpc-extensions-complement-sycl-and-c/]
