Shader validator error with linear broadcast kernel about metal.jl HOT 3 OPEN

declare float @air.sincos.f32(i64)

define void @my_kernel(float addrspace(1)* %a, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %e, i32 %arg2) {
  %tmp4 = bitcast { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %e to float addrspace(1)* addrspace(1)*
  %tmp5 = load float addrspace(1)*, float addrspace(1)* addrspace(1)* %tmp4, align 8
  %b = alloca i32, align 4
  %f = sext i32 %arg2 to i64
  %c = ptrtoint i32* %b to i64
  %1 = call float @air.sincos.f32(i64 %c)
  %d = getelementptr float, float addrspace(1)* %a, i64 %f
  store float 0.000000e+00, float addrspace(1)* %d, align 4
  store float 0.000000e+00, float addrspace(1)* %tmp5, align 4
  ret void
}

!air.kernel = !{!0}
!air.version = !{!6}

!0 = !{void (float addrspace(1)*, { i8 addrspace(1)*, [1 x i64] } addrspace(1)*, i32)* @my_kernel, !1, !2}
!1 = !{}
!2 = !{!3, !4, !5}
!3 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.arg_type_name", !"arr_sin"}
!4 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, !"air.arg_type_name", !"air.arg_name", !"arr_cos"}
!5 = !{i32 2, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint"}
!6 = !{i32 2, i32 6, i32 0}

declare float @air.sincos.f32(float, i64) local_unnamed_addr

define void @my_kernel({ i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0, { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1, i32 %thread_position_in_grid) local_unnamed_addr {
conversion:
  %2 = bitcast { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %0 to float addrspace(1)* addrspace(1)*
  %.unpack8 = load float addrspace(1)*, float addrspace(1)* addrspace(1)* %2, align 8
  %3 = bitcast { i8 addrspace(1)*, [1 x i64] } addrspace(1)* %1 to float addrspace(1)* addrspace(1)*
  %.unpack12 = load float addrspace(1)*, float addrspace(1)* addrspace(1)* %3, align 8
  %4 = alloca i32, align 8
  %5 = sext i32 %thread_position_in_grid to i64
  %6 = getelementptr inbounds float, float addrspace(1)* %.unpack12, i64 %5
  %7 = load float, float addrspace(1)* %6, align 4
  %bitcast_coercion5 = ptrtoint i32* %4 to i64
  %8 = call float @air.sincos.f32(float %7, i64 %bitcast_coercion5)
  %9 = bitcast i32* %4 to float*

  %10 = load float, float* %9, align 8
  %11 = getelementptr inbounds float, float addrspace(1)* %.unpack8, i64 %5
  store float %8, float addrspace(1)* %11, align 4
  store float %10, float addrspace(1)* %6, align 4
  ret void
}

attributes #0 = { argmemonly nocallback nofree nosync nounwind willreturn }

!air.kernel = !{!41}
!air.version = !{!48}
!air.language_version = !{!49}

!9 = distinct !DICompileUnit(language: DW_LANG_Julia, file: !10, producer: "julia", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!10 = !DIFile(filename: "julia", directory: ".")
!41 = !{void ({ i8 addrspace(1)*, [1 x i64] } addrspace(1)*, { i8 addrspace(1)*, [1 x i64] } addrspace(1)*, i32)* @my_kernel, !42, !43}
!42 = !{}
!43 = !{!44, !45, !46}
!44 = !{i32 0, !"air.buffer", !"air.location_index", i32 0, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float32, 1}", !"air.arg_name", !"arr_sin"}
!45 = !{i32 1, !"air.buffer", !"air.location_index", i32 1, i32 1, !"air.read_write", !"air.address_space", i32 1, !"air.arg_type_size", i32 16, !"air.arg_type_align_size", i32 8, !"air.arg_type_name", !"MtlDeviceVector{Float32, 1}", !"air.arg_name", !"arr_cos"}
!46 = !{i32 2, !"air.thread_position_in_grid", !"air.arg_type_name", !"uint"}
!48 = !{i32 2, i32 5, i32 0}
!49 = !{!"Metal", i32 3, i32 1, i32 0}