Esempio n. 1
0
      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 shift(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))
            using (Ocl.Kernel kernel = new Ocl.Kernel())
            using (Ocl.Image2D image2d = new Ocl.Image2D(umat))
            using (Ocl.KernelArg ka = new Ocl.KernelArg(Ocl.KernelArg.Flags.ReadWrite, umatDst))
            {
               float shiftX = 100.5f;
               float shiftY = -50.0f;

               bool success = kernel.Create("shift", ps, buildOpts, errorMsg);
               EmguAssert.IsTrue(success, errorMsg.ToString());
               int idx = 0;
               idx = kernel.Set(idx, image2d);
               idx = kernel.Set(idx, ref shiftX);
               idx = kernel.Set(idx, ref shiftY);
               idx = kernel.Set(idx, ka);
               IntPtr[] globalThreads = new IntPtr[] {new IntPtr(umat.Cols), new IntPtr(umat.Rows), new IntPtr(1) };
               success = kernel.Run(globalThreads, null, true);
               EmguAssert.IsTrue(success, "Failed to run the kernel");
               using (Mat matDst = umatDst.GetMat(AccessType.Read))
               using (Mat saveMat = new Mat())
               {
                  matDst.ConvertTo(saveMat, DepthType.Cv8U, 255.0);
                  saveMat.Save("tmp.jpg");
               }
            }
         }
      }
Esempio n. 2
0
        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 shift(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))
                        using (Ocl.Kernel kernel = new Ocl.Kernel())
                            using (Ocl.Image2D image2d = new Ocl.Image2D(umat))
                                using (Ocl.KernelArg ka = new Ocl.KernelArg(Ocl.KernelArg.Flags.ReadWrite, umatDst))
                                {
                                    float shiftX = 100.5f;
                                    float shiftY = -50.0f;

                                    bool success = kernel.Create("shift", ps, buildOpts, errorMsg);
                                    EmguAssert.IsTrue(success, errorMsg.ToString());
                                    int idx = 0;
                                    idx = kernel.Set(idx, image2d);
                                    idx = kernel.Set(idx, ref shiftX);
                                    idx = kernel.Set(idx, ref shiftY);
                                    idx = kernel.Set(idx, ka);
                                    IntPtr[] globalThreads = new IntPtr[] { new IntPtr(umat.Cols), new IntPtr(umat.Rows), new IntPtr(1) };
                                    success = kernel.Run(globalThreads, null, true);
                                    EmguAssert.IsTrue(success, "Failed to run the kernel");
                                    using (Mat matDst = umatDst.GetMat(AccessType.Read))
                                        using (Mat saveMat = new Mat())
                                        {
                                            matDst.ConvertTo(saveMat, DepthType.Cv8U, 255.0);
                                            saveMat.Save("tmp.jpg");
                                        }
                                }
            }
        }
Esempio n. 3
0
        public void TestOclProgramCompile()
        {
            if (CvInvoke.HaveOpenCL && CvInvoke.UseOpenCL)
            {
                String      sourceStr = @"
__kernel void magnitude_filter_8u(
       __global const uchar* src, int src_step, int src_offset,
       __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
       float scale)
{
   int x = get_global_id(0);
   int y = get_global_id(1);
   if (x < dst_cols && y < dst_rows)
   {
       int dst_idx = y * dst_step + x + dst_offset;
       if (x > 0 && x < dst_cols - 1 && y > 0 && y < dst_rows - 2)
       {
           int src_idx = y * src_step + x + src_offset;
           int dx = (int)src[src_idx]*2 - src[src_idx - 1]          - src[src_idx + 1];
           int dy = (int)src[src_idx]*2 - src[src_idx - 1*src_step] - src[src_idx + 1*src_step];
           dst[dst_idx] = convert_uchar_sat(sqrt((float)(dx*dx + dy*dy)) * scale);
       }
       else
       {
           dst[dst_idx] = 0;
       }
   }
}";
                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;
                        }
            }
        }
Esempio n. 4
0
        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 = @"
__kernel void magnutude_filter_8u(
       __global const uchar* src, int src_step, int src_offset,
       __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
       float scale)
{
   int x = get_global_id(0);
   int y = get_global_id(1);
   if (x < dst_cols && y < dst_rows)
   {
       int dst_idx = y * dst_step + x + dst_offset;
       if (x > 0 && x < dst_cols - 1 && y > 0 && y < dst_rows - 2)
       {
           int src_idx = y * src_step + x + src_offset;
           int dx = (int)src[src_idx]*2 - src[src_idx - 1]          - src[src_idx + 1];
           int dy = (int)src[src_idx]*2 - src[src_idx - 1*src_step] - src[src_idx + 1*src_step];
           dst[dst_idx] = convert_uchar_sat(sqrt((float)(dx*dx + dy*dy)) * scale);
       }
       else
       {
           dst[dst_idx] = 0;
       }
   }
}";

                using (CvString errorMsg = new CvString())
                    using (Ocl.ProgramSource ps = new Ocl.ProgramSource(sourceStr))
                        using (Ocl.Kernel kernel = new Ocl.Kernel())
                            using (Ocl.Image2D image2d = new Ocl.Image2D(umat))
                                using (Ocl.KernelArg ka = new Ocl.KernelArg(Ocl.KernelArg.Flags.ReadWrite, umatDst))
                                {
                                    float shiftX = 100.5f;
                                    float shiftY = -50.0f;

                                    bool success = kernel.Create("myshift", ps, buildOpts, errorMsg);
                                    EmguAssert.IsTrue(success, errorMsg.ToString());
                                    int idx = 0;
                                    idx = kernel.Set(idx, image2d);
                                    idx = kernel.Set(idx, ref shiftX);
                                    idx = kernel.Set(idx, ref shiftY);
                                    idx = kernel.Set(idx, ka);
                                    IntPtr[] globalThreads = new IntPtr[] { new IntPtr(umat.Cols), new IntPtr(umat.Rows), new IntPtr(1) };
                                    success = kernel.Run(globalThreads, null, true);
                                    EmguAssert.IsTrue(success, "Failed to run the kernel");
                                    using (Mat matDst = umatDst.GetMat(AccessType.Read))
                                        using (Mat saveMat = new Mat())
                                        {
                                            matDst.ConvertTo(saveMat, DepthType.Cv8U, 255.0);
                                            saveMat.Save("tmp.jpg");
                                        }
                                }
            }
        }