I found related resources in WWDC. According to https://developer.apple.com/wwdc16/606?time=869 and https://developer.apple.com/wwdc20/10631?time=1481, they are using simdgroup_barrier() if thread group fits in a single SIMD group because threadgroup_barrier() is more expensive than simdgroup_barrier(). simdgroup size is normally 32 (afaik 64 in some AMD gpus) so I think simdgroup_barrier() can be used if thread group size is smaller than 32 or 64.

The sample runs on Monterey beta 7&8! Thanks!!