Investigation: Subgroup Support in WHLSL

[mehmetoguzderin]

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]] / 4	gl_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