This repository was archived by the owner on Dec 8, 2022. It is now read-only.
Description See 0 . Subgroup support in shading languages mostly overlap but SPIR-V has functions without correspondence in other languages. In this article, only mutual functions are inspected. Tables are divided according to feature categories specified in 0 but risks for unified consideration (as SIMD-group in MacOS and quad-group in iOS) of SIMD-group and quad-group in MSL should be further investigated.
Existing Overlap Work
DirectX Shader Compiler which supports SPIR-V backend for HLSL has an overlap work for HLSL to SPIR-V mapping 1 which we can use as a clue for HLSL to GLSL mapping.
Tables
WebGPUSubgroupFeatureBits::BASIC
Proposed Function
HLSL Function 2
MSL Function (SIMD-group / quad-group) 3
GLSL Function 4
u32 subgroupSize(void)uint WaveGetLaneCount(void)[[threads_per_simdgroup]] / 4gl_SubgroupSize
u32 subgroupThreadIndex(void)uint WaveGetLaneIndex(void)[[simdgroup_index_in_threadgroup]] / [[thread_index_in_quadgroup]]gl_SubgroupInvocationID
bool subgroupElect(void)bool WaveIsFirstLane(void)bool simd_is_first(void) / bool quad_is_first(void)bool subgroupElect(void)
WebGPUSubgroupFeatureBits::VOTE
Proposed Function
HLSL Function 2
MSL Function (SIMD-group / quad-group) 3
GLSL Function 4
bool subgroupAll(bool value)bool WaveActiveAllTrue(bool expr)bool simd_all(bool expr) / bool quad_all(bool expr)bool subgroupAll(bool value)
bool subgroupAny(bool value)bool WaveActiveAnyTrue(bool expr)bool simd_any(bool expr) / bool quad_any(bool expr)bool subgroupAny(bool value)
WebGPUSubgroupFeatureBits::BALLOT
Proposed Function
HLSL Function 2
MSL Function (SIMD-group / quad-group) 3
GLSL Function 4
u128 subgroupBallot(bool value)uint4 WaveActiveBallot(bool expr)simd_vote simd_ballot(bool expr) / quad_vote quad_ballot(bool expr)uvec4 subgroupBallot(bool value)
type subgroupBroadcast(type value, u32 threadIndex)<type> WaveReadLaneAt(<type> expr, uint laneIndex)T simd_broadcast(T data,
ushort broadcast_lane_id) / T quad_broadcast(T data,
ushort broadcast_lane_id)genType subgroupBroadcast(genType value, uint id)
type subgroupBroadcastFirst(type value)<type> WaveReadLaneFirst(<type> expr)T simd_broadcast_first(T data) / T quad_broadcast_first(T data)genType subgroupBroadcastFirst(genType value)
WebGPUSubgroupFeatureBits::ARITHMETIC
Proposed Function
HLSL Function 2
MSL Function (SIMD-group / quad-group) 3
GLSL Function 4
type subgroupAdd(type value)<type> WaveActiveSum(<type> expr)T simd_sum(T data) / T quad_sum(T data)genType subgroupAdd(genType value)
type subgroupMul(type value)<type> WaveActiveProduct(<type> expr)T simd_product(T data) / T quad_product(T data)genType subgroupMul(genType value)
type subgroupMin(type value)<type> WaveActiveMin(<type> expr)T simd_min(T data) / T quad_min(T data)genType subgroupMin(genType value)
type subgroupMax(type value)<type> WaveActiveMax(<type> expr)T simd_max(T data) / T quad_max(T data)genType subgroupMax(genType value)
type subgroupAnd(type value)<type> WaveActiveBitAdd(<type> expr)T simd_and(T data) / T quad_and(T data)genType subgroupAnd(genType value)
type subgroupOr(type value)<type> WaveActiveBitOr(<type> expr)T simd_or(T data) / T quad_or(T data)genType subgroupOr(genType value)
type subgroupXor(type value)<type> WaveActiveBitXor(<type> expr)T simd_xor(T data) / T quad_xor(T data)genType subgroupXor(genType value)
type subgroupPrefixExclusiveAdd(type value)<type> WavePrefixSum(<type> expr)T simd_prefix_exclusive_sum(T data) / T quad_ prefix_exclusive_sum(T data)genType subgroupExclusiveAdd(genType value)
type subgroupPrefixExclusiveMul(type value)<type> WavePrefixProduct(<type> expr)T simd_ prefix_exclusive_product(T data) / T quad_ prefix_exclusive_product(T data)genType subgroupExclusiveMul(genType value)
WebGPUSubgroupFeatureBits::QUAD
Proposed Function
HLSL Function 2
MSL Function (SIMD-group / quad-group) 3
GLSL Function 4
type subgroupQuadBroadcast(type value, u32 threadIndex)<type> QuadReadLaneAt(<type> expr, uint laneIndex)T quad_broadcast(T data,
ushort broadcast_lane_id) / T quad_broadcast(T data,
ushort broadcast_lane_id)genType subgroupQuadBroadcast(genType value, uint id)
References Reactions are currently unavailable
