neuron-broken.rs

use ocl::ProQue;

pub fn main() -> ocl::Result<()> {

    let kernel_src = r#"
        float uniform (__global unsigned int* seed)
        {
            *seed = ((*seed) * 16807 ) % 2147483647;
            return  (float)(*seed) * 4.6566129e-10; 
        }

        __kernel void izhikevich(
            __global float* u,
            __global float* v,
            __global float* a,
            __global float* b,
            __global float* c,
            __global float* d,
            __global char* spiking,
            __global unsigned int* seeds
        ) {
            uint const i = get_global_id(0);
            float I = uniform(&seeds[i]) * 5.0;

            v[i] += 0.5 * ((0.04 * v[i] * v[i]) + (5.0 * v[i]) + 140.0 - u[i] + I);
            v[i] += 0.5 * ((0.04 * v[i] * v[i]) + (5.0 * v[i]) + 140.0 - u[i] + I);
            u[i] += a[i] * ((b[i] * v[i]) - u[i]);

            if (v[i] >= 30.0) {
                v[i] = c[i];
                u[i] += d[i];
                spiking[i] = 1;
            } else {
                spiking[i] = 0;
            }
        }
    "#;

    const NUM_NEURONS: usize = 100000;

    let pro_que = ProQue::builder()
        .src(kernel_src)
        .dims(NUM_NEURONS)
        .build()?;

    let v = pro_que.buffer_builder::<f32>().copy_host_slice(&[-65.; NUM_NEURONS]).build()?;
    let u = pro_que.buffer_builder::<f32>().copy_host_slice(&[0.02 * -65.; NUM_NEURONS]).build()?;
    let a = pro_que.buffer_builder::<f32>().copy_host_slice(&[-0.02; NUM_NEURONS]).build()?;
    let b = pro_que.buffer_builder::<f32>().copy_host_slice(&[-0.2; NUM_NEURONS]).build()?;
    let c = pro_que.buffer_builder::<f32>().copy_host_slice(&[-65.; NUM_NEURONS]).build()?;
    let d = pro_que.buffer_builder::<f32>().copy_host_slice(&[-2.; NUM_NEURONS]).build()?;
    let spiking = pro_que.create_buffer::<i8>()?;

    let seeds = pro_que.create_buffer::<u32>()?;
    let seeds_cpu = ocl_extras::scrambled_vec::<u32>((0, std::u32::MAX), NUM_NEURONS);
    seeds.write(&seeds_cpu).enq()?;

    let kernel = pro_que.kernel_builder("izhikevich")
        .arg(&u)
        .arg(&v)
        .arg(&a)
        .arg(&b)
        .arg(&c)
        .arg(&d)
        .arg(&spiking)
        .arg(&seeds)
        .build()?;

    let mut t = 0;

    while t < 10000 {
        t += 1;
        
        unsafe { kernel.enq()?; }
    }

    let mut spiking_cpu = vec![0i8; v.len()];
    spiking.read(&mut spiking_cpu).enq()?;

    let mut v_cpu = vec![0.0f32; v.len()];
    v.read(&mut v_cpu).enq()?;

    println!("{}: {:?}", t, spiking_cpu);

    Ok(())
}

neuron-working.rs

use ocl::ProQue;

pub fn main() -> ocl::Result<()> {

    let kernel_src = r#"
        float uniform (__global unsigned int* seed)
        {
            *seed = ((*seed) * 16807 ) % 2147483647;
            return  (float)(*seed) * 4.6566129e-10; 
        }

        __kernel void izhikevich(
            __global float* u,
            __global float* v,
            __global float* a,
            __global float* b,
            __global float* c,
            __global float* d,
            __global char* spiking,
            __global unsigned int* seeds
        ) {
            uint const i = get_global_id(0);
            float I = uniform(&seeds[i]) * 5.0;

            v[i] += 0.5 * ((0.04 * v[i] * v[i]) + (5.0 * v[i]) + 140.0 - u[i] + I);
            v[i] += 0.5 * ((0.04 * v[i] * v[i]) + (5.0 * v[i]) + 140.0 - u[i] + I);
            u[i] += a[i] * ((b[i] * v[i]) - u[i]);

            if (v[i] >= 30.0) {
                v[i] = c[i];
                u[i] += d[i];
                spiking[i] = 1;
            } else {
                spiking[i] = 0;
            }
        }
    "#;

    const num_neurons: usize = 100000;

    let pro_que = ProQue::builder()
        .src(kernel_src)
        .dims(num_neurons)
        .build()?;

    let v = pro_que.buffer_builder::<f32>().copy_host_slice(&[-65.; num_neurons]).build()?;
    let u = pro_que.buffer_builder::<f32>().copy_host_slice(&[0.02 * -65.; num_neurons]).build()?;
    let a = pro_que.buffer_builder::<f32>().copy_host_slice(&[-0.02; num_neurons]).build()?;
    let b = pro_que.buffer_builder::<f32>().copy_host_slice(&[-0.2; num_neurons]).build()?;
    let c = pro_que.buffer_builder::<f32>().copy_host_slice(&[-65.; num_neurons]).build()?;
    let d = pro_que.buffer_builder::<f32>().copy_host_slice(&[-2.; num_neurons]).build()?;
    let spiking = pro_que.create_buffer::<i8>()?;

    let seeds = pro_que.create_buffer::<u32>()?;
    let seeds_cpu = ocl_extras::scrambled_vec::<u32>((0, std::u32::MAX), num_neurons);
    seeds.write(&seeds_cpu).enq()?;

    let kernel = pro_que.kernel_builder("izhikevich")
        .arg(&u)
        .arg(&v)
        .arg(&a)
        .arg(&b)
        .arg(&c)
        .arg(&d)
        .arg(&spiking)
        .arg(&seeds)
        .build()?;

    let mut t = 0;

    while t < 10000 {
        t += 1;
        
        unsafe { kernel.enq()?; }
    }

    let mut spiking_cpu = vec![0i8; v.len()];
    spiking.read(&mut spiking_cpu).enq()?;

    let mut v_cpu = vec![0.0f32; v.len()];
    v.read(&mut v_cpu).enq()?;

    println!("{}: {:?}", t, spiking_cpu);

    Ok(())
}