2016-03-23 18:34:55 -06:00
|
|
|
__kernel void reduce_min(__global const int* A, __global int* B, __local int* scratch) {
|
2016-04-01 14:22:12 -06:00
|
|
|
//Get local variable data
|
|
|
|
int id = get_global_id(0);
|
|
|
|
int lid = get_local_id(0);
|
|
|
|
int N = get_local_size(0);
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Store valus of global memory into local memory
|
|
|
|
scratch[lid] = A[id];
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Wait for copying to complete
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
for (int i = 1; i < N; i *= 2) {
|
|
|
|
if (!(lid % (i * 2)) && ((lid + i) < N)) {
|
|
|
|
if (scratch[lid] > scratch[lid + i])
|
|
|
|
scratch[lid] = scratch[lid+i];
|
|
|
|
}
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
}
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Store cache in output array
|
|
|
|
if (!lid)
|
|
|
|
atomic_min(&B[0], scratch[lid]);
|
2016-03-23 18:34:55 -06:00
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
__kernel void reduce_max(__global const int* A, __global int* B, __local int* scratch) {
|
2016-04-01 14:22:12 -06:00
|
|
|
//Get local variable data
|
|
|
|
int id = get_global_id(0);
|
|
|
|
int lid = get_local_id(0);
|
|
|
|
int N = get_local_size(0);
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Store valus of global memory into local memory
|
|
|
|
scratch[lid] = A[id];
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Wait for copying to complete
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
for (int i = 1; i < N; i *= 2) {
|
|
|
|
if (!(lid % (i * 2)) && ((lid + i) < N)) {
|
|
|
|
if (scratch[lid] < scratch[lid + i])
|
|
|
|
scratch[lid] = scratch[lid+i];
|
|
|
|
}
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
}
|
2016-03-23 18:34:55 -06:00
|
|
|
|
2016-04-01 14:22:12 -06:00
|
|
|
//Store cache in output array
|
|
|
|
if (!lid)
|
|
|
|
atomic_max(&B[0], scratch[lid]);
|
2016-03-23 18:34:55 -06:00
|
|
|
}
|
|
|
|
|
|
|
|
__kernel void reduce_avg(__global const int* A, __global int* B, __local int* scratch) {
|
2016-04-01 14:22:12 -06:00
|
|
|
//Get local variable data
|
|
|
|
int id = get_global_id(0);
|
|
|
|
int lid = get_local_id(0);
|
|
|
|
int N = get_local_size(0);
|
|
|
|
|
|
|
|
//Store valus of global memory into local memory
|
|
|
|
scratch[lid] = A[id];
|
|
|
|
|
|
|
|
//Wait for copying to complete
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
|
|
|
|
for (int i = 1; i < N; i *= 2) {
|
|
|
|
if (!(lid % (i * 2)) && ((lid + i) < N))
|
|
|
|
{
|
|
|
|
scratch[lid] += scratch[lid+i];
|
|
|
|
}
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
}
|
|
|
|
|
|
|
|
//Store cache in output array
|
|
|
|
if (!lid)
|
|
|
|
atomic_add(&B[0],scratch[lid]);
|
2016-03-23 18:34:55 -06:00
|
|
|
}
|