Untranslated sub-group instructions: OpGroupAll/Any/Broadcast, OpGroupIAdd and other group arithmetic, GroupNonUniformBallot #25

spirv2clc can't translate the sub-group collectives of cl_khr_subgroups, or the ballot of cl_khr_subgroup_ballot. These are what OpenCL.jl's SPIRVIntrinsics emit for KernelInterface's sub-group votes, reductions and scans. Translating a kernel that uses one of them fails:

Presumably the rest of the OpGroup* arithmetic (OpGroupFAdd, OpGroup{S,U,F}{Min,Max}) is missing too. OpGroupNonUniformShuffle/ShuffleXor, OpControlBarrier with sub-group scope, and the sub-group built-in variables already translate.

The OpenCL C equivalents are sub_group_all/any/broadcast, sub_group_reduce_<op>, sub_group_scan_inclusive_<op>/sub_group_scan_exclusive_<op> (cl_khr_subgroups) and sub_group_ballot (cl_khr_subgroup_ballot). The other OpGroupNonUniform* instructions map onto the cl_khr_subgroup_* extensions the same way.

Reproducer, with OpenCL.jl and pocl_jll on the OpenCL C source path:

using OpenCL, pocl_jll
cl.platform!("pocl")
OpenCL.program_backend!(:opencl)
function k(out)
    i = get_global_id(1)
    out[i] = sub_group_reduce_add(out[i])   # or sub_group_any(out[i] > 0), sub_group_ballot(...), ...
    return
end
@opencl global_size=32 local_size=32 k(CLArray(ones(Int32, 32)))

Until this is implemented, OpenCL.jl doesn't report KernelInterface sub-group support for kernels on the OpenCL C path (JuliaGPU/OpenCL.jl#554). The poclc test target and PoCL builds without SPIR-V support (e.g. 'PoCL next' on arm64) use that path.