//----------------------------------------
// 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)));