Random Number Generator

This code requires OpenCL 3.0 or higher.

Rust code

use std::{time::{SystemTime}, mem::MaybeUninit};
use once_cell::sync::Lazy;
use blaze_rs::prelude::*;

#[global_context]
static CONTEXT : SimpleContext = SimpleContext::default();

static CODE : Lazy<String> = Lazy::new(|| {
    let nanos = SystemTime::now().duration_since(SystemTime::UNIX_EPOCH).unwrap();
    format!("#define TIME {}l\n{}", nanos.as_nanos(), include_str!("rng.cl"))
});

#[blaze(Rng)]
#[link = Lazy::force(&CODE)]
pub extern "C" {
    fn next_ints (n: u32, out: *mut MaybeUninit<u32>);
}

#[test]
fn main () -> Result<()> {
    let rng = Rng::new(None)?;
    let mut random = Buffer::<u32>::new_uninit(5, MemAccess::WRITE_ONLY, false)?;
    
    let random = unsafe {
        let _ = rng.next_ints(5, &mut random, [5], None, EMPTY)?.wait()?;
        random.assume_init()  
    };

    println!("{random:?}");
    Ok(())
}

OpenCL C code

#define MUTPILIER 0x5DEECE66Dl
#define ADDEND 0xBl
#define MASK() ((1l << 48) - 1)

global atomic_ulong SEED = ATOMIC_VAR_INIT((8682522807148012l * 1181783497276652981L) ^ TIME);

kernel void next_ints (const uint n, global uint* out) {
    const uint ID = get_global_id(0);
    const uint SIZE = get_global_size(1);
    ulong oldseed, nextseed;

    for (uint i = ID; i < n; i += SIZE) {
        do {
            oldseed = atomic_load(&SEED);
            nextseed = (oldseed * MUTPILIER + ADDEND) & MASK();
        } while (!atomic_compare_exchange_strong(&SEED, &oldseed, nextseed));
        out[i] = nextseed >> 16;
    }
}