commit 6e6d191bda1578aef6137e0aa4bd836b07873946
Author: gfuse <majidgeravand68@gmail.com>
Date:   Fri Jun 18 14:57:59 2021 +0200

    init

diff --git a/CMakeLists.txt b/CMakeLists.txt
new file mode 100755
index 0000000..c00ba93
--- /dev/null
+++ b/CMakeLists.txt
@@ -0,0 +1,14 @@
+cmake_minimum_required(VERSION 3.17)
+project(Vio_OpenCL)
+
+set(CMAKE_CXX_STANDARD 14)
+
+find_package (OpenCL REQUIRED)
+find_package(OpenCV REQUIRED)
+
+include_directories (include ${OpenCL_INCLUDE_DIRS} ${OpenCV_INCLUDE_DIRS} kernel)
+
+add_executable(${PROJECT_NAME} main.cpp include/cl_class.h src/cl_class.cpp)
+target_compile_features(${PROJECT_NAME} PRIVATE cxx_range_for)
+target_link_libraries(${PROJECT_NAME} ${OpenCL_LIBRARIES} ${OpenCV_LIBS})
+ADD_DEFINITIONS(-DTEST_DATA_DIR=\"${PROJECT_SOURCE_DIR}/test/data\")
diff --git a/frame0287.jpg b/frame0287.jpg
new file mode 100755
index 0000000..0c5b7e7
Binary files /dev/null and b/frame0287.jpg differ
diff --git a/include/FastGrayStep.hh b/include/FastGrayStep.hh
new file mode 100644
index 0000000..f76e84a
--- /dev/null
+++ b/include/FastGrayStep.hh
@@ -0,0 +1,80 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#ifndef __CVD_CL_FAST_GRAY_STEP_HH__
+#define __CVD_CL_FAST_GRAY_STEP_HH__
+
+#include <cvd-cl/states/ImageState.hh>
+#include <cvd-cl/states/ListState.hh>
+#include <cvd-cl/worker/WorkerStep.hh>
+
+namespace CVD {
+namespace CL  {
+
+/// \brief WorkerStep to filter a point list by FAST corner criteria.
+class FastGrayStep : public WorkerStep {
+public:
+
+    /// \brief Construct the step.
+    ///
+    /// \param i_image     Input image.
+    /// \param i_points    Input point list.
+    /// \param o_points    Output point list.
+    /// \param threshold   FAST threshold.
+    /// \param ring        Consecutive bits required in FAST ring.
+    explicit FastGrayStep(GrayImageState & i_image, PointListState & i_points, PointListState & o_points, cl_int threshold=40, cl_int ring=9);
+
+    /// \brief De-construct the step.
+    virtual ~FastGrayStep();
+
+    virtual void execute();
+
+protected:
+
+    /// \brief Input image.
+    GrayImageState & i_image;
+
+    /// \brief Input point list.
+    PointListState & i_points;
+
+    /// \brief Output point list.
+    PointListState & o_points;
+
+    /// \brief FAST threshold.
+    cl_int const     threshold;
+
+    /// \brief Consecutive bits required in FAST ring.
+    cl_int const     ring;
+
+    /// \brief OpenCL program.
+    cl::Program      program;
+
+    /// \brief OpenCL kernel.
+    cl::Kernel       kernel;
+};
+
+} // namespace CL
+} // namespace CVD
+
+#endif /* __CVD_CL_FAST_GRAY_STEP_HH__ */
diff --git a/include/PreFastGrayStep.hh b/include/PreFastGrayStep.hh
new file mode 100644
index 0000000..1b8de7b
--- /dev/null
+++ b/include/PreFastGrayStep.hh
@@ -0,0 +1,72 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#ifndef __CVD_CL_PRE_FAST_GRAY_STEP_HH__
+#define __CVD_CL_PRE_FAST_GRAY_STEP_HH__
+
+#include <cvd-cl/worker/WorkerStep.hh>
+#include <cvd-cl/states/ImageState.hh>
+#include <cvd-cl/states/ListState.hh>
+
+namespace CVD {
+namespace CL  {
+
+/// \brief WorkerStep to select points that are reasonably likely to later satisfy a FAST corner test.
+class PreFastGrayStep : public WorkerStep {
+public:
+
+    /// \brief Construct the step.
+    ///
+    /// \param i_image     Input image.
+    /// \param o_points    Output point list.
+    /// \param threshold   FAST threshold.
+    explicit PreFastGrayStep(GrayImageState & i_image, PointListState & o_points, cl_int threshold=40);
+
+    /// \brief De-construct the step.
+    virtual ~PreFastGrayStep();
+
+    virtual void execute();
+
+protected:
+
+    /// \brief Input image.
+    GrayImageState & i_image;
+
+    /// \brief Output point list.
+    PointListState & o_points;
+
+    /// \brief FAST threshold.
+    cl_int const     threshold;
+
+    /// \brief OpenCL program.
+    cl::Program      program;
+
+    /// \brief OpenCL kernel.
+    cl::Kernel       kernel;
+};
+
+} // namespace CL
+} // namespace CVD
+
+#endif /* __CVD_CL_PRE_FAST_GRAY_STEP_HH__ */
diff --git a/include/cl_class.h b/include/cl_class.h
new file mode 100644
index 0000000..b121fa1
--- /dev/null
+++ b/include/cl_class.h
@@ -0,0 +1,189 @@
+//
+// Created by root on 4/27/21.
+//
+
+#ifndef VIO_OPENCL_CL_CLASS_H
+#define VIO_OPENCL_CL_CLASS_H
+#include <CL/cl.hpp>
+#include <iostream>
+#include <vector>
+#include <algorithm>
+#include <fstream>
+#include <exception>
+#include <CL/opencl.h>
+#include <opencv2/opencv.hpp>
+
+class kernel{
+public:
+    kernel(cl::Program* program,std::string name){_kernel = new cl::Kernel(*program,name.c_str());};
+    ~kernel(){
+        _buffers.clear();
+        _images.clear();
+    };
+    template<typename T>
+    int32_t write(size_t id/*buffer ID*/,T* buf,cl::CommandQueue* queue,cl::Context* context,size_t buf_size){
+        _buffers.push_back(std::pair<std::pair<cl::Buffer,size_t>,size_t>(std::pair<cl::Buffer,size_t>(cl::Buffer(*context, CL_MEM_READ_WRITE, sizeof(T) * buf_size),
+                                                                                                       buf_size),id));
+        cl_int err=queue->enqueueWriteBuffer(_buffers.back().first.first,CL_TRUE, 0, sizeof(T) * buf_size, buf);
+        if(err!=0){
+            std::cerr<<"Error C:Kernel, F:write, L:enqueueWriteBuffer\n"<<err<<'\n';
+            return err;
+        }
+        err=_kernel->setArg(id,_buffers.back().first.first);
+        if(err!=0){
+            std::cerr<<"Error C:Kernel, F:write, L3:setArg\n"<<err<<'\n';
+            return err;
+        }
+        return CL_SUCCESS;
+
+    };
+    int32_t write(size_t id/*buffer ID*/,cv::Mat& buf,cl::Context* context){
+        _images.push_back(std::pair<cl::Image2D,size_t>(cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                                                                    cl::ImageFormat(CL_R, CL_UNSIGNED_INT8),
+                                                                    buf.size().width,
+                                                                    buf.size().height,
+                                                                    0,
+                                                                    reinterpret_cast<uchar*>(buf.data)),id));
+
+        cl_int err=_kernel->setArg(id,_images.back().first);
+        if(err!=0){
+            std::cerr<<"Error C:Kernel, F:write image, L:setArg\n"<<err<<'\n';
+            return err;
+        }
+        return CL_SUCCESS;
+    };
+
+    template<typename T>
+    int32_t reload(size_t id,T* buf,cl::CommandQueue* queue){
+        try{
+            if(id>_buffers.size())return CL_MAP_FAILURE;
+            for(auto i:_buffers)if(i.second==id)
+                    queue->enqueueWriteBuffer(i.first.first,CL_TRUE, 0, sizeof(T) * i.first.second, buf);
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    };
+    int32_t reload(size_t id,cv::Mat& buf,cl::Context* context){
+        try{
+            for(auto i:_images)if(i.second==id)
+                    i.first=cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                                        cl::ImageFormat(CL_R, CL_UNSIGNED_INT8),
+                                        buf.size().width,
+                                        buf.size().height,
+                                        0,
+                                        reinterpret_cast<uchar*>(buf.data));
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    };
+    cl::Kernel* _kernel = NULL;
+    std::pair<cl::Buffer,size_t> read(size_t id){
+        for(auto i:_buffers)if(i.second==id)return i.first;
+    }
+    cl::Image2D read_img(size_t id){
+        for(auto i:_images)if(i.second==id)return i.first;
+    }
+private:
+    std::vector<std::pair<std::pair<cl::Buffer,size_t>,size_t>> _buffers;
+    std::vector<std::pair<cl::Image2D,size_t>> _images;
+
+};
+class read_cl{
+public:
+    read_cl(std::string path){
+        src_str=(char *)malloc(0x100000);
+        FILE *fp;
+        fp = fopen(path.c_str(), "r");
+        size = fread(src_str, 1, 0x100000, fp);
+        fclose(fp);
+    }
+    void print(){
+        std::cerr<<src_str<<"\n";
+    }
+    char *src_str = nullptr;
+    size_t size;
+};
+class opencl{
+public:
+    opencl();
+    ~opencl();
+    int32_t make_kernel(std::string name){_kernels.push_back(kernel(program,name));};
+    template<typename T>
+    int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/, int size,T* buf) {
+        try{
+            _kernels.at(id1).write(id2,buf,queue,context,size);
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    }
+    int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,cv::Mat& buf) {
+        try{
+            _kernels.at(id1).write(id2,buf,context);
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    }
+    template<typename T>
+    int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T&& value){
+        _kernels.at(id1)._kernel->setArg(id2,value);
+    };
+    template<typename T>
+    int32_t reload_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T* buf) {
+        try{
+            _kernels.at(id1).reload(id2,buf,queue);
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    }
+    int32_t reload_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,cv::Mat& buf) {
+        try{
+            _kernels.at(id1).reload(id2,buf,context);
+            return CL_SUCCESS;
+        }catch (std::exception& err){
+            std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
+            return CL_MAP_FAILURE;
+        }
+    }
+    int32_t run(size_t id1/*kernal ID*/,std::size_t  x=1,std::size_t y=1,std::size_t z=1) {
+        cl_int err=0;
+        if(z>1 && y>1){
+            err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y,z)/*Global*/, cl::NullRange/*local*/);
+        }else if(z<2 && y>1){
+            err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/);
+        }else{
+            err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x)/*Global*/, cl::NullRange/*local*/);
+        };
+        queue->finish();
+        if(err!=0){
+            std::cerr<<"kernel: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_FUNCTION_NAME>();
+            std::cerr<<" ,Args: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_NUM_ARGS>()<<" ,Run out: "<<err<<'\n';
+            for(int i=0;i<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_NUM_ARGS>();++i)
+                std::cerr<<"argument type "<<i<<":"<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(i)<<'\n';
+        }
+        return err;
+
+    }
+    template<typename T>
+    void read(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/, T* out){
+        queue->enqueueReadBuffer(_kernels.at(id1).read(id2).first, CL_TRUE, 0, sizeof(T) * _kernels.at(id1).read(id2).second, out);
+    }
+    void clear_buf();
+private:
+    std::vector<kernel> _kernels;
+    cl::Context* context = NULL;
+    cl::Device device;
+    cl::Program* program = NULL;
+    cl::CommandQueue* queue=NULL;
+};
+
+#endif //VIO_OPENCL_CL_CLASS_H
diff --git a/kernel/double_add.cl b/kernel/double_add.cl
new file mode 100644
index 0000000..0979f50
--- /dev/null
+++ b/kernel/double_add.cl
@@ -0,0 +1,4 @@
+__kernel void double_add(__global const double* A, __global const double* B, __global double* C) {
+int index = get_global_id(0);
+C[index] = A[index] + B[index];
+};
\ No newline at end of file
diff --git a/kernel/fast-best.cl b/kernel/fast-best.cl
new file mode 100644
index 0000000..38187b3
--- /dev/null
+++ b/kernel/fast-best.cl
@@ -0,0 +1,72 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+
+kernel void fast_best(
+    read_only  image2d_t   scores,
+    global     int2      * corners,
+    global     int2      * filtered,
+    global     int       * icorner
+) {
+
+    // Prepare a suitable OpenCL image sampler.
+    sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+    // Use global work item as 1D offset into corners.
+    int  const ic  = get_global_id(0);
+    int2 const xy  = corners[ic];
+
+    // Read the candidate score.
+    int  const p00 = read_imageui(scores, sampler, xy).x;
+
+    // Read other scores in a tight square around the candidate score.
+    int  const p01 = read_imageui(scores, sampler, xy + (int2)(-1, -1)).x;
+    int  const p02 = read_imageui(scores, sampler, xy + (int2)( 0, -1)).x;
+    int  const p03 = read_imageui(scores, sampler, xy + (int2)( 1, -1)).x;
+    int  const p04 = read_imageui(scores, sampler, xy + (int2)(-1,  0)).x;
+    int  const p05 = read_imageui(scores, sampler, xy + (int2)( 0,  0)).x;
+    int  const p06 = read_imageui(scores, sampler, xy + (int2)( 1,  0)).x;
+    int  const p07 = read_imageui(scores, sampler, xy + (int2)(-1,  1)).x;
+    int  const p08 = read_imageui(scores, sampler, xy + (int2)( 0,  1)).x;
+    int  const p09 = read_imageui(scores, sampler, xy + (int2)( 1,  1)).x;
+
+    // Select the maximum score.
+    int        sco = p00;
+               sco = max(sco, p01);
+               sco = max(sco, p02);
+               sco = max(sco, p03);
+               sco = max(sco, p04);
+               sco = max(sco, p05);
+               sco = max(sco, p06);
+               sco = max(sco, p07);
+               sco = max(sco, p08);
+               sco = max(sco, p09);
+
+    // Keep this score if it is as good as the maximum.
+    if (p00 >= sco) {
+        // Atomically append to filtered corner buffer.
+        filtered[atom_inc(icorner)] = xy;
+    }
+}
+
diff --git a/kernel/fast-gray.cl b/kernel/fast-gray.cl
new file mode 100644
index 0000000..5f45116
--- /dev/null
+++ b/kernel/fast-gray.cl
@@ -0,0 +1,80 @@
+// Copyright (C) 2021  Majid Geravand
+// Copyright (C) 2021  Gfuse
+
+__kernel void fast_gray(
+    __read_only  image2d_t   image,
+    __global     int2      * corners,
+    __global     int       * icorner,
+               int       fast_count
+) {
+
+    // Prepare a suitable OpenCL image sampler.
+    sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+    // Use global work item as 2D image coordinates.
+    int  const x   = get_global_id(0);
+    int  const y   = get_global_id(1);
+    if(x>16 && x<get_image_width(image)-16 && y<get_image_height(image)-16 && y>16){
+        int2 const xy  = (int2)(x, y);
+        // Read the candidate pixel.
+        int  const p00 = read_imageui(image, sampler, xy).x;
+
+        // Read other pixels in a circle around the candidate pixel.
+        int  const p01 = read_imageui(image, sampler, xy + (int2)( 0,  3)).x;
+        int  const p05 = read_imageui(image, sampler, xy + (int2)( 3,  0)).x;
+        int  const p09 = read_imageui(image, sampler, xy + (int2)( 0, -3)).x;
+        int  const p13 = read_imageui(image, sampler, xy + (int2)(-3,  0)).x;
+
+        // Check the absolute difference of each circle pixel.
+        int  const d01 = (abs(p01 - p00) > FAST_THRESH);
+        int  const d05 = (abs(p05 - p00) > FAST_THRESH);
+        int  const d09 = (abs(p09 - p00) > FAST_THRESH);
+        int  const d13 = (abs(p13 - p00) > FAST_THRESH);
+
+        // Check if any two adjacent circle pixels have a high absolute difference.
+
+        if ((d01 && d05) ||(d05 && d09) ||(d09 && d13) ||(d13 && d01)) {
+                // Read other pixels in a circle around the candidate pixel.
+            int  const p02 = read_imageui(image, sampler, xy + (int2)( 1,  3)).x;
+            int  const p03 = read_imageui(image, sampler, xy + (int2)( 2,  2)).x;
+            int  const p04 = read_imageui(image, sampler, xy + (int2)( 3,  1)).x;
+            int  const p06 = read_imageui(image, sampler, xy + (int2)( 3, -1)).x;
+            int  const p07 = read_imageui(image, sampler, xy + (int2)( 2, -2)).x;
+            int  const p08 = read_imageui(image, sampler, xy + (int2)( 1, -3)).x;
+            int  const p10 = read_imageui(image, sampler, xy + (int2)(-1, -3)).x;
+            int  const p11 = read_imageui(image, sampler, xy + (int2)(-2, -2)).x;
+            int  const p12 = read_imageui(image, sampler, xy + (int2)(-3, -1)).x;
+            int  const p14 = read_imageui(image, sampler, xy + (int2)(-3,  1)).x;
+            int  const p15 = read_imageui(image, sampler, xy + (int2)(-2,  2)).x;
+            int  const p16 = read_imageui(image, sampler, xy + (int2)(-1,  3)).x;
+            // Select the maximum score.
+            int     sco = p00;
+                    sco = max(sco, p01);
+                    sco = max(sco, p02);
+                    sco = max(sco, p03);
+                    sco = max(sco, p04);
+                    sco = max(sco, p05);
+                    sco = max(sco, p06);
+                    sco = max(sco, p07);
+                    sco = max(sco, p08);
+                    sco = max(sco, p09);
+                    sco = max(sco, p10);
+                    sco = max(sco, p11);
+                    sco = max(sco, p12);
+                    sco = max(sco, p13);
+                    sco = max(sco, p14);
+                    sco = max(sco, p15);
+
+            // Keep this score if it is as good as the maximum.
+            if (p00 >= sco) {
+                // Atomically append to corner buffer.
+                int const icorn = atom_inc(icorner);
+                if (icorn < fast_count)
+                    corners[icorn] = xy;
+            }
+        }
+
+    }
+
+}
+
diff --git a/kernel/int_add.cl b/kernel/int_add.cl
new file mode 100644
index 0000000..6a7883a
--- /dev/null
+++ b/kernel/int_add.cl
@@ -0,0 +1,4 @@
+__kernel void int_add(__global const int* A, __global const int* B, __global int* C) {
+int index = get_global_id(0);
+C[index] = A[index] + B[index];
+};
\ No newline at end of file
diff --git a/kernel/is_corner_10_greater.cl b/kernel/is_corner_10_greater.cl
new file mode 100644
index 0000000..899434b
--- /dev/null
+++ b/kernel/is_corner_10_greater.cl
@@ -0,0 +1,372 @@
+__kernel void is_corner_10_greater(__global const unsigned char* img,__global const int* x,__global const int* y, int2 image_stride, int wight, short barrier,__global unsigned char* condition) {
+
+/*
+int x_id=x[get_global_id(0)];
+int y_id=y[get_global_id(1)];
+printf("%d, %d \n",x_id,y_id);
+int px=y_id*image_stride+x_id;	   
+int w3 = 3*image_stride;//3*image_stride
+/*
+int t = (img[px] << 8) | img[px+1] - barrier;
+    if (!img[px-1-w3]>t) {
+        if (!img[px+3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]>t) { 
+            if (!img[px-w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+1-w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3+image_stride]>t) { 
+            if (!img[px+3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+1-w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+3]>t) { 
+            if (!img[px-3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+3-image_stride]>t) { 
+            if (!img[px-3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3]>t) { 
+            if (!img[px+2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-2-2*image_stride]>t) { 
+        if (!img[px+3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]>t) { 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+2-2*image_stride]>t) { 
+            if (!img[px-3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+1-w3]>t) { 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-w3]>t) { 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3-image_stride]>t) { 
+        if (!img[px+1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2-2*image_stride]>t) { 
+            if (!img[px+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+1-w3]>t) { 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-w3]>t) { 
+            if (!img[px+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-w3]>t) { 
+        if (!img[px+2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-3+image_stride]>t) { 
+            if (!img[px+1-w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3]>t) { 
+            if (!img[px+3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3]>t) { 
+        if (!img[px+2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2-2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1-w3]>t) { 
+            if (!img[px+1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]>t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+1-w3]>t) { 
+        if (!img[px+1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+2-2*image_stride]>t) { 
+        if (!img[px-3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3+image_stride]>t) { 
+        if (!img[px+3-image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-2+2*image_stride]>t) { 
+        if (!img[px+3-image_stride]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+3-image_stride]>t) { 
+        if (!img[px-1+w3]>t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    }
+    */
+}; 
\ No newline at end of file
diff --git a/kernel/is_corner_10_less.cl b/kernel/is_corner_10_less.cl
new file mode 100644
index 0000000..0815565
--- /dev/null
+++ b/kernel/is_corner_10_less.cl
@@ -0,0 +1,376 @@
+__kernel void is_corner_10_less(__global const uchar* img,__global const int* x,__global const int* y,int image_stride, int wight,int barrier,
+__global uchar* condition) {
+int x_id=x[get_global_id(0)];
+int y_id=y[get_global_id(1)];
+int px=y_id*image_stride+x_id;	   
+int w3 = 3*image_stride;//3*image_stride
+int t = (img[px] << 8) | img[px+1] - barrier;
+if(get_global_id(0)>100){
+    condition[get_global_id(1)*752+get_global_id(0)]=255;
+}else{
+    if(condition[get_global_id(1)*752+get_global_id(0)]!=100)
+    condition[get_global_id(1)*752+get_global_id(0)]=0;
+}
+/*
+    if (!img[px-1-w3]<t) {
+        if (!img[px+3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]<t) { 
+            if (!img[px-w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+1-w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3+image_stride]<t) { 
+            if (!img[px+3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+1-w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+3]<t) { 
+            if (!img[px-3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+3-image_stride]<t) { 
+            if (!img[px-3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3]<t) { 
+            if (!img[px+2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-2-2*image_stride]<t) { 
+        if (!img[px+3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]<t) { 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+2-2*image_stride]<t) { 
+            if (!img[px-3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+1-w3]<t) { 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-w3]<t) { 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3-image_stride]<t) { 
+        if (!img[px+1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2-2*image_stride]<t) { 
+            if (!img[px+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px+1-w3]<t) { 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-w3]<t) { 
+            if (!img[px+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-w3]<t) { 
+        if (!img[px+2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-3+image_stride]<t) { 
+            if (!img[px+1-w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        if (!img[px-3]<t) { 
+            if (!img[px+3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+2-2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3-image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3]<t) { 
+        if (!img[px+2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+2-2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3-image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+1-w3]<t) { 
+            if (!img[px+1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-3+image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-1+w3]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            if (!img[px-2+2*image_stride]<t) { 
+                condition[y_id*wight+x_id]=0;
+            } 
+            condition[y_id*wight+x_id]=255;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+1-w3]<t) { 
+        if (!img[px+1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+2-2*image_stride]<t) { 
+        if (!img[px-3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px-2+2*image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-3+image_stride]<t) { 
+        if (!img[px+3-image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3+image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px-2+2*image_stride]<t) { 
+        if (!img[px+3-image_stride]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        if (!img[px+3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    } 
+    if (!img[px+3-image_stride]<t) { 
+        if (!img[px-1+w3]<t) { 
+            condition[y_id*wight+x_id]=0;
+        } 
+        condition[y_id*wight+x_id]=255;
+    }
+    */
+}; 
\ No newline at end of file
diff --git a/kernel/prefast-gray.cl b/kernel/prefast-gray.cl
new file mode 100644
index 0000000..37ed042
--- /dev/null
+++ b/kernel/prefast-gray.cl
@@ -0,0 +1,103 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+// Enable OpenCL 32-bit integer atomic functions.
+
+__kernel void prefast_gray(
+    __read_only  image2d_t   image,
+    __global     int2      * corners,
+    __global     int       * icorner,
+               int       fast_count
+) {
+    
+    // Prepare a suitable OpenCL image sampler.
+    sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+    // Use global work item as 2D image coordinates.
+    int  const x   = get_global_id(0);
+    int  const y   = get_global_id(1);
+    if(x>16 && x<get_image_width(image)-16 && y<get_image_height(image)-16 && y>16){
+        int2 const xy  = (int2)(x, y);
+        // Read the candidate pixel.
+        int  const p00 = read_imageui(image, sampler, xy).x;
+
+        // Read other pixels in a circle around the candidate pixel.
+        int  const p01 = read_imageui(image, sampler, xy + (int2)( 0,  3)).x;
+        int  const p05 = read_imageui(image, sampler, xy + (int2)( 3,  0)).x;
+        int  const p09 = read_imageui(image, sampler, xy + (int2)( 0, -3)).x;
+        int  const p13 = read_imageui(image, sampler, xy + (int2)(-3,  0)).x;
+
+        // Check the absolute difference of each circle pixel.
+        int  const d01 = (abs(p01 - p00) > FAST_THRESH);
+        int  const d05 = (abs(p05 - p00) > FAST_THRESH);
+        int  const d09 = (abs(p09 - p00) > FAST_THRESH);
+        int  const d13 = (abs(p13 - p00) > FAST_THRESH);
+
+        // Check if any two adjacent circle pixels have a high absolute difference.
+
+        if ((d01 && d05) ||(d05 && d09) ||(d09 && d13) ||(d13 && d01)) {
+                // Read other pixels in a circle around the candidate pixel.
+            int  const p02 = read_imageui(image, sampler, xy + (int2)( 1,  3)).x;
+            int  const p03 = read_imageui(image, sampler, xy + (int2)( 2,  2)).x;
+            int  const p04 = read_imageui(image, sampler, xy + (int2)( 3,  1)).x;
+            int  const p06 = read_imageui(image, sampler, xy + (int2)( 3, -1)).x;
+            int  const p07 = read_imageui(image, sampler, xy + (int2)( 2, -2)).x;
+            int  const p08 = read_imageui(image, sampler, xy + (int2)( 1, -3)).x;
+            int  const p10 = read_imageui(image, sampler, xy + (int2)(-1, -3)).x;
+            int  const p11 = read_imageui(image, sampler, xy + (int2)(-2, -2)).x;
+            int  const p12 = read_imageui(image, sampler, xy + (int2)(-3, -1)).x;
+            int  const p14 = read_imageui(image, sampler, xy + (int2)(-3,  1)).x;
+            int  const p15 = read_imageui(image, sampler, xy + (int2)(-2,  2)).x;
+            int  const p16 = read_imageui(image, sampler, xy + (int2)(-1,  3)).x;
+            // Select the maximum score.
+            int     sco = p00;
+                    sco = max(sco, p01);
+                    sco = max(sco, p02);
+                    sco = max(sco, p03);
+                    sco = max(sco, p04);
+                    sco = max(sco, p05);
+                    sco = max(sco, p06);
+                    sco = max(sco, p07);
+                    sco = max(sco, p08);
+                    sco = max(sco, p09);
+                    sco = max(sco, p10);
+                    sco = max(sco, p11);
+                    sco = max(sco, p12);
+                    sco = max(sco, p13);
+                    sco = max(sco, p14);
+                    sco = max(sco, p15);
+
+            // Keep this score if it is as good as the maximum.
+            if (p00 >= sco) {
+                // Atomically append to corner buffer.
+                int const icorn = atom_inc(icorner);
+                if (icorn < fast_count)
+                    corners[icorn] = xy;
+            }
+        }
+
+    }
+    
+}
+
diff --git a/main.cpp b/main.cpp
new file mode 100755
index 0000000..a0fd4f8
--- /dev/null
+++ b/main.cpp
@@ -0,0 +1,96 @@
+#include <iostream>
+#include "cl_class.h"
+#include <opencv2/opencv.hpp>
+#include <opencv2/core.hpp>
+#include <opencv2/core/mat.hpp>
+// Maximum number of corners.
+#define FAST_COUNT (1 << 18)
+
+int main() {
+    opencl gpu;
+    /*
+    gpu.make_kernel("int_add");
+    int A[] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+    int B[] = { 0, 1, 2, 0, 1, 2, 0, 1, 2, 0 };
+    gpu.write_buf(0,0,10,A);
+    gpu.write_buf(0,1,10,B);
+    int C[10];
+    gpu.run(0,10,1,1,2,10,C);
+    std::cout << " result: \n";
+    for (int i = 0; i < 10; i++){
+        std::cout << C[i] << " ";
+    }
+    std::cout << std::endl;
+    int A1[] = { 10, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+    int B1[] = { 0, 1, 2, 0, 1, 2, 0, 1, 2, 10};
+    gpu.reload_buf(0,0,A1);
+    gpu.reload_buf(0,1,B1);
+    int C1[10];
+    gpu.run(0,10,1,1,2,10,C1);
+    std::cout << " result: \n";
+    for (int i = 0; i < 10; i++){
+        std::cout << C1[i] << " ";
+    }
+    std::cout << std::endl;
+    gpu.make_kernel("double_add");
+    double A2[] = { 0.2, 1.3, 2.5, 3.6, 4.9, 5.12, 6.5, 7.6, 8.4, 9.1 };
+    double B2[] = { 0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0 };
+    gpu.write_buf(1,0,10,A2);
+    gpu.write_buf(1,1,10,B2);
+    double C2[10];
+    gpu.run(1,10,1,1,2,10,C2);
+    std::cout << " result: \n";
+    for (int i = 0; i < 10; i++){
+        std::cout << C2[i] << " ";
+    }
+    std::cout << std::endl;
+     */
+    gpu.make_kernel("fast_gray");
+    cv::Mat img = cv::imread("../frame0287.jpg", CV_8UC1);
+    cv::Mat imageRGBA;
+    cv::Mat downSampled;
+    //cv::imshow("test2",img);
+    //cv::waitKey();
+    cv::resize(img, downSampled, cv::Size(752, 480));
+    cv::cvtColor(downSampled, imageRGBA, CV_GRAY2RGBA);
+    gpu.write_buf(0,0,downSampled);
+    //cv::imshow("test2",downSampled);
+    //cv::waitKey();
+    cl_int2 corners_[740*450];
+    gpu.write_buf(0,1,740*450,corners_);
+    int icorner[1]={0};
+    gpu.write_buf(0,2,1,icorner);
+    int fast_count=FAST_COUNT;
+    gpu.write_buf(0,3,fast_count);
+    printf("\nOpenCL version\n");
+    double time_accumulator = 0;
+    for(int i=0;i<1000;++i){
+        double t = (double)cv::getTickCount();
+        gpu.reload_buf(0,0,downSampled);
+        int icorner[1]={0};
+        gpu.reload_buf(0,2,icorner);
+        gpu.run(0,752,480,1);
+       time_accumulator +=  ((cv::getTickCount() - t) / cv::getTickFrequency());
+    }
+    printf("OpenCL took %f ms (average over %d trials).\n", time_accumulator/((double)1000)*1000.0, 1000 );
+    int icorner_[1]={0};
+    gpu.read(0,2,icorner_);
+    cl_int2 corners[740*450];
+    gpu.read(0,1,corners);
+    //double t = (double)cv::getTickCount();
+    //cv::Mat condition=cv::Mat::ones(cv::Size(752,480),CV_8UC1);
+    //gpu.run(2,752,480,1,6,360960,condition.data);
+
+    int count=0;
+    for(int i=0;i<icorner_[0];++i){
+        if(corners[i].x<0 || corners[i].x>750 || corners[i].y>450 || corners[i].y<0)continue;
+        cv::circle(imageRGBA, cv::Point(corners[i].x,corners[i].y ), 1, cv::Scalar(0, 0,255),1);
+        ++count;
+    }
+    cv::imshow("test3",imageRGBA);
+    std::cerr<<count<<'\n';
+
+    cv::waitKey();
+    gpu.clear_buf();
+    return 0;
+}
diff --git a/readme.md b/readme.md
new file mode 100755
index 0000000..32c6171
--- /dev/null
+++ b/readme.md
@@ -0,0 +1,11 @@
+sudo apt-get install -y gpg-agent wget
+wget -qO - https://repositories.intel.com/graphics/intel-graphics.key |
+  sudo apt-key add -
+sudo apt-add-repository \
+  'deb [arch=amd64] https://repositories.intel.com/graphics/ubuntu bionic main'
+
+sudo apt-get update
+sudo apt install ocl-icd-* opencl-headers intel-level-zero-gpu level-zero intel-opencl intel-igc-opencl-devel level-zero-devel
+
+CodeBlocks build
+cmake .. -G"CodeBlocks - Unix Makefiles"
diff --git a/src/FastGrayStep.cc b/src/FastGrayStep.cc
new file mode 100644
index 0000000..5a2644d
--- /dev/null
+++ b/src/FastGrayStep.cc
@@ -0,0 +1,71 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#include "cvd-cl/steps/FastGrayStep.hh"
+#include "kernels/fast-gray.hh"
+
+namespace CVD {
+namespace CL  {
+
+FastGrayStep::FastGrayStep(GrayImageState & i_image, PointListState & i_points, PointListState & o_points, cl_int threshold, cl_int ring) :
+    WorkerStep (i_image.worker),
+    i_image    (i_image),
+    i_points   (i_points),
+    o_points   (o_points),
+    threshold  (threshold),
+    ring       (ring)
+{
+    char opt[256] = {0,};
+    snprintf(opt, sizeof(opt) - 1, "-DFAST_THRESH=%d -DFAST_RING=%d -DFAST_COUNT=%d", int(threshold), int(ring), int(o_points.size));
+    worker.compile(&program, &kernel, OCL_FAST_GRAY, "fast_gray", opt);
+}
+
+FastGrayStep::~FastGrayStep() {
+    // Do nothing.
+}
+
+void FastGrayStep::execute() {
+    // Read number of input points.
+    size_t const np = i_points.getCount();
+
+    // Assign kernel parameters.
+    kernel.setArg(0, i_image.image);
+    kernel.setArg(1, i_points.buffer);
+    kernel.setArg(2, o_points.buffer);
+    kernel.setArg(3, o_points.count);
+    kernel.setArg(4, (cl_uint) np);
+
+    // Reset number of output points.
+    o_points.setCount(0);
+
+    // Prepare global and local size.
+    cl::NDRange const global(worker.padGlobal(np));
+    cl::NDRange const local(worker.defaultLocal);
+
+    // Queue kernel with global size set to number of input points.
+    worker.queue.enqueueNDRangeKernel(kernel, cl::NullRange, global, local);
+}
+
+} // namespace CL
+} // namespace CVD
diff --git a/src/PreFastGrayStep.cc b/src/PreFastGrayStep.cc
new file mode 100644
index 0000000..4719f70
--- /dev/null
+++ b/src/PreFastGrayStep.cc
@@ -0,0 +1,70 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#include "cvd-cl/steps/PreFastGrayStep.hh"
+#include "kernels/prefast-gray.hh"
+
+namespace CVD {
+namespace CL  {
+
+PreFastGrayStep::PreFastGrayStep(GrayImageState & i_image, PointListState & o_points, cl_int threshold) :
+    WorkerStep (i_image.worker),
+    i_image    (i_image),
+    o_points   (o_points),
+    threshold  (threshold)
+{
+    char opt[256] = {0,};
+    snprintf(opt, sizeof(opt) - 1, "-DFAST_THRESH=%d -DFAST_COUNT=%d", int(threshold), int(o_points.size));
+    worker.compile(&program, &kernel, OCL_PRE_FAST_GRAY, "prefast_gray", opt);
+}
+
+PreFastGrayStep::~PreFastGrayStep() {
+    // Do nothing.
+}
+
+void PreFastGrayStep::execute() {
+    // Reset number of output points.
+    o_points.setCount(0);
+
+    // Assign kernel parameters.
+    kernel.setArg(0, i_image.image);
+    kernel.setArg(1, o_points.buffer);
+    kernel.setArg(2, o_points.count);
+
+    // Read image dimensions.
+    size_t const nx = i_image.nx - 16;
+    size_t const ny = i_image.ny - 16;
+
+    // Construct global, local and offset with a safety boundary.
+    // 16x16 appears to give good performance on most devices.
+    cl::NDRange const global (nx - 16, ny - 16);
+    cl::NDRange const local  (      8,       8);
+    cl::NDRange const offset (      8,       8);
+
+    // Queue kernel.
+    worker.queue.enqueueNDRangeKernel(kernel, offset, global, local);
+}
+
+} // namespace CL
+} // namespace CVD
diff --git a/src/cl_class.cpp b/src/cl_class.cpp
new file mode 100644
index 0000000..fdfbc74
--- /dev/null
+++ b/src/cl_class.cpp
@@ -0,0 +1,48 @@
+//
+// Created by root on 4/27/21.
+//
+#include "cl_class.h"
+opencl::opencl() {
+    std::vector<cl::Platform> all_platforms;
+    cl::Platform::get(&all_platforms);
+    if (all_platforms.size() == 0){
+        std::cerr << " No platforms found. Check OpenCL installation!\n";
+        exit(1);
+    }
+    //get default device of the default platform
+    std::vector<cl::Device> all_devices;
+    for(auto i:all_platforms){
+        std::cout << "Find platform: " << i.getInfo<CL_PLATFORM_NAME>() << "\n";
+        i.getDevices(CL_DEVICE_TYPE_GPU, &all_devices);
+    }
+
+    if (all_devices.size() == 0){
+        std::cout << " No devices found. Check OpenCL installation!\n";
+        exit(1);
+    }
+    device = all_devices[0];
+    std::cout << "CL_DEVICE_NAME: " << device.getInfo<CL_DEVICE_NAME>() <<'\n'
+              << "CL_DEVICE_OPENCL_C_VERSION: " <<device.getInfo<CL_DEVICE_OPENCL_C_VERSION>()<<'\n'
+              << "CL_DEVICE_BUILT_IN_KERNELS: " <<device.getInfo<CL_DEVICE_BUILT_IN_KERNELS>()<<'\n'
+              << "CL_DEVICE_COMPILER_AVAILABLE: " <<device.getInfo<CL_DEVICE_COMPILER_AVAILABLE>()<<'\n'
+              << "CL_DEVICE_LOCAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>()<<'\n'
+              << "CL_DEVICE_GLOBAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>()<<'\n';
+
+    context=new cl::Context({ device });
+    cl::Program::Sources sources;
+    read_cl fast("../kernel/fast-gray.cl");
+    sources.push_back({ fast.src_str, fast.size });
+    program=new cl::Program(*context, sources);
+    if(program->build({ device },"-DFAST_THRESH=40") !=0)
+        std::cout << " Error building: " << program->getBuildInfo<CL_PROGRAM_BUILD_STATUS>(device)<<'\n'
+                  <<program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << '\n';
+    queue=new cl::CommandQueue(*context,device,0,NULL);
+}
+opencl::~opencl() {
+}
+void opencl::clear_buf() {
+    queue->finish();
+    queue->flush();
+    _kernels.clear();
+}
+
diff --git a/src/fastcl.cc b/src/fastcl.cc
new file mode 100644
index 0000000..c561b1e
--- /dev/null
+++ b/src/fastcl.cc
@@ -0,0 +1,135 @@
+// Copyright (C) 2011  Dmitri Nikulin
+// Copyright (C) 2011  Monash University
+//
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+//
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#include "common.h"
+
+#include <cvd-cl/steps/PreFastGrayStep.hh>
+#include <cvd-cl/steps/FastGrayStep.hh>
+
+static void clfast(cl::Device & device, CVD::Image<CVD::byte> const & image) {
+    // Re-interpret image pointer.
+    cl_uchar const * const data = reinterpret_cast<cl_uchar const *>(image.data());
+    int const nx = image.size().x;
+    int const ny = image.size().y;
+
+    CVD::CL::Worker worker(device);
+
+    // Create states.
+    CVD::CL::GrayImageState cl_image(worker, ny, nx, CL_MEM_READ_ONLY);
+    CVD::CL::PointListState cl_corners1(worker, FAST_COUNT);
+    CVD::CL::PointListState cl_corners2(worker, FAST_COUNT);
+
+    // Create steps.
+    CVD::CL::PreFastGrayStep fast1(cl_image, cl_corners1, FAST_THRESH);
+    CVD::CL::FastGrayStep    fast2(cl_image, cl_corners1, cl_corners2, FAST_THRESH, FAST_RING);
+
+    // Write image data.
+    cl_image.set(data);
+
+    // Run kernels for time.
+
+    long const time1 = time(NULL);
+
+    for (int i = 0; i < REPEAT; i++)
+        fast1.execute();
+
+    long const time2 = time(NULL);
+
+    for (int i = 0; i < REPEAT; i++)
+        fast2.execute();
+
+    long const time3 = time(NULL);
+
+    // Read number of corners.
+    int const ncorners1 = cl_corners1.getCount();
+    int const ncorners2 = cl_corners2.getCount();
+
+    // Calculate microseconds per kernel.
+    int const us1 = (((time2 - time1) * 1000000) / REPEAT);
+    int const us2 = (((time3 - time2) * 1000000) / REPEAT);
+
+    // Report timing and number of corners.
+    std::cerr << "OpenCL" << std::endl;
+    std::cerr << std::setw(12) << ncorners1 << " corners 1" << std::endl;
+    std::cerr << std::setw(12) << ncorners2 << " corners 2" << std::endl;
+    std::cerr << std::setw(12) << us1 << " microseconds 1" << std::endl;
+    std::cerr << std::setw(12) << us2 << " microseconds 2" << std::endl;
+    std::cerr << std::endl;
+}
+
+void clfast(CVD::Image<CVD::byte> const & image) {
+    try {
+        // Prepare list for all OpenCL devices on all platforms.
+        std::vector<cl::Device> devices;
+
+        std::vector<cl::Platform> platforms;
+        cl::Platform::get(&platforms);
+
+        std::cerr << "Found " << platforms.size() << " OpenCL platforms" << std::endl;
+
+        for (size_t ip = 0; ip < platforms.size(); ip++) {
+            cl::Platform &pf = platforms.at(ip);
+
+            std::cerr
+                << pf.getInfo<CL_PLATFORM_NAME   >() << " ("
+                << pf.getInfo<CL_PLATFORM_VENDOR >() << ", "
+                << pf.getInfo<CL_PLATFORM_VERSION>() << ")"
+                << std::endl;
+
+            std::vector<cl::Device> newDevices;
+
+            try {
+                pf.getDevices(CL_DEVICE_TYPE_ALL, &newDevices);
+                std::cerr << "  Found " << newDevices.size() << " OpenCL devices" << std::endl;
+            } catch (cl::Error & err) {
+                std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
+            }
+
+            for (size_t id = 0; id < newDevices.size(); id++) {
+                cl::Device &dev = newDevices.at(id);
+
+                std::cerr << "    " << dev.getInfo<CL_DEVICE_NAME>() << std::endl;
+
+                // Add to list of all devices.
+                devices.push_back(dev);
+            }
+        }
+
+        std::cerr << std::endl << std::endl;
+
+        for (size_t id = 0; id < devices.size(); id++) {
+            cl::Device &dev = devices.at(id);
+
+            std::cerr << "Running pipeline for \"" << dev.getInfo<CL_DEVICE_NAME>() << "\"" << std::endl;
+
+            try {
+                clfast(dev, image);
+            } catch (cl::Error & err) {
+                std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
+            }
+        }
+    } catch (cl::Error & err) {
+        std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
+    }
+}
diff --git a/test/data/test1.png b/test/data/test1.png
new file mode 100644
index 0000000..df6935b
Binary files /dev/null and b/test/data/test1.png differ