OneAPI DPC++ (SYCL-based) 簡介

oneAPI Mar 13, 2022

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

Tags

Great! You've successfully subscribed.
Great! Next, complete checkout for full access.
Welcome back! You've successfully signed in.
Success! Your account is fully activated, you now have access to all content.