/*
 * kernel for computational device, to be used in num-int-par-file.c:
 *
 * compute partial sum for computation of pi
 */
__kernel void pi(
        const int          nsteps,
        const float        step_size,
        __local  float*    local_sums,   /* part sums for work items */
        __global float*    partial_sums) /* part sums for work groups */
{
    /* number of work items total */
    int work_items = get_global_size(0);
    /* number of work items in work group */
    int local_work_items  = get_local_size(0);
    /* ID with regard to work group */
    int local_id = get_local_id(0);
    /* workgroup ID */
    int group_id = get_group_id(0);
    /* ID with regard to whole index space */
    int global_id = get_global_id(0);

    float x, sum, accum = 0.0;

    for(int i = global_id; i < nsteps; i += work_items) {
        x = (i+0.5f)*step_size;
        accum += 4.0f/(1.0f+x*x);
    }
    local_sums[local_id] = accum;
    barrier(CLK_LOCAL_MEM_FENCE);
    if (local_id == 0) {
        sum = 0.0f;
        for(int i = 0; i < local_work_items; ++i) {
            sum += local_sums[i];
        }
        partial_sums[group_id] = sum;
    }
}