SYCL ワークサブグループが熱い！？

さっきまでワークサブグループってなんかの単位何だろうな…とあまり意識していませんでした。この感動をどう伝えたらいいのだろうか…。今風の書き方ならまず結果を書くべきなのかな？

vaeの実行時間が、30.35sかかっていたものが、23.66sとか言う数字を叩き出しました。

変更した部分は、昨日まで処理を書き換えていた部分で不要になったコードや引数、コメントを削除して若干変数の取り回しなども確認しながら変更を加えていました。

ついでに、処理のブロック化係数や最大ワークグループはデバイスに依存する数字なので、この値を動的に変更できるかどうか確認していました。

処理のブロック化係数はループの処理回数となっているため、コンパイルでunrollで処理を線型的になるようにコンパイルすることで実行速度を上げています。そのため、コンパイルの時点で実行回数が確定している必要があるわけですが、この処理のブロック化係数もデバイスによって最適な値は変わってくるはずなので、変数にしたい部分ではありますが、定数として、あきらめる必要がありました。

とはいえ、少しづつ書き換えていると、nd_rangeのitemに付随しているgroup, local, sub_groupという単位でいつも悩むわけです。

local_id()は、その処理のlocal_id()というのは直接的な意味なのでとらえやすいのですが、local_range()を考えると…まぁ処理単位の範囲と言うも理解できます。そしてgroup_range()も指定した範囲のはず…ではgroup_id()っていったい何を指しているのだろうか？という事です。ついでにsub_grouptとは？？と謎だらけになってしまいます。

CUDAのコードであれば処理単位を細分化していく考え方なので深く考えなくてもすんなり頭に入るのですが、どうもグローバル、ローカルという言葉を使うと混乱してしまいます。

で、この言葉を整理するために毎回確認しているページがあります。

https://www.isus.jp/wp-content/uploads/dpct/2025/user_guide/reference/compare-prog-models.html#thread-indexing

この部分の表にまとめられているものを毎回確認しながら言葉を整理していたりします(笑)

で、ふと目に留まったのが、その下にあるCUDAのコードをSYCLへの書き換え例のコード…

今までも何度も見ているコードですが、カーネル内で１要素の処理を行っているだけの単純なサンプルコードと思っていましたが…今まで意識したことのなかったsycl::reqd_sub_group_sizeというキーワード。

ここまでさんざん書き換えてきたループが、もしかしてこの1行の指示だけで綺麗に回してくれるんじゃ？？という疑問。

早速、ループ処理で回していた部分を書き換え、呼び出し部分にこのキーワードを追加してコンパイル、実行してみたところ…

30秒を大幅に切る数字が…。

しかもこれの凄いところが、おそらくブロック化係数も変数でも上手く動いてくれそうなところ。

内部的には複雑なブロック処理の結果、ワークアイテムが実行されてゆくのだと思いますが、とりあえず何も考えなくても何とかしてくれそうなこの形…すごいですね…。

ただ、問題も発生しそうなのは、このサブワークの単位も最大数があるので、処理を分割する必要が出てきそうな点は…まぁ後で考えるか(笑)

ちなみに…1024x1024は意味の分からないところでエラーになってしまうので一旦ソースを整える方向で作業を進めています。