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);
+}