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 的差別。如果有任何問題或建議都歡迎在底下留言。