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