CrossTL 3.5.0

Breaking Changes

• CrossGL refuses non-square floatNxM and doubleNxM matrix spellings, such as float2x3, with AmbiguousMatrixSpellingError (project.translate.ambiguous-matrix-spelling): MSL reads float2x3 as 2 columns of 3 rows and HLSL as 2 rows of 3 columns, and the Metal, GLSL, HLSL and SPIR-V generators each took a different one. Write matCxR or dmatCxR. The SPIR-V importer now writes non-square matrices as matCxR. Square float3x3(...) and float3x3(s) now fill columns and build a diagonal matrix in HLSL output, as mat3 does; HLSL output filled rows and repeated s.
• Metal output refuses dFdx, dFdy, fwidth and their ddx/ddy, fine and coarse forms in a compute, ray generation, mesh, object, task or amplification stage, called directly or through a helper, with UnsupportedMetalFeatureError (missing_capabilities metal.kernel-derivatives, metal.mesh-derivatives or metal.object-derivatives). Metal has derivatives only in fragment functions; they used to be written into the kernel, [[mesh]] or [[object]] function, which the Metal compiler rejects.
• Mojo output refuses a vector where a scalar is expected, such as return v; from a function returning float or if (a == b) on two vectors, with MojoScalarConversionError. It used to write the vector unchanged, which Mojo either rejects or, as an if, while or ?: condition, reads as true when any lane is true. Use all, any or a component access to reduce the vector first.
• A type name that is neither built in nor declared by the program, as a structure, enumeration, alias, constant block or generic parameter, is refused with a typed error. Every generator used to copy it into the output.
• WGSL output refuses 64-bit integers, doubles and bfloat16, the vectors and matrices built on them and conversions such as double(x), with WGSLNumericTypeUnsupportedError. It used to narrow double to f32 and long to i32, which changed host buffer layouts, or copy the type unchanged.
• buffer T name[] and buffer T name[N] declare one storage buffer of T elements for every T, structures included, as Metal's device T* does. GLSL buffer block arrays (buffer B { ... } name[N];) keep their meaning as arrays of descriptors.
• CrossGL sizeof of a scalar, a 2- or 4-component vector or a sized array is folded to a constant on Metal, HLSL and GLSL. Other operands are refused with SizeofUnsupportedError; sizeof used to be written out unchanged, so GLSL rejected it and HLSL and Metal measured with their own layout rules.
• The Metal importer refuses sinpi, cospi, tanpi, asinpi, acospi, atanpi and atan2pi, which are exact at integers and half-integers where sin(x * pi) is not, and calls to functions that the source declares but never defines. Both used to reach every target as calls to undeclared functions.
• Metal indirect command buffers (command_buffer, render_command, compute_command) and rasterization rate maps are refused with MetalRuntimeTypeUnsupportedError for every target other than Metal when a stage uses them. HLSL and GLSL output used to copy the type names.
• HLSL output refuses a constant buffer larger than Direct3D's 64 KiB limit, such as a large Metal constant T& argument or GLSL uniform block, with ConstantBufferTooLargeError, and matrix varyings whose semantic indices overlap with InterfaceLocationConflictError. DXC used to reject the generated shader.
• The GLSL importer refuses tensorARM uniforms with GLSLTypeUnsupportedError. It used to drop them, so every target referenced an undeclared name. GLSL output refuses string literals outside #include, #line and debugPrintfEXT formats.
• The Metal importer refuses ray-tracing functions, and the uses of mesh and object functions it cannot import, such as set_indices, with MetalImportRefused. They used to be printed as plain helpers that still called Metal methods.
• HLSL output for the dx10 and dx11 profiles, as selected by backend="dx11", "d3d11" or "directx-11", refuses the WaveSize attribute, from HLSL [WaveSize(...)] or CrossGL @WaveSize(...), with DirectXWaveSizeProfileError, because WaveSize needs Direct3D 12 and Shader Model 6.6 (6.8 for the range forms). It used to be written unchanged, as [WaveSize(32)], which DXC rejects below Shader Model 6.6. The default dx12 profile is unchanged.
• HLSL output for the dx10 and dx11 profiles, as selected by backend="dx11", "d3d11" or "directx-11", refuses mesh and amplification stages (CrossGL mesh, task, amplification and object) with DirectXMeshStageProfileError, because mesh shading needs Direct3D 12 and Shader Model 6.5. They used to be written unchanged, as [shader("mesh")] and [shader("amplification")] entry points that Direct3D 10 and 11 cannot run. The default dx12 profile is unchanged.
• HLSL output for the dx10 and dx11 profiles, as selected by backend="dx11", "d3d11" or "directx-11", refuses ray generation, intersection, any-hit, closest-hit, miss and callable stages with DirectXRayTracingStageProfileError, because DirectX Raytracing needs Direct3D 12 and Shader Model 6.3. They used to be written unchanged, as [shader("raygeneration")] entry points that Direct3D 11 cannot run. The default dx12 profile is unchanged.
• CrossGL reads (a, b) as the C comma operator, which evaluates both operands left to right and has the value of b, and writes a tuple with a trailing comma, (a, b,), as the Mojo and Rust importers now do. The HLSL, Metal and CUDA importers keep comma operators, HLSL, GLSL, Metal, CUDA, HIP and Slang output write them in parentheses, SPIR-V output evaluates the operands in order, WGSL output moves the side effects of the discarded operands before the statement, and Rust and Mojo output refuse them with RustLoweringRefused and MojoLoweringRefused. (a, b) used to be a tuple, so an imported comma expression became an initializer with one element too many, which every compiler rejected, and the side effects of the discarded operands were lost.
• A gl_ attribute that names no GLSL built-in, no built-in of an extension CrossTL reads and none of CrossGL's own, such as @gl_FragBogus, is refused on every target with BuiltinValueUnsupportedError. It used to become a user varying: an interface location on SPIR-V, a user semantic on HLSL and an out variable on GLSL. The aliases gl_IsFrontFace, which the Slang, Rust and Mojo importers write, and gl_FragStencilRefARB are read as gl_FrontFacing and gl_FragStencilRefEXT; they used to become user varyings on SPIR-V, and the ARB stencil reference a user semantic on HLSL and [[attribute(0)]] on Metal.
• A gl_ built-in used in a stage or direction that GLSL does not declare it for, such as a vertex output @gl_SubgroupSize or a fragment input @gl_LaunchIDEXT, is refused with BuiltinValueUnsupportedError. CrossGL's conventions stay accepted: a vertex or fragment input @gl_Position and the fragment inputs @gl_PointSize, @gl_FragColor, @gl_FragData and @gl_SampleMask. HLSL, SPIR-V and GLSL output also refuse a gl_ built-in they have no mapping for, such as the NV per-view vertex outputs on GLSL. Such built-ins used to become user varyings that no stage wrote or read: an interface location on SPIR-V, a : gl_X semantic on HLSL and an out variable on GLSL.
• CUDA and HIP output refuse a program that contains a construct they cannot lower, such as an image atomic, a depth comparison or a missing resource query, with CudaLoweringRefused or HipLoweringRefused, which list every such construct. They used to return the program with an /* unsupported ... */ comment in the construct's place, often followed by a placeholder value such as 0, so the output did not compile or computed something else.
• A CrossGL storage buffer without a @buffer_layout attribute holds its elements at std430 offsets on Metal and HLSL, as on every other target and as the host lays them out. Metal output declares such elements with packed vectors and padding, so struct { vec3 a; float b; } is 16 bytes with b at byte 12; MSL's own rules used to put b at byte 16 in a 32-byte element. HLSL output reads and writes such buffers as 32-bit words at their std430 offsets, so struct { float a; vec3 b; } keeps b at byte 16 in a 32-byte element and vec3 v[] a 16-byte stride; HLSL's structured-buffer rules used to put b at byte 4 and give vec3 a 12-byte stride. Buffers whose layouts already agree are unchanged.
• Metal output samples a combined sampler, such as GLSL sampler2D, samplerCube or sampler2DShadow, through a sampler <texture>Sampler [[sampler(n)]] entry argument at the texture's binding n, which the host must bind, and a helper that samples one takes the sampler as a parameter after the texture. It used to sample through an inline sampler(mag_filter::linear, min_filter::linear), which ignored the host's sampler state: mipmapped textures were read at level 0, every address mode was clamp_to_edge and depth comparisons had no compare function.
• Sampler feedback textures, FeedbackTexture2D<KIND> and FeedbackTexture2DArray<KIND> and the HLSL importer's feedbackTexture2D spellings, are refused with FeedbackTextureUnsupportedError for every target other than Direct3D, and on Direct3D when KIND is not SAMPLER_FEEDBACK_MIN_MIP or SAMPLER_FEEDBACK_MIP_REGION_USED. The other generators used to copy the type into their output, which no other target compiler accepts. The undeclared-type check reads the two feedback kinds as values rather than type names, and HLSL output accepts the float2 gradients that DXC declares for WriteSamplerFeedbackGrad on a FeedbackTexture2DArray, where it used to demand float3 (#2798).
• The OpenCL importer refuses the OpenCL C 2.0 device-side enqueue event functions create_user_event, retain_event, release_event, set_user_event_status, is_valid_event and capture_event_profiling_info with OpenCLTargetUnsupportedError for every target, because no target has the clk_event_t objects they create and signal. They used to reach every target as calls to undeclared functions. A function the program defines under one of these names is still called.
• An inout vertex entry parameter bound to SV_Position is refused with VertexInputPositionUnsupportedError on every target other than HLSL: its input is the Position vertex attribute and its output the position built-in, which one declaration cannot name. GLSL and SPIR-V output used to read gl_Position for it, Metal output read zero and WGSL output put @builtin(position) on a vertex input.
• Mojo output declares file-scope constants, enumeration constants and an enumeration's underlying type with comptime instead of alias, which Mojo 1.2 rejects at file scope with expressions must not appear at file scope. Mojo 24.x and 25.x, which do not know comptime, reject the new output.
• Vertex inputs and varyings take their location from one rule on GLSL, SPIR-V, WGSL and Metal: an explicit location, then the classic semantic table (POSITION 0, NORMAL 1, TANGENT 2, BINORMAL 3, TEXCOORD and TEXCOORD0 5, TEXCOORD1 to TEXCOORD7 6 to 12, COLOR and COLOR0 to COLOR7 13 to 20), then the first free location in declaration order. Each target used to number them its own way, so locations change: an HLSL COLOR vertex input moves from attribute 1 to 13 on Metal and a TEXCOORD input from location 4 to 5 on GLSL, and host vertex layouts written for the old numbers need updating. GLSL output declares vertex inputs without a classic semantic with their location instead of leaving it to the linker, reads a vertex input member with a built-in semantic such as SV_VertexID from the built-in instead of an unbound attribute, and reserves every location a vertex input array spans. Fragment outputs named COLORn write colour attachment n on every target; a structure that carries them and is also another interface is refused with FragmentColourOutputUnsupportedError.

Added

