ocl bug report

kernel.cl

typedef struct {
    float mass;
    float rx;
    float ry;
    float vx;
    float vy;
} Particle_t;

__kernel void time_to_collision(
    __global Particle_t parts[PART_NUM],
    __global float result[PART_NUM][PART_NUM])
{
    result[0][get_global_id(0)] = parts[get_global_id(0)].vx;
}

main.rs

use arraylib::Array;

#[repr(C)]
#[derive(Clone, Copy, Debug, Default, PartialEq)]
struct Particle {
    mass: f32,
    rx: f32,
    ry: f32,
    vx: f32,
    vy: f32
}

unsafe impl ocl::traits::OclPrm for Particle {}

fn main() -> ocl::Result<()> {
    const PART_NUM: usize = 3;

    let pro_que = ocl::ProQue::builder()
        .device(
            ocl::Device::list(ocl::Platform::default(), Some(ocl::DeviceType::new().gpu()))
                .unwrap()
                .get(0)
                .expect("Get at least 1 GPU"),
        )
        .prog_bldr({
            let mut prog = ocl::Program::builder();
            prog.src_file("src/kernel.cl")
                .cmplr_def("PART_NUM", PART_NUM as i32);
            prog
        })
        .dims(PART_NUM)
        .build()?;

    let parts = <[_; PART_NUM]>::from_fn(|i| Particle {
        mass: (i*i) as f32,
        rx: i as f32,
        ry: 1000.0 - i as f32,
        vx: 1.0,
        vy: -1.0
    });

    let buffer = pro_que.buffer_builder()
        .copy_host_slice(&parts[..])
        .build()?;

    let result = pro_que.buffer_builder::<f32>()
        .len(PART_NUM)
        .build()?;

    let kernel = pro_que
        .kernel_builder("time_to_collision")
        .arg(&buffer)
        .arg(&result)
        .build()?;

    let mut event = ocl::Event::empty();
    unsafe { kernel.cmd().enew(&mut event).enq()?; }
    let mut answer = [0.0f32; PART_NUM];
    result.read(&mut answer[..]).ewait(&event).enq()?;
    assert_eq!(
        answer,
        <[_; PART_NUM]>::from_iter(
            parts.iter().map(|p| p.vx)).unwrap()
    );

    Ok(())
}