Intel's OpenCL driver (NEO) stores a wrong `Bool` field in a struct followed by a tuple

[maleadt]

A kernel that stores a struct with a Bool field followed by a tuple field loses the Bool on Intel's OpenCL driver: it reads back as false, while the tuple is right. With the fields in the other order, or on PoCL, the result is correct.

using OpenCL
OpenCL.cl.platform!(only(filter(p -> occursin("Intel(R) OpenCL Graphics", p.name), OpenCL.cl.platforms())))

struct FlagFirst; valid::Bool; value::Tuple{Float32, Int32}; end
struct ValueFirst; value::Tuple{Float32, Int32}; valid::Bool; end

function kernel(out, x, ::Type{S}) where {S}
    i = get_global_id()
    v = (x[i], Int32(i))
    @inbounds out[i] = S === FlagFirst ? FlagFirst(x[i] >= 0, v) : ValueFirst(v, x[i] >= 0)
    return
end

x = Float32[0.5, 1.5, 2.5, 3.5]
for S in (FlagFirst, ValueFirst)
    out = CLArray{S}(undef, 4)
    @opencl global_size=4 kernel(out, CLArray(x), S)
    println(S, ": ", Array(out))
end

On an Intel Iris Xe (device 0x9a49, "Intel(R) OpenCL Graphics"):

FlagFirst: FlagFirst[FlagFirst(false, (0.5f0, 1)), FlagFirst(false, (1.5f0, 2)), FlagFirst(false, (2.5f0, 3)), FlagFirst(false, (3.5f0, 4))]
ValueFirst: ValueFirst[ValueFirst((0.5f0, 1), true), ValueFirst((1.5f0, 2), true), ValueFirst((2.5f0, 3), true), ValueFirst((3.5f0, 4), true)]

On PoCL (pocl_jll 7.2, x86_64), both are true. OpenCL.jl 0.10.11, Julia 1.13.0-rc4. With a constant flag (FlagFirst(true, v)) the kernel does not compile at all: llc fails with "unable to translate instruction: call llvm.spv.const.composite.i32".

Since PoCL runs the same SPIR-V correctly, the miscompilation is probably in Intel's graphics compiler; the SPIR-V may need forwarding to intel/compute-runtime or intel/intel-graphics-compiler.

AcceleratedKernels.jl works around it by putting the value first in its (value, valid) partial results (WORKAROUND(Intel NEO) in src/reduce/utilities.jl on JuliaGPU/AcceleratedKernels.jl#133); with the flag first, its findmin-style tuple reductions and scans gave wrong results on NEO.