Merge pull request #872 from MichaelHancock/master
Added opencl language support
This commit is contained in:
commit
083c485b4d
|
@ -0,0 +1,77 @@
|
||||||
|
__kernel void reduce_min(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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)){
|
||||||
|
if (scratch[lid] > scratch[lid + i])
|
||||||
|
scratch[lid] = scratch[lid+i];
|
||||||
|
}
|
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE);
|
||||||
|
}
|
||||||
|
|
||||||
|
//Store cache in output array
|
||||||
|
if (!lid)
|
||||||
|
atomic_min(&B[0], scratch[lid]);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__kernel void reduce_max(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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)){
|
||||||
|
if (scratch[lid] < scratch[lid + i])
|
||||||
|
scratch[lid] = scratch[lid+i];
|
||||||
|
}
|
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE);
|
||||||
|
}
|
||||||
|
|
||||||
|
//Store cache in output array
|
||||||
|
if (!lid)
|
||||||
|
atomic_max(&B[0], scratch[lid]);
|
||||||
|
}
|
||||||
|
|
||||||
|
__kernel void reduce_avg(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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]);
|
||||||
|
}
|
|
@ -0,0 +1,77 @@
|
||||||
|
__kernel void reduce_min(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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)){
|
||||||
|
if (scratch[lid] > scratch[lid + i])
|
||||||
|
scratch[lid] = scratch[lid+i];
|
||||||
|
}
|
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE);
|
||||||
|
}
|
||||||
|
|
||||||
|
//Store cache in output array
|
||||||
|
if (!lid)
|
||||||
|
atomic_min(&B[0], scratch[lid]);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__kernel void reduce_max(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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)){
|
||||||
|
if (scratch[lid] < scratch[lid + i])
|
||||||
|
scratch[lid] = scratch[lid+i];
|
||||||
|
}
|
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE);
|
||||||
|
}
|
||||||
|
|
||||||
|
//Store cache in output array
|
||||||
|
if (!lid)
|
||||||
|
atomic_max(&B[0], scratch[lid]);
|
||||||
|
}
|
||||||
|
|
||||||
|
__kernel void reduce_avg(__global const int* A, __global int* B, __local int* scratch) {
|
||||||
|
//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]);
|
||||||
|
}
|
|
@ -8,6 +8,7 @@ module.exports = {
|
||||||
###
|
###
|
||||||
grammars: [
|
grammars: [
|
||||||
"C"
|
"C"
|
||||||
|
"opencl"
|
||||||
]
|
]
|
||||||
|
|
||||||
###
|
###
|
||||||
|
@ -15,6 +16,7 @@ module.exports = {
|
||||||
###
|
###
|
||||||
extensions: [
|
extensions: [
|
||||||
"c"
|
"c"
|
||||||
|
"cl"
|
||||||
]
|
]
|
||||||
|
|
||||||
options:
|
options:
|
||||||
|
|
Loading…
Reference in New Issue