gistfile1.txt

require 'opencl_ruby_ffi'
require 'benchmark'

source = <<EOF
__kernel void addition( float2 alpha, __global const float *x, __global float *y) {
  size_t ig = get_global_id(0);
  y[ig] = (alpha.s0 + alpha.s1 + x[ig])*0.3333333333333333333f;
}
EOF

n_times = 2**20

Benchmark.bm do |x|
  OpenCL::platforms.each { |platform|
    platform.devices.each { |device|
      x.report(device.name+":") do
        context = OpenCL::create_context(device)
        queue = context.create_command_queue(device, :properties => OpenCL::CommandQueue::PROFILING_ENABLE)
        prog = context.create_program_with_source( source )
        prog.build
        a_in = NArray.sfloat(n_times).random(1.0)
        a_out = NArray.sfloat(n_times)
        float = OpenCL::Float2::new(3.0,2.0)
        b_in = context.create_buffer(a_in.size * a_in.element_size, :flags => OpenCL::Mem::COPY_HOST_PTR, :host_ptr => a_in)
        b_out = context.create_buffer(a_out.size * a_out.element_size)
        event = prog.addition(queue, [n_times], float, b_in, b_out, :local_work_size => [128])
        queue.enqueue_read_buffer(b_out, a_out, :event_wait_list => [event])
        queue.finish
        diff = (a_in - a_out*3.0)
        n_times.times { |i|
          raise "Computation error #{i} : #{diff[i]+float.s0+float.s1}" if (diff[i]+float.s0+float.s1).abs > 0.00001
        }
      end
    }
  }
end