Ⅰ DPC++簡介
DPC++是Data Parallel C++（資料並行C++）的首字母縮寫，它是Intel為了將SYCL引入LLVM和oneAPI所開發的開源專案。SYCL是為了提高各種加速裝置上的程式設計效率而開發的一種高階別的程式設計模型，簡單來說它是一種跨平臺的抽象層，使用者不需要關心底層的加速器具體是什麼，按照標準編寫統一的程式碼就可以在各種平臺上執行。可以說SYCL大大提高了編寫異構計算程式碼的可移植性和程式設計效率，已經成為了異構計算的行業標準。值得一提的是SYCL並不是由多個單詞的首字母的縮寫。DPC++正是建立在SYCL和現代C++語言之上，具體來說是建立在C++17標準之上的。
寫本篇文章的目是為了討論現代C++語言在DPC++中的應用，算是對《現代C++語言核心特性解析》一書的補充，而不是要探究異構計算的原理，因為這是一個龐大的話題，需要資深專家才好駕馭。
關於實驗環境，我選擇的是本地安裝Intel oneApi Toolkit，因為本地工具用起來還是更加方便一些。不過，如果讀者朋友們的硬體條件不允許，那麼我們可以註冊使用DevCloud。DevCloud是Intel公司提供的遠端開發環境，包含了最新的Intel 硬體和軟體叢集。
Ⅱ DPC++背景
1．什麼是資料並行程式設計
資料並行程式設計既可以被描述為一種思維方式，也可以被描述為一種程式設計方式。 資料由一組並行的處理單元進行操作。 每個處理單元都是能夠對資料進行計算的硬體裝置。這些處理單元可能存在於單個裝置上，也可能存在於我們計算機系統中的多個裝置上。 我們可以指定程式碼以核心的形式處理我們的資料。
核心是資料並行程式設計中一個重要的概念，它的功能是讓裝置上的處理單元執行計算。這個術語在SYCL、OpenCL、CUDA 和 DPC++都有使用到。
2．什麼是異構系統
異構系統是包含多種型別的計算裝置的任何系統。 例如，同時具有CPU和GPU的系統就是異構系統。現在已經有很多中這樣的計算裝置了，包括 CPU、GPU、FPGA、DSP、ASIC和AI 晶片。異構系統的出現帶來了一個很大的挑戰，就是剛剛提到的這些裝置，每一種都具有不同的架構，也具有不同的特性，這就導致對每個裝置有不同程式設計和優化需求，而DPC++開發一個動機就是幫助解決這樣的挑戰。
3．為什麼需要異構系統
因為異構計算很重要，一直以來計算機架構師致力於限制功耗、減少延遲和提高吞吐量的工作。從1990年到2006年，由於處理器效能每兩到三年翻一番（主要是因為時脈頻率每兩年翻一番），導致那個時候應用程式的效能都跟著有所提升。這種情況在2006年左右結束，一個多核和多核處理器的新時代出現了。由於架構向並行處理的轉變為多工系統帶來了效能提升，但是在不改變程式設計程式碼的情況下，並沒有為大多數現有的單個應用程式帶來效能提升。在這個新時代，GPU等加速器因為能夠更高效的加速應用程式變得比以往任何時候都流行。這催生了一個異構計算時代，誕生了大量的具有自己的專業處理能力的加速器以及許多不同的程式設計模型。它們通過更加專業化的加速器設計可以在特定問題上提供更高效能的計算，因為它們不必去處理所有問題。這是一個經典的計算機架構權衡。它通常意味著加速器只能支援為處理器設計的程式語言的子集。事實上，在DPC++中，只有在核心中編寫的程式碼才能在加速器中執行。
加速器架構可以分為幾大類，這些類別會影響我們對程式設計模型、演算法以及如何高效使用加速器的決策。例如，CPU是通用程式碼的最佳選擇，包括標量和決策程式碼，並且通常內建向量加速器。GPU則是尋求加速向量和密切相關的張量。DSP尋求是以低延遲加速特定數學運算，通常用於處理手機的模擬訊號等。AI加速器通常用於加速矩陣運算，儘管有些加速器也可能加速圖。FPGA和ASIC特別適用於加速計算空間問題。
4．為什麼使用DPC++
一方面因為DPC++具有可移植性、高階性和非專有性，同時滿足現代異構計算機體系結構的要求。另一方面，它可以讓跨主機和計算裝置的程式碼使用相同的程式設計環境，即現代C++的程式設計環境。最後，計算機體系結構的未來包括跨越標量、向量、矩陣和空間 (SVMS) 操作的加速器，需要對包括 SVMS 功能在內的異構機器的支援。並且這種支援應該涵蓋高度複雜的可程式設計裝置，以及可程式設計性較低的固定功能或專用的裝置。
Ⅲ 初探DPC++
在開始討論現代C++語言在DPC++中的應用之前，讓我們先看一遍完整的程式碼，順便測試我們的實驗環境：
#include <CL/sycl.hpp>
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<info::device::name>();
const std::string DeviceVendor = Device.get_info<info::device::vendor>();
return Device.is_gpu() && (DeviceName.find("Intel") != std::string::npos) ? 100 : 0;
}
};
int main() {
IntelGPUSelector d;
queue q(d);
int* data = malloc_shared<int>(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);
}
編譯執行上面的程式碼，如果沒有問題應該輸出：
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
簡單解釋一下這段程式碼，sycl是DPC++的實體的名稱空間，用using namespace sycl;開啟名稱空間可以簡化後續程式碼。IntelGPUSelector是一個繼承了device_selector的裝置選擇器，其中device_selector是純虛類，它有個純虛擬函式int operator()(const device& Device) const需要派生類來實現，該函式會遍歷計算機上的計算裝置，並且返回使用裝置的優先順序，返回數字越高優先順序越高，這裡選擇Intel的GPU作為首選的計算裝置，注意這個函式使用了override來說明其目的是覆蓋虛擬函式。queue的目的是指定工作的目標位置，這裡設定的是Intel的GPU。函式模板malloc_shared分配了可在裝置上使用的工作記憶體。成員函式parallel_for執行平行計算。值得注意的是free呼叫的是sycl::free而不是C執行時庫的free。在這段程式碼中，比較明顯使用了現在C++語法的地方是函式parallel_for的實參，
[=](auto i) { data[i] = i; }
這是一個lambda表示式。
Ⅳ DPC++和lambda表示式
如果要選出一個對DPC++最重要的現代C++語言特性，我覺得lambda表示式應該可以被選上。因為在DPC++的程式碼中，核心程式碼一般都是以lambda表示式的形式出現。比如上面的例子就是將lambda表示式作為物件傳入到Intel的GPU裝置上然後進行計算的。在這個lambda表示式中，[=]是捕獲列表，它可以捕獲當前定義作用域內的變數的值，這也是它可以在函式體內使用data[i]的原因。捕獲列表[=]之後的是形參列表(auto i)，注意這裡的形參型別使用的是auto佔位符，也就是說，我們將形參型別的確認工作交給了編譯器。我們一般稱這種lambda表示式為泛型lambda表示式。當然，如果在編譯時選擇C++20標準，我們還可以將其改為模板語法的泛型lambda表示式：
[=]<typename T>(T i) { data[i] = i; }
lambda表示式的捕獲列表功能非常強大，除了捕獲值以外，還可以捕獲引用，例如：
[&](auto i) { data[i] = i; }
以上程式碼會捕獲當前定義作用域內的變數的引用，不過值得注意的是，由於這裡的程式碼會交給加速核心執行，捕獲引用並不是一個正確的做法，會導致編譯出錯。另外一般來說，我們並不推薦直接捕獲所有可捕獲的物件，而是有選擇的捕獲，例如：
[data](auto i) { data[i] = i; }
當然，除了使用lambda表示式，我們也可以選擇其他形式的程式碼來執行裝置，比如使用仿函式：
struct AssginTest {
void operator()(auto i) const { data_[i] = i; }
int* data_;
};
AssginTest functor{data};
q.parallel_for(N, functor).wait();
但是很明顯，這種方法沒有使用lambda表示式來的簡單直接。
Ⅴ DPC++和泛型能力
之所以能夠讓parallel_for這麼靈活的接受各種形式的實參，是因為parallel_for本身是一個成員函式模板：
template <typename KernelName = detail::auto_name, typename KernelType>
event parallel_for(range<1> NumWorkItems,
_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
_CODELOCARG(&CodeLoc);
return parallel_for_impl<KernelName>(NumWorkItems, KernelFunc, CodeLoc);
}
其中KernelFunc就是傳入的lambda表示式或者仿函式，KernelType是KernelFunc的型別。
如果從這裡的程式碼一路執行跟蹤下去，會發現它們都是用模板傳遞實參型別，直到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
這是因為sycld.dll是一個二進位制模組，它無法以模板的形式提供程式碼，所有的型別必須確定下來，為了解決這個問題，cl::sycl::queue::submit_impl使用了std::function：
event submit_impl(function_class<void(handler &)> CGH,
const detail::code_location &CodeLoc);
函式模板cl::sycl::queue::parallel_for_impl將KernelFunc封裝到另外一個lambda表示式物件中，並且通過function_class<void(handler &)>來傳遞整個lambda表示式：
template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
event parallel_for_impl(
range<Dims> NumWorkItems, KernelType KernelFunc,
const detail::code_location &CodeLoc = detail::code_location::current()) {
return submit(
[&](handler &CGH) {
CGH.template parallel_for<KernelName, KernelType>(NumWorkItems,
KernelFunc);
},
CodeLoc);
}
其中function_class就是std::function。注意這裡CGH.template parallel_for需要說明符template否則尖括號會解析出錯。DPC++通過這樣一系列的操作，最大限度的保留了使用者程式設計的靈活性。
Ⅵ DPC++和模板推導
DPC++程式碼中大量的運用了C++17標準才引入的模板推導特性，關於這些特性我們還是從一個DPC++的小例子開始：
int main() {
IntelGPUSelector d;
queue q(d);
std::vector<int> v1(N);
std::array<int, N> 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] << " ";
}
這段程式碼沒有使用malloc_shared分配記憶體，取而代之的是使用buffer和accessor，其中buffer用於封裝資料，accessor用於訪問資料。這裡以buffer為例解析DPC++對模板推導的使用。
首先觀察buffer的兩個例項，它們的建構函式的實參分別是std::vector<int>和std::array<int, N>型別。之所以能夠這樣呼叫建構函式，並不是因為buffer為這兩個型別過載了它的建構函式，而是因為其建構函式使用了模板。這裡涉及到一個C++17標準新特性——類别範本的模板實參推導。在以往，類别範本的例項化必須是顯式傳入模板實參，否則會造成編譯出錯。在新的標準中，類别範本的模板實參已經可以根據建構函式來推導了。來看一下buffer的建構函式：
template <typename T, int dimensions = 1,
typename AllocatorT = cl::sycl::buffer_allocator,
typename = typename detail::enable_if_t<(dimensions > 0) &&
(dimensions <= 3)>>
class buffer {
public:
...
template <class Container, int N = dimensions,
typename = EnableIfOneDimension<N>,
typename = EnableIfContiguous<Container>>
buffer(Container &container, AllocatorT allocator,
const property_list &propList = {})
: Range(range<1>(container.size())) {
impl = std::make_shared<detail::buffer_impl>(
container.data(), get_count() * sizeof(T),
detail::getNextPowerOfTwo(sizeof(T)), propList,
make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT>>(
allocator));
}
template <class Container, int N = dimensions,
typename = EnableIfOneDimension<N>,
typename = EnableIfContiguous<Container>>
buffer(Container &container, const property_list &propList = {})
: buffer(container, {}, propList) {}
...
};
程式碼buffer buf1(v1);會執行
buffer(Container &container, const property_list &propList = {})
這條建構函式，值得注意的是該建構函式並沒有實際的實現程式碼，而是通過委託建構函式的方法呼叫了
buffer(Container &container, AllocatorT allocator, const property_list &propList = {})
委託建構函式是C++11引入的特性，它可以讓某個建構函式將構造的執行權交給另外的建構函式。回到模板推導，這裡通過建構函式會推匯出Container是std::vector<int>，dimensions的推導結果是1，而後面兩個模板引數是用來檢查前兩個模板引數是否正確的，這裡大量的使用了模板超程式設計的技巧：
template <int dims>
using EnableIfOneDimension = typename detail::enable_if_t<1 == dims>;
template <class Container>
using EnableIfContiguous =
detail::void_t<detail::enable_if_t<std::is_convertible<
detail::remove_pointer_t<decltype(
std::declval<Container>().data())> (*)[],
const T (*)[]>::value>,
decltype(std::declval<Container>().size())>;
首先它們都是使用using定義的別名模板，它們的目的分別是檢查dims是否為1和Container是否為連續的。第一個別名模板很簡單，直接檢查dims是否為1，detail::enable_if_t就是std::enable_if_t。第二個檢查連續性的方法稍微麻煩一些，簡單來說就是檢查容器物件的成員函式data()返回值的型別的陣列指標是否能和const T (*)[]轉換，這裡主要檢查兩點，第一容器具有data()成員函式，第二返回型別的指標和T const T (*)[]轉換。事實上，在標準容器中，只有連續容器有data()成員函式，其他的都會因為沒有data()而報錯，例如：
no member named 'data' in 'std::list<int>'
仔細閱讀上面程式碼的朋友應該會發現另外一個問題，那就是沒有任何地方可以幫助編譯器推匯出buffer的類别範本形參T。這就不得不說DPC++將C++17關於模板推導的新特性用的淋漓盡致了。實際上在程式碼中，有這樣一句使用者自定義推導指引的程式碼：
template <class Container>
buffer(Container &, const property_list & = {})
->buffer<typename Container::value_type, 1>;
使用者自定義推導指引是指程式設計師可以指導編譯器如何通過函式實參推導模板形參的型別。最後在這個例子中，需要注意一下，buffer在析構的時候才會將快取的資料寫到v1和v2，所以這裡用了單獨的作用域。
~buffer_impl() {
try {
BaseT::updateHostMemory();
} catch (...) {
}
}
Ⅶ 總結
本篇文章從幾個簡單的DPC++的例子展開，逐步探究了DPC++對於現代C++語言特性的運用，其中比較重要的包括lambda表示式、泛型和模板推導，當然DPC++運用的新特性遠不止這些。從另一方面來看，這些新特性的加入確實的幫助DPC++完成了過去無法完成的工作，這也是近幾年C++的發展趨勢，越來越多的程式碼庫開始引入新的特性，並且有一些非常”神奇“的程式碼也孕育而生。DPC++就是其中之一，光是閱讀DPC++中使用新特性的程式碼就已經足夠讓人拍案叫絕了，更何況還有程式碼的組織架構、底層的抽象等等。我知道，單單一篇文章並不能討論清楚DPC++中現代C++語言的特性，所以王婆賣瓜的推薦自己寫的書《現代C++語言核心特性解析》和盛格塾課程《現代C++42講》，相信看完這本書或者經過課程訓練後朋友們會對現代C++語言的特性有一個比較深入的理解。
參考文獻
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/17382710/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/]