From b69c00c9ff81c740479e13d7e9f99d766b193513 Mon Sep 17 00:00:00 2001
From: Shahzaib Gill <shahzaibgill@gmail.com>
Date: Sat, 8 Oct 2016 15:07:55 -0400
Subject: [PATCH] Kernel and helper function
---
examples/median-filter.py | 61 +++++++++++++++++++++++++++++++++++++++
1 file changed, 61 insertions(+)
diff --git a/examples/median-filter.py b/examples/median-filter.py
index 860f09c9..59e842a0 100644
--- a/examples/median-filter.py
+++ b/examples/median-filter.py
@@ -22,3 +22,64 @@ else:
# Create queue for each kernel execution
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
+
+
+# Kernel function
+src = '''
+void sort(int *a, int *b, int *c) {
+ int swap;
+ if(*a > *b) {
+ swap = *a;
+ *a = *b;
+ *b = swap;
+ }
+ if(*a > *c) {
+ swap = *a;
+ *a = *c;
+ *c = swap;
+ }
+ if(*b > *c) {
+ swap = *b;
+ *b = *c;
+ *c = swap;
+ }
+}
+__kernel void medianFilter(__global float *img, __global float *res, __global int *width, __global int *height)
+{
+ int w = *width;
+ int h = *height;
+ int posx = get_global_id(1);
+ int posy = get_global_id(0);
+ int i = w*posy + posx;
+ // Keeping the edge pixels the same
+ if( posx == 0 || posy == 0 || posx == w-1 || posy == h-1 )
+ {
+ res[ i ] = img[ i ];
+ }
+ else
+ {
+ int pixel00, pixel01, pixel02, pixel10, pixel11, pixel12, pixel20, pixel21, pixel22;
+ pixel00 = img[i - 1 - w];
+ pixel01 = img[i- w];
+ pixel02 = img[i + 1 - w];
+ pixel10 = img[i - 1];
+ pixel11 = img[i];
+ pixel12 = img[i + 1];
+ pixel20 = img[i - 1 + w];
+ pixel21 = img[i + w];
+ pixel22 = img[i + 1 + w];
+ //sort the rows
+ sort( &(pixel00), &(pixel01), &(pixel02) );
+ sort( &(pixel10), &(pixel11), &(pixel12) );
+ sort( &(pixel20), &(pixel21), &(pixel22) );
+ //sort the columns
+ sort( &(pixel00), &(pixel10), &(pixel20) );
+ sort( &(pixel01), &(pixel11), &(pixel21) );
+ sort( &(pixel02), &(pixel12), &(pixel22) );
+ //sort the diagonal
+ sort( &(pixel00), &(pixel11), &(pixel22) );
+ // median is the the middle value of the diagonal
+ res [ i ] = pixel11;
+ }
+}
+'''
\ No newline at end of file
--
GitLab