Translations

This document attempts to document how WGSL translates into the various backends for the cases where the translation is not a direct mapping.

Access Control

HLSL

  • ReadOnly -> ByteAddressBuffer
  • ReadWrite -> RWByteAddressBuffer

MSL

  • ReadOnly -> const

SPIR-V

There are two ways this can be achieved in SPIR-V. Either the variable can be decorated with NonWritable or each member of the struct can be decorated with NonWritable. We chose to go the struct member route.

  • The read-only becomes part of the type in this case. Otherwise, you are treating the readonly type information as part of the variable which is confusing.
  • Treating the readonly as part of the variable means we should be deduplicating the types behind the access control, which causes confusing with the type_names and various tracking systems within Tint.

Builtin Decorations

NameSPIR-VMSLHLSL
positionSpvBuiltInPositionpositionSV_Position
vertex_indexSpvBuiltInVertexIndexvertex_idSV_VertexID
instance_indexSpvBuiltInInstanceIndexinstance_idSV_InstanceID
front_facingSpvBuiltInFrontFacingfront_facingSV_IsFrontFacing
frag_coordSpvBuiltInFragCoordpositionSV_Position
frag_depthSpvBuiltInFragDepthdepth(any)SV_Depth
local_invocation_idSpvBuiltInLocalInvocationIdthread_position_in_threadgroupSV_GroupThreadID
local_invocation_indexSpvBuiltInLocalInvocationIndexthread_index_in_threadgroupSV_GroupIndex
global_invocation_idSpvBuiltInGlobalInvocationIdthread_position_in_gridSV_DispatchThreadID

Builtins Methods

NameSPIR-VMSLHLSL
absGLSLstd450FAbs or GLSLstd450SAbsfabs or absabs
acosGLSLstd450Acosacosacos
allSpvOpAllallall
anySpvOpAnyanyany
arrayLengthSpvOpArrayLength
asinGLSLstd450Asinasinasin
atanGLSLstd450Atanatanatan
atan2GLSLstd450Atan2atan2atan2
ceilGLSLstd450Ceilceilceil
clampGLSLstd450NClamp or GLSLstd450UClamp or GLSLstd450SClampclampclamp
cosGLSLstd450Coscoscos
coshGLSLstd450Coshcoshcosh
countOneBitsSpvOpBitCountpopcountcountbits
crossGLSLstd450Crosscrosscross
determinantGLSLstd450Determinantdeterminantdeterminant
distanceGLSLstd450Distancedistancedistance
dotSpOpDotdotdot
dpdxSpvOpDPdxdpdxddx
dpdxCoarseSpvOpDPdxCoarsedpdxddx_coarse
dpdxFineSpvOpDPdxFinedpdxddx_fine
dpdySpvOpDPdydpdyddy
dpdyCoarseSpvOpDPdyCoarsedpdyddy_coarse
dpdyFineSpvOpDPdyFinedpdyddy_fine
expGLSLstd450Expexpexp
exp2GLSLstd450Exp2exp2exp2
faceForwardGLSLstd450FaceForwardfaceforwardfaceforward
floorGLSLstd450Floorfloorfloor
fmaGLSLstd450Fmafmafma
fractGLSLstd450Fractfractfrac
frexpGLSLstd450Frexp
fwidthSpvOpFwidthfwidthfwidth
fwidthCoarseSpvOpFwidthCoarsefwidthfwidth
fwidthFineSpvOpFwidthFinefwidthfwidth
inverseSqrtGLSLstd450InverseSqrtrsqrtrsqrt
ldexpGLSLstd450Ldexp
lengthGLSLstd450Lengthlengthlength
logGLSLstd450Logloglog
log2GLSLstd450Log2log2log2
maxGLSLstd450NMax or GLSLstd450SMax or GLSLstd450UMaxfmax or maxmax
minGLSLstd450NMin or GLSLstd450SMin or GLSLstd450UMinfmin or minmin
mixGLSLstd450FMixmixmix
modfGLSLstd450Modf
normalizeGLSLstd450Normalizenormalizenormalize
powGLSLstd450Powpowpow
reflectGLSLstd450Reflectreflectreflect
reverseBitsSpvOpBitReversereverse_bitsreversebits
roundGLSLstd450Roundroundround
selectSpvOpSelectselect
signGLSLstd450FSignsignsign
sinGLSLstd450Sinsinsin
sinhGLSLstd450Sinhsinhsinh
smoothStepGLSLstd450SmoothStepsmoothstepsmoothstep
sqrtGLSLstd450Sqrtsqrtsqrt
stepGLSLstd450Stepstepstep
tanGLSLstd450Tantantan
tanhGLSLstd450Tanhtanhtanh
truncGLSLstd450Trunctrunctrunc