• crosstl doctor lists, for each backend, the validation tools CrossTL looks for and which of them are on PATH. A validation that cannot run now names the missing tool.
• CrossGL textureResidency(read, texel) and textureResident(code) query sparse texture residency on GLSL, Metal, HLSL and SPIR-V.
• metal_sizeof(x) measures a type with Metal's layout rules. The Metal importer writes it for operands it cannot size, such as template parameters.
• Metal [[mesh]] and [[object]] functions import as CrossGL mesh and task stages.
• gl_InstanceIndex translates to OpenGL as gl_InstanceID + gl_BaseInstance, to OpenGL ES as gl_InstanceID, to HLSL as SV_InstanceID and to SPIR-V. As with gl_VertexIndex, HLSL's value does not include the base instance.
• Metal fixed sampler state (constexpr sampler) translates to OpenGL through host sampler objects. GLSL output starts with a // crosstl-fixed-samplers: line that lists each state with its uniforms and texture units. WebGL still refuses it.
• The SPIR-V importer reads subpass inputs as subpassLoad. MSL clz and ctz import through findMSB and findLSB for 8-, 16- and 32-bit operands, and Metal output lowers findMSB and findLSB instead of copying them.
• CrossGL structures may declare member functions. Each becomes a free function Owner_name that takes the receiver first, as in when the body cannot write it and inout when it may, and recv.name(args), Owner::name(args), this and unqualified member names in the body are rewritten to match. Methods of generic structures, writing methods called on a value that is not an lvalue and calls whose receiver type cannot be inferred are refused with StructMethodUnsupportedError. The parser used to skip a member function up to its first ; and read the rest of its body as module-scope code, so every target called an undeclared method.
• CrossGL, GLSL, HLSL, Slang and CUDA input accept a structure declared in a function body, as in struct S { float a; } s;. Every target receives it at stage or program scope as <function>_S, with its uses renamed; a structure whose members depend on function-local names is refused with LocalStructUnsupportedError. Such programs used to fail to parse or be refused for an undeclared type. Slang output also declares the structures of a stage block, which it used to drop.
• A project can state that read-only entry storage is immutable for a dispatch with immutable_storage entries naming a source, entryPoint and parameter. The DirectX and OpenGL software-subgroup uniformity proofs then treat loads from that storage at uniform offsets as uniform, the artifacts carry the contract in a crosstl-immutable-storage header, and reflection, packaging and native dispatch refuse a dispatch in which a writable binding can reach that storage.
• Project translation for DirectX and OpenGL declares a compute workgroup whose shape the target cannot declare but whose thread count it can, such as a Metal kernel dispatched in (1, 2, 256) threadgroups, as (X * Y * Z, 1, 1), and rebuilds the local and global invocation IDs, the workgroup size and the threads per grid from the flattened local index, the workgroup ID and the workgroup count, so they keep their values. The execution record, runtime manifests and native descriptors carry the declared size as targetWorkgroupSize. Direct3D 12 and OpenGL cap a workgroup's z dimension at 64, so such a kernel used to fail with project.translate.workgroup-size-invalid although its 512 threads fit. Programs that read those built-ins outside entry parameters, and executions that depend on the subgroup width or a host-dispatch contract, are still refused.
• GLSL mesh and task shaders (GL_EXT_mesh_shader) translate to HLSL and Metal. The module-scope outputs, gl_MeshVerticesEXT, per-vertex and perprimitiveEXT out arrays and blocks and gl_PrimitiveTriangleIndicesEXT and its point and line forms, become vertex and primitive structures and an index array passed to the entry, and taskPayloadSharedEXT becomes the mesh entry's payload parameter and the groupshared payload a task entry passes to DispatchMesh. HLSL output used to refuse the module-scope outputs and Metal output the global stores. Built-ins neither target carries yet, such as gl_PointSize and gl_CullPrimitiveEXT, reads of an output, helpers that use the mesh interface and a payload shared by a task and a mesh stage are refused with GlslMeshStageUnsupportedError.
• GLSL output reads a specialization constant that has no default, such as an MSL [[function_constant]] declared without a value, from a CROSSGL_SPECIALIZE_<name> macro that the host defines after #version when the driver compiles the source, as OpenGL programs select variants, and the metadata comment records the macro as source_macro. Compiled from source, such a shader used to stop at an #error, because only SPIR-V compilation has constant_id; without the macro the #error now names it.
• The GLSL importer reads the GL_ARB_sparse_texture2 and GL_ARB_sparse_texture_clamp functions, such as sparseTextureARB and sparseTexelsResidentARB, and the Metal importer reads a local initialised from sparse_sample, sparse_read, sparse_gather or another sparse_* texture member, with its value() and resident(), as CrossGL textureResidency and textureResident queries. GLSL input used to pass the calls through, so every other target called functions it does not declare, and Metal input refused every sparse_* member. Other uses of sparse_color, such as a parameter, a structure member or a copy, are refused.

Improved

• Mojo output converts scalar boolean values to float and half through Int, so casts from boolean variables and comparison expressions compile.
• WGSL output monomorphizes generic functions, passes inout and out parameters as pointers, copies written parameters into locals, binds every entry parameter with @location or @builtin and gives integer varyings @interpolate(flat). Memory barriers, matrix-scalar +, - and /, matrixCompMult, outerProduct, dot4add_i8packed and dot4add_u8packed are lowered to WGSL instead of being copied.
• Inline ray queries have one CrossGL spelling, the DXR method names, and translate between HLSL, GLSL and SPIR-V.
• The Metal importer reads more of the MSL that real headers use: function templates deduce their arguments before applying defaults, partial specializations constrained with enable_if are selected, alias templates work as constructors, a local alias may reuse its template's name, simd type aliases work as constructors, an unscoped enumerator is found in the enclosing scope as in C++, and powr and exp10 are lowered for other targets.
• Metal sources keep their argument table indices as SPIR-V bindings, and their fixed sampler state becomes host samplers on Direct3D and Vulkan. Metal output binds arrays of HLSL structured buffers and GLSL storage blocks as one entry argument per element, as MSL requires. Vulkan subpass inputs become texture reads on targets without input attachments.
• Generic functions deduce a type parameter from a vector argument, as in pick(ivec2(1, 2)) for T pick<T>(vec2<T> v), and generically spelled vectors such as vec2<u16> map to the right type on GLSL, HLSL and MSL.
• The HLSL importer imports [[vk::shader_record_ext]] constant buffers as shader record blocks. The Slang importer keeps groupshared on globals.
• Translating large sources is faster: the lexers pick token patterns by first character, and the Metal preprocessor and several lowering passes walk the program fewer times.
• Metal output renames identifiers that MSL reads as keywords, such as kernel, device, thread and half, gives gl_InstanceIndex an [[instance_id]] input, binds a ray generation stage's launch index and size to the kernel grid and spells its ray flags as constants, and passes resources and module-scope groupshared values to mesh and object functions.
• The Metal importer types matrix arithmetic, the saturate family and the results of texture methods, imports depth texture reads as the depth value, and keeps [[attribute(n)]] vertex inputs at location n on every target instead of renaming them to HLSL semantics.
• SPIR-V output folds module-scope constant expressions exactly, rounding binary32 results after each operation, instead of emitting instructions that module scope cannot hold. Shading-rate built-ins, per-primitive mesh outputs and arrays of constant buffers are declared with their SPIR-V built-ins and capabilities.
• INFINITY, NAN, FLT_MAX, FLT_MIN, FLT_EPSILON and the other standard float limits are written as exact binary32 values on HLSL and GLSL.
• Arrays of HLSL structured buffers keep HLSL's tight element layout on GLSL, SPIR-V and Metal.
• Texel lookup offsets that a program keeps in constant locals or computes in short constant loops become constant expressions on GLSL, HLSL, WGSL, SPIR-V and Slang. SPIR-V output refuses a non-constant offset outside gathers instead of writing an invalid module.
• Metal half matrices translate to SPIR-V and WGSL as float matrices that hold binary16 values, rounded after each operation, and half operands in mixed expressions are typed. The Metal importer maps the simd:: vector and matrix types like the Metal ones.
• GLSL length() of a vector or matrix becomes its component or column count on other targets.
• Structure member extents named by constants, such as float data[Total]; with constant constexpr size_t Total = X * Y * Z; declared after the structure, are folded to integers on every target.
• The Metal importer resolves the <simd/simd.h> aliases such as simd::float4, simd_float4x4 and vector_float3, declares the <metal_stdlib> plain-data structures a source uses without declaring, such as the tessellation factor and indirect argument records, and imports atomic_fetch_sub_explicit as an atomic add of the negated value instead of copying the call.
• HLSL implicit vector truncation in assignments and initializers becomes a leading-component swizzle on every target.
• SPIR-V output lowers Metal structures with device, constant, threadgroup or thread pointer members to resource handles, as HLSL and GLSL output do. On all three targets such a member may hold a null pointer that is overwritten before every use, and a thread pointer member may point into a fixed-size local array.
• The Metal importer reads more of the C++ in MLX's kernels: decltype of locals and decltype(e)::member in template arguments, conditional_t selections once their condition folds, constexpr locals initialized by type traits such as is_same_v, Owner::member static constants, arithmetic on integral_constant values, partial specializations spelled with an alias of their argument and the overload an explicit instantiation selects, and it skips static_assert in a discarded if constexpr branch. Two struct specializations whose arguments differ only in punctuation, such as !flag and flag, get distinct names; the second used to reuse the first one's members. DirectX and OpenGL software subgroups lower simd_shuffle_up, simd_shuffle_xor and quad_sum and prove more loops and early returns uniform, and HLSL output renames a type alias named float16_t or vector together with the types that use it. Project translation reuses entry discovery, template expansions and the source host interface while a source's include closure is unchanged.
• HLSL and GLSL output lower 8- and 16-bit arrays held in 32-bit union words, such as the byte lanes a compare-and-swap loop updates, to lane get and set helpers, and lower scalar pointer views &local.member of a structure member or of a union's 32-bit storage word. The Metal importer specializes template unions per argument, folds sizeof in their member extents and folds value variable templates such as template <typename T> constexpr int n = ...; at their uses. Such kernels used to be refused with member-shape-unsupported or missing-view-backing. HLSL output refuses a call that passes one scalar as two inout views, f(&x, &x), whose copy-back used to lose one of the writes.
• Translating large Metal sources with many template specializations is faster: the Metal preprocessor finds structure declarations, using namespace directives and static constexpr calls with one scan instead of one per name, and project translation rebuilds the constrained-call type context only when it reaches a call in a newly materialized body. PyTorch's BinaryKernel.metal, 2.97 MB after preprocessing, takes 57.8 s instead of 155.7 s.
• The Metal importer resolves a plain type name that no alias, typedef or template binding names without walking its alias tables, so large Metal units translate faster: MLX's gemv.metal to HLSL takes 66 s instead of 87 s.
• Metal output converts the int operand of a comparison with a uint to uint explicitly, as GLSL and HLSL do implicitly, so the output compiles with -Wsign-compare warnings treated as errors. Non-negative literals are left as written.

Fixed

