Bug 264373

//---------------------------------------- // Function: fused_fused_decode1_take_kernel //---------------------------------------- @group(0) @binding(0) var<storage, read_write> T_take : array<f32>; @group(0) @binding(1) var<storage, read> lv : array<u32>; @group(0) @binding(2) var<storage, read> lv_1 : array<i32>; @group(0) @binding(3) var<storage, read> lv1 : array<u32>; struct PODArgs { n: i32, vocab_size: i32, packGridDimX: u32 } @group(0) @binding(4) var<uniform> podArgs : PODArgs; @compute @workgroup_size(256, 1, 1) fn fused_fused_decode1_take_kernel( @builtin(workgroup_id) blockIdx : vec3<u32>, @builtin(num_workgroups) gridDim : vec3<u32>, @builtin(local_invocation_id) threadIdx : vec3<u32> ) { if (blockIdx.z * gridDim.x + blockIdx.x > podArgs.packGridDimX) { return; } let v__1 : i32 = i32(blockIdx.z * gridDim.x + blockIdx.x); T_take[((v__1 * 256i) + i32(threadIdx.x))] = fma(f32(((lv[(((lv_1[(v__1>>4u)] * 512i) + ((v__1 & 15i) * 32i)) + (i32(threadIdx.x)>>3u))]>>u32(((i32(threadIdx.x) & 7i) * 4i))) & 15u)), bitcast<f32>(((lv1[(((lv_1[(v__1>>4u)] * 128i) + ((v__1 & 15i) * 8i)) + (i32(threadIdx.x)>>5u))] & 65535u)<<16u)), bitcast<f32>((((lv1[(((lv_1[(v__1>>4u)] * 128i) + ((v__1 & 15i) * 8i)) + (i32(threadIdx.x)>>5u))]>>16u) & 65535u)<<16u))); } generates ('global2' is not declared): ... ABs are omitted ... [[kernel]] void function0(vec<unsigned, 3> parameter0 [[threadgroup_position_in_grid]], vec<unsigned, 3> parameter1 [[threadgroups_per_grid]], vec<unsigned, 3> parameter2 [[thread_position_in_threadgroup]], constant __ArgumentBufferT_0& __ArgumentBuffer_0 [[buffer(0)]]) { const constant type0& global4 = __ArgumentBuffer_0.global4; const device array<unsigned, 1>& global3 = __ArgumentBuffer_0.global3; device array<float, 1>& global0 = __ArgumentBuffer_0.global0; const device array<unsigned, 1>& global1 = __ArgumentBuffer_0.global1; if ((((parameter0.z * parameter1.x) + parameter0.x) > global4.field2)) { return; } int local0 = int(((parameter0.z * parameter1.x) + parameter0.x)); global0[((local0 * 256) + int(parameter2.x))] = fma(float(((global1[(((global2[(local0 >> 4u)] * 512) + ((local0 & 15) * 32)) + (int(parameter2.x) >> 3u))] >> unsigned(((int(parameter2.x) & 7) * 4))) & 15u)), as_type<float>(((global3[(((global2[(local0 >> 4u)] * 128) + ((local0 & 15) * 8)) + (int(parameter2.x) >> 5u))] & 65535u) << 16u)), as_type<float>((((global3[(((global2[(local0 >> 4u)] * 128) + ((local0 & 15) * 8)) + (int(parameter2.x) >> 5u))] >> 16u) & 65535u) << 16u))); } MSL compilation error: Error Domain=MTLLibraryErrorDomain Code=3 "program_source:47:76: error: use of undeclared identifier 'global2' global0[((local0 * 256) + int(parameter2.x))] = fma(float(((global1[(((global2[(local0 >> 4u)] * 512) + ((local0 & 15) * 32)) + (int(parameter2.x) >> 3u))] >> unsigned(((int(parameter2.x) & 7) * 4))) & 15u)), as_type<float>(((global3[(((global2[(local0 >> 4u)] * 128) + ((local0 & 15) * 8)) + (int(parameter2.x) >> 5u))] & 65535u) << 16u)), as_type<float>((((global3[(((global2[(local0 >> 4u)] * 128) + ((local0 & 15) * 8)) + (int(parameter2.x) >> 5u))] >> 16u) & 65535u) << 16u)));