Added opencl language support

OpenCL language support is achieved here simply by attaching '.CL'
extension to the 'C' language module. OpenCL is C based however beautify
will not support beautification despite having C language support, due
to the different file extension.
This commit is contained in:
Michael Hancock 2016-03-24 00:34:55 +00:00
parent 96e8cde969
commit 18ae07eaff
3 changed files with 156 additions and 0 deletions

View File

@ -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]);
}

View File

@ -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]);
}

View File

@ -8,6 +8,7 @@ module.exports = {
###
grammars: [
"C"
"opencl"
]
###
@ -15,6 +16,7 @@ module.exports = {
###
extensions: [
"c"
"cl"
]
options: