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"); } } } }
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"); } } } }
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"); } } } }