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 裡使用。