Added Ocl Context and Ocl Program.

Emgu.CV.Extern/core/ocl_c.cpp

@@ -93,6 +93,55 @@ void* oclDeviceGetPtr(cv::ocl::Device* device)
 }
 
 
+//----------------------------------------------------------------------------
+//
+//  OclContext
+//
+//----------------------------------------------------------------------------
+cv::ocl::Context* oclContextCreate()
+{
+	return new cv::ocl::Context();
+}
+const cv::ocl::Context* oclContextGetDefault(bool initialize)
+{
+	return &cv::ocl::Context::getDefault(initialize);
+}
+void oclContextRelease(cv::ocl::Context** context)
+{
+	delete* context;
+	*context = 0;
+}
+const cv::ocl::Program* oclContextGetProg(
+	cv::ocl::Context* context,
+	cv::ocl::ProgramSource* prog,
+	cv::String* buildopt,
+	cv::String* errmsg)
+{
+	cv::String tmp;
+	errmsg = errmsg ? errmsg : &tmp;
+	cv::ocl::Program program = context->getProg(*prog, *buildopt, *errmsg);
+	cv::ocl::Program* p = new cv::ocl::Program(program);
+	return p;
+}
+
+//----------------------------------------------------------------------------
+//
+//  OclProgram
+//
+//----------------------------------------------------------------------------
+cv::ocl::Program* oclProgramCreate()
+{
+	return new cv::ocl::Program();
+}
+void oclProgramRelease(cv::ocl::Program** program)
+{
+	delete* program;
+	*program = 0;
+}
+void oclProgramGetBinary(cv::ocl::Program* program, std::vector<char>* binary)
+{
+	program->getBinary(*binary);
+}
 
 //----------------------------------------------------------------------------
 //
@@ -124,6 +173,7 @@ cv::ocl::Kernel* oclKernelCreateDefault()
 }
 bool oclKernelCreate(cv::ocl::Kernel* kernel, cv::String* kname, cv::ocl::ProgramSource* source, cv::String* buildOpts, cv::String* errmsg)
 {
+
 	return kernel->create(kname->c_str(), *source, *buildOpts, errmsg);
 }
 void oclKernelRelease(cv::ocl::Kernel** kernel)

Emgu.CV.Extern/core/ocl_c.h

@@ -58,6 +58,29 @@ CVAPI(const cv::ocl::Device*) oclDeviceGetDefault();
 CVAPI(void) oclDeviceRelease(cv::ocl::Device** device);
 CVAPI(void*) oclDeviceGetPtr(cv::ocl::Device* device);
 
+//----------------------------------------------------------------------------
+//
+//  OclContext
+//
+//----------------------------------------------------------------------------
+CVAPI(cv::ocl::Context*) oclContextCreate();
+CVAPI(const cv::ocl::Context*) oclContextGetDefault(bool initialize);
+CVAPI(void) oclContextRelease(cv::ocl::Context** context);
+CVAPI(const cv::ocl::Program*) oclContextGetProg(
+    cv::ocl::Context* context,
+    cv::ocl::ProgramSource* prog,
+    cv::String* buildopt, 
+    cv::String* errmsg);
+
+//----------------------------------------------------------------------------
+//
+//  OclProgram
+//
+//----------------------------------------------------------------------------
+CVAPI(cv::ocl::Program*) oclProgramCreate();
+CVAPI(void) oclProgramRelease(cv::ocl::Program** program);
+CVAPI(void) oclProgramGetBinary(cv::ocl::Program* program, std::vector<char>* binary);
+
 //----------------------------------------------------------------------------
 //
 //  OclProgramSource

Emgu.CV.Test/AutoTestOpenCL.cs

@@ -6,6 +6,7 @@ using System;
 using System.Collections.Generic;
 using System.Diagnostics;
 using System.Drawing;
+using System.IO;
 using System.Runtime.InteropServices;
 using System.Text;
 using Emgu.CV.CvEnum;
@@ -13,6 +14,7 @@ using Emgu.CV.OCR;
 using Emgu.CV.Structure;
 using Emgu.CV.Util;
 using Emgu.CV.Features2D;
+using Emgu.CV.Ocl;
 using Emgu.CV.XFeatures2D;
 
 
@@ -40,24 +42,10 @@ namespace Emgu.CV.Test
         }
 
         [Test]