Types

Sampler Types

WGSLSPIR-VMSLHLSL
samplerOpTypeSamplersamplerSamplerState
sampler_comparisonOpTypeSamplersamplerSamplerComparisonState

Texture Types

WGSLSPIR-VMSLHLSL
texture_1d<type>OpTypeImage 1D Sampled=1texture1d<type, access::sample>Texture1D
texture_2d<type>OpTypeImage 2D Sampled=1texture2d<type, access::sample>Texture2D
texture_2d_array<type>OpTypeImage 2D Arrayed=1 Sampled=1texture2d_array<type, access::sample>Texture2DArray
texture_3d<type>OpTypeImage 3D Sampled=1texture3d<type, access::sample>Texture3D
texture_cube<type>OpTypeImage Cube Sampled=1texturecube<type, access::sample>TextureCube
texture_cube_array<type>OpTypeImage Cube Arrayed=1 Sampled=1texturecube_array<type, access::sample>TextureCubeArray
texture_multisampled_2d<type>OpTypeImage 2D MS=1 Sampled=1texture2d_ms<type, access::sample>Texture2D
texture_depth_2dOpTypeImage 2D Depth=1 Sampled=1depth2d<float, access::sample>Texture2D
texture_depth_2d_arrayOpTypeImage 2D Depth=1 Arrayed=1 Sampled=1depth2d_array<float, access::sample>Texture2DArray
texture_depth_cubeOpTypeImage Cube Depth=1 Sampled=1depthcube<float, access::sample>TextureCube
texture_depth_cube_arrayOpTypeImage Cube Depth=1 Arrayed=1 Sampled=1depthcube_array<float, access::sample>TextureCubeArray
texture_depth_multisampled_2dOpTypeImage 2D Depth=1 MS=1 Sampled=1depth2d<float, access::sample>Texture2DMSArray
texture_storage_1d<image_storage_type>OpTypeImage 1D Sampled=2texture1d<type, access::read>RWTexture1D
texture_storage_2d<image_storage_type>OpTypeImage 2D Sampled=2texture2d<type, access::read>RWTexture2D
texture_storage_2d_array<image_storage_type>OpTypeImage 2D Arrayed=1 Sampled=2texture2d_array<type, access::read>RWTexture2DArray
texture_storage_3d<image_storage_type>OpTypeImage 3D Sampled=2texture3d<type, access::read>RWTexture3D
texture_storage_1d<image_storage_type>OpTypeImage 1D Sampled=2texture1d<type, access::write>RWTexture1D
texture_storage_2d<image_storage_type>OpTypeImage 2D Sampled=1texture2d<type, access::write>RWTexture2D
texture_storage_2d_array<image_storage_type>OpTypeImage 2D Arrayed=1 Sampled=2texture2d_array<type, access::write>RWTexture2DArray
texture_storage_3d<image_storage_type>OpTypeImage 3D Sampled=2texture3d<type, access::write>RWTexture3D

Short-circuting

HLSL

TODO(dsinclair): Nested if's

SPIR-V

TODO(dsinclair): Nested if's

Storage classes

TODO(dsinclair): do ...

Storage buffers

HLSL

TODO(dsinclair): Rewriting of accessors to loads

Loop blocks

HLSL

TODO(dsinclair): Rewrite with bools

MSL

TODO(dsinclair): Rewrite with bools

Input / Output storage class

HLSL

TODO(dsinclair): Structs and params

MSL

TODO(dsinclair): Structs and params

Discard

HLSL

  • discard

MSL

  • discard_fragment()

Specialization constants

HLSL

#ifndef WGSL_SPEC_CONSTANT_<id>
-- if default provided
#define WGSL_SPEC_CONSTANT_<id> default value
-- else
#error spec constant required for constant id
--
#endif
static const <type> <name> = WGSL_SPEC_CONSTANT_<id>

MSL

@function_constant(<id>)