[writer/hlsl,msl] Correctly generate inf and nan literals
Also add missing msl macros to the renamer.
Bug: tint:951
Change-Id: I543e6eae885c979596ca63f9d6c7378dd5425e8a
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56941
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/transform/renamer.cc b/src/transform/renamer.cc
index b6f6b78..ac975b3 100644
--- a/src/transform/renamer.cc
+++ b/src/transform/renamer.cc
@@ -32,807 +32,844 @@
namespace {
// This list is used for a binary search and must be kept in sorted order.
-const char* kReservedKeywordsHLSL[] = {"AddressU",
- "AddressV",
- "AddressW",
- "AllMemoryBarrier",
- "AllMemoryBarrierWithGroupSync",
- "AppendStructuredBuffer",
- "BINORMAL",
- "BLENDINDICES",
- "BLENDWEIGHT",
- "BlendState",
- "BorderColor",
- "Buffer",
- "ByteAddressBuffer",
- "COLOR",
- "CheckAccessFullyMapped",
- "ComparisonFunc",
- "CompileShader",
- "ComputeShader",
- "ConsumeStructuredBuffer",
- "D3DCOLORtoUBYTE4",
- "DEPTH",
- "DepthStencilState",
- "DepthStencilView",
- "DeviceMemoryBarrier",
- "DeviceMemroyBarrierWithGroupSync",
- "DomainShader",
- "EvaluateAttributeAtCentroid",
- "EvaluateAttributeAtSample",
- "EvaluateAttributeSnapped",
- "FOG",
- "Filter",
- "GeometryShader",
- "GetRenderTargetSampleCount",
- "GetRenderTargetSamplePosition",
- "GroupMemoryBarrier",
- "GroupMemroyBarrierWithGroupSync",
- "Hullshader",
- "InputPatch",
- "InterlockedAdd",
- "InterlockedAnd",
- "InterlockedCompareExchange",
- "InterlockedCompareStore",
- "InterlockedExchange",
- "InterlockedMax",
- "InterlockedMin",
- "InterlockedOr",
- "InterlockedXor",
- "LineStream",
- "MaxAnisotropy",
- "MaxLOD",
- "MinLOD",
- "MipLODBias",
- "NORMAL",
- "NULL",
- "Normal",
- "OutputPatch",
- "POSITION",
- "POSITIONT",
- "PSIZE",
- "PixelShader",
- "PointStream",
- "Process2DQuadTessFactorsAvg",
- "Process2DQuadTessFactorsMax",
- "Process2DQuadTessFactorsMin",
- "ProcessIsolineTessFactors",
- "ProcessQuadTessFactorsAvg",
- "ProcessQuadTessFactorsMax",
- "ProcessQuadTessFactorsMin",
- "ProcessTriTessFactorsAvg",
- "ProcessTriTessFactorsMax",
- "ProcessTriTessFactorsMin",
- "RWBuffer",
- "RWByteAddressBuffer",
- "RWStructuredBuffer",
- "RWTexture1D",
- "RWTexture2D",
- "RWTexture2DArray",
- "RWTexture3D",
- "RasterizerState",
- "RenderTargetView",
- "SV_ClipDistance",
- "SV_Coverage",
- "SV_CullDistance",
- "SV_Depth",
- "SV_DepthGreaterEqual",
- "SV_DepthLessEqual",
- "SV_DispatchThreadID",
- "SV_DomainLocation",
- "SV_GSInstanceID",
- "SV_GroupID",
- "SV_GroupIndex",
- "SV_GroupThreadID",
- "SV_InnerCoverage",
- "SV_InsideTessFactor",
- "SV_InstanceID",
- "SV_IsFrontFace",
- "SV_OutputControlPointID",
- "SV_Position",
- "SV_PrimitiveID",
- "SV_RenderTargetArrayIndex",
- "SV_SampleIndex",
- "SV_StencilRef",
- "SV_Target",
- "SV_TessFactor",
- "SV_VertexArrayIndex",
- "SV_VertexID",
- "Sampler",
- "Sampler1D",
- "Sampler2D",
- "Sampler3D",
- "SamplerCUBE",
- "StructuredBuffer",
- "TANGENT",
- "TESSFACTOR",
- "TEXCOORD",
- "Texcoord",
- "Texture",
- "Texture1D",
- "Texture2D",
- "Texture2DArray",
- "Texture2DMS",
- "Texture2DMSArray",
- "Texture3D",
- "TextureCube",
- "TextureCubeArray",
- "TriangleStream",
- "VFACE",
- "VPOS",
- "VertexShader",
- "abort",
- "allow_uav_condition",
- "asdouble",
- "asfloat",
- "asint",
- "asm",
- "asm_fragment",
- "asuint",
- "auto",
- "bool",
- "bool1",
- "bool1x1",
- "bool1x2",
- "bool1x3",
- "bool1x4",
- "bool2",
- "bool2x1",
- "bool2x2",
- "bool2x3",
- "bool2x4",
- "bool3",
- "bool3x1",
- "bool3x2",
- "bool3x3",
- "bool3x4",
- "bool4",
- "bool4x1",
- "bool4x2",
- "bool4x3",
- "bool4x4",
- "branch",
- "break",
- "call",
- "case",
- "catch",
- "cbuffer",
- "centroid",
- "char",
- "class",
- "clip",
- "column_major",
- "compile_fragment",
- "const",
- "const_cast",
- "continue",
- "countbits",
- "ddx",
- "ddx_coarse",
- "ddx_fine",
- "ddy",
- "ddy_coarse",
- "ddy_fine",
- "degrees",
- "delete",
- "discard",
- "do",
- "double",
- "double1",
- "double1x1",
- "double1x2",
- "double1x3",
- "double1x4",
- "double2",
- "double2x1",
- "double2x2",
- "double2x3",
- "double2x4",
- "double3",
- "double3x1",
- "double3x2",
- "double3x3",
- "double3x4",
- "double4",
- "double4x1",
- "double4x2",
- "double4x3",
- "double4x4",
- "dst",
- "dword",
- "dword1",
- "dword1x1",
- "dword1x2",
- "dword1x3",
- "dword1x4",
- "dword2",
- "dword2x1",
- "dword2x2",
- "dword2x3",
- "dword2x4",
- "dword3",
- "dword3x1",
- "dword3x2",
- "dword3x3",
- "dword3x4",
- "dword4",
- "dword4x1",
- "dword4x2",
- "dword4x3",
- "dword4x4",
- "dynamic_cast",
- "else",
- "enum",
- "errorf",
- "explicit",
- "export",
- "extern",
- "f16to32",
- "f32tof16",
- "false",
- "fastopt",
- "firstbithigh",
- "firstbitlow",
- "flatten",
- "float",
- "float1",
- "float1x1",
- "float1x2",
- "float1x3",
- "float1x4",
- "float2",
- "float2x1",
- "float2x2",
- "float2x3",
- "float2x4",
- "float3",
- "float3x1",
- "float3x2",
- "float3x3",
- "float3x4",
- "float4",
- "float4x1",
- "float4x2",
- "float4x3",
- "float4x4",
- "fmod",
- "for",
- "forcecase",
- "frac",
- "friend",
- "fxgroup",
- "goto",
- "groupshared",
- "half",
- "half1",
- "half1x1",
- "half1x2",
- "half1x3",
- "half1x4",
- "half2",
- "half2x1",
- "half2x2",
- "half2x3",
- "half2x4",
- "half3",
- "half3x1",
- "half3x2",
- "half3x3",
- "half3x4",
- "half4",
- "half4x1",
- "half4x2",
- "half4x3",
- "half4x4",
- "if",
- "in",
- "inline",
- "inout",
- "int",
- "int1",
- "int1x1",
- "int1x2",
- "int1x3",
- "int1x4",
- "int2",
- "int2x1",
- "int2x2",
- "int2x3",
- "int2x4",
- "int3",
- "int3x1",
- "int3x2",
- "int3x3",
- "int3x4",
- "int4",
- "int4x1",
- "int4x2",
- "int4x3",
- "int4x4",
- "interface",
- "isfinite",
- "isinf",
- "isnan",
- "lerp",
- "lineadj",
- "linear",
- "lit",
- "log10",
- "long",
- "loop",
- "mad",
- "matrix",
- "min10float",
- "min10float1",
- "min10float1x1",
- "min10float1x2",
- "min10float1x3",
- "min10float1x4",
- "min10float2",
- "min10float2x1",
- "min10float2x2",
- "min10float2x3",
- "min10float2x4",
- "min10float3",
- "min10float3x1",
- "min10float3x2",
- "min10float3x3",
- "min10float3x4",
- "min10float4",
- "min10float4x1",
- "min10float4x2",
- "min10float4x3",
- "min10float4x4",
- "min12int",
- "min12int1",
- "min12int1x1",
- "min12int1x2",
- "min12int1x3",
- "min12int1x4",
- "min12int2",
- "min12int2x1",
- "min12int2x2",
- "min12int2x3",
- "min12int2x4",
- "min12int3",
- "min12int3x1",
- "min12int3x2",
- "min12int3x3",
- "min12int3x4",
- "min12int4",
- "min12int4x1",
- "min12int4x2",
- "min12int4x3",
- "min12int4x4",
- "min16float",
- "min16float1",
- "min16float1x1",
- "min16float1x2",
- "min16float1x3",
- "min16float1x4",
- "min16float2",
- "min16float2x1",
- "min16float2x2",
- "min16float2x3",
- "min16float2x4",
- "min16float3",
- "min16float3x1",
- "min16float3x2",
- "min16float3x3",
- "min16float3x4",
- "min16float4",
- "min16float4x1",
- "min16float4x2",
- "min16float4x3",
- "min16float4x4",
- "min16int",
- "min16int1",
- "min16int1x1",
- "min16int1x2",
- "min16int1x3",
- "min16int1x4",
- "min16int2",
- "min16int2x1",
- "min16int2x2",
- "min16int2x3",
- "min16int2x4",
- "min16int3",
- "min16int3x1",
- "min16int3x2",
- "min16int3x3",
- "min16int3x4",
- "min16int4",
- "min16int4x1",
- "min16int4x2",
- "min16int4x3",
- "min16int4x4",
- "min16uint",
- "min16uint1",
- "min16uint1x1",
- "min16uint1x2",
- "min16uint1x3",
- "min16uint1x4",
- "min16uint2",
- "min16uint2x1",
- "min16uint2x2",
- "min16uint2x3",
- "min16uint2x4",
- "min16uint3",
- "min16uint3x1",
- "min16uint3x2",
- "min16uint3x3",
- "min16uint3x4",
- "min16uint4",
- "min16uint4x1",
- "min16uint4x2",
- "min16uint4x3",
- "min16uint4x4",
- "msad4",
- "mul",
- "mutable",
- "namespace",
- "new",
- "nointerpolation",
- "noise",
- "noperspective",
- "numthreads",
- "operator",
- "out",
- "packoffset",
- "pass",
- "pixelfragment",
- "pixelshader",
- "point",
- "precise",
- "printf",
- "private",
- "protected",
- "public",
- "radians",
- "rcp",
- "refract",
- "register",
- "reinterpret_cast",
- "return",
- "row_major",
- "rsqrt",
- "sample",
- "sampler",
- "sampler1D",
- "sampler2D",
- "sampler3D",
- "samplerCUBE",
- "sampler_state",
- "saturate",
- "shared",
- "short",
- "signed",
- "sincos",
- "sizeof",
- "snorm",
- "stateblock",
- "stateblock_state",
- "static",
- "static_cast",
- "string",
- "struct",
- "switch",
- "tbuffer",
- "technique",
- "technique10",
- "technique11",
- "template",
- "tex1D",
- "tex1Dbias",
- "tex1Dgrad",
- "tex1Dlod",
- "tex1Dproj",
- "tex2D",
- "tex2Dbias",
- "tex2Dgrad",
- "tex2Dlod",
- "tex2Dproj",
- "tex3D",
- "tex3Dbias",
- "tex3Dgrad",
- "tex3Dlod",
- "tex3Dproj",
- "texCUBE",
- "texCUBEbias",
- "texCUBEgrad",
- "texCUBElod",
- "texCUBEproj",
- "texture",
- "texture1D",
- "texture1DArray",
- "texture2D",
- "texture2DArray",
- "texture2DMS",
- "texture2DMSArray",
- "texture3D",
- "textureCube",
- "textureCubeArray",
- "this",
- "throw",
- "transpose",
- "triangle",
- "triangleadj",
- "true",
- "try",
- "typedef",
- "typename",
- "uint",
- "uint1",
- "uint1x1",
- "uint1x2",
- "uint1x3",
- "uint1x4",
- "uint2",
- "uint2x1",
- "uint2x2",
- "uint2x3",
- "uint2x4",
- "uint3",
- "uint3x1",
- "uint3x2",
- "uint3x3",
- "uint3x4",
- "uint4",
- "uint4x1",
- "uint4x2",
- "uint4x3",
- "uint4x4",
- "uniform",
- "union",
- "unorm",
- "unroll",
- "unsigned",
- "using",
- "vector",
- "vertexfragment",
- "vertexshader",
- "virtual",
- "void",
- "volatile",
- "while"};
+const char* kReservedKeywordsHLSL[] = {
+ "AddressU",
+ "AddressV",
+ "AddressW",
+ "AllMemoryBarrier",
+ "AllMemoryBarrierWithGroupSync",
+ "AppendStructuredBuffer",
+ "BINORMAL",
+ "BLENDINDICES",
+ "BLENDWEIGHT",
+ "BlendState",
+ "BorderColor",
+ "Buffer",
+ "ByteAddressBuffer",
+ "COLOR",
+ "CheckAccessFullyMapped",
+ "ComparisonFunc",
+ "CompileShader",
+ "ComputeShader",
+ "ConsumeStructuredBuffer",
+ "D3DCOLORtoUBYTE4",
+ "DEPTH",
+ "DepthStencilState",
+ "DepthStencilView",
+ "DeviceMemoryBarrier",
+ "DeviceMemroyBarrierWithGroupSync",
+ "DomainShader",
+ "EvaluateAttributeAtCentroid",
+ "EvaluateAttributeAtSample",
+ "EvaluateAttributeSnapped",
+ "FOG",
+ "Filter",
+ "GeometryShader",
+ "GetRenderTargetSampleCount",
+ "GetRenderTargetSamplePosition",
+ "GroupMemoryBarrier",
+ "GroupMemroyBarrierWithGroupSync",
+ "Hullshader",
+ "InputPatch",
+ "InterlockedAdd",
+ "InterlockedAnd",
+ "InterlockedCompareExchange",
+ "InterlockedCompareStore",
+ "InterlockedExchange",
+ "InterlockedMax",
+ "InterlockedMin",
+ "InterlockedOr",
+ "InterlockedXor",
+ "LineStream",
+ "MaxAnisotropy",
+ "MaxLOD",
+ "MinLOD",
+ "MipLODBias",
+ "NORMAL",
+ "NULL",
+ "Normal",
+ "OutputPatch",
+ "POSITION",
+ "POSITIONT",
+ "PSIZE",
+ "PixelShader",
+ "PointStream",
+ "Process2DQuadTessFactorsAvg",
+ "Process2DQuadTessFactorsMax",
+ "Process2DQuadTessFactorsMin",
+ "ProcessIsolineTessFactors",
+ "ProcessQuadTessFactorsAvg",
+ "ProcessQuadTessFactorsMax",
+ "ProcessQuadTessFactorsMin",
+ "ProcessTriTessFactorsAvg",
+ "ProcessTriTessFactorsMax",
+ "ProcessTriTessFactorsMin",
+ "RWBuffer",
+ "RWByteAddressBuffer",
+ "RWStructuredBuffer",
+ "RWTexture1D",
+ "RWTexture2D",
+ "RWTexture2DArray",
+ "RWTexture3D",
+ "RasterizerState",
+ "RenderTargetView",
+ "SV_ClipDistance",
+ "SV_Coverage",
+ "SV_CullDistance",
+ "SV_Depth",
+ "SV_DepthGreaterEqual",
+ "SV_DepthLessEqual",
+ "SV_DispatchThreadID",
+ "SV_DomainLocation",
+ "SV_GSInstanceID",
+ "SV_GroupID",
+ "SV_GroupIndex",
+ "SV_GroupThreadID",
+ "SV_InnerCoverage",
+ "SV_InsideTessFactor",
+ "SV_InstanceID",
+ "SV_IsFrontFace",
+ "SV_OutputControlPointID",
+ "SV_Position",
+ "SV_PrimitiveID",
+ "SV_RenderTargetArrayIndex",
+ "SV_SampleIndex",
+ "SV_StencilRef",
+ "SV_Target",
+ "SV_TessFactor",
+ "SV_VertexArrayIndex",
+ "SV_VertexID",
+ "Sampler",
+ "Sampler1D",
+ "Sampler2D",
+ "Sampler3D",
+ "SamplerCUBE",
+ "StructuredBuffer",
+ "TANGENT",
+ "TESSFACTOR",
+ "TEXCOORD",
+ "Texcoord",
+ "Texture",
+ "Texture1D",
+ "Texture2D",
+ "Texture2DArray",
+ "Texture2DMS",
+ "Texture2DMSArray",
+ "Texture3D",
+ "TextureCube",
+ "TextureCubeArray",
+ "TriangleStream",
+ "VFACE",
+ "VPOS",
+ "VertexShader",
+ "abort",
+ "allow_uav_condition",
+ "asdouble",
+ "asfloat",
+ "asint",
+ "asm",
+ "asm_fragment",
+ "asuint",
+ "auto",
+ "bool",
+ "bool1",
+ "bool1x1",
+ "bool1x2",
+ "bool1x3",
+ "bool1x4",
+ "bool2",
+ "bool2x1",
+ "bool2x2",
+ "bool2x3",
+ "bool2x4",
+ "bool3",
+ "bool3x1",
+ "bool3x2",
+ "bool3x3",
+ "bool3x4",
+ "bool4",
+ "bool4x1",
+ "bool4x2",
+ "bool4x3",
+ "bool4x4",
+ "branch",
+ "break",
+ "call",
+ "case",
+ "catch",
+ "cbuffer",
+ "centroid",
+ "char",
+ "class",
+ "clip",
+ "column_major",
+ "compile_fragment",
+ "const",
+ "const_cast",
+ "continue",
+ "countbits",
+ "ddx",
+ "ddx_coarse",
+ "ddx_fine",
+ "ddy",
+ "ddy_coarse",
+ "ddy_fine",
+ "degrees",
+ "delete",
+ "discard",
+ "do",
+ "double",
+ "double1",
+ "double1x1",
+ "double1x2",
+ "double1x3",
+ "double1x4",
+ "double2",
+ "double2x1",
+ "double2x2",
+ "double2x3",
+ "double2x4",
+ "double3",
+ "double3x1",
+ "double3x2",
+ "double3x3",
+ "double3x4",
+ "double4",
+ "double4x1",
+ "double4x2",
+ "double4x3",
+ "double4x4",
+ "dst",
+ "dword",
+ "dword1",
+ "dword1x1",
+ "dword1x2",
+ "dword1x3",
+ "dword1x4",
+ "dword2",
+ "dword2x1",
+ "dword2x2",
+ "dword2x3",
+ "dword2x4",
+ "dword3",
+ "dword3x1",
+ "dword3x2",
+ "dword3x3",
+ "dword3x4",
+ "dword4",
+ "dword4x1",
+ "dword4x2",
+ "dword4x3",
+ "dword4x4",
+ "dynamic_cast",
+ "else",
+ "enum",
+ "errorf",
+ "explicit",
+ "export",
+ "extern",
+ "f16to32",
+ "f32tof16",
+ "false",
+ "fastopt",
+ "firstbithigh",
+ "firstbitlow",
+ "flatten",
+ "float",
+ "float1",
+ "float1x1",
+ "float1x2",
+ "float1x3",
+ "float1x4",
+ "float2",
+ "float2x1",
+ "float2x2",
+ "float2x3",
+ "float2x4",
+ "float3",
+ "float3x1",
+ "float3x2",
+ "float3x3",
+ "float3x4",
+ "float4",
+ "float4x1",
+ "float4x2",
+ "float4x3",
+ "float4x4",
+ "fmod",
+ "for",
+ "forcecase",
+ "frac",
+ "friend",
+ "fxgroup",
+ "goto",
+ "groupshared",
+ "half",
+ "half1",
+ "half1x1",
+ "half1x2",
+ "half1x3",
+ "half1x4",
+ "half2",
+ "half2x1",
+ "half2x2",
+ "half2x3",
+ "half2x4",
+ "half3",
+ "half3x1",
+ "half3x2",
+ "half3x3",
+ "half3x4",
+ "half4",
+ "half4x1",
+ "half4x2",
+ "half4x3",
+ "half4x4",
+ "if",
+ "in",
+ "inline",
+ "inout",
+ "int",
+ "int1",
+ "int1x1",
+ "int1x2",
+ "int1x3",
+ "int1x4",
+ "int2",
+ "int2x1",
+ "int2x2",
+ "int2x3",
+ "int2x4",
+ "int3",
+ "int3x1",
+ "int3x2",
+ "int3x3",
+ "int3x4",
+ "int4",
+ "int4x1",
+ "int4x2",
+ "int4x3",
+ "int4x4",
+ "interface",
+ "isfinite",
+ "isinf",
+ "isnan",
+ "lerp",
+ "lineadj",
+ "linear",
+ "lit",
+ "log10",
+ "long",
+ "loop",
+ "mad",
+ "matrix",
+ "min10float",
+ "min10float1",
+ "min10float1x1",
+ "min10float1x2",
+ "min10float1x3",
+ "min10float1x4",
+ "min10float2",
+ "min10float2x1",
+ "min10float2x2",
+ "min10float2x3",
+ "min10float2x4",
+ "min10float3",
+ "min10float3x1",
+ "min10float3x2",
+ "min10float3x3",
+ "min10float3x4",
+ "min10float4",
+ "min10float4x1",
+ "min10float4x2",
+ "min10float4x3",
+ "min10float4x4",
+ "min12int",
+ "min12int1",
+ "min12int1x1",
+ "min12int1x2",
+ "min12int1x3",
+ "min12int1x4",
+ "min12int2",
+ "min12int2x1",
+ "min12int2x2",
+ "min12int2x3",
+ "min12int2x4",
+ "min12int3",
+ "min12int3x1",
+ "min12int3x2",
+ "min12int3x3",
+ "min12int3x4",
+ "min12int4",
+ "min12int4x1",
+ "min12int4x2",
+ "min12int4x3",
+ "min12int4x4",
+ "min16float",
+ "min16float1",
+ "min16float1x1",
+ "min16float1x2",
+ "min16float1x3",
+ "min16float1x4",
+ "min16float2",
+ "min16float2x1",
+ "min16float2x2",
+ "min16float2x3",
+ "min16float2x4",
+ "min16float3",
+ "min16float3x1",
+ "min16float3x2",
+ "min16float3x3",
+ "min16float3x4",
+ "min16float4",
+ "min16float4x1",
+ "min16float4x2",
+ "min16float4x3",
+ "min16float4x4",
+ "min16int",
+ "min16int1",
+ "min16int1x1",
+ "min16int1x2",
+ "min16int1x3",
+ "min16int1x4",
+ "min16int2",
+ "min16int2x1",
+ "min16int2x2",
+ "min16int2x3",
+ "min16int2x4",
+ "min16int3",
+ "min16int3x1",
+ "min16int3x2",
+ "min16int3x3",
+ "min16int3x4",
+ "min16int4",
+ "min16int4x1",
+ "min16int4x2",
+ "min16int4x3",
+ "min16int4x4",
+ "min16uint",
+ "min16uint1",
+ "min16uint1x1",
+ "min16uint1x2",
+ "min16uint1x3",
+ "min16uint1x4",
+ "min16uint2",
+ "min16uint2x1",
+ "min16uint2x2",
+ "min16uint2x3",
+ "min16uint2x4",
+ "min16uint3",
+ "min16uint3x1",
+ "min16uint3x2",
+ "min16uint3x3",
+ "min16uint3x4",
+ "min16uint4",
+ "min16uint4x1",
+ "min16uint4x2",
+ "min16uint4x3",
+ "min16uint4x4",
+ "msad4",
+ "mul",
+ "mutable",
+ "namespace",
+ "new",
+ "nointerpolation",
+ "noise",
+ "noperspective",
+ "numthreads",
+ "operator",
+ "out",
+ "packoffset",
+ "pass",
+ "pixelfragment",
+ "pixelshader",
+ "point",
+ "precise",
+ "printf",
+ "private",
+ "protected",
+ "public",
+ "radians",
+ "rcp",
+ "refract",
+ "register",
+ "reinterpret_cast",
+ "return",
+ "row_major",
+ "rsqrt",
+ "sample",
+ "sampler",
+ "sampler1D",
+ "sampler2D",
+ "sampler3D",
+ "samplerCUBE",
+ "sampler_state",
+ "saturate",
+ "shared",
+ "short",
+ "signed",
+ "sincos",
+ "sizeof",
+ "snorm",
+ "stateblock",
+ "stateblock_state",
+ "static",
+ "static_cast",
+ "string",
+ "struct",
+ "switch",
+ "tbuffer",
+ "technique",
+ "technique10",
+ "technique11",
+ "template",
+ "tex1D",
+ "tex1Dbias",
+ "tex1Dgrad",
+ "tex1Dlod",
+ "tex1Dproj",
+ "tex2D",
+ "tex2Dbias",
+ "tex2Dgrad",
+ "tex2Dlod",
+ "tex2Dproj",
+ "tex3D",
+ "tex3Dbias",
+ "tex3Dgrad",
+ "tex3Dlod",
+ "tex3Dproj",
+ "texCUBE",
+ "texCUBEbias",
+ "texCUBEgrad",
+ "texCUBElod",
+ "texCUBEproj",
+ "texture",
+ "texture1D",
+ "texture1DArray",
+ "texture2D",
+ "texture2DArray",
+ "texture2DMS",
+ "texture2DMSArray",
+ "texture3D",
+ "textureCube",
+ "textureCubeArray",
+ "this",
+ "throw",
+ "transpose",
+ "triangle",
+ "triangleadj",
+ "true",
+ "try",
+ "typedef",
+ "typename",
+ "uint",
+ "uint1",
+ "uint1x1",
+ "uint1x2",
+ "uint1x3",
+ "uint1x4",
+ "uint2",
+ "uint2x1",
+ "uint2x2",
+ "uint2x3",
+ "uint2x4",
+ "uint3",
+ "uint3x1",
+ "uint3x2",
+ "uint3x3",
+ "uint3x4",
+ "uint4",
+ "uint4x1",
+ "uint4x2",
+ "uint4x3",
+ "uint4x4",
+ "uniform",
+ "union",
+ "unorm",
+ "unroll",
+ "unsigned",
+ "using",
+ "vector",
+ "vertexfragment",
+ "vertexshader",
+ "virtual",
+ "void",
+ "volatile",
+ "while",
+};
// This list is used for a binary search and must be kept in sorted order.
-const char* kReservedKeywordsMSL[] = {"access",
- "alignas",
- "alignof",
- "and",
- "and_eq",
- "array",
- "array_ref",
- "as_type",
- "asm",
- "atomic",
- "atomic_bool",
- "atomic_int",
- "atomic_uint",
- "auto",
- "bitand",
- "bitor",
- "bool",
- "bool2",
- "bool3",
- "bool4",
- "break",
- "buffer",
- "case",
- "catch",
- "char",
- "char16_t",
- "char2",
- "char3",
- "char32_t",
- "char4",
- "class",
- "compl",
- "const",
- "const_cast",
- "const_reference",
- "constant",
- "constexpr",
- "continue",
- "decltype",
- "default",
- "delete",
- "depth2d",
- "depth2d_array",
- "depth2d_ms",
- "depth2d_ms_array",
- "depthcube",
- "depthcube_array",
- "device",
- "discard_fragment",
- "do",
- "double",
- "dynamic_cast",
- "else",
- "enum",
- "explicit",
- "extern",
- "false",
- "final",
- "float",
- "float2",
- "float2x2",
- "float2x3",
- "float2x4",
- "float3",
- "float3x2",
- "float3x3",
- "float3x4",
- "float4",
- "float4x2",
- "float4x3",
- "float4x4",
- "for",
- "fragment",
- "friend",
- "goto",
- "half",
- "half2",
- "half2x2",
- "half2x3",
- "half2x4",
- "half3",
- "half3x2",
- "half3x3",
- "half3x4",
- "half4",
- "half4x2",
- "half4x3",
- "half4x4",
- "if",
- "imageblock",
- "inline",
- "int",
- "int16_t",
- "int2",
- "int3",
- "int32_t",
- "int4",
- "int64_t",
- "int8_t",
- "kernel",
- "long",
- "long2",
- "long3",
- "long4",
- "main",
- "metal",
- "mutable",
- "namespace",
- "new",
- "noexcept",
- "not",
- "not_eq",
- "nullptr",
- "operator",
- "or",
- "or_eq",
- "override",
- "packed_bool2",
- "packed_bool3",
- "packed_bool4",
- "packed_char2",
- "packed_char3",
- "packed_char4",
- "packed_float2",
- "packed_float3",
- "packed_float4",
- "packed_half2",
- "packed_half3",
- "packed_half4",
- "packed_int2",
- "packed_int3",
- "packed_int4",
- "packed_short2",
- "packed_short3",
- "packed_short4",
- "packed_uchar2",
- "packed_uchar3",
- "packed_uchar4",
- "packed_uint2",
- "packed_uint3",
- "packed_uint4",
- "packed_ushort2",
- "packed_ushort3",
- "packed_ushort4",
- "patch_control_point",
- "private",
- "protected",
- "ptrdiff_t",
- "public",
- "r16snorm",
- "r16unorm",
- "r8unorm",
- "reference",
- "register",
- "reinterpret_cast",
- "return",
- "rg11b10f",
- "rg16snorm",
- "rg16unorm",
- "rg8snorm",
- "rg8unorm",
- "rgb10a2",
- "rgb9e5",
- "rgba16snorm",
- "rgba16unorm",
- "rgba8snorm",
- "rgba8unorm",
- "sampler",
- "short",
- "short2",
- "short3",
- "short4",
- "signed",
- "size_t",
- "sizeof",
- "srgba8unorm",
- "static",
- "static_assert",
- "static_cast",
- "struct",
- "switch",
- "template",
- "texture",
- "texture1d",
- "texture1d_array",
- "texture2d",
- "texture2d_array",
- "texture2d_ms",
- "texture2d_ms_array",
- "texture3d",
- "texture_buffer",
- "texturecube",
- "texturecube_array",
- "this",
- "thread",
- "thread_local",
- "threadgroup",
- "threadgroup_imageblock",
- "throw",
- "true",
- "try",
- "typedef",
- "typeid",
- "typename",
- "uchar",
- "uchar2",
- "uchar3",
- "uchar4",
- "uint",
- "uint16_t",
- "uint2",
- "uint3",
- "uint32_t",
- "uint4",
- "uint64_t",
- "uint8_t",
- "ulong2",
- "ulong3",
- "ulong4",
- "uniform",
- "union",
- "unsigned",
- "ushort",
- "ushort2",
- "ushort3",
- "ushort4",
- "using",
- "vec",
- "vertex",
- "virtual",
- "void",
- "volatile",
- "wchar_t",
- "while",
- "xor",
- "xor_eq"};
+const char* kReservedKeywordsMSL[] = {
+ "HUGE_VALF",
+ "HUGE_VALH",
+ "INFINITY",
+ "MAXFLOAT",
+ "MAXHALF",
+ "M_1_PI_F",
+ "M_1_PI_H",
+ "M_2_PI_F",
+ "M_2_PI_H",
+ "M_2_SQRTPI_F",
+ "M_2_SQRTPI_H",
+ "M_E_F",
+ "M_E_H",
+ "M_LN10_F",
+ "M_LN10_H",
+ "M_LN2_F",
+ "M_LN2_H",
+ "M_LOG10E_F",
+ "M_LOG10E_H",
+ "M_LOG2E_F",
+ "M_LOG2E_H",
+ "M_PI_2_F",
+ "M_PI_2_H",
+ "M_PI_4_F",
+ "M_PI_4_H",
+ "M_PI_F",
+ "M_PI_H",
+ "M_SQRT1_2_F",
+ "M_SQRT1_2_H",
+ "M_SQRT2_F",
+ "M_SQRT2_H",
+ "NAN",
+ "access",
+ "alignas",
+ "alignof",
+ "and",
+ "and_eq",
+ "array",
+ "array_ref",
+ "as_type",
+ "asm",
+ "atomic",
+ "atomic_bool",
+ "atomic_int",
+ "atomic_uint",
+ "auto",
+ "bitand",
+ "bitor",
+ "bool",
+ "bool2",
+ "bool3",
+ "bool4",
+ "break",
+ "buffer",
+ "case",
+ "catch",
+ "char",
+ "char16_t",
+ "char2",
+ "char3",
+ "char32_t",
+ "char4",
+ "class",
+ "compl",
+ "const",
+ "const_cast",
+ "const_reference",
+ "constant",
+ "constexpr",
+ "continue",
+ "decltype",
+ "default",
+ "delete",
+ "depth2d",
+ "depth2d_array",
+ "depth2d_ms",
+ "depth2d_ms_array",
+ "depthcube",
+ "depthcube_array",
+ "device",
+ "discard_fragment",
+ "do",
+ "double",
+ "dynamic_cast",
+ "else",
+ "enum",
+ "explicit",
+ "extern",
+ "false",
+ "final",
+ "float",
+ "float2",
+ "float2x2",
+ "float2x3",
+ "float2x4",
+ "float3",
+ "float3x2",
+ "float3x3",
+ "float3x4",
+ "float4",
+ "float4x2",
+ "float4x3",
+ "float4x4",
+ "for",
+ "fragment",
+ "friend",
+ "goto",
+ "half",
+ "half2",
+ "half2x2",
+ "half2x3",
+ "half2x4",
+ "half3",
+ "half3x2",
+ "half3x3",
+ "half3x4",
+ "half4",
+ "half4x2",
+ "half4x3",
+ "half4x4",
+ "if",
+ "imageblock",
+ "infinity",
+ "inline",
+ "int",
+ "int16_t",
+ "int2",
+ "int3",
+ "int32_t",
+ "int4",
+ "int64_t",
+ "int8_t",
+ "kernel",
+ "long",
+ "long2",
+ "long3",
+ "long4",
+ "main",
+ "metal",
+ "mutable",
+ "namespace",
+ "new",
+ "noexcept",
+ "not",
+ "not_eq",
+ "nullptr",
+ "operator",
+ "or",
+ "or_eq",
+ "override",
+ "packed_bool2",
+ "packed_bool3",
+ "packed_bool4",
+ "packed_char2",
+ "packed_char3",
+ "packed_char4",
+ "packed_float2",
+ "packed_float3",
+ "packed_float4",
+ "packed_half2",
+ "packed_half3",
+ "packed_half4",
+ "packed_int2",
+ "packed_int3",
+ "packed_int4",
+ "packed_short2",
+ "packed_short3",
+ "packed_short4",
+ "packed_uchar2",
+ "packed_uchar3",
+ "packed_uchar4",
+ "packed_uint2",
+ "packed_uint3",
+ "packed_uint4",
+ "packed_ushort2",
+ "packed_ushort3",
+ "packed_ushort4",
+ "patch_control_point",
+ "private",
+ "protected",
+ "ptrdiff_t",
+ "public",
+ "r16snorm",
+ "r16unorm",
+ "r8unorm",
+ "reference",
+ "register",
+ "reinterpret_cast",
+ "return",
+ "rg11b10f",
+ "rg16snorm",
+ "rg16unorm",
+ "rg8snorm",
+ "rg8unorm",
+ "rgb10a2",
+ "rgb9e5",
+ "rgba16snorm",
+ "rgba16unorm",
+ "rgba8snorm",
+ "rgba8unorm",
+ "sampler",
+ "short",
+ "short2",
+ "short3",
+ "short4",
+ "signed",
+ "size_t",
+ "sizeof",
+ "srgba8unorm",
+ "static",
+ "static_assert",
+ "static_cast",
+ "struct",
+ "switch",
+ "template",
+ "texture",
+ "texture1d",
+ "texture1d_array",
+ "texture2d",
+ "texture2d_array",
+ "texture2d_ms",
+ "texture2d_ms_array",
+ "texture3d",
+ "texture_buffer",
+ "texturecube",
+ "texturecube_array",
+ "this",
+ "thread",
+ "thread_local",
+ "threadgroup",
+ "threadgroup_imageblock",
+ "throw",
+ "true",
+ "try",
+ "typedef",
+ "typeid",
+ "typename",
+ "uchar",
+ "uchar2",
+ "uchar3",
+ "uchar4",
+ "uint",
+ "uint16_t",
+ "uint2",
+ "uint3",
+ "uint32_t",
+ "uint4",
+ "uint64_t",
+ "uint8_t",
+ "ulong2",
+ "ulong3",
+ "ulong4",
+ "uniform",
+ "union",
+ "unsigned",
+ "ushort",
+ "ushort2",
+ "ushort3",
+ "ushort4",
+ "using",
+ "vec",
+ "vertex",
+ "virtual",
+ "void",
+ "volatile",
+ "wchar_t",
+ "while",
+ "xor",
+ "xor_eq",
+};
} // namespace
diff --git a/src/transform/renamer_test.cc b/src/transform/renamer_test.cc
index 1c57b4a..9db7e3b 100644
--- a/src/transform/renamer_test.cc
+++ b/src/transform/renamer_test.cc
@@ -847,248 +847,286 @@
"volatile"));
// "while" // WGSL reserved keyword
-INSTANTIATE_TEST_SUITE_P(RenamerTestMsl,
- RenamerTestMsl,
- testing::Values(
- // c++14 spec
- "alignas",
- "alignof",
- "and",
- "and_eq",
- // "asm", // Also reserved in WGSL
- "auto",
- "bitand",
- "bitor",
- // "bool", // Also used in WGSL
- // "break", // Also used in WGSL
- // "case", // Also used in WGSL
- "catch",
- "char",
- "char16_t",
- "char32_t",
- "class",
- "compl",
- // "const", // Also used in WGSL
- "const_cast",
- "constexpr",
- // "continue", // Also used in WGSL
- "decltype",
- // "default", // Also used in WGSL
- "delete",
- // "do", // Also used in WGSL
- "double",
- "dynamic_cast",
- // "else", // Also used in WGSL
- // "enum", // Also used in WGSL
- "explicit",
- "extern",
- // "false", // Also used in WGSL
- "final",
- "float",
- // "for", // Also used in WGSL
- "friend",
- "goto",
- // "if", // Also used in WGSL
- "inline",
- "int",
- "long",
- "mutable",
- "namespace",
- "new",
- "noexcept",
- "not",
- "not_eq",
- "nullptr",
- "operator",
- "or",
- "or_eq",
- "override",
- // "private", // Also used in WGSL
- "protected",
- "public",
- "register",
- "reinterpret_cast",
- // "return", // Also used in WGSL
- "short",
- "signed",
- "sizeof",
- "static",
- "static_assert",
- "static_cast",
- // "struct", // Also used in WGSL
- // "switch", // Also used in WGSL
- "template",
- "this",
- "thread_local",
- "throw",
- // "true", // Also used in WGSL
- "try",
- // "typedef", // Also used in WGSL
- "typeid",
- "typename",
- "union",
- "unsigned",
- // "using", // WGSL reserved keyword
- "virtual",
- // "void", // Also used in WGSL
- "volatile",
- "wchar_t",
- // "while", // WGSL reserved keyword
- "xor",
- "xor_eq",
+INSTANTIATE_TEST_SUITE_P(
+ RenamerTestMsl,
+ RenamerTestMsl,
+ testing::Values(
+ // c++14 spec
+ "alignas",
+ "alignof",
+ "and",
+ "and_eq",
+ // "asm", // Also reserved in WGSL
+ "auto",
+ "bitand",
+ "bitor",
+ // "bool", // Also used in WGSL
+ // "break", // Also used in WGSL
+ // "case", // Also used in WGSL
+ "catch",
+ "char",
+ "char16_t",
+ "char32_t",
+ "class",
+ "compl",
+ // "const", // Also used in WGSL
+ "const_cast",
+ "constexpr",
+ // "continue", // Also used in WGSL
+ "decltype",
+ // "default", // Also used in WGSL
+ "delete",
+ // "do", // Also used in WGSL
+ "double",
+ "dynamic_cast",
+ // "else", // Also used in WGSL
+ // "enum", // Also used in WGSL
+ "explicit",
+ "extern",
+ // "false", // Also used in WGSL
+ "final",
+ "float",
+ // "for", // Also used in WGSL
+ "friend",
+ "goto",
+ // "if", // Also used in WGSL
+ "inline",
+ "int",
+ "long",
+ "mutable",
+ "namespace",
+ "new",
+ "noexcept",
+ "not",
+ "not_eq",
+ "nullptr",
+ "operator",
+ "or",
+ "or_eq",
+ "override",
+ // "private", // Also used in WGSL
+ "protected",
+ "public",
+ "register",
+ "reinterpret_cast",
+ // "return", // Also used in WGSL
+ "short",
+ "signed",
+ "sizeof",
+ "static",
+ "static_assert",
+ "static_cast",
+ // "struct", // Also used in WGSL
+ // "switch", // Also used in WGSL
+ "template",
+ "this",
+ "thread_local",
+ "throw",
+ // "true", // Also used in WGSL
+ "try",
+ // "typedef", // Also used in WGSL
+ "typeid",
+ "typename",
+ "union",
+ "unsigned",
+ // "using", // WGSL reserved keyword
+ "virtual",
+ // "void", // Also used in WGSL
+ "volatile",
+ "wchar_t",
+ // "while", // WGSL reserved keyword
+ "xor",
+ "xor_eq",
- // Metal Spec
- "access",
- // "array", // Also used in WGSL
- "array_ref",
- "as_type",
- // "atomic", // Also used in WGSL
- "atomic_bool",
- "atomic_int",
- "atomic_uint",
- "bool2",
- "bool3",
- "bool4",
- "buffer",
- "char2",
- "char3",
- "char4",
- "const_reference",
- "constant",
- "depth2d",
- "depth2d_array",
- "depth2d_ms",
- "depth2d_ms_array",
- "depthcube",
- "depthcube_array",
- "device",
- "discard_fragment",
- "float2",
- "float2x2",
- "float2x3",
- "float2x4",
- "float3",
- "float3x2",
- "float3x3",
- "float3x4",
- "float4",
- "float4x2",
- "float4x3",
- "float4x4",
- "fragment",
- "half",
- "half2",
- "half2x2",
- "half2x3",
- "half2x4",
- "half3",
- "half3x2",
- "half3x3",
- "half3x4",
- "half4",
- "half4x2",
- "half4x3",
- "half4x4",
- "imageblock",
- "int16_t",
- "int2",
- "int3",
- "int32_t",
- "int4",
- "int64_t",
- "int8_t",
- "kernel",
- "long2",
- "long3",
- "long4",
- "main", // No functions called main
- "metal", // The namespace
- "packed_bool2",
- "packed_bool3",
- "packed_bool4",
- "packed_char2",
- "packed_char3",
- "packed_char4",
- "packed_float2",
- "packed_float3",
- "packed_float4",
- "packed_half2",
- "packed_half3",
- "packed_half4",
- "packed_int2",
- "packed_int3",
- "packed_int4",
- "packed_short2",
- "packed_short3",
- "packed_short4",
- "packed_uchar2",
- "packed_uchar3",
- "packed_uchar4",
- "packed_uint2",
- "packed_uint3",
- "packed_uint4",
- "packed_ushort2",
- "packed_ushort3",
- "packed_ushort4",
- "patch_control_point",
- "ptrdiff_t",
- "r16snorm",
- "r16unorm",
- // "r8unorm", // Also used in WGSL
- "reference",
- "rg11b10f",
- "rg16snorm",
- "rg16unorm",
- // "rg8snorm", // Also used in WGSL
- // "rg8unorm", // Also used in WGSL
- "rgb10a2",
- "rgb9e5",
- "rgba16snorm",
- "rgba16unorm",
- // "rgba8snorm", // Also used in WGSL
- // "rgba8unorm", // Also used in WGSL
- // "sampler", // Also used in WGSL
- "short2",
- "short3",
- "short4",
- "size_t",
- "srgba8unorm",
- "texture",
- "texture1d",
- "texture1d_array",
- "texture2d",
- "texture2d_array",
- "texture2d_ms",
- "texture2d_ms_array",
- "texture3d",
- "texture_buffer",
- "texturecube",
- "texturecube_array",
- "thread",
- "threadgroup",
- "threadgroup_imageblock",
- "uchar",
- "uchar2",
- "uchar3",
- "uchar4",
- "uint",
- "uint16_t",
- "uint2",
- "uint3",
- "uint32_t",
- "uint4",
- "uint64_t",
- "uint8_t",
- "ulong2",
- "ulong3",
- "ulong4",
- // "uniform", // Also used in WGSL
- "ushort",
- "ushort2",
- "ushort3",
- "ushort4",
- // "vec", // WGSL reserved keyword
- "vertex"));
+ // Metal Spec
+ "access",
+ // "array", // Also used in WGSL
+ "array_ref",
+ "as_type",
+ // "atomic", // Also used in WGSL
+ "atomic_bool",
+ "atomic_int",
+ "atomic_uint",
+ "bool2",
+ "bool3",
+ "bool4",
+ "buffer",
+ "char2",
+ "char3",
+ "char4",
+ "const_reference",
+ "constant",
+ "depth2d",
+ "depth2d_array",
+ "depth2d_ms",
+ "depth2d_ms_array",
+ "depthcube",
+ "depthcube_array",
+ "device",
+ "discard_fragment",
+ "float2",
+ "float2x2",
+ "float2x3",
+ "float2x4",
+ "float3",
+ "float3x2",
+ "float3x3",
+ "float3x4",
+ "float4",
+ "float4x2",
+ "float4x3",
+ "float4x4",
+ "fragment",
+ "half",
+ "half2",
+ "half2x2",
+ "half2x3",
+ "half2x4",
+ "half3",
+ "half3x2",
+ "half3x3",
+ "half3x4",
+ "half4",
+ "half4x2",
+ "half4x3",
+ "half4x4",
+ "imageblock",
+ "int16_t",
+ "int2",
+ "int3",
+ "int32_t",
+ "int4",
+ "int64_t",
+ "int8_t",
+ "kernel",
+ "long2",
+ "long3",
+ "long4",
+ "main", // No functions called main
+ "metal", // The namespace
+ "packed_bool2",
+ "packed_bool3",
+ "packed_bool4",
+ "packed_char2",
+ "packed_char3",
+ "packed_char4",
+ "packed_float2",
+ "packed_float3",
+ "packed_float4",
+ "packed_half2",
+ "packed_half3",
+ "packed_half4",
+ "packed_int2",
+ "packed_int3",
+ "packed_int4",
+ "packed_short2",
+ "packed_short3",
+ "packed_short4",
+ "packed_uchar2",
+ "packed_uchar3",
+ "packed_uchar4",
+ "packed_uint2",
+ "packed_uint3",
+ "packed_uint4",
+ "packed_ushort2",
+ "packed_ushort3",
+ "packed_ushort4",
+ "patch_control_point",
+ "ptrdiff_t",
+ "r16snorm",
+ "r16unorm",
+ // "r8unorm", // Also used in WGSL
+ "reference",
+ "rg11b10f",
+ "rg16snorm",
+ "rg16unorm",
+ // "rg8snorm", // Also used in WGSL
+ // "rg8unorm", // Also used in WGSL
+ "rgb10a2",
+ "rgb9e5",
+ "rgba16snorm",
+ "rgba16unorm",
+ // "rgba8snorm", // Also used in WGSL
+ // "rgba8unorm", // Also used in WGSL
+ // "sampler", // Also used in WGSL
+ "short2",
+ "short3",
+ "short4",
+ "size_t",
+ "srgba8unorm",
+ "texture",
+ "texture1d",
+ "texture1d_array",
+ "texture2d",
+ "texture2d_array",
+ "texture2d_ms",
+ "texture2d_ms_array",
+ "texture3d",
+ "texture_buffer",
+ "texturecube",
+ "texturecube_array",
+ "thread",
+ "threadgroup",
+ "threadgroup_imageblock",
+ "uchar",
+ "uchar2",
+ "uchar3",
+ "uchar4",
+ "uint",
+ "uint16_t",
+ "uint2",
+ "uint3",
+ "uint32_t",
+ "uint4",
+ "uint64_t",
+ "uint8_t",
+ "ulong2",
+ "ulong3",
+ "ulong4",
+ // "uniform", // Also used in WGSL
+ "ushort",
+ "ushort2",
+ "ushort3",
+ "ushort4",
+ // "vec", // WGSL reserved keyword
+ "vertex",
+
+ // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+ // Table 6.5. Constants for single-precision floating-point math
+ // functions
+ "MAXFLOAT",
+ "HUGE_VALF",
+ "INFINITY",
+ "infinity",
+ "NAN",
+ "M_E_F",
+ "M_LOG2E_F",
+ "M_LOG10E_F",
+ "M_LN2_F",
+ "M_LN10_F",
+ "M_PI_F",
+ "M_PI_2_F",
+ "M_PI_4_F",
+ "M_1_PI_F",
+ "M_2_PI_F",
+ "M_2_SQRTPI_F",
+ "M_SQRT2_F",
+ "M_SQRT1_2_F",
+ "MAXHALF",
+ "HUGE_VALH",
+ "M_E_H",
+ "M_LOG2E_H",
+ "M_LOG10E_H",
+ "M_LN2_H",
+ "M_LN10_H",
+ "M_PI_H",
+ "M_PI_2_H",
+ "M_PI_4_H",
+ "M_1_PI_H",
+ "M_2_PI_H",
+ "M_2_SQRTPI_H",
+ "M_SQRT2_H",
+ "M_SQRT1_2_H"));
} // namespace
} // namespace transform
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 06d9221..c93caa7 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -15,6 +15,7 @@
#include "src/writer/hlsl/generator_impl.h"
#include <algorithm>
+#include <cmath>
#include <iomanip>
#include <set>
#include <utility>
@@ -2474,7 +2475,14 @@
if (auto* l = lit->As<ast::BoolLiteral>()) {
out << (l->IsTrue() ? "true" : "false");
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
- out << FloatToString(fl->value()) << "f";
+ if (std::isinf(fl->value())) {
+ out << (fl->value() >= 0 ? "asfloat(0x7f800000u)"
+ : "asfloat(0xff800000u)");
+ } else if (std::isnan(fl->value())) {
+ out << "asfloat(0x7fc00000u)";
+ } else {
+ out << FloatToString(fl->value()) << "f";
+ }
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
out << sl->value();
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 67e4da7..6601bdf 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -15,6 +15,7 @@
#include "src/writer/msl/generator_impl.h"
#include <algorithm>
+#include <cmath>
#include <iomanip>
#include <utility>
#include <vector>
@@ -1086,7 +1087,13 @@
if (auto* l = lit->As<ast::BoolLiteral>()) {
out << (l->IsTrue() ? "true" : "false");
} else if (auto* fl = lit->As<ast::FloatLiteral>()) {
- out << FloatToString(fl->value()) << "f";
+ if (std::isinf(fl->value())) {
+ out << (fl->value() >= 0 ? "INFINITY" : "-INFINITY");
+ } else if (std::isnan(fl->value())) {
+ out << "NAN";
+ } else {
+ out << FloatToString(fl->value()) << "f";
+ }
} else if (auto* sl = lit->As<ast::SintLiteral>()) {
out << sl->value();
} else if (auto* ul = lit->As<ast::UintLiteral>()) {
diff --git a/test/bug/tint/943.spvasm.expected.msl b/test/bug/tint/943.spvasm.expected.msl
index 551f911..31a04fe 100644
--- a/test/bug/tint/943.spvasm.expected.msl
+++ b/test/bug/tint/943.spvasm.expected.msl
@@ -1,32 +1,506 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct Uniforms {
+ /* 0x0000 */ float tint_symbol;
+ /* 0x0004 */ int8_t tint_pad[12];
+ /* 0x0010 */ packed_int3 aShape;
+ /* 0x001c */ int8_t tint_pad_1[4];
+ /* 0x0020 */ packed_int3 bShape;
+ /* 0x002c */ int8_t tint_pad_2[4];
+ /* 0x0030 */ packed_int3 outShape;
+ /* 0x003c */ int8_t tint_pad_3[4];
+ /* 0x0040 */ packed_int2 outShapeStrides;
+};
+struct ssbOut {
+ /* 0x0000 */ float result[1];
+};
+struct ssbA {
+ /* 0x0000 */ float A[1];
+};
+struct ssbB {
+ /* 0x0000 */ float B[1];
+};
+struct tint_array_wrapper_1 {
+ float arr[64];
+};
+struct tint_array_wrapper {
+ tint_array_wrapper_1 arr[64];
+};
+struct tint_array_wrapper_3 {
+ float arr[1];
+};
+struct tint_array_wrapper_2 {
+ tint_array_wrapper_3 arr[64];
+};
+struct tint_array_wrapper_4 {
+ tint_array_wrapper_3 arr[1];
+};
+bool coordsInBounds_vi2_vi2_(thread int2* const coord, thread int2* const shape) {
+ bool x_87 = false;
+ bool x_88_phi = false;
+ int2 const x_76 = *(coord);
+ bool const x_81 = all((x_76 >= int2(0, 0)));
+ x_88_phi = x_81;
+ if (x_81) {
+ int2 const x_84 = *(coord);
+ int2 const x_85 = *(shape);
+ x_87 = all((x_84 < x_85));
+ x_88_phi = x_87;
+ }
+ bool const x_88 = x_88_phi;
+ return x_88;
+}
-Validation Failure:
+float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread int* const row, thread int* const col, thread int* const tint_symbol_5, thread int* const tint_symbol_6, thread int* const tint_symbol_7) {
+ int batchASize = 0;
+ int2 param_10 = 0;
+ int2 param_11 = 0;
+ float x_430 = 0.0f;
+ int const x_417 = x_48.aShape.y;
+ int const x_419 = x_48.aShape.z;
+ batchASize = (x_417 * x_419);
+ int const x_421 = *(row);
+ int const x_422 = *(col);
+ int const x_424 = *(tint_symbol_5);
+ int const x_425 = *(tint_symbol_6);
+ param_10 = int2(x_421, x_422);
+ param_11 = int2(x_424, x_425);
+ bool const x_429 = coordsInBounds_vi2_vi2_(&(param_10), &(param_11));
+ if (x_429) {
+ int const x_438 = *(tint_symbol_7);
+ int const x_439 = batchASize;
+ int const x_441 = *(row);
+ int const x_442 = *(tint_symbol_6);
+ int const x_445 = *(col);
+ float const x_448 = x_165.A[(((x_438 * x_439) + (x_441 * x_442)) + x_445)];
+ x_430 = x_448;
+ } else {
+ x_430 = 0.0f;
+ }
+ float const x_450 = x_430;
+ return x_450;
+}
-Compilation failed:
+float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread int* const row_1, thread int* const col_1, thread int* const tint_symbol_8, thread int* const tint_symbol_9, thread int* const tint_symbol_10) {
+ int batchBSize = 0;
+ int2 param_12 = 0;
+ int2 param_13 = 0;
+ float x_468 = 0.0f;
+ int const x_455 = x_48.bShape.y;
+ int const x_457 = x_48.bShape.z;
+ batchBSize = (x_455 * x_457);
+ int const x_459 = *(row_1);
+ int const x_460 = *(col_1);
+ int const x_462 = *(tint_symbol_8);
+ int const x_463 = *(tint_symbol_9);
+ param_12 = int2(x_459, x_460);
+ param_13 = int2(x_462, x_463);
+ bool const x_467 = coordsInBounds_vi2_vi2_(&(param_12), &(param_13));
+ if (x_467) {
+ int const x_475 = *(tint_symbol_10);
+ int const x_476 = batchBSize;
+ int const x_478 = *(row_1);
+ int const x_479 = *(tint_symbol_9);
+ int const x_482 = *(col_1);
+ float const x_485 = x_185.B[(((x_475 * x_476) + (x_478 * x_479)) + x_482)];
+ x_468 = x_485;
+ } else {
+ x_468 = 0.0f;
+ }
+ float const x_487 = x_468;
+ return x_487;
+}
-program_source:5:22: error: expected parameter declarator
- /* 0x0000 */ float NAN;
- ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:28: note: expanded from macro 'NAN'
-#define NAN __builtin_nanf("")
- ^
-program_source:5:22: error: expected ')'
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:28: note: expanded from macro 'NAN'
-#define NAN __builtin_nanf("")
- ^
-program_source:5:22: note: to match this '('
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_types:214:27: note: expanded from macro 'NAN'
-#define NAN __builtin_nanf("")
- ^
-program_source:494:31: warning: equality comparison with extraneous parentheses
+int getOutputFlatIndex_vi3_(constant Uniforms& x_48, thread int3* const coords) {
+ int3 const x_99 = *(coords);
+ int const x_105 = x_48.outShapeStrides.x;
+ int const x_107 = x_48.outShapeStrides.y;
+ return int(dot(float3(x_99), float3(int3(x_105, x_107, 1))));
+}
+
+void setOutput_i1_f1_(device ssbOut& x_54, thread int* const flatIndex, thread float* const value) {
+ int const x_95 = *(flatIndex);
+ float const x_96 = *(value);
+ x_54.result[x_95] = x_96;
+ return;
+}
+
+void setOutput_i1_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const d0, thread int* const d1, thread int* const d2, thread float* const value_1) {
+ int flatIndex_1 = 0;
+ int3 param = 0;
+ int param_1 = 0;
+ float param_2 = 0.0f;
+ int const x_115 = *(d0);
+ int const x_116 = *(d1);
+ int const x_117 = *(d2);
+ param = int3(x_115, x_116, x_117);
+ int const x_120 = getOutputFlatIndex_vi3_(x_48, &(param));
+ flatIndex_1 = x_120;
+ int const x_122 = flatIndex_1;
+ param_1 = x_122;
+ float const x_124 = *(value_1);
+ param_2 = x_124;
+ setOutput_i1_f1_(x_54, &(param_1), &(param_2));
+ return;
+}
+
+void mm_write_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const row_2, thread int* const col_2, thread float* const value_2, thread int* const tint_symbol_11) {
+ int3 outCoord = 0;
+ int param_14 = 0;
+ int param_15 = 0;
+ int param_16 = 0;
+ float param_17 = 0.0f;
+ int const x_491 = *(tint_symbol_11);
+ int const x_492 = *(row_2);
+ int const x_493 = *(col_2);
+ outCoord = int3(x_491, x_492, x_493);
+ int const x_496 = *(tint_symbol_11);
+ param_14 = x_496;
+ int const x_498 = *(row_2);
+ param_15 = x_498;
+ int const x_500 = *(col_2);
+ param_16 = x_500;
+ float const x_502 = *(value_2);
+ param_17 = x_502;
+ setOutput_i1_i1_i1_f1_(x_48, x_54, &(param_14), &(param_15), &(param_16), &(param_17));
+ return;
+}
+
+void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const dimAOuter, thread int* const dimInner, thread int* const dimBOuter, thread uint3* const tint_symbol_12, thread uint3* const tint_symbol_13, thread int* const tint_symbol_14, thread int* const tint_symbol_15, thread int* const tint_symbol_16, threadgroup tint_array_wrapper* const tint_symbol_17, thread int* const tint_symbol_18, threadgroup tint_array_wrapper_2* const tint_symbol_19) {
+ int tileRow = 0;
+ int tileCol = 0;
+ int globalRow = 0;
+ int globalCol = 0;
+ int numTiles = 0;
+ int innerRow = 0;
+ int innerCol = 0;
+ tint_array_wrapper_4 acc = {};
+ int tileColA = 0;
+ int tileRowB = 0;
+ int t = 0;
+ int innerRow_1 = 0;
+ int innerCol_1 = 0;
+ int inputRow = 0;
+ int inputCol = 0;
+ int param_3 = 0;
+ int param_4 = 0;
+ int innerRow_2 = 0;
+ int innerCol_2 = 0;
+ int inputRow_1 = 0;
+ int inputCol_1 = 0;
+ int param_5 = 0;
+ int param_6 = 0;
+ int k = 0;
+ int inner = 0;
+ tint_array_wrapper_3 BCached = {};
+ int innerRow_3 = 0;
+ float ACached = 0.0f;
+ int innerCol_3 = 0;
+ int innerRow_4 = 0;
+ int innerCol_4 = 0;
+ int param_7 = 0;
+ int param_8 = 0;
+ float param_9 = 0.0f;
+ uint const x_132 = (*(tint_symbol_12)).y;
+ tileRow = (as_type<int>(x_132) * 1);
+ uint const x_137 = (*(tint_symbol_12)).x;
+ tileCol = (as_type<int>(x_137) * 1);
+ uint const x_143 = (*(tint_symbol_13)).y;
+ globalRow = (as_type<int>(x_143) * 1);
+ uint const x_148 = (*(tint_symbol_13)).x;
+ globalCol = (as_type<int>(x_148) * 1);
+ int const x_152 = *(dimInner);
+ numTiles = (((x_152 - 1) / 64) + 1);
+ innerRow = 0;
+ while (true) {
+ int const x_163 = innerRow;
+ if ((x_163 < 1)) {
+ } else {
+ break;
+ }
+ innerCol = 0;
+ while (true) {
+ int const x_171 = innerCol;
+ if ((x_171 < 1)) {
+ } else {
+ break;
+ }
+ int const x_177 = innerRow;
+ int const x_178 = innerCol;
+ acc.arr[x_177].arr[x_178] = 0.0f;
+ {
+ int const x_181 = innerCol;
+ innerCol = (x_181 + 1);
+ }
+ }
+ {
+ int const x_183 = innerRow;
+ innerRow = (x_183 + 1);
+ }
+ }
+ uint const x_187 = (*(tint_symbol_12)).x;
+ tileColA = (as_type<int>(x_187) * 64);
+ uint const x_192 = (*(tint_symbol_12)).y;
+ tileRowB = (as_type<int>(x_192) * 1);
+ t = 0;
+ while (true) {
+ int const x_201 = t;
+ int const x_202 = numTiles;
+ if ((x_201 < x_202)) {
+ } else {
+ break;
+ }
+ innerRow_1 = 0;
+ while (true) {
+ int const x_210 = innerRow_1;
+ if ((x_210 < 1)) {
+ } else {
+ break;
+ }
+ innerCol_1 = 0;
+ while (true) {
+ int const x_218 = innerCol_1;
+ if ((x_218 < 64)) {
+ } else {
+ break;
+ }
+ int const x_221 = tileRow;
+ int const x_222 = innerRow_1;
+ inputRow = (x_221 + x_222);
+ int const x_225 = tileColA;
+ int const x_226 = innerCol_1;
+ inputCol = (x_225 + x_226);
+ int const x_233 = inputRow;
+ int const x_234 = inputCol;
+ int const x_235 = globalRow;
+ int const x_236 = innerRow_1;
+ int const x_238 = t;
+ int const x_240 = inputCol;
+ param_3 = (x_235 + x_236);
+ param_4 = ((x_238 * 64) + x_240);
+ float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_14, tint_symbol_15, tint_symbol_16);
+ (*(tint_symbol_17)).arr[x_233].arr[x_234] = x_244;
+ {
+ int const x_247 = innerCol_1;
+ innerCol_1 = (x_247 + 1);
+ }
+ }
+ {
+ int const x_249 = innerRow_1;
+ innerRow_1 = (x_249 + 1);
+ }
+ }
+ innerRow_2 = 0;
+ while (true) {
+ int const x_257 = innerRow_2;
+ if ((x_257 < 1)) {
+ } else {
+ break;
+ }
+ innerCol_2 = 0;
+ while (true) {
+ int const x_265 = innerCol_2;
+ if ((x_265 < 1)) {
+ } else {
+ break;
+ }
+ int const x_268 = tileRowB;
+ int const x_269 = innerRow_2;
+ inputRow_1 = (x_268 + x_269);
+ int const x_272 = tileCol;
+ int const x_273 = innerCol_2;
+ inputCol_1 = (x_272 + x_273);
+ int const x_278 = inputRow_1;
+ int const x_279 = inputCol_1;
+ int const x_280 = t;
+ int const x_282 = inputRow_1;
+ int const x_284 = globalCol;
+ int const x_285 = innerCol_2;
+ param_5 = ((x_280 * 64) + x_282);
+ param_6 = (x_284 + x_285);
+ float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_15, tint_symbol_18, tint_symbol_16);
+ (*(tint_symbol_19)).arr[x_278].arr[x_279] = x_289;
+ {
+ int const x_291 = innerCol_2;
+ innerCol_2 = (x_291 + 1);
+ }
+ }
+ {
+ int const x_293 = innerRow_2;
+ innerRow_2 = (x_293 + 1);
+ }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ k = 0;
+ while (true) {
+ int const x_302 = k;
+ if ((x_302 < 64)) {
+ } else {
+ break;
+ }
+ inner = 0;
+ while (true) {
+ int const x_310 = inner;
+ if ((x_310 < 1)) {
+ } else {
+ break;
+ }
+ int const x_314 = inner;
+ int const x_315 = k;
+ int const x_316 = tileCol;
+ int const x_317 = inner;
+ float const x_320 = (*(tint_symbol_19)).arr[x_315].arr[(x_316 + x_317)];
+ BCached.arr[x_314] = x_320;
+ {
+ int const x_322 = inner;
+ inner = (x_322 + 1);
+ }
+ }
+ innerRow_3 = 0;
+ while (true) {
+ int const x_330 = innerRow_3;
+ if ((x_330 < 1)) {
+ } else {
+ break;
+ }
+ int const x_333 = tileRow;
+ int const x_334 = innerRow_3;
+ int const x_336 = k;
+ float const x_338 = (*(tint_symbol_17)).arr[(x_333 + x_334)].arr[x_336];
+ ACached = x_338;
+ innerCol_3 = 0;
+ while (true) {
+ int const x_345 = innerCol_3;
+ if ((x_345 < 1)) {
+ } else {
+ break;
+ }
+ int const x_347 = innerRow_3;
+ int const x_348 = innerCol_3;
+ float const x_349 = ACached;
+ int const x_350 = innerCol_3;
+ float const x_352 = BCached.arr[x_350];
+ float const x_355 = acc.arr[x_347].arr[x_348];
+ acc.arr[x_347].arr[x_348] = (x_355 + (x_349 * x_352));
+ {
+ int const x_358 = innerCol_3;
+ innerCol_3 = (x_358 + 1);
+ }
+ }
+ {
+ int const x_360 = innerRow_3;
+ innerRow_3 = (x_360 + 1);
+ }
+ }
+ {
+ int const x_362 = k;
+ k = (x_362 + 1);
+ }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ {
+ int const x_364 = t;
+ t = (x_364 + 1);
+ }
+ }
+ innerRow_4 = 0;
+ while (true) {
+ int const x_372 = innerRow_4;
+ if ((x_372 < 1)) {
+ } else {
+ break;
+ }
+ innerCol_4 = 0;
+ while (true) {
+ bool x_393 = false;
+ bool x_394_phi = false;
+ int const x_380 = innerCol_4;
+ if ((x_380 < 1)) {
+ } else {
+ break;
+ }
+ int const x_382 = globalCol;
+ int const x_383 = innerCol_4;
+ int const x_385 = *(dimBOuter);
+ bool const x_386 = ((x_382 + x_383) < x_385);
+ x_394_phi = x_386;
+ if (x_386) {
+ int const x_389 = globalRow;
+ int const x_390 = innerRow_4;
+ int const x_392 = *(dimAOuter);
+ x_393 = ((x_389 + x_390) < x_392);
+ x_394_phi = x_393;
+ }
+ bool const x_394 = x_394_phi;
+ if (x_394) {
+ int const x_397 = globalRow;
+ int const x_398 = innerRow_4;
+ int const x_400 = globalCol;
+ int const x_401 = innerCol_4;
+ int const x_403 = innerRow_4;
+ int const x_404 = innerCol_4;
+ param_7 = (x_397 + x_398);
+ param_8 = (x_400 + x_401);
+ float const x_409 = acc.arr[x_403].arr[x_404];
+ param_9 = x_409;
+ mm_write_i1_i1_f1_(x_48, x_54, &(param_7), &(param_8), &(param_9), tint_symbol_16);
+ }
+ {
+ int const x_411 = innerCol_4;
+ innerCol_4 = (x_411 + 1);
+ }
+ }
+ {
+ int const x_413 = innerRow_4;
+ innerRow_4 = (x_413 + 1);
+ }
+ }
+ return;
+}
+
+void main_1(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const tint_symbol_20, thread int* const tint_symbol_21, thread int* const tint_symbol_22, thread uint3* const tint_symbol_23, thread int* const tint_symbol_24, thread uint3* const tint_symbol_25, threadgroup tint_array_wrapper* const tint_symbol_26, threadgroup tint_array_wrapper_2* const tint_symbol_27) {
+ int param_18 = 0;
+ int param_19 = 0;
+ int param_20 = 0;
+ int const x_67 = x_48.aShape.y;
+ *(tint_symbol_20) = x_67;
+ int const x_71 = x_48.aShape.z;
+ *(tint_symbol_21) = x_71;
+ int const x_75 = x_48.bShape.z;
+ *(tint_symbol_22) = x_75;
+ uint const x_505 = (*(tint_symbol_23)).z;
+ *(tint_symbol_24) = as_type<int>(x_505);
+ int const x_508 = *(tint_symbol_20);
+ param_18 = x_508;
+ int const x_510 = *(tint_symbol_21);
+ param_19 = x_510;
+ int const x_512 = *(tint_symbol_22);
+ param_20 = x_512;
+ mm_matMul_i1_i1_i1_(x_48, x_165, x_185, x_54, &(param_18), &(param_19), &(param_20), tint_symbol_25, tint_symbol_23, tint_symbol_20, tint_symbol_21, tint_symbol_24, tint_symbol_26, tint_symbol_22, tint_symbol_27);
+ return;
+}
+
+kernel void tint_symbol_1(uint3 gl_LocalInvocationID_param [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& x_48 [[buffer(3)]], const device ssbA& x_165 [[buffer(1)]], const device ssbB& x_185 [[buffer(2)]], device ssbOut& x_54 [[buffer(0)]]) {
+ threadgroup tint_array_wrapper tint_symbol_28;
+ threadgroup tint_array_wrapper_2 tint_symbol_29;
+ thread uint3 tint_symbol_30 = 0u;
+ thread uint3 tint_symbol_31 = 0u;
+ thread int tint_symbol_32 = 0;
+ thread int tint_symbol_33 = 0;
+ thread int tint_symbol_34 = 0;
+ thread int tint_symbol_35 = 0;
if ((local_invocation_index == 0u)) {
- ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:494:31: note: remove extraneous parentheses around the comparison to silence this warning
- if ((local_invocation_index == 0u)) {
- ~ ^ ~
-program_source:494:31: note: use '=' to turn this equality comparison into an assignment
- if ((local_invocation_index == 0u)) {
- ^~
- =
+ tint_array_wrapper const tint_symbol_3 = {.arr={}};
+ tint_symbol_28 = tint_symbol_3;
+ tint_array_wrapper_2 const tint_symbol_4 = {.arr={}};
+ tint_symbol_29 = tint_symbol_4;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_symbol_30 = gl_LocalInvocationID_param;
+ tint_symbol_31 = gl_GlobalInvocationID_param;
+ main_1(x_48, x_165, x_185, x_54, &(tint_symbol_32), &(tint_symbol_33), &(tint_symbol_34), &(tint_symbol_31), &(tint_symbol_35), &(tint_symbol_30), &(tint_symbol_28), &(tint_symbol_29));
+ return;
+}
+
diff --git a/test/bug/tint/951.spvasm b/test/bug/tint/951.spvasm
new file mode 100644
index 0000000..e5119f1
--- /dev/null
+++ b/test/bug/tint/951.spvasm
@@ -0,0 +1,158 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 10
+; Bound: 86
+; Schema: 0
+ OpCapability Shader
+ %1 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
+ OpExecutionMode %main LocalSize 128 1 1
+ OpSource GLSL 450
+ OpName %main "main"
+ OpName %setOutput_i1_f1_ "setOutput(i1;f1;"
+ OpName %flatIndex "flatIndex"
+ OpName %value "value"
+ OpName %getAAtOutCoords_ "getAAtOutCoords("
+ OpName %unaryOperation_f1_ "unaryOperation(f1;"
+ OpName %a "a"
+ OpName %ssbOut "ssbOut"
+ OpMemberName %ssbOut 0 "result"
+ OpName %_ ""
+ OpName %ssbA "ssbA"
+ OpMemberName %ssbA 0 "A"
+ OpName %__0 ""
+ OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
+ OpName %index "index"
+ OpName %Uniforms "Uniforms"
+ OpMemberName %Uniforms 0 "NAN"
+ OpMemberName %Uniforms 1 "aShape"
+ OpMemberName %Uniforms 2 "outShape"
+ OpMemberName %Uniforms 3 "outShapeStrides"
+ OpMemberName %Uniforms 4 "size"
+ OpName %__1 ""
+ OpName %a_0 "a"
+ OpName %param "param"
+ OpName %param_0 "param"
+ OpName %param_1 "param"
+ OpDecorate %_runtimearr_float ArrayStride 4
+ OpMemberDecorate %ssbOut 0 NonReadable
+ OpMemberDecorate %ssbOut 0 Offset 0
+ OpDecorate %ssbOut BufferBlock
+ OpDecorate %_ DescriptorSet 0
+ OpDecorate %_ Binding 0
+ OpDecorate %_runtimearr_float_0 ArrayStride 4
+ OpMemberDecorate %ssbA 0 NonWritable
+ OpMemberDecorate %ssbA 0 Offset 0
+ OpDecorate %ssbA BufferBlock
+ OpDecorate %__0 DescriptorSet 0
+ OpDecorate %__0 Binding 1
+ OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+ OpMemberDecorate %Uniforms 0 Offset 0
+ OpMemberDecorate %Uniforms 1 Offset 4
+ OpMemberDecorate %Uniforms 2 Offset 8
+ OpMemberDecorate %Uniforms 3 Offset 12
+ OpMemberDecorate %Uniforms 4 Offset 16
+ OpDecorate %Uniforms Block
+ OpDecorate %__1 DescriptorSet 0
+ OpDecorate %__1 Binding 2
+ OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %_ptr_Function_int = OpTypePointer Function %int
+ %float = OpTypeFloat 32
+ %_ptr_Function_float = OpTypePointer Function %float
+ %10 = OpTypeFunction %void %_ptr_Function_int %_ptr_Function_float
+ %15 = OpTypeFunction %float
+ %18 = OpTypeFunction %float %_ptr_Function_float
+ %_runtimearr_float = OpTypeRuntimeArray %float
+ %ssbOut = OpTypeStruct %_runtimearr_float
+ %_ptr_Uniform_ssbOut = OpTypePointer Uniform %ssbOut
+ %_ = OpVariable %_ptr_Uniform_ssbOut Uniform
+ %int_0 = OpConstant %int 0
+ %_ptr_Uniform_float = OpTypePointer Uniform %float
+ %_runtimearr_float_0 = OpTypeRuntimeArray %float
+ %ssbA = OpTypeStruct %_runtimearr_float_0
+ %_ptr_Uniform_ssbA = OpTypePointer Uniform %ssbA
+ %__0 = OpVariable %_ptr_Uniform_ssbA Uniform
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %uint_0 = OpConstant %uint 0
+ %_ptr_Input_uint = OpTypePointer Input %uint
+ %float_0 = OpConstant %float 0
+ %bool = OpTypeBool
+ %float_0x1p_128 = OpConstant %float 0x1p+128
+ %Uniforms = OpTypeStruct %float %int %int %int %int
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+ %__1 = OpVariable %_ptr_Uniform_Uniforms Uniform
+ %int_4 = OpConstant %int 4
+ %_ptr_Uniform_int = OpTypePointer Uniform %int
+ %uint_128 = OpConstant %uint 128
+ %uint_1 = OpConstant %uint 1
+ %gl_WorkGroupSize = OpConstantComposite %v3uint %uint_128 %uint_1 %uint_1
+ %main = OpFunction %void None %3
+ %5 = OpLabel
+ %index = OpVariable %_ptr_Function_int Function
+ %a_0 = OpVariable %_ptr_Function_float Function
+ %param = OpVariable %_ptr_Function_float Function
+ %param_0 = OpVariable %_ptr_Function_int Function
+ %param_1 = OpVariable %_ptr_Function_float Function
+ %60 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %61 = OpLoad %uint %60
+ %62 = OpBitcast %int %61
+ OpStore %index %62
+ %63 = OpLoad %int %index
+ %69 = OpAccessChain %_ptr_Uniform_int %__1 %int_4
+ %70 = OpLoad %int %69
+ %71 = OpSLessThan %bool %63 %70
+ OpSelectionMerge %73 None
+ OpBranchConditional %71 %72 %73
+ %72 = OpLabel
+ %75 = OpFunctionCall %float %getAAtOutCoords_
+ OpStore %a_0 %75
+ %77 = OpLoad %float %a_0
+ OpStore %param %77
+ %78 = OpFunctionCall %float %unaryOperation_f1_ %param
+ %80 = OpLoad %int %index
+ OpStore %param_0 %80
+ OpStore %param_1 %78
+ %82 = OpFunctionCall %void %setOutput_i1_f1_ %param_0 %param_1
+ OpBranch %73
+ %73 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ %setOutput_i1_f1_ = OpFunction %void None %10
+ %flatIndex = OpFunctionParameter %_ptr_Function_int
+ %value = OpFunctionParameter %_ptr_Function_float
+ %14 = OpLabel
+ %27 = OpLoad %int %flatIndex
+ %28 = OpLoad %float %value
+ %30 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %27
+ OpStore %30 %28
+ OpReturn
+ OpFunctionEnd
+ %getAAtOutCoords_ = OpFunction %float None %15
+ %17 = OpLabel
+ %41 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %42 = OpLoad %uint %41
+ %43 = OpAccessChain %_ptr_Uniform_float %__0 %int_0 %42
+ %44 = OpLoad %float %43
+ OpReturnValue %44
+ OpFunctionEnd
+ %unaryOperation_f1_ = OpFunction %float None %18
+ %a = OpFunctionParameter %_ptr_Function_float
+ %21 = OpLabel
+ %47 = OpLoad %float %a
+ %50 = OpFOrdLessThan %bool %47 %float_0
+ OpSelectionMerge %52 None
+ OpBranchConditional %50 %51 %52
+ %51 = OpLabel
+ OpReturnValue %float_0x1p_128
+ %52 = OpLabel
+ %55 = OpLoad %float %a
+ %56 = OpExtInst %float %1 Log %55
+ OpReturnValue %56
+ OpFunctionEnd
diff --git a/test/bug/tint/951.spvasm.expected.hlsl b/test/bug/tint/951.spvasm.expected.hlsl
new file mode 100644
index 0000000..24ddc28
--- /dev/null
+++ b/test/bug/tint/951.spvasm.expected.hlsl
@@ -0,0 +1,63 @@
+RWByteAddressBuffer x_16 : register(u0, space0);
+ByteAddressBuffer x_20 : register(t1, space0);
+static uint3 gl_GlobalInvocationID = uint3(0u, 0u, 0u);
+cbuffer cbuffer_x_24 : register(b2, space0) {
+ uint4 x_24[2];
+};
+
+float getAAtOutCoords_() {
+ const uint x_42 = gl_GlobalInvocationID.x;
+ const float x_44 = asfloat(x_20.Load((4u * x_42)));
+ return x_44;
+}
+
+float unaryOperation_f1_(inout float a) {
+ const float x_47 = a;
+ if ((x_47 < 0.0f)) {
+ return asfloat(0x7f800000u);
+ }
+ const float x_55 = a;
+ return log(x_55);
+}
+
+void setOutput_i1_f1_(inout int flatIndex, inout float value) {
+ const int x_27 = flatIndex;
+ const float x_28 = value;
+ x_16.Store((4u * uint(x_27)), asuint(x_28));
+ return;
+}
+
+void main_1() {
+ int index = 0;
+ float a_1 = 0.0f;
+ float param = 0.0f;
+ int param_1 = 0;
+ float param_2 = 0.0f;
+ const uint x_61 = gl_GlobalInvocationID.x;
+ index = asint(x_61);
+ const int x_63 = index;
+ const uint scalar_offset = (16u) / 4;
+ const int x_70 = asint(x_24[scalar_offset / 4][scalar_offset % 4]);
+ if ((x_63 < x_70)) {
+ const float x_75 = getAAtOutCoords_();
+ a_1 = x_75;
+ param = a_1;
+ const float x_78 = unaryOperation_f1_(param);
+ param_1 = index;
+ param_2 = x_78;
+ setOutput_i1_f1_(param_1, param_2);
+ }
+ return;
+}
+
+struct tint_symbol_1 {
+ uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID;
+};
+
+[numthreads(128, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+ const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
+ gl_GlobalInvocationID = gl_GlobalInvocationID_param;
+ main_1();
+ return;
+}
diff --git a/test/bug/tint/951.spvasm.expected.msl b/test/bug/tint/951.spvasm.expected.msl
new file mode 100644
index 0000000..ea5f0f8
--- /dev/null
+++ b/test/bug/tint/951.spvasm.expected.msl
@@ -0,0 +1,70 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct ssbOut {
+ /* 0x0000 */ float result[1];
+};
+struct ssbA {
+ /* 0x0000 */ float A[1];
+};
+struct Uniforms {
+ /* 0x0000 */ float tint_symbol;
+ /* 0x0004 */ int aShape;
+ /* 0x0008 */ int outShape;
+ /* 0x000c */ int outShapeStrides;
+ /* 0x0010 */ int size;
+};
+
+float getAAtOutCoords_(const device ssbA& x_20, thread uint3* const tint_symbol_3) {
+ uint const x_42 = (*(tint_symbol_3)).x;
+ float const x_44 = x_20.A[x_42];
+ return x_44;
+}
+
+float unaryOperation_f1_(thread float* const a) {
+ float const x_47 = *(a);
+ if ((x_47 < 0.0f)) {
+ return INFINITY;
+ }
+ float const x_55 = *(a);
+ return log(x_55);
+}
+
+void setOutput_i1_f1_(device ssbOut& x_16, thread int* const flatIndex, thread float* const value) {
+ int const x_27 = *(flatIndex);
+ float const x_28 = *(value);
+ x_16.result[x_27] = x_28;
+ return;
+}
+
+void main_1(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_16, thread uint3* const tint_symbol_4) {
+ int index = 0;
+ float a_1 = 0.0f;
+ float param = 0.0f;
+ int param_1 = 0;
+ float param_2 = 0.0f;
+ uint const x_61 = (*(tint_symbol_4)).x;
+ index = as_type<int>(x_61);
+ int const x_63 = index;
+ int const x_70 = x_24.size;
+ if ((x_63 < x_70)) {
+ float const x_75 = getAAtOutCoords_(x_20, tint_symbol_4);
+ a_1 = x_75;
+ float const x_77 = a_1;
+ param = x_77;
+ float const x_78 = unaryOperation_f1_(&(param));
+ int const x_80 = index;
+ param_1 = x_80;
+ param_2 = x_78;
+ setOutput_i1_f1_(x_16, &(param_1), &(param_2));
+ }
+ return;
+}
+
+kernel void tint_symbol_1(uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], constant Uniforms& x_24 [[buffer(2)]], const device ssbA& x_20 [[buffer(1)]], device ssbOut& x_16 [[buffer(0)]]) {
+ thread uint3 tint_symbol_5 = 0u;
+ tint_symbol_5 = gl_GlobalInvocationID_param;
+ main_1(x_24, x_20, x_16, &(tint_symbol_5));
+ return;
+}
+
diff --git a/test/bug/tint/951.spvasm.expected.spvasm b/test/bug/tint/951.spvasm.expected.spvasm
new file mode 100644
index 0000000..608d251
--- /dev/null
+++ b/test/bug/tint/951.spvasm.expected.spvasm
@@ -0,0 +1,164 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 92
+; Schema: 0
+ OpCapability Shader
+ %46 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %tint_symbol
+ OpExecutionMode %main LocalSize 128 1 1
+ OpName %ssbOut "ssbOut"
+ OpMemberName %ssbOut 0 "result"
+ OpName %x_16 "x_16"
+ OpName %ssbA "ssbA"
+ OpMemberName %ssbA 0 "A"
+ OpName %x_20 "x_20"
+ OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
+ OpName %Uniforms "Uniforms"
+ OpMemberName %Uniforms 0 "NAN"
+ OpMemberName %Uniforms 1 "aShape"
+ OpMemberName %Uniforms 2 "outShape"
+ OpMemberName %Uniforms 3 "outShapeStrides"
+ OpMemberName %Uniforms 4 "size"
+ OpName %x_24 "x_24"
+ OpName %tint_symbol "tint_symbol"
+ OpName %getAAtOutCoords_ "getAAtOutCoords_"
+ OpName %unaryOperation_f1_ "unaryOperation_f1_"
+ OpName %a "a"
+ OpName %setOutput_i1_f1_ "setOutput_i1_f1_"
+ OpName %flatIndex "flatIndex"
+ OpName %value "value"
+ OpName %main_1 "main_1"
+ OpName %index "index"
+ OpName %a_1 "a_1"
+ OpName %param "param"
+ OpName %param_1 "param_1"
+ OpName %param_2 "param_2"
+ OpName %main "main"
+ OpDecorate %ssbOut Block
+ OpMemberDecorate %ssbOut 0 Offset 0
+ OpDecorate %_runtimearr_float ArrayStride 4
+ OpDecorate %x_16 DescriptorSet 0
+ OpDecorate %x_16 Binding 0
+ OpDecorate %ssbA Block
+ OpMemberDecorate %ssbA 0 Offset 0
+ OpDecorate %x_20 NonWritable
+ OpDecorate %x_20 DescriptorSet 0
+ OpDecorate %x_20 Binding 1
+ OpDecorate %Uniforms Block
+ OpMemberDecorate %Uniforms 0 Offset 0
+ OpMemberDecorate %Uniforms 1 Offset 4
+ OpMemberDecorate %Uniforms 2 Offset 8
+ OpMemberDecorate %Uniforms 3 Offset 12
+ OpMemberDecorate %Uniforms 4 Offset 16
+ OpDecorate %x_24 NonWritable
+ OpDecorate %x_24 DescriptorSet 0
+ OpDecorate %x_24 Binding 2
+ OpDecorate %tint_symbol BuiltIn GlobalInvocationId
+ %float = OpTypeFloat 32
+%_runtimearr_float = OpTypeRuntimeArray %float
+ %ssbOut = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_ssbOut = OpTypePointer StorageBuffer %ssbOut
+ %x_16 = OpVariable %_ptr_StorageBuffer_ssbOut StorageBuffer
+ %ssbA = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_ssbA = OpTypePointer StorageBuffer %ssbA
+ %x_20 = OpVariable %_ptr_StorageBuffer_ssbA StorageBuffer
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+ %13 = OpConstantNull %v3uint
+%gl_GlobalInvocationID = OpVariable %_ptr_Private_v3uint Private %13
+ %int = OpTypeInt 32 1
+ %Uniforms = OpTypeStruct %float %int %int %int %int
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+ %x_24 = OpVariable %_ptr_Uniform_Uniforms Uniform
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%tint_symbol = OpVariable %_ptr_Input_v3uint Input
+ %20 = OpTypeFunction %float
+ %uint_0 = OpConstant %uint 0
+%_ptr_Private_uint = OpTypePointer Private %uint
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_ptr_Function_float = OpTypePointer Function %float
+ %30 = OpTypeFunction %float %_ptr_Function_float
+ %float_0 = OpConstant %float 0
+ %bool = OpTypeBool
+%float_0x1p_128 = OpConstant %float 0x1p+128
+ %void = OpTypeVoid
+%_ptr_Function_int = OpTypePointer Function %int
+ %47 = OpTypeFunction %void %_ptr_Function_int %_ptr_Function_float
+ %59 = OpTypeFunction %void
+ %63 = OpConstantNull %int
+ %65 = OpConstantNull %float
+ %uint_4 = OpConstant %uint 4
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%getAAtOutCoords_ = OpFunction %float None %20
+ %22 = OpLabel
+ %25 = OpAccessChain %_ptr_Private_uint %gl_GlobalInvocationID %uint_0
+ %26 = OpLoad %uint %25
+ %28 = OpAccessChain %_ptr_StorageBuffer_float %x_20 %uint_0 %26
+ %29 = OpLoad %float %28
+ OpReturnValue %29
+ OpFunctionEnd
+%unaryOperation_f1_ = OpFunction %float None %30
+ %a = OpFunctionParameter %_ptr_Function_float
+ %34 = OpLabel
+ %36 = OpLoad %float %a
+ %38 = OpFOrdLessThan %bool %36 %float_0
+ OpSelectionMerge %40 None
+ OpBranchConditional %38 %41 %40
+ %41 = OpLabel
+ OpReturnValue %float_0x1p_128
+ %40 = OpLabel
+ %44 = OpLoad %float %a
+ %45 = OpExtInst %float %46 Log %44
+ OpReturnValue %45
+ OpFunctionEnd
+%setOutput_i1_f1_ = OpFunction %void None %47
+ %flatIndex = OpFunctionParameter %_ptr_Function_int
+ %value = OpFunctionParameter %_ptr_Function_float
+ %53 = OpLabel
+ %55 = OpLoad %int %flatIndex
+ %57 = OpLoad %float %value
+ %58 = OpAccessChain %_ptr_StorageBuffer_float %x_16 %uint_0 %55
+ OpStore %58 %57
+ OpReturn
+ OpFunctionEnd
+ %main_1 = OpFunction %void None %59
+ %61 = OpLabel
+ %index = OpVariable %_ptr_Function_int Function %63
+ %a_1 = OpVariable %_ptr_Function_float Function %65
+ %param = OpVariable %_ptr_Function_float Function %65
+ %param_1 = OpVariable %_ptr_Function_int Function %63
+ %param_2 = OpVariable %_ptr_Function_float Function %65
+ %69 = OpAccessChain %_ptr_Private_uint %gl_GlobalInvocationID %uint_0
+ %70 = OpLoad %uint %69
+ %71 = OpBitcast %int %70
+ OpStore %index %71
+ %72 = OpLoad %int %index
+ %75 = OpAccessChain %_ptr_Uniform_int %x_24 %uint_4
+ %76 = OpLoad %int %75
+ %77 = OpSLessThan %bool %72 %76
+ OpSelectionMerge %78 None
+ OpBranchConditional %77 %79 %78
+ %79 = OpLabel
+ %80 = OpFunctionCall %float %getAAtOutCoords_
+ OpStore %a_1 %80
+ %81 = OpLoad %float %a_1
+ OpStore %param %81
+ %82 = OpFunctionCall %float %unaryOperation_f1_ %param
+ %84 = OpLoad %int %index
+ OpStore %param_1 %84
+ OpStore %param_2 %82
+ %85 = OpFunctionCall %void %setOutput_i1_f1_ %param_1 %param_2
+ OpBranch %78
+ %78 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %59
+ %89 = OpLabel
+ %90 = OpLoad %v3uint %tint_symbol
+ OpStore %gl_GlobalInvocationID %90
+ %91 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/bug/tint/951.spvasm.expected.wgsl b/test/bug/tint/951.spvasm.expected.wgsl
new file mode 100644
index 0000000..b6f5fdd
--- /dev/null
+++ b/test/bug/tint/951.spvasm.expected.wgsl
@@ -0,0 +1,82 @@
+type RTArr = [[stride(4)]] array<f32>;
+
+type RTArr_1 = [[stride(4)]] array<f32>;
+
+[[block]]
+struct ssbOut {
+ result : RTArr_1;
+};
+
+[[block]]
+struct ssbA {
+ A : RTArr_1;
+};
+
+[[block]]
+struct Uniforms {
+ NAN : f32;
+ aShape : i32;
+ outShape : i32;
+ outShapeStrides : i32;
+ size : i32;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> x_16 : ssbOut;
+
+[[group(0), binding(1)]] var<storage, read> x_20 : ssbA;
+
+var<private> gl_GlobalInvocationID : vec3<u32>;
+
+[[group(0), binding(2)]] var<uniform> x_24 : Uniforms;
+
+fn getAAtOutCoords_() -> f32 {
+ let x_42 : u32 = gl_GlobalInvocationID.x;
+ let x_44 : f32 = x_20.A[x_42];
+ return x_44;
+}
+
+fn unaryOperation_f1_(a : ptr<function, f32>) -> f32 {
+ let x_47 : f32 = *(a);
+ if ((x_47 < 0.0)) {
+ return 0x1p+128;
+ }
+ let x_55 : f32 = *(a);
+ return log(x_55);
+}
+
+fn setOutput_i1_f1_(flatIndex : ptr<function, i32>, value : ptr<function, f32>) {
+ let x_27 : i32 = *(flatIndex);
+ let x_28 : f32 = *(value);
+ x_16.result[x_27] = x_28;
+ return;
+}
+
+fn main_1() {
+ var index : i32;
+ var a_1 : f32;
+ var param : f32;
+ var param_1 : i32;
+ var param_2 : f32;
+ let x_61 : u32 = gl_GlobalInvocationID.x;
+ index = bitcast<i32>(x_61);
+ let x_63 : i32 = index;
+ let x_70 : i32 = x_24.size;
+ if ((x_63 < x_70)) {
+ let x_75 : f32 = getAAtOutCoords_();
+ a_1 = x_75;
+ let x_77 : f32 = a_1;
+ param = x_77;
+ let x_78 : f32 = unaryOperation_f1_(&(param));
+ let x_80 : i32 = index;
+ param_1 = x_80;
+ param_2 = x_78;
+ setOutput_i1_f1_(&(param_1), &(param_2));
+ }
+ return;
+}
+
+[[stage(compute), workgroup_size(128, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] gl_GlobalInvocationID_param : vec3<u32>) {
+ gl_GlobalInvocationID = gl_GlobalInvocationID_param;
+ main_1();
+}
diff --git a/test/expressions/literals/-inf.spvasm b/test/expressions/literals/-inf.spvasm
new file mode 100644
index 0000000..2cb1527
--- /dev/null
+++ b/test/expressions/literals/-inf.spvasm
@@ -0,0 +1,21 @@
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
+ OpExecutionMode %main OriginUpperLeft
+ OpSource HLSL 600
+ OpName %out_var_SV_TARGET "out.var.SV_TARGET"
+ OpName %main "main"
+ OpDecorate %out_var_SV_TARGET Location 0
+ %float = OpTypeFloat 32
+ %float_0x1p_128 = OpConstant %float -0x1p+128
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
+ %main = OpFunction %void None %9
+ %10 = OpLabel
+ %12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
+ OpStore %out_var_SV_TARGET %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/-inf.spvasm.expected.hlsl b/test/expressions/literals/-inf.spvasm.expected.hlsl
new file mode 100644
index 0000000..3dc5854
--- /dev/null
+++ b/test/expressions/literals/-inf.spvasm.expected.hlsl
@@ -0,0 +1,20 @@
+static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void main_1() {
+ out_var_SV_TARGET = float4(asfloat(0xff800000u), asfloat(0xff800000u), asfloat(0xff800000u), asfloat(0xff800000u));
+ return;
+}
+
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol {
+ float4 out_var_SV_TARGET_1 : SV_Target0;
+};
+
+tint_symbol main() {
+ main_1();
+ const main_out tint_symbol_1 = {out_var_SV_TARGET};
+ const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
+ return tint_symbol_2;
+}
diff --git a/test/expressions/literals/-inf.spvasm.expected.msl b/test/expressions/literals/-inf.spvasm.expected.msl
new file mode 100644
index 0000000..fbaa889
--- /dev/null
+++ b/test/expressions/literals/-inf.spvasm.expected.msl
@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol_1 {
+ float4 out_var_SV_TARGET_1 [[color(0)]];
+};
+
+void main_1(thread float4* const tint_symbol_4) {
+ *(tint_symbol_4) = float4(-INFINITY, -INFINITY, -INFINITY, -INFINITY);
+ return;
+}
+
+fragment tint_symbol_1 tint_symbol() {
+ thread float4 tint_symbol_5 = 0.0f;
+ main_1(&(tint_symbol_5));
+ main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
+ tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
+ return tint_symbol_3;
+}
+
diff --git a/test/expressions/literals/-inf.spvasm.expected.spvasm b/test/expressions/literals/-inf.spvasm.expected.spvasm
new file mode 100644
index 0000000..db2fddd
--- /dev/null
+++ b/test/expressions/literals/-inf.spvasm.expected.spvasm
@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %tint_symbol_1
+ OpExecutionMode %main OriginUpperLeft
+ OpName %out_var_SV_TARGET "out_var_SV_TARGET"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %main_1 "main_1"
+ OpName %main_out "main_out"
+ OpMemberName %main_out 0 "out_var_SV_TARGET_1"
+ OpName %tint_symbol_2 "tint_symbol_2"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %tint_symbol_1 Location 0
+ OpMemberDecorate %main_out 0 Offset 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+ %5 = OpConstantNull %v4float
+%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+%float_n0x1p_128 = OpConstant %float -0x1p+128
+ %13 = OpConstantComposite %v4float %float_n0x1p_128 %float_n0x1p_128 %float_n0x1p_128 %float_n0x1p_128
+ %main_out = OpTypeStruct %v4float
+ %14 = OpTypeFunction %void %main_out
+ %main_1 = OpFunction %void None %8
+ %11 = OpLabel
+ OpStore %out_var_SV_TARGET %13
+ OpReturn
+ OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %14
+%tint_symbol = OpFunctionParameter %main_out
+ %18 = OpLabel
+ %19 = OpCompositeExtract %v4float %tint_symbol 0
+ OpStore %tint_symbol_1 %19
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %8
+ %21 = OpLabel
+ %22 = OpFunctionCall %void %main_1
+ %24 = OpLoad %v4float %out_var_SV_TARGET
+ %25 = OpCompositeConstruct %main_out %24
+ %23 = OpFunctionCall %void %tint_symbol_2 %25
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/-inf.spvasm.expected.wgsl b/test/expressions/literals/-inf.spvasm.expected.wgsl
new file mode 100644
index 0000000..3e916f6
--- /dev/null
+++ b/test/expressions/literals/-inf.spvasm.expected.wgsl
@@ -0,0 +1,17 @@
+var<private> out_var_SV_TARGET : vec4<f32>;
+
+fn main_1() {
+ out_var_SV_TARGET = vec4<f32>(-0x1p+128, -0x1p+128, -0x1p+128, -0x1p+128);
+ return;
+}
+
+struct main_out {
+ [[location(0)]]
+ out_var_SV_TARGET_1 : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main() -> main_out {
+ main_1();
+ return main_out(out_var_SV_TARGET);
+}
diff --git a/test/expressions/literals/inf.spvasm b/test/expressions/literals/inf.spvasm
new file mode 100644
index 0000000..d958091
--- /dev/null
+++ b/test/expressions/literals/inf.spvasm
@@ -0,0 +1,21 @@
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
+ OpExecutionMode %main OriginUpperLeft
+ OpSource HLSL 600
+ OpName %out_var_SV_TARGET "out.var.SV_TARGET"
+ OpName %main "main"
+ OpDecorate %out_var_SV_TARGET Location 0
+ %float = OpTypeFloat 32
+ %float_0x1p_128 = OpConstant %float 0x1p+128
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
+ %main = OpFunction %void None %9
+ %10 = OpLabel
+ %12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
+ OpStore %out_var_SV_TARGET %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/inf.spvasm.expected.hlsl b/test/expressions/literals/inf.spvasm.expected.hlsl
new file mode 100644
index 0000000..d973da7
--- /dev/null
+++ b/test/expressions/literals/inf.spvasm.expected.hlsl
@@ -0,0 +1,20 @@
+static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void main_1() {
+ out_var_SV_TARGET = float4(asfloat(0x7f800000u), asfloat(0x7f800000u), asfloat(0x7f800000u), asfloat(0x7f800000u));
+ return;
+}
+
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol {
+ float4 out_var_SV_TARGET_1 : SV_Target0;
+};
+
+tint_symbol main() {
+ main_1();
+ const main_out tint_symbol_1 = {out_var_SV_TARGET};
+ const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
+ return tint_symbol_2;
+}
diff --git a/test/expressions/literals/inf.spvasm.expected.msl b/test/expressions/literals/inf.spvasm.expected.msl
new file mode 100644
index 0000000..37e0e9c
--- /dev/null
+++ b/test/expressions/literals/inf.spvasm.expected.msl
@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol_1 {
+ float4 out_var_SV_TARGET_1 [[color(0)]];
+};
+
+void main_1(thread float4* const tint_symbol_4) {
+ *(tint_symbol_4) = float4(INFINITY, INFINITY, INFINITY, INFINITY);
+ return;
+}
+
+fragment tint_symbol_1 tint_symbol() {
+ thread float4 tint_symbol_5 = 0.0f;
+ main_1(&(tint_symbol_5));
+ main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
+ tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
+ return tint_symbol_3;
+}
+
diff --git a/test/expressions/literals/inf.spvasm.expected.spvasm b/test/expressions/literals/inf.spvasm.expected.spvasm
new file mode 100644
index 0000000..eb416fd
--- /dev/null
+++ b/test/expressions/literals/inf.spvasm.expected.spvasm
@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %tint_symbol_1
+ OpExecutionMode %main OriginUpperLeft
+ OpName %out_var_SV_TARGET "out_var_SV_TARGET"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %main_1 "main_1"
+ OpName %main_out "main_out"
+ OpMemberName %main_out 0 "out_var_SV_TARGET_1"
+ OpName %tint_symbol_2 "tint_symbol_2"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %tint_symbol_1 Location 0
+ OpMemberDecorate %main_out 0 Offset 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+ %5 = OpConstantNull %v4float
+%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+%float_0x1p_128 = OpConstant %float 0x1p+128
+ %13 = OpConstantComposite %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
+ %main_out = OpTypeStruct %v4float
+ %14 = OpTypeFunction %void %main_out
+ %main_1 = OpFunction %void None %8
+ %11 = OpLabel
+ OpStore %out_var_SV_TARGET %13
+ OpReturn
+ OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %14
+%tint_symbol = OpFunctionParameter %main_out
+ %18 = OpLabel
+ %19 = OpCompositeExtract %v4float %tint_symbol 0
+ OpStore %tint_symbol_1 %19
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %8
+ %21 = OpLabel
+ %22 = OpFunctionCall %void %main_1
+ %24 = OpLoad %v4float %out_var_SV_TARGET
+ %25 = OpCompositeConstruct %main_out %24
+ %23 = OpFunctionCall %void %tint_symbol_2 %25
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/inf.spvasm.expected.wgsl b/test/expressions/literals/inf.spvasm.expected.wgsl
new file mode 100644
index 0000000..c4dd750
--- /dev/null
+++ b/test/expressions/literals/inf.spvasm.expected.wgsl
@@ -0,0 +1,17 @@
+var<private> out_var_SV_TARGET : vec4<f32>;
+
+fn main_1() {
+ out_var_SV_TARGET = vec4<f32>(0x1p+128, 0x1p+128, 0x1p+128, 0x1p+128);
+ return;
+}
+
+struct main_out {
+ [[location(0)]]
+ out_var_SV_TARGET_1 : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main() -> main_out {
+ main_1();
+ return main_out(out_var_SV_TARGET);
+}
diff --git a/test/expressions/literals/nan.spvasm b/test/expressions/literals/nan.spvasm
new file mode 100644
index 0000000..95df6dc
--- /dev/null
+++ b/test/expressions/literals/nan.spvasm
@@ -0,0 +1,21 @@
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %out_var_SV_TARGET
+ OpExecutionMode %main OriginUpperLeft
+ OpSource HLSL 600
+ OpName %out_var_SV_TARGET "out.var.SV_TARGET"
+ OpName %main "main"
+ OpDecorate %out_var_SV_TARGET Location 0
+ %float = OpTypeFloat 32
+ %float_0x1p_128 = OpConstant %float 0x1.1p+128
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %out_var_SV_TARGET = OpVariable %_ptr_Output_v4float Output
+ %main = OpFunction %void None %9
+ %10 = OpLabel
+ %12 = OpCompositeConstruct %v4float %float_0x1p_128 %float_0x1p_128 %float_0x1p_128 %float_0x1p_128
+ OpStore %out_var_SV_TARGET %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/nan.spvasm.expected.hlsl b/test/expressions/literals/nan.spvasm.expected.hlsl
new file mode 100644
index 0000000..f3c1e53
--- /dev/null
+++ b/test/expressions/literals/nan.spvasm.expected.hlsl
@@ -0,0 +1,20 @@
+static float4 out_var_SV_TARGET = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void main_1() {
+ out_var_SV_TARGET = float4(asfloat(0x7fc00000u), asfloat(0x7fc00000u), asfloat(0x7fc00000u), asfloat(0x7fc00000u));
+ return;
+}
+
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol {
+ float4 out_var_SV_TARGET_1 : SV_Target0;
+};
+
+tint_symbol main() {
+ main_1();
+ const main_out tint_symbol_1 = {out_var_SV_TARGET};
+ const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
+ return tint_symbol_2;
+}
diff --git a/test/expressions/literals/nan.spvasm.expected.msl b/test/expressions/literals/nan.spvasm.expected.msl
new file mode 100644
index 0000000..df5d951
--- /dev/null
+++ b/test/expressions/literals/nan.spvasm.expected.msl
@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct main_out {
+ float4 out_var_SV_TARGET_1;
+};
+struct tint_symbol_1 {
+ float4 out_var_SV_TARGET_1 [[color(0)]];
+};
+
+void main_1(thread float4* const tint_symbol_4) {
+ *(tint_symbol_4) = float4(NAN, NAN, NAN, NAN);
+ return;
+}
+
+fragment tint_symbol_1 tint_symbol() {
+ thread float4 tint_symbol_5 = 0.0f;
+ main_1(&(tint_symbol_5));
+ main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
+ tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
+ return tint_symbol_3;
+}
+
diff --git a/test/expressions/literals/nan.spvasm.expected.spvasm b/test/expressions/literals/nan.spvasm.expected.spvasm
new file mode 100644
index 0000000..b4b90b7
--- /dev/null
+++ b/test/expressions/literals/nan.spvasm.expected.spvasm
@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %tint_symbol_1
+ OpExecutionMode %main OriginUpperLeft
+ OpName %out_var_SV_TARGET "out_var_SV_TARGET"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %main_1 "main_1"
+ OpName %main_out "main_out"
+ OpMemberName %main_out 0 "out_var_SV_TARGET_1"
+ OpName %tint_symbol_2 "tint_symbol_2"
+ OpName %tint_symbol "tint_symbol"
+ OpName %main "main"
+ OpDecorate %tint_symbol_1 Location 0
+ OpMemberDecorate %main_out 0 Offset 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+ %5 = OpConstantNull %v4float
+%out_var_SV_TARGET = OpVariable %_ptr_Private_v4float Private %5
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void
+%float_0x1_1p_128 = OpConstant %float 0x1.1p+128
+ %13 = OpConstantComposite %v4float %float_0x1_1p_128 %float_0x1_1p_128 %float_0x1_1p_128 %float_0x1_1p_128
+ %main_out = OpTypeStruct %v4float
+ %14 = OpTypeFunction %void %main_out
+ %main_1 = OpFunction %void None %8
+ %11 = OpLabel
+ OpStore %out_var_SV_TARGET %13
+ OpReturn
+ OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %14
+%tint_symbol = OpFunctionParameter %main_out
+ %18 = OpLabel
+ %19 = OpCompositeExtract %v4float %tint_symbol 0
+ OpStore %tint_symbol_1 %19
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %8
+ %21 = OpLabel
+ %22 = OpFunctionCall %void %main_1
+ %24 = OpLoad %v4float %out_var_SV_TARGET
+ %25 = OpCompositeConstruct %main_out %24
+ %23 = OpFunctionCall %void %tint_symbol_2 %25
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/literals/nan.spvasm.expected.wgsl b/test/expressions/literals/nan.spvasm.expected.wgsl
new file mode 100644
index 0000000..a72149c
--- /dev/null
+++ b/test/expressions/literals/nan.spvasm.expected.wgsl
@@ -0,0 +1,17 @@
+var<private> out_var_SV_TARGET : vec4<f32>;
+
+fn main_1() {
+ out_var_SV_TARGET = vec4<f32>(0x1.1p+128, 0x1.1p+128, 0x1.1p+128, 0x1.1p+128);
+ return;
+}
+
+struct main_out {
+ [[location(0)]]
+ out_var_SV_TARGET_1 : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main() -> main_out {
+ main_1();
+ return main_out(out_var_SV_TARGET);
+}
