import sys
import pyopencl as cl
import numpy

KERNEL = """
__kernel void test(__global const int *a,
                   __global const int *b,
                   __global int *c
                  )
{
       int agid = get_global_id(0);
       for (int g=0; g < %(bsize)d; g++)
           c[agid] = popcount(a[agid] ^ b[g]);
}
"""
for ASIZE in [29120, 29122]:
    BSIZE = 29097
    a = numpy.random.randint(-1595081346, 1595081346, ASIZE).astype(numpy.int32)
    b = numpy.random.randint(-1595081346, 1595081346, BSIZE).astype(numpy.int32)
    c_result = numpy.zeros(ASIZE, numpy.int32)
    for platform in cl.get_platforms():
        for device in platform.get_devices():
            ctx = cl.Context([device])
            queue = cl.CommandQueue(ctx,
                                    properties=cl.command_queue_properties.PROFILING_ENABLE)

            mf = cl.mem_flags
            a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
            b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
            dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY | mf.COPY_HOST_PTR, hostbuf=c_result)
            prg = cl.Program(ctx, KERNEL % {'bsize' : BSIZE}).build()
            exec_evt = prg.test(queue, a.shape, None, a_buf, b_buf, dest_buf)
            exec_evt.wait()
            elapsed = 1e-9*(exec_evt.profile.end - exec_evt.profile.start)
            print "Time for ASIZE: %s [%s]: %g s" % (ASIZE,
                                                     cl.device_type.to_string(device.type),
                                                     elapsed)
            cl.enqueue_read_buffer(queue, dest_buf, c_result).wait()