OneAPI DPC++ (SYCL-based) 簡介
Intel 前一陣子開始都有再釋出一些 Intel GPU 的消息，如 DG1 等。但如何使用他們呢? 有別於 NVIDIA 的 CUDA 以及 AMD 的 HIP，他們是鑑於 SYCL 這套語言往上加一些額外功能，而這些功能有些就會回流到 SYCL 上面 (例如 SYCL2020 中就有幾項被加過去，例如 subgroup 等)
SYCL
SYCL 是一個可跨平台的語言，一開始應該更多是為了 OpenCL 發展，在一些地方都還看得到 OpenCL 的影子，又或者目前的 Backend 大多都是先有 OpenCL 版本，但現在已經不限於 OpenCL，例如 Intel 他們的 Level_Zero 。SYCL 基於 C++ 之上，在他們的教學或說明中，可以看到利用 lambda 或 template的情形。但背後還得靠各廠商基於他們的硬體，來讓編譯器來編出他們所能用的；但使用者方面可以只需寫一份然後編譯出能跑在各種硬體上的程式 (理想上)
來源: SYCL Overview - The Khronos Group Inc
這圖可以看到 CUDA/HIP 正在開發中，以及要介紹的 OneAPI
OneAPI
隨著 Intel 公布的顯示卡計畫，也公布了開發者要怎麼使用他們的硬體的方式，統稱為 OneAPI，裡頭有著編譯器 DPC++、數學函式庫 OneMKL、又或者平行函數OneDPL 等。當然跟之前 AMD HIP 很像，由於 CUDA 的開發者眾多，所以 Intel 也有一款 DPCT 幫助從 CUDA 變成 DPC++ 的程式，來吸引開發者又或者降低跨過來的門檻。這篇主要是簡介 DPC++，DPCT 轉 CUDA 的部分所會遇到的問題等會在另一篇介紹。
DPC++
由於 DPC++ 是基於 SYCL 的，那跟 CUDA/HIP 的預設觀點就不太一樣。 SYCL 是利用 buffer 來管理記憶體，然後程式利用各個函數對於資料的使用順序，自動安排執行過程，類似於 OpenMP 的 task depend，這跟 CUDA/HIP 預設是用 default stream 的流程不太一樣，CUDA/HIP 執行完一個函數才會在執行另一個，如果需要同時執行或者決定之間的順序的話，則需要用不同 stream 或者 graph/event 達成。
OneAPI 可以讓你在創建 queue (所有東西都是由 queue 去執行，類似於 stream) 時，指定 in_order，這樣在執行上，就會執行完一個再執行第二個，更接近 CUDA/HIP；還有可以不用 buffer 的方式， malloc_shared (Unified shared memory) 以及 malloc_device (device only)，如果不是很在意效能的話，建議使用 malloc_shared，因為有蠻多 OneAPI 目前還只支援 Unified shared memory。
CUDA/HIP: dim3 - SYCL: range
另外一個差異蠻大的點是在順序，CUDA/HIP 的 dim3 順序是 x, y, z ，對應到 OneAPI 又或者說是 SYCL， range(0, 1, 2) 並非 0 <-> x 這樣，而是 range(z, y, x) ，這跟執行的安排有關。 CUDA 中 dim3(x) -> dim3(x, y) 時，threads 的順序還是沿著 x 方向，也就是說 thread0->x0y0, thread1->x1y0, ..., thread32->x0y1 這樣，但在 SYCL range(r1)->range(r1, r2) 時，threads (SYCL 稱之 work items) 的順序就會先沿著 r2 而非 r1 了 ，更像多層 for 迴圈，第一層 r1 第二層 r2 這樣，thread0->r1(0)r2(0), thread1->r1(0)r2(1), ..., thread32->r1(1)r2(0) 因此，如果要符合 CUDA 的順序的話，無論原先是只有 x 或者 x, y, z ，都會弄成 3d range(z, y, x) 的形式
CUDA/HIP: block/grid - SYCL: workgoup/nd_range
CUDA 執行時是用cuda_kernel<<<grid, block>>>(...) 而 SYCL 則是用 queue->parallel_for(nd_range(global_size, local_size), lambda-function...) ，如果要用CUDA 形式來對應的話 local_size = range<3>(block.z, block.y, block.x), num_group = range<3>(grid.z, grid.y, grid.x), global_size = local_size * num_group ，local_size 也就是指定 workgroup 大小的，那跟前面所講的一樣要用三維倒序的方式，但 global_size 並不直接等於 grid，要用 grid 倒序成 range 之後在乘上剛剛的 local_size 這才是他所要的 global_size。敘述上可能有點差異，CUDA 是我有 grid 個 block，而 SYCL 則是，我有 global_size 這麼多，每 local_size 框成一組。
簡單範例
這邊就以我要對每個陣列元素都加 1 來做為範例
// CUDA
__global__ void add_one(int num, float* A) {
auto tid = BlockIdx.x * blockDim.x + threadIdx.x;
if (tid < num) {
A[tid] += 1;
}
}
// execute
add_one<<<ngrid, 256>>>(num, A);// CUDA->SYCL
void add_one(int num, float* A, nd_item<3> item_ct1) {
auto tid = item_ct1.get_group(2) * item_ct1.get_local_range().get(2) + item_ct1.get_local_id(2);
// or auto tid = item_ct1.get_global_linear_id();
if (tid < num) {
A[tid] += 1;
}
}
// execute
queue->parallel_for(nd_range<3>(range<3>(1, 1, ngrid) * range<3>(1, 1, 256), range<3>(1, 1, 256)), [=](sycl::nd_item<3> item_ct1){
add_one(num, A, item_ct1);
});如果從 CUDA 轉到 SYCL，會發現變長很多，這還是縮減掉 submit 的版本，用 DPCT 來轉大致上也會長這樣，所以我們實驗室其實有在額外做一些讓兩者間更靠近一點，這之後會再介紹。那當然如果像這樣的程式其實如果直接用 SYCL 來寫，就會變短很多
// SYCL
queue->parallel_for(range<1>(num), [=](item<1> idx){
auto tid = idx.get_linear_id();
A[tid] += 1;
});SYCL 可以直接弄 global_size，會執行函數的也都只會在指定範圍之下，所以不像前面兩種，要先確認是否 tid 有沒有超過範圍。因此只是要 element-wise 的這類操作的話，直接用 SYCL 的方式寫會簡單許多。
總結
我其實是從 CUDA 轉 DPC++ 的這種角度去學習的，在敘述上會比較著重於這兩者的差異，有時候會覺得 SYCL 需要繞一下才能做到跟 CUDA 一樣的事情，但有的時候 SYCL 反而可以用更直覺的方式來寫。這篇大概簡單介紹了一下 DPC++ (或者說 SYCL)，以及會最先遇到跟 CUDA/HIP 的差別。如果有任何問題或建議都歡迎在底下留言。