Basic scalar rand() support

[dnadlinger]

For use with https://github.com/SciML/PSOGPU.jl, it would be great to have basic support of on-device rand(). A simple MWE, adapted from https://github.com/JuliaGPU/Metal.jl/issues/406, would be

function kernel(state)
    idx = thread_position_in_grid_1d()
    state[idx] = rand(Float32)
    return
end
n = 10
state = Metal.ones(n)
@metal threads=n kernel(state)

which currently fails with

julia> @metal threads=n kernel(state)
ERROR: InvalidIRError: compiling MethodInstance for kernel(::MtlDeviceVector{Float32, 1}) resulted in invalid LLVM IR
Reason: unsupported call to an unknown function (call to julia.get_pgcstack)

Would a solution such as in https://github.com/JuliaGPU/CUDA.jl/blob/master/src/device/random.jl be appropriate? What would be the most appropriate way to store the state in Metal?

[maleadt]

Would a solution such as in https://github.com/JuliaGPU/CUDA.jl/blob/master/src/device/random.jl be appropriate?

Yes, although ideally it should be implemented in GPUArrays.jl now that it's based on KernelAbstractions.jl. This may require some improvements to KA.jl though.