SYCL(DPC++) sub_group 簡介

SYCL sub_group 類似於 CUDA warp 的概念,threads 在同一個 subgroup 中是同時執行,且不用透過 shared memory 就可以交換資料 (shuffle) 等。 subgroup 一開始是由 Intel 所額外加上去的,後來被納入進 SYCL 2020 中。

Shuffle

sub_group 也有許多跟 cuda cooperative group shuffle 功能一樣的函數,但是名字有點不一樣。大致上是這樣對應 shfl_XX (cuda) <-> shuffle_XX (SYCL) 。但有一些不同的地方, shfl_up/shfl_down 在 CUDA 中,當該 thread 是存取無效的位置時,是回傳原本資料,但在 SYCL shuffle_up/shuffle_down 則是說明回傳值並不特定,如果有要處理這方面的值時,需要注意。

另外在 SYCL 是只有提供 free function 但是在 DPC++ 這些可以直接於 sub_group 的成員函數使用

// SYCL <-> DPCPP
sycl::shift_group_left(sg, x, offset); <-> sg.shuffle_down(x, offset);
sycl::shift_group_right(sg, x, offset); <-> sg.shuffle_up(x, offset);
sycl::select_from_group(sg, x, id); <-> sg.shuffle(x, id);
sycl::permute_group_by_xor(sg, x, mask); <-> sg.shuffle_xor(x, mask);

大小

sub_group 的大小會根據硬體有所改變,Nvidia 的就固定在 32,但 Intel CPU 大多是 4, 8, 16, 32, 64 (CPU 可以用 32, 64,是在 DPC++ 2021 某一版後才加入) ,而 Intel GPU 是 8, 16, 32。目前大致上是這樣,但我也沒有找到 Intel 有說他們之後的也都會給這些 sub_group 還是會更動,因此在開發時,都要因此做出一定彈性或者取捨。要用 reqd_sub_group_size 決定大小,並在 parallel_for 時給出,沒有指定的話,編譯器會自動使用預設值。當然,在 kernel 中是不能改動大小的。

範例

利用 reqd_sub_group_size 指定大小,並利用 shuffle_xor 做資料交換

h.parallel_for(nd_range<1>(N, B), 
  [=](nd_item<1> item) [[sycl::reqd_sub_group_size(16)]] {
    auto sg = item.get_sub_group();
    auto i = item.get_global_id(0);
    auto local = data[i];
    for (int mask = 0; mask < sg.size(); mask <<= 1) {
      auto remote = sg.shuffle_xor(local, mask);
      local += remote;
    }
    data[i] = local;
});

也可以直接用 group algorithm 裡的 reduce 完成

h.parallel_for(nd_range<1>(N, B), 
  [=](nd_item<1> item) [[sycl::reqd_sub_group_size(16)]] {
    auto sg = item.get_sub_group();
    auto i = item.get_global_id(0);
    data[i] = reduce(sg, data[i], ONEAPI::plus<>());
});

結語

雖然 SYCL 允許不同的 sub_group (跟硬體有關),可以根據需求先給定好大小,但不同硬體參數也導致開發上需要做取捨或者要個別處理;另外並沒有辦法像 CUDA subwarp 那樣使用,有些函式也因此沒辦法直接轉換過去。並介紹在 CUDA warp 上很常用的 shuffle 要如何對應到 SYCL 裡使用。