-      public void TestOclKernel()
+        public void TestOclProgramCompile()
         {
             if (CvInvoke.HaveOpenCL && CvInvoke.UseOpenCL)
             {
-
-            Ocl.Device defaultDevice = Ocl.Device.Default;
-
-            Mat img = EmguAssert.LoadMat("lena.jpg");
-            Mat imgGray = new Mat();
-            CvInvoke.CvtColor(img, imgGray, ColorConversion.Bgr2Gray);
-            Mat imgFloat = new Mat();
-            imgGray.ConvertTo(imgFloat, DepthType.Cv32F, 1.0/255);
-            UMat umat = imgFloat.GetUMat(AccessType.Read, UMat.Usage.AllocateDeviceMemory);
-            UMat umatDst = new UMat();
-            umatDst.Create(umat.Rows, umat.Cols, DepthType.Cv32F, umat.NumberOfChannels, UMat.Usage.AllocateDeviceMemory);
-            
-            String buildOpts = String.Format("-D dstT={0}", Ocl.OclInvoke.TypeToString(umat.Depth));
-    
                 String sourceStr = @"
 __kernel void magnitude_filter_8u(
        __global const uchar* src, int src_step, int src_offset,
@@ -82,6 +70,49 @@ __kernel void magnitude_filter_8u(
        }
    }
 }";
+                Ocl.Context context = Context.Default;
+                String buildOpts = String.Format("-D dstT={0}", Ocl.OclInvoke.TypeToString(DepthType.Cv8U));
+                using (Ocl.ProgramSource ps = new Ocl.ProgramSource(sourceStr))
+                using (CvString errorMsg = new CvString())
+                using (Program p = context.GetProgram(ps, buildOpts, errorMsg))
+                {
+                    byte[] binary = p.Binary;
+                }
+            }
+        }
+
+        [Test]
+        public void TestOclKernel()
+        {
+            if (CvInvoke.HaveOpenCL && CvInvoke.UseOpenCL)
+            {
+                Ocl.Device defaultDevice = Ocl.Device.Default;
+
+                Mat img = EmguAssert.LoadMat("lena.jpg");
+                Mat imgGray = new Mat();
+                CvInvoke.CvtColor(img, imgGray, ColorConversion.Bgr2Gray);
+                Mat imgFloat = new Mat();
+                imgGray.ConvertTo(imgFloat, DepthType.Cv32F, 1.0 / 255);
+                UMat umat = imgFloat.GetUMat(AccessType.Read, UMat.Usage.AllocateDeviceMemory);
+                UMat umatDst = new UMat();
+                umatDst.Create(umat.Rows, umat.Cols, DepthType.Cv32F, umat.NumberOfChannels,
+                    UMat.Usage.AllocateDeviceMemory);
+
+                String buildOpts = String.Format("-D dstT={0}", Ocl.OclInvoke.TypeToString(umat.Depth));
+
+                String sourceStr = @"
+__constant sampler_t samplerLN = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
+__kernel void myshift(const image2d_t src, float shift_x, float shift_y, __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+{
+   int x = get_global_id(0);
+   int y = get_global_id(1);
+   if (x >= dst_cols) return;
+   int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset));
+   __global dstT *dstf = (__global dstT *)(dst + dst_index);
+   float2 coord = (float2)((float)x+0.5f+shift_x, (float)y+0.5f+shift_y);
+   dstf[0] = (dstT)read_imagef(src, samplerLN, coord).x;
+}";
+
 
                 using (CvString errorMsg = new CvString())
                 using (Ocl.ProgramSource ps = new Ocl.ProgramSource(sourceStr))
@@ -106,7 +137,8 @@ __kernel void magnitude_filter_8u(
                     using (Mat saveMat = new Mat())
                     {
                         matDst.ConvertTo(saveMat, DepthType.Cv8U, 255.0);
-                  saveMat.Save("tmp.jpg");
+                        FileInfo fi = new FileInfo("tmp.jpg");
+                        saveMat.Save(fi.FullName);
                     }
                 }
             }

Emgu.CV/Ocl/Context.cs

