/perf/kseta

__kernel void multiply(__global float *res, __global float *a, __global float *b, unsigned long local_size, __global float *scratch, __local float *shmem) {

    int i;

    float sum = 0;

 

    int threads = get_global_size(0);

    int groups = get_num_groups(0);

    int group = get_group_id(0);

    int pos = get_global_id(0);

    int tid = get_local_id(0);

 

    for(i = 0; i < local_size; i++) {

        sum  += i;//a[i * threads + pos] * b[i * threads + pos];

    }

 

    shmem[tid] = sum;

 

    barrier(CLK_LOCAL_MEM_FENCE);

 

    for(i = get_local_size(0)/2; i>0; i >>= 1) {

        if (tid < i) shmem[tid] += shmem[tid + i];

        barrier(CLK_LOCAL_MEM_FENCE);

    }

 

    if (tid == 0) { 

        scratch[group] = shmem[0];

    }

 

    barrier(CLK_GLOBAL_MEM_FENCE);

 

    if (pos == 0) {

        sum = 0;

        for (i = 0; i < groups; i++)

            sum += scratch[i];

 

        *res = sum;

    }

}