• HIP output lowers ptr<T> helper parameters imported from HIP or CUDA to native pointer types instead of emitting the CrossGL spelling unchanged.
• CrossGL accepts an empty statement, a ; on its own, which it used to refuse with CrossGLFunctionBodyParseError, so every target refused a function containing one. A switch whose last case or default label is followed by no statement now leaves the switch there with break;; Metal and GLSL ES output used to end the switch on the bare label, which the Metal compiler and GLSL ES 3.00 refuse.
• The Slang importer writes a Slang floatRxC as the same CrossGL matrix, matCxR, as the HLSL importer does. Constructors and initializer lists fill rows, a single scalar fills every element, m[i][j] and the _mRC and _RC swizzles name element (i, j), m[i] is row i, and * of two matrices is componentwise. It used to keep the arguments and subscripts of the transposed matrix and still write mul(a, b) as a * b, so every mul with a matrix operand computed a different product; row_major and column_major members are now kept as @ row_major and @ column_major. A row write through a multi-component swizzle and a conversion from a value of unknown type are refused with SlangImportRefused.
• GLSL input names the instance it gives a buffer block without one, such as buffer Data { float values[]; };, so that no local, parameter or global of the program has the same name: with a local float data the members are read through data_2. The name used to be checked against global declarations only, so the local shadowed the instance and data.values read a member of a float, which does not compile.
• CUDA and HIP output apply the scalar math built-ins, such as sin, exp2, inversesqrt, pow and floor, to native half2 (f16vec2) values through per-component helpers that widen each component to float and pack the result with __floats2half2_rn, as in __floats2half2_rn(sinf(__half2float(value.x)), sinf(__half2float(value.y))). They used to pass the native half2 straight to the scalar float function, as in sinf(h), which does not compile because no overload accepts half2.
• CUDA and HIP output apply the scalar math built-ins, such as sin, exp, inversesqrt and pow, to f16vec3 and f16vec4 values through per-component helpers that widen each __half to float, as in __float2half(sinf(__half2float(value.x))). They used to pass the half vector straight to the float function, as in sinf(h), which does not compile because no overload accepts the generated cgl_half3 and cgl_half4 structures.
• CUDA and HIP output map int64 and uint64 locals, shared variables and constructors to native 64-bit integer types instead of undeclared CrossGL names (#3189).
• CUDA and HIP output lower scalar length, distance and normalize to standard math calls: fabsf, fabs, and a float round trip for half. normalize of a call or other compound operand, such as normalize(f(x)), goes through a cgl_normalize_float style helper so the operand is evaluated once. They used to fall through to bare calls like length(x), distance(a, b) and normalize(x), which the CUDA and HIP compilers reject because they provide no scalar overloads for these builtins.
• CUDA and HIP output expand a multi-component swizzle of a half vector, such as e.xyz or t.xxyz, through cgl_make_half3 and cgl_make_half4, and a two-component one, such as e.wy, into a native half2 through __floats2half2_rn. The swizzle is typed as a half vector, so r + e.wzy uses the cgl_cgl_half3_add helper, and an assignment to a swizzle, such as o.xy = p or o.wzy += r, writes each component. It used to be written as a member access, as in e.xyz, which does not compile because the generated cgl_half3 and cgl_half4 structures only have the members x, y, z and w.
• CUDA and HIP output expand a multi-component swizzle of a native half2 (f16vec2), such as a.yx, into __floats2half2_rn(a.y, a.x), write an assignment to one, such as t.yx = b or t.xy += b, component by component, and pass a half2 operand of a half vector constructor as its two components, as in cgl_make_half4(a.x, a.y, b.x, b.y) for f16vec4(a, b), through a single-evaluation helper when the operand is a call. They used to keep a.yx as a member access and pass a as one argument, which does not compile because __half2 only has the members x and y and cgl_make_half3 and cgl_make_half4 take one value per component.
• GLSL input keeps an assignment or compound assignment that is an operand in parentheses, so (b = b + 1.0) * 2.0 imports as ((b = (b + 1.0)) * 2.0), -(b -= 0.5) as (-(b -= 0.5)) and the arms of ?: likewise. It used to write (b = (b + 1.0) * 2.0), which binds the surrounding operator to the right-hand side of the assignment, so every target computed something else or did not compile.
• CUDA and HIP output compute % and mod on half vectors per component, as for float vectors, on the components converted to float and rounded back to __half: % with fmodf and mod with the floored quotient. The helpers used to apply % to the __half components, which is not a floating-point remainder and does not compile because __half defines no % operator.
• CUDA and HIP output lower length, normalize, distance and cross on cgl_half3 and cgl_half4 vectors to generated helpers that compute in float through __half2float, such as cgl_cgl_half3_length(s), and distance on float and double vectors to a cgl_float3_distance style helper. They used to write bare calls such as length(s) and distance(a, b), which CUDA and HIP do not declare.
• CUDA and HIP output apply mix, fract, step, atan with two arguments and fma to f16vec3 and f16vec4 values through per-component helpers that widen each __half to float and narrow the result back, as in __float2half(fmaf(__half2float(arg0.x), __half2float(arg1.x), __half2float(arg2.x))). They used to write the calls unchanged, as in lerp(v, w, s) or fmaf(v, w, v), which do not compile because no overload accepts the generated cgl_half3 and cgl_half4 structures.
• HLSL Shader Model 6.9 dx::MaybeReorderThread(hint, bits) is refused with ShaderExecutionReorderingUnsupportedError for every target other than Direct3D, which writes the call back unchanged. The Metal, GLSL, Slang, CUDA, HIP, Rust and Mojo generators copied the HLSL spelling into their output, where each target compiler rejects it.
• CUDA and HIP reference parameters that the function writes, T& and T (&v)[N], import as inout, so the writes reach the caller; a reference the function only reads and every const T& import by value. They all used to import by value, so writes were lost. A call that binds such a parameter to shared, global or pointed-to memory, to a local whose address is taken, or to two overlapping arguments raises ReferenceParameterError, as do written T&& and T*& parameters. HIP input parses T (&v)[N] parameters.
• The CUDA and HIP importers write a namespace-scope constexpr or const variable whose initializer is a constant expression, such as constexpr int N = 3; or const float S = N * 0.5f;, as a CrossGL constant, const i32 N = 3;, which CUDA and HIP output declare const and HLSL output static const (#3328). It used to import as a mutable global, var N: i32 = 3;, so CUDA and HIP output wrote a host variable that the kernel cannot read, and HLSL output a uniform without its value. A const whose initializer calls a function or reads a variable, such as const float R = sqrtf(2.0f);, pointers, arrays, __constant__, __device__, __shared__ and extern declarations and function locals keep their lowering.
• CUDA and HIP output combine a vector with a scalar constructor operand such as float(i) or int(x) under +, -, *, / and % through the existing vector-scalar helpers, as in cgl_float3_mul_scalar(v, float(i)). They used to keep the native operator, as in (v * float(i)), which does not compile because float3 and the other vector types define no arithmetic operators. Scalar constructor calls now have the type they construct.
• CUDA and HIP output combine a half vector with a half scalar under +, -, * and / through the existing vector-scalar helpers, such as cgl_cgl_half3_mul_scalar(v, h). They used to keep the native operator, as in (v * h), which does not compile because the generated cgl_half3 and cgl_half4 structures define no arithmetic operators.
• CrossGL array<T, N> with a nonzero extent is the array T name[N] in parameters, locals, struct members, globals and return types, so CUDA and HIP input, which spells C++ arrays that way, reaches GLSL, HLSL, Slang and SPIR-V as arrays. GLSL, HLSL and Slang used to write array<i32, 2> verbatim and SPIR-V refused it. array<T, 0> and array<T> keep the standard-array type. SPIR-V output passes a by-value array held in a struct member or an array element through its own Function variable, since OpFunctionCall takes only memory object declarations; spirv-val rejected the access chain it used to pass.
• CUDA and HIP output splat a half vector constructor whose only argument is a component of a native half2, such as f16vec3(b.x) for an f16vec2 b, into cgl_make_half3(b.x, b.x, b.x). It used to be written as cgl_make_half3(b.x), which does not compile because cgl_make_half3 takes three components.
• WGSL output drops function prototypes before matching definitions, preserving the real function body when a prototype appears first (#3207).
• WGSL output discards the value of a call statement to a @must_use built-in such as min(a, 4u); or a structure constructor such as Pair(a, b); with a phony assignment, _ = min(a, 4u);. The call used to be written as a statement, which WGSL rejects. Atomics, barriers, textureStore and functions the shader declares are written unchanged.
• WGSL output lowers the two-argument frexp(x, e) and modf(x, i) through a generated helper per value and written type, such as crosstl_frexp_vec3_f32_vec3_i32(a0, &e), which calls WGSL's one-argument built-in, stores its exp or whole member through the pointer and returns fract; a written variable of another type, such as the float exponent of HLSL's frexp, receives the member converted to its type. The calls used to be copied unchanged, which naga rejects with too many arguments passed to frexp because WGSL's frexp and modf take one argument and return a structure. A written argument that WGSL cannot pass as a ptr<function, T>, such as e.x, a structure member, an array or storage buffer element or a module-scope variable, is refused, as it is for out parameters (#2520).
• OpenGL and WebGL output accept vals[k++] = 1.0; vals[k] = 2.0;. ASTNode.child_nodes yielded a child kept under two field names, such as ArrayAccessNode.index and index_expr, once per name, so the GLSL interval analysis applied the increment twice, took k to be 2 and refused the in-range vals[k] as a constant index out of range. Each child is now yielded once (#3233).
• WGSL output writes a one-scalar matrix constructor such as mat2(s) or mat2x3(s) column by column with s on the diagonal and zeros elsewhere, as in mat2x2<f32>(vec2<f32>(s, 0.0), vec2<f32>(0.0, s)), including a negated argument such as mat2(-1.0), and converts an integer argument to the element type. It used to be copied as mat2x2<f32>(s), which WGSL rejects because it has no diagonal matrix constructor. An argument with side effects, which would be evaluated once per diagonal element, is refused with a message asking for a local variable.
• CUDA and HIP output splat a vector constructor whose only argument is a scalar cast, such as ivec3(int(x)), through the cgl_int3_splat helper. It used to be written as make_int3(int(x)), which does not compile because make_int3 takes one argument per component.
• The Slang importer keeps out and inout on the parameters of every function, not only of entry points, so a helper's writes to an inout int k reach the caller (#3215). A constref parameter imports as the value it reads. A ref parameter, which names the caller's storage, is refused with SlangImportRefused; it used to be read as a type named ref and produced invalid CrossGL.
• The Slang importer reads a cast of a literal zero to an array type, such as (float[3])0 or (float4[2])0, the idiom for a zero-filled array, as float[3](0), as the HLSL importer does. It used to stop with Unexpected token in declaration: NUMBER (#3348). A cast of any other operand to an array type is refused with SlangImportRefused, because Slang fills only the first element from a scalar, as in {x, 0, 0} for (float[3])x, where HLSL fills every element.
• CUDA and HIP output keep sinpi and cospi, which both math libraries provide exactly at integers and half-integers: sinpif and cospif for float, sinpi and cospi for double, per component for vectors. They used to be refused with MetalMathFunctionUnsupportedError, whose message also named MSL when the call came from CUDA or HIP source.
• Mojo output of a component-wise vector comparison such as a < b on vec2 operands is a SIMD[DType.bool, 2] mask written with a.lt(b), and any and all of a boolean vector become reduce_or() and reduce_and(), ignoring the padding lane of three-component vectors. A bvec2 initialized from a comparison used to be refused.
• OpenGL and WebGL output copy an unqualified array parameter in, as GLSL and every other target do, so a helper that writes its array parameter no longer changes the caller's array (#3201). The generator gave every written fixed-array parameter inout, because asking whether a helper writes its array parameter also marked it for writeback. Writeback now stays with callers that pass caller-visible storage such as &local, and with the helpers they forward it to.
• Mojo output splats a bool into a boolean vector with the fill keyword, as in SIMD[DType.bool, 2](fill=s) for bvec2(s) and for a bool compared with a bvec2 on either side. It used to write SIMD[DType.bool, 2](s), or pass s unchanged on the right of a comparison, which Mojo rejects because Bool is not Scalar[DType.bool]. A comparison or && / || of scalars is now typed bool rather than its operand type, so bvec2(x < 1.0) is written SIMD[DType.bool, 2](fill=(x < 1.0)) instead of calling .cast[DType.bool]() on a Bool, a && (x < 1.0) on a bvec2 a splats the comparison like a bool operand, and mix(a, b, x < 0.5) on scalars selects b or a instead of passing a Bool to lerp.
• Mojo output writes ! of a boolean vector, such as !(a == b), as the lane-wise ~ and types the result as a mask of the same width, so any, all and mix read it per lane and bool r = !(a == b); is refused with MojoScalarConversionError. It used to write not, which Mojo rejects where a SIMD[DType.bool, N] mask is expected and, assigned to a bool, computes !any(a == b).
• Mojo output writes a float or integer vector built from a bool, such as vec2(s), vec2(s, x) or ivec2(s), as SIMD[DType.float32, 2](Float32(Int(s))), SIMD[DType.float32, 2](Float32(Int(s)), x) and SIMD[DType.int32, 2](Int32(Int(s))), and bool vector components the same way, as in Float32(Int(b[0])). It used to write (s).cast[DType.float32](), which Mojo 1.2.0 rejects because Bool has no cast method.
• Mojo output combines boolean vector masks under && and || per component with & and |, as in (x).lt(y) & (y).lt(z). It used to write and and or, which Mojo accepts but evaluates as a whole operand chosen by whether any lane of the left mask is true, so a && b with a = (true, false) and b = (true, true) gave (true, true). A scalar bool operand of a mask is splatted to its lanes, as in (a & SIMD[DType.bool, 2](fill=s)), and two scalar bool operands keep and and or.
• CUDA and HIP output declare a CrossGL out or inout parameter as a C++ reference, float& x or float (&v)[2], so the callee's writes reach the caller (#3159). They used to declare it by value and drop the writes. A call that passes one variable to an out or inout parameter and to another parameter that reaches the caller, where references and copy-out differ, raises CudaLoweringRefused or HipLoweringRefused.
• CUDA and HIP output refuse a swizzle of several components passed to an out or inout parameter, such as set2(v.xy), with CudaLoweringRefused or HipLoweringRefused. They build the swizzle as a new vector, as in set2(make_float2(v.x, v.y)), which the parameter's reference cannot bind to, so the output did not compile.
• CUDA and HIP output refuse a call to a generic function that passes a swizzle of several components to an out or inout parameter, such as bump<vec2>(v.xy, step), or one variable to two parameters that reach the caller, such as both<float>(a, a), with CudaLoweringRefused or HipLoweringRefused, as they refuse such calls to other functions (#3400). They used to write the call to the specialization: bump_vec2(make_float2(v.x, v.y), step) did not compile, and both_float(a, a) bound both references to a.
• WGSL and Slang output keep the named fields of a braced structure initializer such as Pair { b: y, a: x }, in member order, and refuse a missing, unknown or repeated field. They used to construct the structure with no arguments, which zeroed every field.
• WGSL output gives imageSize the shape GLSL defines: i32 for 1D images, vec3<i32> for 3D images, and vec3<i32> with the layer count last for 2D array images, or the number of cubes for cube array images. It used to write vec2<i32>(textureDimensions(image)) for every image, which naga rejects for 1D, 3D and array images.
• A for-loop initializer that declares several variables, for (int a = 0, b = n; a < b; a++, b--), reaches GLSL, WebGL, MSL, Slang, CUDA and HIP output as one declaration and SPIR-V output as both declarations. GLSL, WebGL, MSL and Slang output used to write the parser's list of declarations into the loop header, CUDA output wrote the list of updates, HIP output ran the updates once before the loop, and SPIR-V output refused the program. The GLSL interval analysis also kept the declared variables at their initial values inside such a loop and dropped branches it took to be unreachable. Declarations of different types raise ForInitializerError (#3252, #3253, #3255, #3257, #3258, #3261, #3278).
• CUDA and HIP output writes min and max of integer operands as the integer intrinsics instead of fminf and fmaxf, which converted the operands to float and rounded values above 2^24, and writes double operands as fmin and fmax. Operands of different integer types are cast to the type they promote to, as min(static_cast<unsigned int>(i), 32u).
• WGSL output gives resources without an explicit binding the lowest free slots in their bind group, as GLSL and HLSL output do. Explicit bindings in any group used to push them higher: next to cbuffer Params @binding(3), two storage buffers without a binding took slots 4 and 5 instead of 0 and 1.
• Slang, CUDA and HIP output write a nested { ... } block, such as a block in a function body or a braced case, as a block. Slang output used to write BlockNode(statements=2);, which slangc rejects, and CUDA and HIP output dropped the braces, so declarations in the block redeclared outer names. A declaration in the block ends with it, so a shadowed outer variable keeps its own type after the block (#3259, #3270, #3272).
• Metal output of a mixed-signedness vector operation declared with an integer vector type, such as int4 r = a + b on an int4 a and a uint4 b, keeps both operands in the common type, as in int4 r = int4(uint4(a) + b);. It used to convert the right operand back to the declared type, as in int4(uint4(a) + int4(b)), which metal rejects because the operands are mixed again. int4 q = a + g on a float4 g is likewise written int4(float4(a) + g) instead of int4(float4(a) + int4(g)).
• Metal output of matrix shape conversions such as (float3x3)m and of fragment colour attachment inputs compiles, and HLSL GetDimensions assigns its stride output on Metal.
• CUDA, HIP and OpenCL C input keep a nested { ... } block and a braced case body as a CrossGL block. The parsers spliced it into the enclosing block, so a declaration in it redeclared an outer name, or one in another case, and a use after the block read the inner variable (#3256, #3273, #3274, #3276). A HIP union shadowed in a nested block, which was refused for that reason, now translates.
• HLSL output of Metal sources gives half fragment returns SV_Target, declares half textures without a stray array size and drops <TargetConditionals.h>.
• The runtime loader manifest's targetProfiles for HLSL artifacts, from directx_target_profiles_for_source, is ["directx-12"] when the output has ray tracing stages (raygeneration, intersection, anyhit, closesthit, miss, callable), mesh or amplification stages, or wave and quad intrinsics such as WaveActiveSum and QuadReadAcrossX. It used to also list directx-11, which cannot run them and whose dx11 HLSL profile refuses them.
• GLSL, Slang, CUDA, HIP and OpenCL C input write a C-family octal literal such as 017 as the CrossGL octal 0o17. They copied it unchanged and CrossGL read it as decimal, so 017 + 010 computed 27 instead of 23 in every target (#3282). A leading-zero integer with an 8 or a 9 raises IntegerLiteralError.
• CUDA, HIP and OpenCL C input rename a struct named let, match or var, which CrossGL reads as keywords, and in CUDA input one named kernel, at its declaration and every use, including constructors such as let{1.0f}, as HLSL input does (#3386). Enums, type aliases, functions and variables with those names are renamed the same way, and a rename skips every name the source uses, so let becomes let__ where the source has a let_. These importers used to keep the name, or rename only a struct's declaration, so the CrossGL they wrote did not parse (SyntaxError: Expected binding name, got LBRACE).
• An early return; in a GLSL main that writes stage outputs returns those outputs on every target.
• CUDA output writes CrossGL scalar casts such as f64(x) with C type names.
• CUDA and HIP output write an HLSL Interlocked* call that returns the original value through its last argument, on a groupshared variable or a member of a structured-buffer element, as an assignment, as in orig = atomicAdd(&counter, 1u);, and InterlockedCompareStore as atomicCAS. The call used to be replaced with /* unsupported ... */ 0u;, so the operation never ran and the original value was never assigned, and InterlockedCompareStore was copied unchanged, which does not compile.
• Metal output converts the result of a vector expression to the declared type rather than each operand, so float4 x = a / int4(2, 2, 2, 2) stays an integer division and int4 x = b / 2u an unsigned one (#3166). A vector whose component type ranks below a scalar operand's is converted first, float4(a) * 0.5, because MSL would convert the scalar to the vector's component type instead; the result then has the common type, so int4 x = a & 4u converts the uint4 result back. A floating literal keeps a half vector's precision: h * 2.0 stays a half4 product.
• CUDA output passes the address of the target to block- and system-scoped atomics such as atomicAdd_block and atomicInc_system, as it already did for atomicAdd and atomicInc.
• CUDA and HIP output copy an in array parameter that the function writes, such as an HLSL float v[2] parameter, into a local array at function entry, so the caller's array keeps its values as it does in HLSL. The parameter used to be emitted as a C array, which decays to a pointer, so writes in the callee changed the caller's array. A parameter counts as written when the function assigns, increments or takes the address of an element, writes through it with *v, v->a or v + 1, or passes the array or an element to an out or inout parameter, a non-const pointer parameter, an unqualified array parameter or frexp, modf or sincos, including through an explicitly specialised generic call such as fill<int>(v, 7). Unwritten in arrays, out and inout arrays, arrays without a qualifier and arrays of pointers, references, structured buffers or other resources are unchanged; a struct the program declares is copied whatever its name, such as Texture2DParams.
• Metal output renames a user function named like a Metal Standard Library function, such as fwidth, min or simd_sum, with every overload and every call to it (fwidth_fn). using namespace metal; made a call to it ambiguous, which the Metal compiler rejects (#3156).
• Metal output renames a shader-scope constant spelled like a Metal keyword, such as device, half, thread or constant, together with its uses, as it already did for program-scope variables: const float device = 2.0; becomes constant float device_ = 2.0;, and a type constructor such as half(1.0) keeps its name. The constant used to be emitted verbatim, which the Metal compiler rejects (#3385).
• Metal output lowers length, dot, distance and normalize on scalars to abs(x), x * y, abs(x - y) and sign(x). Metal declares them only for vectors, so the scalar calls used to fail to compile with call to 'length' is ambiguous.
• A program with several stages of one kind, such as a Metal library with several vertex functions, has every stage rewritten, not only the first.
• Translating a malformed shader no longer takes time that grows much faster than its length.
• SPIR-V output evaluates a subscript index that writes, such as k--, next() or (j = k + 1), once when a declaration's initializer reads an element (float a = t[k--];, S s = data[k++];, float x = s[k++].x;) and when an assignment is used as a value (a = (t[k++] = x);). The emitter first checked whether the initializer aliased a resource, or stored through the assignment target, and then lowered the access again, so the side effect ran twice and the read used the wrong element (#3294).
• The Metal check that refuses a derivative a kernel, mesh or object function reaches through a helper follows every overload a call can select: those that accept its argument count, narrowed by the argument types when they decide. It followed only the last overload of the name, so a derivative in an earlier overload reached the Metal compiler, which rejected the kernel (#3168).
• An assignment used as a value keeps its own parentheses: int f = (e = 4) + 1; used to become e = 4 + 1 on several targets, which assigns 5.
• CUDA, HIP and Metal output keep the parentheses around the operand of a postfix ++ or -- unless it is a name, member access, index or call, so (*p)++ stays (*p)++. It used to be written *p++, which increments the pointer and discards the value because postfix ++ binds tighter than unary *, and does not compile for an array, as in (*v)++ on a float v[2].
• CUDA and HIP output keep the parentheses of a nested prefix operator that would otherwise merge with the one before it, so -(-a) stays -(-a) and -(--a) stays -(--a). It used to be written --a and ---a: a pre-decrement where the source has none, which changes the variable, and output that does not compile.
• SPIR-V output short-circuits &&, || and ?:. The right operand and both arms used to run unconditionally.
• WGSL output writes a conditional on a structure, array or matrix that is the whole initializer of a local variable or the whole value of an assignment, such as Pair r = c ? p : q;, as var r: Pair; and an if/else that assigns the chosen branch, so only that branch is evaluated. It used to write select(q, p, c), which naga rejected because select takes scalars and vectors only. Such a conditional anywhere else, such as (c ? p : q).a, a return value or a module-scope constant, and an initializer that reads an outer variable of the name it declares, are refused with a ValueError.
• GLSL output declares long, ulong, size_t and ptrdiff_t as 64-bit integers instead of 32-bit ones, which changed arithmetic and block layouts; GLSL ES and WebGL refuse them. SPIR-V output converts size_t(x) and ptrdiff_t(x) and widens uint without sign extension.
• 8- and 16-bit integer vector arithmetic wraps at its own width on GLSL and Metal, and Metal widens native ushort2 and similar fields and buffer elements when it reads them.
• Storage buffer arrays translate to GLSL, SPIR-V, HLSL, WGSL, Slang and Metal. Sized arrays on HLSL used to be static arrays whose stores never reached memory, and unsized arrays of structures were refused on HLSL and WGSL.
• WGSL output writes the length() of a buffer T* storage buffer, such as outp.length() for buffer float* outp, as i32(arrayLength(&outp)), the number of T elements the bound buffer holds, and passes a buffer T* helper parameter, which is already a pointer, as arrayLength(values). It used to copy outp.length() unchanged, which naga rejects (#2396).
• A GLSL entry point that writes one of its inputs, built-ins included, writes a local copy instead of the read-only input.
• HLSL entry points that write stage outputs through out and inout parameters translate to GLSL, HLSL, Metal, SPIR-V, WGSL and Slang, and SV_Target<N> reaches colour attachment N on GLSL, WGSL, Slang and Metal. Multi-stage shaders written this way used to fail, and GLSL and WGSL numbered attachments in declaration order.
• WGSL output of a switch has one default arm and keeps fallthrough by merging case bodies, writes discard;, and converts scalar operands of vector stores, shifts and bitwise operators, which naga used to reject.
• WGSL output declares the integer and unsigned 1D, 3D, cube and cube array samplers, such as isampler3D or usamplerCube, as texture_3d<i32>, texture_cube<u32> and so on. It used to declare them under their CrossGL names, which naga rejects. sampler2DRect, samplerBuffer and the other samplers with no WGSL texture type are refused with a diagnostic, as sampler2DMS already was.
• An HLSL vertex input with SV_Position is a regular vertex attribute on GLSL, Metal, WGSL and SPIR-V. Metal and WGSL rejected it, and SPIR-V read the gl_Position output instead.
• WGSL output writes the flat, noperspective, centroid and sample qualifiers of floating-point vertex outputs and fragment inputs, from CrossGL, GLSL and the HLSL nointerpolation, noperspective, centroid and sample modifiers, as @interpolate(flat), @interpolate(linear), @interpolate(perspective, centroid), @interpolate(perspective, sample) and their linear combinations, and gives a vertex output the interpolation of the fragment input at its location, which WebGPU requires to match. They used to be dropped, so every such value was interpolated with perspective correction at the pixel centre. flat with centroid or sample becomes @interpolate(flat), its meaning in GLSL, and integer IO with noperspective or smooth is refused with WGSLStageIOUnsupportedError (#3483, #3519, #3520, #3521, #3522).
• WGSL output refuses texture(), textureLod(), textureGrad(), textureOffset() and their offset forms on integer samplers such as isampler2D or usampler2DArray with WGSLIntegerTextureSampleUnsupportedError. It used to write textureSample, which WGSL only defines for f32 textures, so naga rejected the output. texelFetch() and textureSize() on these samplers still translate.
• Function-scope static variables keep their value across calls on every target, and HLSL output keeps the initializers of static globals.
• HIP output keeps static and const on variable declarations, as CUDA output does, so an HLSL static const int N = 3; stays a constant that kernels can read and use as an array size. It used to be written as int N = 3;, a mutable host variable that hipcc rejects in device code and as the size of float a[N].
• WGSL output writes the bias form of texture(), texture(s, coords, bias), as textureSampleBias(s, s_sampler, coords, bias) with the texture's split sampler, passing the layer of a 2D array or cube array texture as its own argument and converting an integer bias with f32(). It used to write textureSample(s, coords, bias), which passes the coordinate as the sampler and the bias as the coordinate, so naga rejected the output (#2462). A bias on a 1D texture, which textureSampleBias has no overload for, or on an element of a sampled texture array selected at run time is refused with WGSLTextureBiasUnsupportedError, and a bias in a compute stage, which has no derivatives, with ImplicitLodStageUnsupportedError, as in a vertex stage.
• WGSL output of texture(), textureLod(), texelFetch() and texelFetchOffset() on 2D array and cube array textures passes the layer as its own array_index argument, as WGSL requires. It used to leave the layer inside the coordinate, so naga rejected the output. Sampled lookups round the layer with floor(layer + 0.5), as GLSL does.
• WGSL output of textureGrad(), textureOffset(), textureLodOffset(), textureGradOffset() and the shadow lookups on 2D array and cube array textures, sampler2DArrayShadow and samplerCubeArrayShadow included, passes the layer as its own array_index argument too. They used to leave it inside the coordinate, so naga rejected the output. The offset forms on cube array textures, which WGSL does not have, are refused with WGSLCubeTextureOffsetUnsupportedError.
• Vector && and || combine per component on every target. HLSL 2021 and WGSL output used to be invalid, and GLSL refused it.
• Metal output of HLSL vertex inputs that mix attributes with built-ins such as SV_VertexID compiles, and an int SV_VertexID is accepted.
• Metal fragment parameters whose members carry [[color(n)]] read colour attachments on every target, or are refused where the target has no framebuffer fetch. They used to become ordinary inputs. A default-constructed Metal sampler gets MSL's default sampler state.
• The Metal importer reads m.columns[i] of a matrix as the column m[i]. The columns array used as a value is refused with MetalMatrixColumnsError.
• HLSL output stores bfloat16 vectors as per-lane uint payloads, writes fwidthFine and fwidthCoarse with fine and coarse derivatives, and the HLSL importer maps packed 8-bit structured buffer elements.
• The HLSL importer reads round as CrossGL roundEven, because HLSL rounds halfway cases to the nearest even integer. Metal output used to write round and CUDA and HIP output roundf, which round them away from zero, and GLSL output round, which leaves them to the implementation, so round(2.5) gave 3 instead of 2. Metal output now writes rint, converting an integer operand to float because rint has only float and half overloads, as in rint(float(i)); CUDA and HIP output write rintf, or rint for a double operand, and round float and double vectors per component; GLSL output writes roundEven; Slang output writes rint, which rounds halfway cases to even on every Slang target, where Slang round leaves them to the target; Rust output calls round_even instead of round; and Mojo output calls its round_even helper, which now has overloads for Float16, Float32 and Float64 scalars and 2- and 4-wide vectors, converting an integer operand to Float32 and spelling a literal operand as Float32(2.5). A function the shader defines named round is still called when the call's argument count fits its parameters.
• Rust output writes a call with explicit type arguments with the turbofish, as in subgroup_broadcast_const::<_, 0>(a0). It used to write subgroup_broadcast_const<_, 0>(a0), which rustc reads as two comparisons and does not parse (#2342).
• A scalar swizzle such as s.xxx becomes a vector constructor on Metal, WGSL and GLSL ES, which do not accept it.
• Slang output reads gl_NumWorkGroups as WorkgroupCount(), and Metal output infers the type of a shader-scope let from its initializer.
• WGSL output declares a shader-scope let or var with no written type with the type of its initializer, as in var<private> scale: f32 = 2.0;, and refuses one whose initializer has no type it can infer. It used to declare it as var<private> scale: void = 2.0;, which naga rejects because void is not a WGSL type (#2444).
• Metal device T& and const device T& arguments that reference a structure read the buffer on every target. HLSL, GLSL and SPIR-V used to declare them as stage inputs, so the shader read vertex attributes or nothing.
• Metal [[vertex_id]], [[instance_id]] and [[primitive_id]] arguments keep their uint type, so arithmetic on them wraps as in MSL. They used to import as int.
• OpenCL C barrier(flags) and work_group_barrier(flags[, scope]) order the memory their flags name. CLK_GLOBAL_MEM_FENCE and CLK_IMAGE_MEM_FENCE put the CrossGL buffer or image fence on each side of the execution barrier, as in threadgroup_barrier(mem_flags::mem_device) for Metal, DeviceMemoryBarrierWithGroupSync() for HLSL and storageBarrier() around workgroupBarrier() for WGSL. They used to lower to workgroupBarrier() whatever the flags said, which orders work-group memory only, so a write to a global buffer before the barrier could be read stale after it. Flags that are not only CLK_ names, such as a variable, order all memory. A memory_scope wider than the device is refused with OpenCLFenceScopeUnsupportedError.
• The OpenCL importer reads a vector constructor whose type is a typedef of an OpenCL vector type, such as (duo)(x, 3.0f) for typedef float2 duo;, as a vector constructor, resolving the alias to its target type. It used to read it as a cast of a comma expression, so duo b = (duo)(x, 3.0f); lowered to vec2<f32>((x , 3.0f)), which keeps only the last operand; the DirectX generator then refused the initializer and the Metal generator wrote float2({x, 3.0}).
• A GLSL local whose initializer reads an outer variable of the same name, as in float scale = scale * 3.0;, is renamed on HLSL, Metal, CUDA, HIP and Slang, where the initializer would read the new, uninitialized variable. The GLSL importer no longer reads a local that shadows a member of an anonymous buffer block, or of a uniform block it flattens, as that member.
• Row-major matrix struct members (@ row_major, from GLSL layout(row_major) block members, HLSL and Slang row_major and SPIR-V RowMajor) keep their bytes on GLSL, WebGL, SPIR-V, Metal, WGSL, CUDA, HIP, Rust and Mojo. Those targets declare the member as its column-major transpose and rewrite every read and write; HLSL and Slang do so for GLSL buffer blocks only. The member used to be declared column-major, so reads, writes and the offsets of later members changed. Uses the transposed storage cannot keep, such as writing one column, raise RowMajorMemberError.
• GLSL output zero-initializes arrays with one constructor argument per element, and refuses an array whose size it cannot evaluate with OpenGLZeroValueError. HLSL output writes a one-argument vector constructor of unknown width as a cast, which splats a scalar.
• Locals named after built-ins that GLSL and HLSL output call, such as texture, textureSize, mod and imageSize, are renamed so they no longer hide those functions.
• HLSL output adds no second SV_Position parameter to a fragment entry whose input structure has one, declares int members that carry SV_RenderTargetArrayIndex, SV_ViewportArrayIndex, SV_PrimitiveID or an index system value as uint, and gives members named view, layer or viewport no system-value semantic. A Metal varying float3 view used to carry the view index.
• WGSL output binds a fragment input that carries gl_PrimitiveID, SV_PrimitiveID or MSL [[primitive_id]] to @builtin(primitive_index) as u32 and starts the module with enable primitive_index;, and reads an int input through i32(...), so the shader computes in the type it declared. It used to refuse the input with WGSLStageIOUnsupportedError, and wrote a CrossGL @ primitive_id input as a user @location. WGSL has the primitive index only as a fragment input, so it is refused in other stages, as an output and with a non-integer type, as is a write to an int structure member that carries it (#3505).
• Early fragment tests reach HLSL as [earlydepthstencil], SPIR-V and Metal output, and the HLSL importer reads [earlydepthstencil]. Only GLSL output kept them.
• The HLSL importer keeps a value-returning function marked [numthreads] as a helper, as DXC does, and GLSL output converts a scalar returned by a float4 pixel entry to the output's type.
• The Slang importer keeps an entry point's [WaveSize(...)] attribute, in the single and the range forms, as CrossGL @ WaveSize(...), as the HLSL importer does. It used to drop it without a diagnostic, so DirectX output ran with whatever wave size the driver picked. Slang and SPIR-V output now refuse the attribute, as they already did for HLSL input.
• SPIR-V output declares immutable scalar, vector and matrix globals with constant initializers as constants, so constant uint T = W * 3; is valid, and flattens structure members of vertex outputs and fragment inputs that carry no location, such as the nested structures of a Metal mesh pipeline's [[stage_in]]. They used to overlap the next member's location.
• GLSL output of a multi-stage file emits each helper only under the macros of the stages that call it, so a fragment-only helper that calls fwidth no longer breaks the vertex stage. Written half-precision stage inputs are copied into locals, and varyings no longer take the names of entry resource parameters.
• Depth comparisons with 16-bit float reference values translate.
• WGSL output converts a signed shift amount to u32 whatever its expression form, so a bitwise amount such as 1 << (n & 3) joins the arithmetic, name and element spellings that were already converted, and a ~ amount or the select a ternary lowers to are converted the same way. A bitwise amount was left signed, which naga rejects with automatic conversions cannot convert elements of i32 to u32. A signed bitwise scalar beside a vector bitwise operand is now replicated into a vector, as the other scalar spellings are.
• An int or uint image load assigned to a float vector, such as a uimage2D imageLoad assigned to a vec4, and an int load assigned to a uint vector convert the value as GLSL does: Metal output writes float4(tex.read(...)) and HLSL output converts through the __crossgl_uint_to_float4 helpers. Metal, HLSL and GLSL output used to refuse the shader as a result-kind mismatch. A float load into an integer vector, which GLSL rejects, is still refused.
• HLSL and Metal output pass gl_BaryCoordEXT, gl_BaryCoordNoPerspEXT, their NV spellings and gl_ShadingRateEXT, which GLSL reads as globals, to the fragment entry point as parameters carrying the built-in and on to the helpers that read them, and every target but GLSL replaces the gl_ShadingRateFlag2VerticalPixelsEXT style flags with their values 1, 2, 4 and 8. Both used to write the gl_ names unchanged, which neither compiler declares. HLSL output gives a no-perspective barycentric input SV_Barycentrics1, so a shader that reads both no longer declares two SV_Barycentrics inputs, and Metal output gives a no-perspective parameter [[center_no_perspective]] instead of interpolating it with perspective. Metal output refuses a fragment stage that reads both interpolations with MetalBarycentricsUnsupportedError and still refuses gl_ShadingRateEXT.
• Metal output accepts a scalar or non-float4 vector colour return, such as float main(...) : SV_Target, and refuses only bool; it used to require float4. A fragment function whose [[stage_in]] structure carries SV_Position and that also reads the fragment coordinate, as a lowered Vulkan SubpassLoad does, reuses that member instead of declaring a second [[position]] input, which MSL rejects. DXR ray-state intrinsics such as WorldRayDirection() and RayTCurrent(), and GLSL ray built-ins such as gl_HitTEXT, in closest-hit, any-hit, miss, callable and intersection stages are refused with MetalRayDispatchUnsupportedError; they used to be written unchanged into [[visible]] functions, which do not receive them.
• Metal output renames a user function, program-scope constant or structure named like a type, enumeration or alias that using namespace metal; brings into scope, such as bias, level, filter, coord, access or quad, together with every call, reference and constructor, as in constant float bias_ = 0.005;. Such a name used to fail to compile with reference to 'bias' is ambiguous or redefinition of 'quad' as different kind of symbol. Parameters, locals and members keep their names.
• HLSL and Metal output declare the built-in outputs that a vertex or fragment stage writes by name, such as gl_Position = p; or gl_FragDepth = d;, and the stage's in and out variables in the entry point's input and output structures. They used to write them as undeclared identifiers, and Metal output declared the stage variables as uninitialized locals that read and wrote nothing. HLSL output maps gl_SampleID to SV_SampleIndex, the NV barycentrics to SV_Barycentrics and the shading rates to uint SV_ShadingRate, and reads the fragment subgroup, helper-invocation and sample-count built-ins through intrinsics. SPIR-V output declares gl_HelperInvocation, gl_FullyCoveredEXT, the AMD barycentrics, gl_DeviceIndex and the NV shader-SM built-ins with their capabilities and extensions, and reads gl_BaryCoordNV as the KHR built-in. GLSL output writes gl_Layer and gl_ViewportIndex from vertex and tessellation evaluation stages under GL_ARB_shader_viewport_layer_array.
• @invariant on a gl_Position output, whether a structure member, a stage variable or the entry point's return, is kept: GLSL and GLSL ES output redeclare invariant gl_Position;, Metal output writes [[position]] [[invariant]] and SPIR-V output decorates the Position built-in Invariant. GLSL output used to drop the qualifier, and Metal output read @invariant @gl_Position as a user varying and wrote a zero position. Metal output refuses @invariant on any other output with UnsupportedMetalFeatureError, because MSL allows [[invariant]] only on the position.
• SPIR-V output declares DepthReplacing and StencilRefReplacingEXT for a fragment stage that writes the depth or stencil reference through a stage-level output, as it did for one that returns a structure. The module used to omit the execution mode, which spirv-val rejects.
• Rust output passes out and inout parameters as &mut T, copies an inout value into a local at entry and back before every return and at the end of the body, and passes &mut place at call sites. It used to pass them by value, so the callee's writes never reached the caller. An argument that cannot be borrowed mutably, such as a swizzle, an expression that is not a place or one place passed to two such parameters, is refused with RustLoweringRefused. The Rust importer reads &mut T and &mut self parameters as inout instead of by value.
• CrossGL cbuffer P { ... } p; and cbuffer P { ... } p[N]; declare the instance, as a structure P and a uniform ConstantBuffer<P> p, and a @buffer_layout on a ConstantBuffer<T> places the members of T at the declared offsets on every target whose own rules differ. The parser used to drop the instance name, so every target read an undeclared name, and every generator laid T out by its own rules.
• GLSL output declares a ConstantBuffer<T> that a stage declares without uniform as a uniform block, as it does a module-scope one. It used to drop the declaration and leave every read of it undeclared.
• Metal output declares a @ row_major matrix member of a constant block without a @buffer_layout as its column-major transpose and rewrites every read, as it does for structure members. It used to declare the member column-major, so every read saw the transpose, and a non-square matrix took the size of the other shape, which moved the members after it.
• CUDA and HIP output rename a user function named like a device library function, such as rsqrt, fabs, sinf, atomicInc or make_float2, with every call to it, as rsqrt_fn; nvcc and hipcc rejected the redefinition or the ambiguous overload. WGSL output calls a user function named like an HLSL built-in, such as lerp(x, 1.0), instead of rewriting the call to mix. Entry points keep their names.
• CUDA and HIP output declare 64-bit integer vectors, i64vec2 to u64vec4 and vec2<i64>, as longlong2 to ulonglong4 with their make_ constructors, buffer pointers included, and keep all 64 bits in component reads, swizzles, scalar operands, bitwise operators, abs and sign. They used to write u64vec2, vec2<i64> or i642, which name no CUDA or HIP type, so a CUDA kernel using longlong2 did not translate back to CUDA.
• CUDA and HIP output declare file-scope storage buffers, textures, samplers and images, with their length, counter and query companions, as __device__ variables that the host binds by name, and constant arrays and vectors as __device__ const variables with brace initializers. They used to be plain globals, which nvcc and hipcc reject in device code (reference to __host__ variable).
• CUDA and HIP output pass the out arguments of frexp, modf, sincos, sincospi and remquo by address, as in frexpf(x, &(e)), and lower vector calls, calls whose out arguments have other types and vector ldexp through per-component helpers; a float exponent to ldexp computes x * exp2(e), as DXC lowers it. The calls used to be written as in CrossGL, which passes an int where an int* is expected and has no vector overloads. Half values, mismatched component counts and out arguments that are not one assignable value are refused.
• OpenGL, WebGL, WGSL and Metal output convert a numeric vector operand of any and all to the boolean vector of its width, as in any(bvec3(mask)) for an ivec3 mask, and write the reduction of a numeric scalar as the scalar converted to bool. These targets reduce only boolean values, so glslangValidator rejected the output with 'any' : no matching overloaded function found.
• WGSL output writes ^ of two boolean operands, such as p ^ q on bool or vec2<bool> values, as (p != q), and a boolean t ^= v as t = (t != v). WGSL defines ^ only for integers, so naga rejected the output. A for-loop update of this form whose target calls a function is refused.
• WGSL output writes a numeric unary plus, +x, as x. WGSL has no unary +, so naga rejected the output. Unary plus on a boolean is refused.
• HLSL output declares a shader-scope static const matrix, an array of them or a constant structure with matrix fields, and writes the product of such a matrix and a vector with mul. It used to refuse the declaration and write the product with *, which HLSL computes component by component. The HLSL importer reads an array of matrices whose brace initializer lists one matrix constructor per element.
• The Metal importer deduces a member template's arguments when the call names a struct template specialization through a type alias, as in a device Box<float16_t>* argument with typedef half float16_t;. It used to compare Box<half> with the unexpanded Box<float16_t> and refuse the call as not binding consistently.
• Slang, CUDA and HIP output compute floating % as the truncated remainder x - y * trunc(x / y), as HLSL, MSL, C, WGSL and SPIR-V do: Slang output keeps its %, and CUDA and HIP output call fmodf, or fmod for double, per component for vectors. They used to floor the quotient as GLSL mod does, so -0.5 % 2.0 gave 1.5 instead of -0.5; mod() keeps the floored form. Metal, OpenGL, WebGL, CUDA and HIP output write a floating x %= y as x = x % y, and refuse it with FloatRemainderAssignmentError when the target has side effects. Metal output used to write %=, which the Metal compiler rejects, and OpenGL output refused it.
• A structure constructor converts each argument of another scalar or vector type with the same lane count to its member type, as GLSL does, so Pair(x, k) with an int k is written Pair(x, float(k)), and CUDA and HIP output write structure constructors as aggregate initializers, Pair{x, float(k)}. CUDA and HIP output used to call a constructor the structure does not declare, Metal's aggregate initialization rejected the narrowing, and WGSL, SPIR-V and GLSL ES output, which have no implicit conversions, did not compile.
• CUDA and HIP output reduce all and any of a vector through generated cgl_<type>_all and cgl_<type>_any helpers that test each component, and of a scalar as ((x) != 0). Neither has a device all or any for vectors, and nvcc's legacy all and any are warp votes on a bool, so the output did not compile. Other operand types are refused.
• CUDA, HIP and OpenCL C input keep an assignment that is the operand of a binary, unary or conditional expression in parentheses, so (t -= 3) > 0 keeps its meaning. It used to import as t -= 3 > 0, which CrossGL reads as t -= (3 > 0).
• WGSL output writes every CrossGL array constructor as a WGSL array value, such as array<f32, 3>(...), including one that is returned, assigned, subscripted or passed as an argument. Only a constructor that initialized a declaration used to be converted; the others kept the CrossGL spelling float[3](...), which naga rejects. An unsized constructor takes its count from its arguments, and a size that does not match the argument count is refused.
• HIP output converts the argument of __ffsll, including the one findLSB of a 64-bit value writes, to unsigned long long unless it already is long long or unsigned long long, calls normf and rnormf for float arrays and norm and rnorm for double ones, and passes atomicAddNoRet the address of its target. hipcc rejected the ambiguous __ffsll call, the float array passed to norm and the value passed to atomicAddNoRet. A norm or rnorm of another element type is refused with HipLoweringRefused.
• CUDA and HIP output write GLSL's lessThan, lessThanEqual, greaterThan, greaterThanEqual, equal, notEqual and not as operators, as HLSL, MSL, WGSL and Slang output do. They used to write calls that nvcc and hipcc reject as undeclared. A user function of the same name is still called.
• CUDA and HIP output declare the structures a stage declares, such as the block type of a GLSL buffer Data { ... } data; written inside a compute stage, at module scope after the shader's structures. They used to write only the shader's structures, so Data data; named an undeclared type. Two stages that declare different structures of one name are refused.
• A helper function that reads a compute or vertex built-in, such as gl_GlobalInvocationID, gl_LocalInvocationIndex, gl_VertexID or gl_InstanceID, receives it from the entry point as a trailing parameter on HLSL, Metal, Slang and WGSL, as fragment helpers already did on some of them, and the vertex and instance index are converted from uint to int once in the entry point. Slang and WGSL fragment helpers receive their built-ins the same way. The helper used to read an identifier those targets do not declare.
• Metal, Slang, Rust, CUDA, HIP and SPIR-V output declare a constant block written inside a stage, such as uniform Params { mat4 m; vec4 bias; }; in a compute stage, as they declare a module-scope one. Metal, Slang, Rust, CUDA and HIP output used to drop the block and read its members undeclared, and SPIR-V output refused the program. A block with the name, members and attributes of a module-scope block is that block, and one that would change what a name means once it moves to module scope is refused with StageConstantBlockConflictError (project.translate.stage-constant-block-conflict). Slang output writes mat * vec on a block member with mul().
• HLSL and GLSL output lower a CrossGL buffer T* p helper parameter as they lower an MSL device T* one: HLSL output passes an RWStructuredBuffer<T> and an element offset, and GLSL output specializes the helper per buffer with an element offset. HLSL output used to write T *p and GLSL output float values[], which DXC and glslang reject. A read-only const buffer T* or MSL const device T* parameter that a call passes a writable buffer gets an RWStructuredBuffer<T> overload of the helper, through helper chains; DXC used to reject the call with no matching function. A helper that would need more than 16 overloads, or whose overloads do not settle within 16 passes, raises DirectXResourcePointerParameterError.
• HLSL output passes a constant or device array parameter that a caller passes a view of, such as strides + 2, as a buffer and an element offset; it used to pass the whole buffer. HLSL and GLSL output write a member or swizzle store through a pointer held in a private aggregate, such as row[i].x = value, through a helper that writes that member in place. They used to assign to the result of the load helper, which DXC and glslang reject. A store to a member of a storage view's element is refused.
• WGSL, SPIR-V and every target other than Metal, HLSL and GLSL replace a CrossGL typedef name with the type it names wherever it is used as a type, following chains of aliases and the scoping of block-scope aliases; Metal, HLSL and GLSL output keep program-scope aliases. WGSL output used to declare var x: real for typedef float real;, which naga rejects. An alias with a qualifier or attribute, and a constructor through an alias of a type without a constructor spelling, are refused with TypeAliasUnsupportedError.
• The HLSL, Slang and CUDA importers write typedef and typealias declarations, block-scope ones included, as CrossGL typedefs. The HLSL importer used to write type real = float;, which no generator reads, the Slang importer dropped block-scope aliases and the CUDA importer wrote a comment, so the imported program used an undeclared type. The CUDA importer writes typedef struct X Y as typedef X Y and drops typedef struct X X, and the HIP importer writes a functional cast through an alias as the aliased type, f32(x) for real(x). The HLSL importer refuses vector<T, N> and matrix<T, R, C> with more than four lanes, which it used to read as vec4 and mat4.
• CUDA and HIP output write CrossGL float literals with the f suffix, as in 2.0f and 1e-3f, so float arithmetic stays in single precision and float2 v = {1.0f, 2.0f} does not narrow from double. float y = x * 2.0; used to compute in double precision, and the initializer failed hipcc -Werror=c++11-narrowing.
• == and != on structures and constant-length arrays compare member by member on every target other than GLSL and GLSL ES, which have the operators: p == q becomes the && of its scalar comparisons, all(a == b) of its vectors and of its matrix columns, in member order, and != the || of != comparisons. MSL, HLSL, WGSL, SPIR-V and Slang output used to write the operators unchanged, which their compilers reject, and CUDA and HIP output compared arrays as pointers. Operands with side effects are refused with AggregateEqualityOperandError, and samplers, textures, pointers, runtime-length arrays, operands of unknown type and comparisons with too many leaves with AggregateEqualityUnsupportedError.
• CrossGL reads t[j = k + 1] on a declared variable as a subscript whose index is an assignment, in initializers, assignment values and call arguments. It used to read the brackets as named specialization arguments and keep t[j = k + 1] as one identifier, which GLSL, HLSL and Metal output copied and SPIR-V output refused. A comma inside the brackets, t[i, j], is a parse error instead of identifier text.
• WGSL output negates u32 scalars and vecN<u32> vectors as (~x + 1u) and folds a negated unsigned literal, -2u to 4294967294u. WGSL has unary - only for signed and floating values, so naga rejected -a on a u32.
• The Metal importer replaces a class template member with an explicit specialization, such as template <> void Owner<float>::put(const device float2* w), when the two match by type, so a member declared through an owner alias such as const device pair_T*, through another template's member alias such as typename Traits<T>::scalar_T, or as vec<float, 2> instead of float2 now matches. The importer compared the spellings, so such a specialization was ignored and the primary body was emitted for the specialized owner without a diagnostic. A specialization whose name matches a member but whose signature matches none of its overloads is refused with MetalTemplateSpecializationError.
• Project translation of Metal sources matches a template helper parameter spelled metal::vec<T, N> or ::vec<T, N> with an unqualified vec<float, N> argument, so a call such as helper<T>(...) in a member of a template structure is materialized. It used to fail with project.translate.template-materialization-unsupported on DirectX and OpenGL. Other namespaces still have to match.
• GLSL, WebGL, HLSL, Slang and Metal output declare a mutable stage-scope variable without a storage, interface or resource qualifier, which has one copy per invocation shared by every function of the stage, at module scope, renaming it when it collides with another name and moving an initializer that is not a constant expression to the start of the entry point. GLSL and HLSL output used to drop the declaration, Slang output declared it as a uniform, and Metal output made it a local of the entry point that the stage's helpers could not see. Slang output also keeps static on module-scope values, which it dropped for HLSL and Slang input.
• HLSL and Slang output write CrossGL m *= n and v *= m with a matrix right operand, and Metal output m *= n, as the product m = m * n or v = v * m, through mul() where the target needs it. HLSL and Slang define *= on matrices component-wise and MSL has no *= for two matrices, so the output computed a different value or did not compile. A target with side effects, which the product would evaluate twice, is refused with MatrixCompoundAssignmentError (#3211).
• The HLSL importer reads a do { ... } while (cond); loop as a CrossGL do/while loop, so a continue in the body reaches the condition. It used to write while (true) { body; if (!cond) break; }, where continue skipped the test, so the loop did not end (#3204).
• HLSL output of an array assignment from a brace list whose elements read the destination or may have side effects, such as a = {a[1], a[0]}, fills a temporary of the array type and then assigns it whole. It used to store the elements one by one while evaluating the list, so a = {a[1], a[0]} left both elements equal to the old a[1]. A list with a call keeps the element stores when the destination is a local or parameter, which the call cannot see.
• HLSL and Metal output of a SPIR-V mesh shader declare its outputs: the index array it writes becomes the indices parameter, and each member of a Block output structure becomes a vertex member at the structure's location plus the locations of the members before it. Both generators used to drop those declarations and keep writes to an undeclared vertexOutput. A structure output that mixes per-vertex and per-primitive members is refused.
• HLSL output accepts a mixed-width integer compound assignment whose right operand reads buffer elements, such as acc += a[i] * b[j] on an int64_t acc with uint64_t and int64_t buffers from a Metal source. The side-effect check took the buffer loads that Metal pointer reads import as for arbitrary calls and refused the assignment with DirectXContextualConversionError. An index with a side effect, such as a[i++], is still refused.
• The SPIR-V importer reads an array constant, such as the OpConstantComposite that glslang writes for an array constructor of constants, as an initializer list and a structure constant as a constructor of its structure. It used to write both as calls to spirv_constant_composite(...), which Vulkan and Metal output refused and GLSL and HLSL output copied unchanged. A composite constant of another type without a constructor is refused with SpirvInstructionUnsupportedError. Metal output stores a list assigned to an array element by element when no element has side effects or reads that array, where it used to refuse every such list (#3060).
• The HLSL importer keeps the namespace of dx::HitObject, so a Shader Model 6.9 program that uses hit objects translates to HLSL, which declares its TraceRay and Invoke payload structures [raypayload]. The importer used to rename the type dx_HitObject, which DXC does not declare. A dx::HitObject type or dx::HitObject:: call is refused with HitObjectUnsupportedError for every target other than Direct3D; those generators copied it into their output, which Metal and glslang reject with use of undeclared identifier 'dx'.
• Slang output keeps the parentheses around the operand of a member access or subscript, such as (v + v).x or (m * 2.0)[1], types the results of matrix +, -, * and /, transpose, inverse, matrixCompMult, outerProduct and ?:, so a subscript of a matrix-valued expression reads a CrossGL column as m[i] does, and lowers matrixCompMult to * and outerProduct to a generated helper per shape. It used to drop those parentheses, so the access applied to the last operand only, and to write matrixCompMult and outerProduct unchanged, which slangc rejects with undefined identifier. An outerProduct whose result is not a float or double matrix and a matrixCompMult without a matrix operand are refused with SlangLoweringRefused.
• The OpenCL importer writes atomic_sub, atomic_fetch_sub, atomic_fetch_sub_explicit and atomic_dec as atomicAdd of the negated operand, which returns the same previous value, negating as 0u - x on uint memory because WGSL has no unary minus on u32. It used to write atomicSub, which CrossGL does not define.
• The OpenCL importer drops the memory_order and memory_scope arguments of an *_explicit atomic, such as atomic_fetch_add_explicit(p, 1, memory_order_relaxed, memory_scope_device), when every order is memory_order_relaxed and the scope is no wider than memory_scope_device, which is what CrossGL atomics provide. It used to pass them on as extra arguments of the CrossGL atomic. Any other order or scope, an order that is not a literal and a wrong argument count are refused with OpenCLAtomicOrderUnsupportedError.
• The OpenCL importer lowers read_imagef(image, sampler, coord) to textureLod(image, sampler, coord, 0.0) and the form without a sampler to texelFetch, or imageLoad on a read_write image, imports read_only images, the OpenCL default, as sampled textures, turns a sampler_t constant into a CrossGL fixed sampler with the same CLK_* state, and binds image and sampler kernel arguments as resources. It used to copy read_imagef unchanged, import every image as a storage image, keep a sampler's CLK_* bitmask as its value and put image and sampler arguments in the k_Args constant buffer. Unnormalized coordinates (CLK_NORMALIZED_COORDS_FALSE) translate to Metal and are refused by the other targets' fixed-sampler checks; mip-level and gradient reads, depth and buffer images and the other reads with no exact lowering are refused with OpenCLImageReadUnsupportedError.
• The OpenCL importer lowers write_imagef, write_imagei and write_imageui to imageStore, declares an image written with write_imagei or write_imageui as an iimage or uimage so its element type matches the stored value, and imports image1d_buffer_t as imageBuffer. It used to copy the calls unchanged, so every target called an undeclared function. Writes with a mip level, writes to depth images, an image written by more than one of the functions and integer writes through a helper parameter are refused with OpenCLImageWriteUnsupportedError.
• The OpenCL importer lowers mem_fence, read_mem_fence, write_mem_fence and atomic_work_item_fence to the CrossGL fence over the memory their flag names, memoryBarrierShared() for CLK_LOCAL_MEM_FENCE, memoryBarrierBuffer() for CLK_GLOBAL_MEM_FENCE and memoryBarrierImage() for CLK_IMAGE_MEM_FENCE, or memoryBarrier() for several flags or a flag that is not a literal, and sub_group_barrier to subgroupBarrier(). mem_fence used to become a comment, which dropped the fence. An atomic_work_item_fence with a scope wider than the device and a sub_group_barrier with a scope wider than the sub-group are refused with OpenCLFenceScopeUnsupportedError.
• The OpenCL importer drops prefetch(p, n), which only hints that memory will be read soon, and keeps the side effects of its arguments, such as an i++, as statements of their own. It used to copy the call, usually with a pointer cast, which no target declares.
• Metal output lowers the fine and coarse derivatives, dFdxFine, dFdyCoarse, fwidthFine and the others in their GLSL and HLSL spellings, to dfdx, dfdy and fwidth, because MSL does not let a shader choose either, and HLSL and WGSL output accept the fwidth_fine and fwidth_coarse spellings that the GLSL importer writes for fwidthFine and fwidthCoarse, refusing them in HLSL vertex and compute stages as they refuse fwidthFine. These calls used to reach the output unchanged, which the target compilers reject. WGSL output calls a function the shader defines under a derivative name, such as ddx_fine or dFdx, instead of the WGSL built-in.
• HIP output passes the address of the target to the system-scoped atomics, such as atomicAdd_system, atomicExch_system and atomicCAS_system, on structured buffer elements and buffer pointers, as CUDA output does. It used to pass the element by value, which does not compile because HIP declares them with a pointer first argument (#3092).
• CUDA output writes CrossGL norm and rnorm of a float array or pointer as normf and rnormf and of a double one as norm and rnorm, and refuses other or unresolved element types with CudaLoweringRefused. It used to write norm and rnorm, which take const double*, for float arrays too, which nvcc rejects (#2656).
• CUDA and HIP output lower doubleBitsToLong and longBitsToDouble to __double_as_longlong and __longlong_as_double and type their results as long long and double. They used to write the CrossGL names, which neither compiler declares. A vector or matrix operand is refused with CudaLoweringRefused or HipLoweringRefused. A typed refusal in an arm of a CUDA match now stops the translation; it used to replace the whole match, valid arms included, with an unsupported CUDA match comment.
• CUDA and HIP output lower countLeadingZeros to __clz, or __clzll for a 64-bit operand, taking the width from a sized cast such as countLeadingZeros(u64(x)). They used to write the CrossGL name, which neither compiler declares. A function the program defines with that name is still called.
• WGSL, Metal and GLSL output lower the HLSL 8-bit pack intrinsics pack_u8, pack_s8, pack_clamp_u8, pack_clamp_s8, unpack_u8u32 and unpack_s8s32: WGSL output to pack4xU8, pack4xI8, pack4xU8Clamp, pack4xI8Clamp, unpack4xU8 and unpack4xI8, Metal output, which also lowers unpack_u8u16 and unpack_s8s16, to byte vectors reinterpreted with as_type, as in as_type<uint>(uchar4(uint4(v))), and GLSL output to packUnorm4x8 and lane shifts. The results are typed, so storing one in a variable of another signedness or width gets a conversion. They used to be written unchanged, which none of the three compilers declares.
• Metal output lowers dot4add_i8packed and dot4add_u8packed to the dot product of the four byte lanes of each operand added to the accumulator, taking the signed sum in uint so it wraps as HLSL defines, where signed overflow in MSL is undefined. They used to be written unchanged, which the Metal compiler rejects.
• GLSL and WebGL output wrap 8- and 16-bit scalar integer +, -, *, / and % at their own width when the result is used without being stored, so float(v * v) on an int8_t v = 13 is -87.0. They wrapped only results stored in a narrow variable, so the expression was evaluated in 32 bits and gave 169.0. Unary - and ~ and shifts by an operand of the same narrow type wrap as well, and scalar results from Metal, CUDA, HIP and OpenCL C sources, which promote narrow operands to int, stay 32-bit: uchar x = 200; float(x * x) is 40000 (#3032, #3033).
• The Mojo importer no longer declares a comptime or alias that names a DType, such as comptime float_dtype = DType.float32, as a CrossGL global; it already resolved such aliases where they name a type. It used to also write let float_dtype = DType.float32;, an untyped global that no target can declare. Other comptime values are still declared.
• Formatting CUDA, HIP and the other C-family outputs with clang-format keeps the #include order the generator wrote. clang-format used to sort the includes, moving <hip/device_functions.h> and <hip/hip_fp16.h> ahead of <hip/hip_runtime.h>, which hipcc rejects with use of undeclared identifier '__ockl_sdot2' (#3039).
• The Metal importer reads make_filled_matrix, which fills every element of a matrix with one value, as a matrix of column splats, and refuses one with bfloat elements, more than four columns or rows, or a size that is not a literal with MetalImportRefused. It used to pass the call to every target as a call to an undeclared function.
• WGSL output binds a Slang fragment result or structure member with SV_Target or SV_TargetN, which the Slang importer writes as Out_Color or Out_ColorN, to @location(N). It used to have no binding, which naga rejects with Entry point arguments and return values must all have bindings (#2912).
• Slang output lowers bitCount, findMSB, findLSB and bitfieldReverse to countbits, firstbithigh, firstbitlow and reversebits, converting the uint result of the first three to the int of the operand's width so a scan that finds no bit still gives -1, and the Slang importer maps the four intrinsics back. Slang output used to write the CrossGL names, which slangc rejects with undefined identifier. An operand whose integer type is not known is refused with SlangLoweringRefused.
• Slang output maps the 8- and 64-bit integer spellings int8, uint8, int64, uint64, i64 and u64, the i8vecN, u8vecN, i64vecN and u64vecN vectors and generic vectors such as vec2<i64> to int8_t, uint8_t, int64_t, uint64_t and vector<int64_t, 2> and the like, and int32 and uint32 to int and uint. It used to write them unchanged, which slangc rejects.
• Slang output gives an entry parameter or result with @builtin(name) the Slang system value of that built-in, as in @builtin(global_invocation_id) to SV_DispatchThreadID and @builtin(sample_mask) to SV_Coverage, and declares a compute parameter bound to gl_NumWorkGroups or @builtin(num_workgroups) as a local read from WorkgroupCount(). It used to write : builtin and : gl_NumWorkGroups, which slangc rejects. An unknown @builtin name and a gl_NumWorkGroups parameter outside a compute stage are refused with SlangLoweringRefused (#2910).
• Slang output declares a CrossGL storage buffer buffer T* name as RWStructuredBuffer<T> name, or StructuredBuffer<T> name when it is readonly, with the binding the other targets give it, and reads *name as name[0]. It used to declare a global pointer, T* name, which slangc turns into a uniform block of buffer device addresses for SPIR-V and a Ptr parameter that DXC rejects for HLSL (#2921, #3003).
• The Slang importer reads SV_Position on a vertex output as gl_Position, as the HLSL importer does, while a fragment input stays gl_FragCoord and a vertex input an ordinary attribute. It used to map SV_Position to Out_Position, which no generator treats as the position built-in, so OpenGL output wrote a user varying and never set gl_Position. An SV_Position that is both a vertex input and a stage output, such as an inout parameter, is refused with SlangImportRefused (#2295).
• WGSL output writes the two-argument atan(y, x) as atan2(y, x). WGSL's atan takes one argument, so naga rejected the call with too many arguments passed to atan. A function the shader defines named atan is still called.
• WGSL output of textureSize returns the signed i32 or vecN<i32> that GLSL defines, or the type the result is stored in, and adds the layer count from textureNumLayers as the last component for array textures. It used to write the unsigned textureDimensions(t) result, which naga rejects where an int or ivec is expected, without the layer count.
• WGSL output writes mix(x, y, a) with a bool or bvecN selector as select(x, y, a), which takes y where a is true, as mix does. WGSL's mix takes only a floating-point selector, so naga rejected the call. A selector whose width differs from the known width of the values is refused.
• HLSL output reads gl_HelperInvocation as the Shader Model 6.6 intrinsic IsHelperLane(), and the dx10 and dx11 profiles refuse it with DirectXLoweringRefused. It used to write gl_HelperInvocation, which DXC does not declare. A local or global the program declares with that name is left alone.
• Slang output writes fwidthCoarse and fwidthFine as fwidth_coarse and fwidth_fine. It used to copy the CrossGL names, which slangc rejects as undefined identifiers, so Slang input that used them did not survive the round trip.
• CUDA output includes <cstdint>, so the fixed-width types it writes for CrossGL's 8-, 16- and 64-bit integers, such as uint8_t and int64_t, are declared. nvcc used to stop with identifier "uint64_t" is undefined.
• GLSL output lowers HLSL dot4add_i8packed(a, b, acc) and dot4add_u8packed(a, b, acc) to acc plus a dot() of the four 8-bit lanes of a and b, byte 0 first, sign-extending the lanes of the signed form. It used to copy the HLSL calls, which glslang rejects with no matching overloaded function found.
• The OpenCL importer reads get_local_linear_id() as the local invocation index, gl_LocalInvocationIndex, which HLSL output writes as SV_GroupIndex and Metal output as thread_index_in_threadgroup. It used to reach every target as a call to an undeclared function.
• The Slang importer maps the HLSL barrier intrinsics as the HLSL importer does: GroupMemoryBarrierWithGroupSync() to workgroupBarrier(), GroupMemoryBarrier() to memoryBarrierShared(), DeviceMemoryBarrier() and AllMemoryBarrier() to deviceMemoryBarrier() and allMemoryBarrier(), and DeviceMemoryBarrierWithGroupSync() and AllMemoryBarrierWithGroupSync() to a fence on each side of an execution barrier. They used to pass through unchanged, so Metal and GLSL output did not compile and SPIR-V output raised SPIRVBuiltinLoweringError.
• The Slang importer reads int8_t4_packed and uint8_t4_packed, also as structured buffer elements, as uint, as the HLSL importer does. It used to copy the type names, which CrossGL refused as undeclared types.
• Slang output writes the 16-bit integer spellings int16, uint16, i16vec2 to i16vec4 and u16vec2 to u16vec4 as int16_t, uint16_t, vector<int16_t, N> and vector<uint16_t, N>. It used to copy them, which slangc rejects as undefined identifiers.
• Componentwise matrix arithmetic, +, - and / of a matrix and a scalar on either side, / of two matrices of one shape and -m, translates to HLSL, OpenGL, Metal and SPIR-V, and HLSL output types matrixCompMult and outerProduct results as matrices. HLSL output typed m / 2.0 and m - 1.0 as vectors, so (m / 2.0)[1] read a row and (m - 1.0) * r used * instead of mul(); OpenGL output typed m - 1.0 as a float and refused (m - 1.0) * r; Metal output wrote operators MSL does not have and now calls one helper per shape, such as __crossgl_sub_float2x2_float(m, 1.0); SPIR-V output refused these expressions and wrote an invalid OpFNegate on a matrix. A Metal compound update such as m += 1.0 whose target has side effects is refused with UnsupportedMetalFeatureError.
• WebGL output accepts a float local whose only value is an integer sum compared for equality, such as float d = float(n - 8); tested with d == 0.0. GLSL output declared such a local precise, which GLSL ES 3.00 does not have, so WebGL refused the shader; precise is now added only when one of the sums is floating or of unknown type.
• SPIR-V output declares a GLSL interface block instance, such as out Block { vec2 uv; } vout[]; in a geometry, tessellation or mesh stage, as one stage interface variable with the block's direction, patch, perprimitiveEXT and interpolation qualifiers, location and array extents. It used to declare no variable, so every shader that used one was refused with does not name storage or Unknown variable. A block whose members carry their own location or component, or with transform feedback or stream layouts, is refused with InterfaceBlockInstanceUnsupportedError.
• Stage layout counts written as constant expressions, such as layout(max_vertices = 4 * groupSizeX * groupSizeY) out; over const uint declarations, and the max_primitives, vertices and invocations counts and their HLSL attribute forms such as [maxvertexcount(N)], are folded to literals on every target. The generators accepted only a literal or a single constant name and refused such shaders.
• GLSL output declares vertex outputs and fragment inputs of structure type after the structures they name; they used to come first, so such a shader never compiled. In a program with a mesh stage, user varyings without a location take the locations SPIR-V gives them, vertex members first and then primitive members, because GL_EXT_mesh_shader compiles only to SPIR-V and glslang rejected the unlocated ones. A structure varying with a built-in field, such as a Metal nested [[stage_in]] with [[position]], is refused with GLSLNestedStageBuiltinError instead of becoming a user varying.
• GLSL and SPIR-V output declare a task payload in task-payload storage, taskPayloadSharedEXT and TaskPayloadWorkgroupEXT: the groupshared variable, local or module-scope, that a task entry passes to DispatchMesh, as HLSL amplification shaders write it, and a mesh entry's @mesh_payload parameter. A local task payload used to be declared as shared workgroup storage, which EmitMeshTasksEXT does not pass to the mesh stage, and a module-scope one was refused; GLSL output dropped the mesh parameter, so the entry read an undeclared name, and SPIR-V output flattened it into vertex inputs that no stage writes.
• GLSL and SPIR-V output take a mesh stage's max_vertices and max_primitives from the sizes of its @vertices, @primitives and @indices arrays when no attribute gives them; SPIR-V output used to declare one primitive. HLSL output gives the members of mesh vertex and primitive structures that have no semantic the TEXCOORD<n> semantics vertex outputs default to, numbering primitive members after vertex members so they match the pixel input, instead of refusing every such member. The threadgroup size a host chooses for a Metal mesh or object function, given with source_options={'threadgroup_size': (x, y, z)} or a project workgroup-size rule, is applied as it is for kernels; it used to be dropped.
• WGSL output receives global_invocation_id, local_invocation_id, workgroup_id and num_workgroups as vec3<u32>, and vertex_index, instance_index, local_invocation_index and sample_index as u32, then copies them into a local of the type the source declares, such as an MSL uint2 thread position or the int of gl_VertexID. It used to declare the built-in with the source's type, which naga rejects.
• WGSL output lowers texelFetchOffset(t, p, lod, offset) to textureLoad(t, (p + offset), lod). It used to copy the GLSL call, which WGSL does not declare.
• The HLSL, GLSL and Metal importers keep a default label where the source puts it among the case labels and keep an empty default:, and the Metal importer keeps the source's break statements instead of ending every clause with one, so a clause without a jump falls through to the next one in source order. A default in the middle of a switch used to move last, an empty default: was dropped and Metal fallthrough was lost, so every target ran a different clause for some selectors.
• Slang output keeps the parentheses around a member, index or postfix operand and around the operand of a prefix operator, so (p + q).yx stays (p + q).yx and -(-a) stays -(-a), and WGSL output does the same for prefix operators. Slang output used to write p + q.yx, which swizzles only q, --a, which decrements a, and p * 2.0.y, which slangc rejects; WGSL output wrote --a, which naga rejects.
• WGSL output writes an increment, decrement or assignment used as a value, such as a[i++], int a = i++; or x = y = 7, as statements before the expression in left-to-right order, keeping the old value in a let for the postfix forms. A while or for condition with one is checked at the top of the loop body, and an else if condition moves into its else block. It used to write a[i += 1] and var a: i32 = i += 1, which naga rejects and which read the updated value. A side effect where it may not run, such as the right operand of && or ||, a ?: arm or a for update, and an l-value that would call a function twice are refused.
• HLSL output shifts 16-bit integer vectors such as u16vec2 and i16vec2 at their own width and negates and complements a uint16_t or int16_t scalar as an int, as in -int(a), following the MSL rules CrossGL uses for narrow integers. It used to shift on int2 operands, so (n << 4).x could reach 1048560, and to negate a uint16_t natively, so -a of 1 gave 65535 instead of -1.
• HLSL, GLSL, Metal and SPIR-V output translate a call to a CrossGL generic function with explicit type arguments, such as zero_of<float>(), to the specialization it selects: the type arguments bind the leading generic parameters, the others are deduced from the value arguments, and a default applies only to a parameter that cannot be deduced, so scale_by<int>(float16(2), 3) for generic<T, U = float> T scale_by(U factor, T value) calls scale_by_int_float16. The call used to be written unchanged, as zero_of<float>(), which names no function in the target. A call that passes more type arguments than the function has parameters, or leaves a parameter unbound, is refused.
• Metal sources that pass a structure template with a short capitalised name, such as MD, RMS or Vec, as a template argument, as in Holder<T, MD<T>>, translate to HLSL, GLSL and SPIR-V. They used to be refused with Holder missing MD, because the check for unresolved template parameters took such names for parameters.
• HLSL output declares a shader-scope let or var with no written type as a static global with the type of its initializer and keeps the initializer, as in static float scale = 2.0;, and refuses one whose type it cannot infer with DirectXLoweringRefused. It used to declare it with the type None and no value in the implicit $Globals constant buffer, which DXC rejects (#2444).
• Slang output writes CrossGL matCxR as Slang floatRxC, as HLSL output does: mat4x3 becomes float3x4, a constructor fills columns, as in transpose(float2x2(a, b, c, d)), a single scalar fills the diagonal, as in float2x2(2.0, 0.0, 0.0, 2.0) for mat2(2.0), m[i] and m[i][j] name column i, as in m._m00_m10 for m[0], and a matrix times a vector read from a buffer is written mul(m, v) (#2767). It used to write mat4x3 as float4x3, fill constructors by rows, splat a scalar into every element and index rows, so non-square matrices had the transposed shape and products and subscripts computed other values.
• WGSL output passes workgroupUniformLoad a pointer to its workgroup variable, as in workgroupUniformLoad(&a0). It used to pass the variable itself, which naga rejects because the built-in takes a workgroup pointer (#2634).
• The HLSL importer converts a floating-point texel location or mip level passed to Load or the texture subscript operator to an integer, as HLSL does implicitly, so velocity[pos.xy] imports as texelFetch(velocity, ivec2(pos.xy), 0). It used to pass the float coordinates to texelFetch, which GLSL refuses. Metal output gives a texel fetch the texel type of its texture, as it already did for samples, so a fetched texel assigned to a half4 is converted explicitly, as MSL requires between vector types.
• Metal output applies the ray flags of TraceRay to the intersector, as in set_triangle_cull_mode(triangle_cull_mode::back) for RAY_FLAG_CULL_BACK_FACING_TRIANGLES and accept_any_intersection(true) for RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH, and selects each mode at run time when the flags are not a constant. It used to drop the flags, so a trace ran without the culling, forced opacity or first-hit termination it asked for. Constant flags that DXR and SPIR-V leave undefined, such as both face culling bits, are refused with MetalLoweringRefused.
• The Metal importer keeps long, ulong and their vector types 64-bit as as_type targets, so as_type<long>(words) on a uint2 becomes packInt2x32(ivec2(words)) on GLSL, as as_type<int64_t> already did (#2183). GLSL output used to read the target as the 32-bit int and refuse the bitcast with as_type<int> cannot bitcast from uvec2: total widths are 32 and 64 bits.
• HLSL and Metal output pass the fragment built-ins that GLSL reads as globals, gl_FrontFacing, gl_SampleID, gl_PrimitiveID, gl_FragCoord and, on Metal, gl_PointCoord, to the entry point as parameters, as in bool front [[front_facing]] or bool front : SV_IsFrontFace, and on to every helper that reads them. Only gl_FragCoord used to be bound, on Metal only in the entry point, so every other read was an undeclared identifier. HLSL output refuses gl_PointCoord with BuiltinValueUnsupportedError.
• Binding, set, buffer, texture, sampler, location, colour, index and specialization-constant id attributes that name an integer constant, an enumerator or an integer expression of them, such as Metal [[buffer(BufferIndexUniforms)]] with an enumerator from a header shared with the host or GLSL layout(location = BASE + 1), take that value on every target, and an index that is not such a constant is refused with BindingIndexUnresolvedError. HLSL output used to assign the next free register, Metal output dropped the attribute and GLSL output refused it or printed the syntax tree, so the binding moved without a diagnostic. The Metal importer keeps a [[color(RenderTargetAlbedo)]] output at its attachment instead of moving it to render target 0, and folds an enumerator [[function_constant(...)]] index, which Metal and SPIR-V output used to refuse.
• HLSL and SPIR-V output accept a fragment colour output of any numeric scalar or 1- to 4-component vector type, including float16_t vectors, so HLSL float main(PSIn input) : SV_TARGET translates back to HLSL and to SPIR-V, and HLSL output gives a fragment output member with @location(n) the semantic SV_TARGETn. Such outputs used to be refused, and a @location(n) member was written with the semantic : location, which DXC rejects.
• An array extent that names an enumerator, such as float w[Mode::Count], is folded to the enumerator's value, and a conversion to an enumeration, such as Mode(x), becomes a conversion to its underlying type, on every target. Both used to reach the targets as written. The Metal importer keeps a scoped Channel::Green instead of mangling it to Channel_u3a_u3aGreen.
• HLSL output initializes a structure constant that has an array member from an initializer list, as in static const P S1 = {2.0, {6.0, 7.0}};, and Metal output copies a named array constant into an array member element by element, as in {WS[0], WS[1]}. HLSL output used to pass the array literal {6.0, 7.0} to the generated constructor helper and Metal output wrote P{1.0, WS}, which neither compiler accepts. HLSL output refuses an array literal member outside a declaration initializer.
• The Metal importer accepts a C-style cast of a parenthesised expression after a statement keyword, such as return (float)(1.0 / x);, and a parenthesised expression statement after a condition, such as if (c) (y) = x;, which it used to refuse with MetalStructMethodError as unresolved temporary function objects. A typedef of xhalf, xhalf2 or another xhalf name decides its type: typedef half2 xhalf2; used to fail to parse, and after typedef float xhalf; an xhalf value was imported as half. HLSL and GLSL output of Metal sources drop #include <simd/simd.h> and the other <simd/...> and <metal_...> system headers, which neither compiler can resolve.
• Metal output accepts a reference to a constant buffer value, such as constant float& bias [[buffer(0)]], as a texture call operand, as in t.sample(s, uv, level(bias)), which it used to refuse because float& is not a scalar. A write to some components of a vector in a byte-addressed buffer block, such as p.vel.y *= 0.5 or p.color[i] = v, stores each written component on its own; it used to use the load as the target, as in float4(...).x = v, which does not compile.
• The HLSL importer reads SubpassInput<T> and SubpassInputMS<T> as CrossGL subpassInput, isubpassInput or usubpassInput of their element type, with an MS suffix for the multisampled form, so SPIR-V output declares them as DXC does and other targets read them as textures. SubpassInput used to import as sampler2d, which lost the element type and left subpassLoad unresolved, and SubpassInputMS was not recognised.
• SPIR-V output writes a constant splat texel offset such as ivec3(1) as ConstOffset. It used to write it as a non-constant Offset, which Vulkan accepts only on gathers.