✨🔥🎨🔨Add Image2D

Author: ShoupingShan

CL/ImageFilter2D.cl

@@ -0,0 +1,35 @@
+
+// Edge filter of image
+
+__kernel void edge_filter(__read_only image2d_t srcImg,
+                              __write_only image2d_t dstImg,
+                              sampler_t sampler,
+                              int width, int height)
+{
+    
+    float kernelWeights[9] = { -1.0f, 0.0f, -1.0f,
+                               0.0f  , 4.0f, 0.0f,
+                               -1.0f, 0.0f, -1.0f };
+
+    int2 startImageCoord = (int2) (get_global_id(0) - 1, get_global_id(1) - 1);
+    int2 endImageCoord   = (int2) (get_global_id(0) + 1, get_global_id(1) + 1);
+    int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));
+
+    if (outImageCoord.x < width && outImageCoord.y < height)
+    {
+        int weight = 0;
+        float4 outColor = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+        for( int y = startImageCoord.y; y <= endImageCoord.y; y++)
+        {
+            for( int x = startImageCoord.x; x <= endImageCoord.x; x++)
+            {
+                //outColor:(R,G,B,A)
+                outColor += (read_imagef(srcImg, sampler, (int2)(x, y)) * kernelWeights[weight]);
+                weight += 1;
+            }
+        }
+
+        // Write the output value to image
+        write_imagef(dstImg, outImageCoord, outColor);
+    }
+}

GPU/image2D.png

@@ -10,9 +10,18 @@ filter2D is studied as well.
     Inter(R) Corporation Intel(R) Core(TM) i7-4710MQ CPU@2.5GHZ
     NVIDIA Corporation GeForce GTX 850M
 ## How to Run
+Buffer Method
 ``` shell
 $ git clone git@github.com:ShoupingShan/OpenCL-Image-Filter.git
-$ cd OpenCL-Image-Filter
+$ cd OpenCL-Image-Filter/src/buffer_method
+$ make clean
+$ make
+$ ./pic
+```
+Image2D Method
+``` shell
+$ git clone git@github.com:ShoupingShan/OpenCL-Image-Filter.git
+$ cd OpenCL-Image-Filter/src/bimage2D_method
 $ make clean
 $ make
 $ ./pic
@@ -35,29 +44,49 @@ $ ./pic
 #### 2.OpenCV
 ![OpenCV](CPU/OpenCV.jpg)
 
-### OpenCL_CPU and OpenCL_GPU results
+### OpenCL_CPU, OpenCL_GPU and OpenCL_Image2D results
 #### 1.OpenCL_CPU
 ![OpenCL_CPU](CPU/OpenCL_cpu.jpg)
 #### 2.OpenCL_GPU
 ![OpenCL_GPU](GPU/OpenCL_gpu.jpg)
+#### 3.OpenCL_Image2D
+![OpenCL_GPU](GPU/image2D.png)
 
 ### Run time compare
-| `Method`/`Time(/s)`  | Normal  | OpenCV  | OpenCL(CPU)  | OpenCL(GPU)   |
-|:--------------------:|:-------:|:-------:|:------------:|:-------------:|
-| 1  | 0.015371  | 0.00033 |0.005138 |0.000637 |
-| 2  | 0.015318  | 0.000313|0.005108 |0.00064  |
-| 3  | 0.01417   | 0.000312|0.005323 |0.000651 |
-| 4  | 0.014164  | 0.000331|0.005341 |0.000621 |
-| 5  | 0.014439  | 0.000308|0.005168 |0.000623 |
-| 6  | 0.014289  | 0.00031 |0.006278 |0.000624 |
-| 7  | 0.014349  | 0.000311|0.005307 |0.000625 |
-| 8  | 0.014133  | 0.000307|0.00513  |0.000624 |
-| 9  | 0.014122  | 0.000315|0.005227 |0.000624 |
-|10  | 0.014315  | 0.00033 |0.005174 |0.000629 |
-|Average Time/s  | 0.014467| 0.0003145|0.0053194 |0.0006298 |
+| `Method`/`Time(/s)`  | Normal  | OpenCV  | OpenCL(CPU)  | OpenCL(GPU)   |Image2D|
+|:--------------------:|:-------:|:-------:|:------------:|:-------------:|:--:|
+| 1  | 0.015371  | 0.00033 |0.005138 |0.000637 | 0.000375|
+| 2  | 0.015318  | 0.000313|0.005108 |0.00064  |0.000367|
+| 3  | 0.01417   | 0.000312|0.005323 |0.000651 |0.000377|
+| 4  | 0.014164  | 0.000331|0.005341 |0.000621 |0.00038|
+| 5  | 0.014439  | 0.000308|0.005168 |0.000623 |0.00037|
+| 6  | 0.014289  | 0.00031 |0.006278 |0.000624 |0.000364|
+| 7  | 0.014349  | 0.000311|0.005307 |0.000625 |0.000374|
+| 8  | 0.014133  | 0.000307|0.00513  |0.000624 |0.000379|
+| 9  | 0.014122  | 0.000315|0.005227 |0.000624 |0.000377|
+|10  | 0.014315  | 0.00033 |0.005174 |0.000629 |0.000372|
+|Average Time/s  | `0.014467`| `0.0003145`|`0.0053194` |`0.0006298` |`0.000374`|
 
 
 ![runtime](image/runtime.png)
 
+### Run time changed by image size
+
+| `Method`/`Time(/s)`  | Normal  | OpenCV  | OpenCL(CPU)  | OpenCL(GPU)   | Image2D|
+|:--------------------:|:-------:|:-------:|:------------:|:-------------:|:---:|
+| 256*256  | 0.003827  | 0.000123 |0.001939 |0.000408 | 0.000351 |
+| 512*512  | 0.014761  | 0.0003444|0.005273 |0.000589  | 0.000375|
+| 1280*720  | 0.0050005   | 0.001282|0.017817 |0.001345 | 0.000485|
+| 1920*1080  | 0.110346  | 0.002831|0.037708 |0.002273 | 0.000926|
+| 2560*1440  | 0.195084 | 0.005242|0.06523|0.003523 |0.001367 |
+| 2880*2560  | 0.390584  | 0.09822 |0.128243 |0.009293 | 0.0022|
+| 3840*2160  | 0.441809  | 0.0118741|0.143425 |0.010045 |0.002367 |
+| 7680*4320 | 1.7679  | 0.04676|0.56525  |0.029638 |0.007554 |
+
+
+![dimensiontime](image/dimension_time.png)
+
+### OpenCV & OpenCL_GPU & OpenCL_Image2D
+![OOO](image/image2dopencvopencl.png)
 ## Contact me
 **Email**: shp395210@outlook.com

README.md

@@ -7,7 +7,7 @@
 #include <CL/cl.h>
 #include <sys/types.h>
 #include<fcntl.h>
-#define COMPUTE_KERNEL_FILENAME ("./CL/oclFilter.cl") //  path to file
+#define COMPUTE_KERNEL_FILENAME ("../../CL/oclFilter.cl") //  path to file
 using namespace std;
 int width = 512,height = 512;       // image resolution, change it to yours one
 cl_int err;                         // error output
@@ -237,7 +237,7 @@ int main()
     fout1.open("./result/OpenCL.txt",ios::out);
     fout2.open("./result/Normal.txt",ios::out);
     fout3.open("./result/OpenCV.txt",ios::out);
-    Mat rgbimage = imread("./data/Lena.jpg");
+    Mat rgbimage = imread("../../data/Lena.jpg");
     if(!rgbimage.data)//judge weather the image load successfully
     {
         cout << "Fail to load image!" << endl;