From 2c8140bbdfbe6e4cf1894e0eb725955dbd2234da Mon Sep 17 00:00:00 2001
From: Shahzaib Gill <shahzaibgill@gmail.com>
Date: Sat, 8 Oct 2016 16:15:43 -0400
Subject: [PATCH 1/4] Initial commit

---
 examples/median-filter.py | 0
 1 file changed, 0 insertions(+), 0 deletions(-)
 create mode 100644 examples/median-filter.py

diff --git a/examples/median-filter.py b/examples/median-filter.py
new file mode 100644
index 00000000..e69de29b

From a809d1682f389cf7e61dcd588dd3f7aa4246f2d8 Mon Sep 17 00:00:00 2001
From: Shahzaib Gill <shahzaibgill@gmail.com>
Date: Sat, 8 Oct 2016 18:03:02 -0400
Subject: [PATCH 2/4] Read in grayscale image

---
 examples/median-filter.py | 13 +++++++++++++
 1 file changed, 13 insertions(+)

diff --git a/examples/median-filter.py b/examples/median-filter.py
index e69de29b..89a27618 100644
--- a/examples/median-filter.py
+++ b/examples/median-filter.py
@@ -0,0 +1,13 @@
+import numpy as np
+from scipy.misc import imread, imsave
+import pycuda.autoinit
+import pycuda.driver as drv
+from pycuda.compiler import SourceModule
+
+# Maximum thread size for GPU is dependent on GPU, but normally 512.
+# Threads per block should be a multiple of 32.
+# Block and Grid Size is dependent on the image.
+# This example uses a 256x256 pixel image. A 2D block (16x16) and a 1D grid (256,1) is used
+
+#Read in image
+img = imread('noisyImage.jpg', flatten=True).astype(np.float32)

From 67c39a9b1c9e8cc49de0845ed11ce309c6c2f082 Mon Sep 17 00:00:00 2001
From: Shahzaib Gill <shahzaibgill@gmail.com>
Date: Sat, 8 Oct 2016 18:03:19 -0400
Subject: [PATCH 3/4] Kernel and helper function

---
 examples/median-filter.py | 53 +++++++++++++++++++++++++++++++++++++++
 1 file changed, 53 insertions(+)

diff --git a/examples/median-filter.py b/examples/median-filter.py
index 89a27618..236cd56f 100644
--- a/examples/median-filter.py
+++ b/examples/median-filter.py
@@ -11,3 +11,56 @@
 
 #Read in image
 img = imread('noisyImage.jpg', flatten=True).astype(np.float32)
+
+
+mod = SourceModule('''
+__host__ __device__ 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;
+    }
+}
+__global__ void medianFilter(float *result, float *img, int w, int h) {
+    //2D Blocks, 1D Grid. Finding respective index
+    int i = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
+    // Keeping the edge pixels the same
+    if (i < w || i > w * (h-1)-1 || i % (w-1) == 0 ) {
+        result[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
+        result[i] = pixel11;
+    }
+}''')
\ No newline at end of file

From a49f15166d89e63b04b1e18d24bddcdee519429d Mon Sep 17 00:00:00 2001
From: Shahzaib Gill <shahzaibgill@gmail.com>
Date: Sat, 8 Oct 2016 18:03:52 -0400
Subject: [PATCH 4/4] Calling kernel to be executed in parallel

---
 examples/median-filter.py | 16 ++++++++++++++--
 1 file changed, 14 insertions(+), 2 deletions(-)

diff --git a/examples/median-filter.py b/examples/median-filter.py
index 236cd56f..fc5ba0b7 100644
--- a/examples/median-filter.py
+++ b/examples/median-filter.py
@@ -12,7 +12,6 @@
 #Read in image
 img = imread('noisyImage.jpg', flatten=True).astype(np.float32)
 
-
 mod = SourceModule('''
 __host__ __device__ void sort(int *a, int *b, int *c) {
     int swap;
@@ -63,4 +62,17 @@
         // median is the the middle value of the diagonal
         result[i] = pixel11;
     }
-}''')
\ No newline at end of file
+}''')
+
+medianFilter = mod.get_function("medianFilter")
+
+#This will need tweaking based on your image
+blockWidth = np.int32(img.shape[1]/16)
+blockHeight = np.int32(img.shape[0]/16)
+gridSize = np.int32((img.shape[0] * img.shape[1])/(blockWidth * blockHeight))
+
+#Empty array for computation output
+result = np.zeros_like(img)
+#Kernel execution
+medianFilter(drv.Out(result), drv.In(img),np.int32(img.shape[1]),np.int32(img.shape[0]), block=(blockWidth,blockHeight,1), grid=(gridSize,1))
+imsave('medianFilter-CUDA.jpg',result)
\ No newline at end of file