@@ -0,0 +1,101 @@
+//----------------------------------------------------------------------------
+//  Copyright (C) 2004-2020 by EMGU Corporation. All rights reserved.       
+//----------------------------------------------------------------------------
+
+using System;
+using System.Collections.Generic;
+using System.Text;
+using Emgu.CV;
+using Emgu.Util;
+using System.Runtime.InteropServices;
+
+namespace Emgu.CV.Ocl
+{
+    /// <summary>
+    /// This class contains ocl context information
+    /// </summary>
+    public partial class Context : UnmanagedObject
+    {
+        private static Context _defaultContext = new Context(OclInvoke.oclContextGetDefault(), false);
+
+        private bool _needDispose;
+
+        /// <summary>
+        /// Create a empty OclContext object
+        /// </summary>
+        public Context()
+           : this(OclInvoke.oclContextCreate(), true)
+        {
+        }
+
+        /// <summary>
+        /// Get the default OclContext. Do not dispose this context.
+        /// </summary>
+        public static Context Default
+        {
+            get
+            {
+                return _defaultContext;
+            }
+        }
+
+        internal Context(IntPtr ptr, bool needDispose)
+        {
+            _ptr = ptr;
+            _needDispose = needDispose;
+        }
+
+        /// <summary>
+        /// Release all the unmanaged memory associated with this OclContext
+        /// </summary>
+        protected override void DisposeObject()
+        {
+            if (_needDispose)
+            {
+                if (_ptr != IntPtr.Zero)
+                {
+                    OclInvoke.oclContextRelease(ref _ptr);
+                }
+            }
+        }
+
+        /// <summary>
+        /// Compile the program 
+        /// </summary>
+        /// <param name="prog">The program source</param>
+        /// <param name="buildOpt">The build option</param>
+        /// <param name="errMsg">Error message</param>
+        /// <returns>The compiled program</returns>
+        public Program GetProgram(ProgramSource prog, String buildOpt, CvString errMsg)
+        {
+            using (CvString csBuildOpt = new CvString(buildOpt))
+            {
+                return new Program(OclInvoke.oclContextGetProg(_ptr, prog, csBuildOpt, errMsg));
+            }
+        }
+
+    }
+
+    /// <summary>
+    /// Class that contains ocl functions
+    /// </summary>
+    public static partial class OclInvoke
+    {
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern IntPtr oclContextCreate();
+
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern IntPtr oclContextGetDefault();
+
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern void oclContextRelease(ref IntPtr oclContext);
+
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern IntPtr oclContextGetProg(
+            IntPtr context,
+            IntPtr prog,
+            IntPtr buildopt,
+            IntPtr errmsg);
+
+    }
+}

Emgu.CV/Ocl/Device.cs

@@ -12,7 +12,7 @@ using System.Runtime.InteropServices;
 namespace Emgu.CV.Ocl
 {
     /// <summary>
-   /// This class contains ocl runtime information
+    /// This class contains ocl device information
     /// </summary>
     public partial class Device : UnmanagedObject
     {
@@ -46,7 +46,7 @@ namespace Emgu.CV.Ocl
         }
 
         /// <summary>
-      /// Release all the unmanaged memory associated with this OclInfo
+        /// Release all the unmanaged memory associated with this OclDevice
         /// </summary>
         protected override void DisposeObject()
         {

Emgu.CV/Ocl/Program.cs

@@ -0,0 +1,79 @@
+//----------------------------------------------------------------------------
+//  Copyright (C) 2004-2020 by EMGU Corporation. All rights reserved.       
+//----------------------------------------------------------------------------
+
+using System;
+using System.Collections.Generic;
+using System.Text;
+using Emgu.CV;
+using Emgu.Util;
+using System.Runtime.InteropServices;
+using Emgu.CV.Util;
+
+namespace Emgu.CV.Ocl
+{
+    /// <summary>
+    /// This class contains ocl Program information
+    /// </summary>
+    public partial class Program : UnmanagedObject
+    {
+        /// <summary>
+        /// Create a empty OclProgram object
+        /// </summary>
+        public Program()
+           : this(OclInvoke.oclProgramCreate())
+        {
+        }
+
+        internal Program(IntPtr ptr)
+        {
+            _ptr = ptr;
+        }
+
+        
+        /// <summary>
+        /// Get the program binary
+        /// </summary>
+        public byte[] Binary
+        {
+            get
+            {
+                if (_ptr == IntPtr.Zero)
+                    return null;
+
+                using (VectorOfByte vb = new VectorOfByte())
+                {
+                    OclInvoke.oclProgramGetBinary(_ptr, vb);
+                    return vb.ToArray();
+                }
+            }
+        }
+
+        /// <summary>
+        /// Release all the unmanaged memory associated with this OclProgram
+        /// </summary>
+        protected override void DisposeObject()
+        {
+            if (_ptr != IntPtr.Zero)
+            {
+                OclInvoke.oclProgramRelease(ref _ptr);
+            }
+        }
+
+    }
+
+    /// <summary>
+    /// Class that contains ocl functions
+    /// </summary>
+    public static partial class OclInvoke
+    {
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern IntPtr oclProgramCreate();
+
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern void oclProgramRelease(ref IntPtr oclProgram);
+
+        [DllImport(CvInvoke.ExternLibrary, CallingConvention = CvInvoke.CvCallingConvention)]
+        internal static extern void oclProgramGetBinary(IntPtr program, IntPtr binary);
+    }
+}