I'm new to opencl but I have some experience using HLSL. In HLSL multiple passes are used when you need to finish a computation before moving on to the next step.
I would like to know how this sort of thing is done in opencl.
I would like to do something like
__kernel void imagingTest(__read_only image2d_t srcImg, __write_only image2d_t dstImg)
{
//DO Some stuff here
// write the image
write_imageui(dstImg, coord, bgra);
}
I would like to apply this kernel multiple times
EDIT
At the moment the C# code I have to call the kernel is
using System;
using System.Collections;
using System.Collections.Generic;
using System.Drawing;
using System.Drawing.Imaging;
using System.IO;
using System.Runtime.InteropServices;
using Emgu.CV;
using Emgu.Util;
using Emgu;
using Emgu.CV.Structure;
using OpenCL.Net;
namespace HLSLTest
{
public class Computations
{
private Cl.Context _context;
private Cl.Device _device;
private Cl.Kernel kernel;
private Cl.Kernel kernel2;
private void CheckErr(Cl.ErrorCode err, string name)
{
if (err != Cl.ErrorCode.Success)
{
Console.WriteLine("ERROR: " + name + " (" + err.ToString() + ")");
}
}
private void ContextNotify(string errInfo, byte[] data, IntPtr cb, IntPtr userData)
{
Console.WriteLine("OpenCL Notification: " + errInfo);
}
public void Setup()
{
Cl.ErrorCode error;
Cl.Platform[] platforms = Cl.GetPlatformIDs(out error);
List<Cl.Device> devicesList = new List<Cl.Device>();
CheckErr(error, "Cl.GetPlatformIDs");
foreach (Cl.Platform platform in platforms)
{
string platformName = Cl.GetPlatformInfo(platform, Cl.PlatformInfo.Name, out error).ToString();
Console.WriteLine("Platform: " + platformName);
CheckErr(error, "Cl.GetPlatformInfo");
//We will be looking only for GPU devices
foreach (Cl.Device device in Cl.GetDeviceIDs(platform, Cl.DeviceType.Gpu, out error))
{
CheckErr(error, "Cl.GetDeviceIDs");
Console.WriteLine("Device: " + device.ToString());
devicesList.Add(device);
}
}
if (devicesList.Count <= 0)
{
Console.WriteLine("No devices found.");
return;
}
_device = devicesList[0];
if (Cl.GetDeviceInfo(_device, Cl.DeviceInfo.ImageSupport, out error).CastTo<Cl.Bool>() == Cl.Bool.False)
{
Console.WriteLine("No image support.");
return;
}
_context = Cl.CreateContext(null, 1, new[] { _device }, ContextNotify, IntPtr.Zero, out error); //Second parameter is amount of devices
CheckErr(error, "Cl.CreateContext");
//Load and compile kernel source code.
string programPath = Environment.CurrentDirectory + "/../../../ImagingTest.cl"; //The path to the source file may vary
if (!System.IO.File.Exists(programPath))
{
Console.WriteLine("Program doesn't exist at path " + programPath);
return;
}
string programSource = System.IO.File.ReadAllText(programPath);
using (Cl.Program program = Cl.CreateProgramWithSource(_context, 1, new[] { programSource }, null, out error))
{
CheckErr(error, "Cl.CreateProgramWithSource");
//Compile kernel source
error = Cl.BuildProgram(program, 1, new[] { _device }, string.Empty, null, IntPtr.Zero);
CheckErr(error, "Cl.BuildProgram");
//Check for any compilation errors
if
(
Cl.GetProgramBuildInfo
(
program,
_device,
Cl.ProgramBuildInfo.Status,
out error
).CastTo<Cl.BuildStatus>() != Cl.BuildStatus.Success
)
{
CheckErr(error, "Cl.GetProgramBuildInfo");
Console.WriteLine("Cl.GetProgramBuildInfo != Success");
Console.WriteLine(Cl.GetProgramBuildInfo(program, _device, Cl.ProgramBuildInfo.Log, out error));
return;
}
//Create the required kernel (entry function)
kernel = Cl.CreateKernel(program, "imagingTest", out error);
kernel2 = Cl.CreateKernel(program, "imagingTest", out error);
CheckErr(error, "Cl.CreateKernel");
}
}
public void ImagingTest(Image<Gray, Single> InputImage, out Image<Gray, Single> outputImage)
{
Cl.ErrorCode error;
int intPtrSize = 0;
intPtrSize = Marshal.SizeOf(typeof(IntPtr));
//Image's RGBA data converted to an unmanaged[] array
byte[] inputByteArray;
//OpenCL memory buffer that will keep our image's byte[] data.
Cl.Mem inputImage2DBuffer;
Cl.ImageFormat clImageFormat = new Cl.ImageFormat(Cl.ChannelOrder.RGBA, Cl.ChannelType.Unsigned_Int8);
int inputImgWidth, inputImgHeight;
int inputImgBytesSize;
int inputImgStride;
inputImgWidth = InputImage.Width;
inputImgHeight = InputImage.Height;
System.Drawing.Bitmap bmpImage = InputImage.ToBitmap();
//Get raw pixel data of the bitmap
//The format should match the format of clImageFormat
BitmapData bitmapData = bmpImage.LockBits
(
new Rectangle(0, 0, bmpImage.Width, bmpImage.Height),
ImageLockMode.ReadOnly,
PixelFormat.Format32bppArgb
);
inputImgStride = bitmapData.Stride;
inputImgBytesSize = bitmapData.Stride * bitmapData.Height;
//Copy the raw bitmap data to an unmanaged byte[] array
inputByteArray = new byte[inputImgBytesSize];
Marshal.Copy(bitmapData.Scan0, inputByteArray, 0, inputImgBytesSize);
//Allocate OpenCL image memory buffer
inputImage2DBuffer = Cl.CreateImage2D
(
_context,
Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadOnly,
clImageFormat,
(IntPtr)bitmapData.Width,
(IntPtr)bitmapData.Height,
(IntPtr)0,
inputByteArray,
out error
);
CheckErr(error, "Cl.CreateImage2D input");
//Unmanaged output image's raw RGBA byte[] array
byte[] outputByteArray = new byte[inputImgBytesSize];
//Allocate OpenCL image memory buffer
Cl.Mem outputImage2DBuffer = Cl.CreateImage2D
(
_context,
Cl.MemFlags.CopyHostPtr | Cl.MemFlags.WriteOnly,
clImageFormat,
(IntPtr)inputImgWidth,
(IntPtr)inputImgHeight,
(IntPtr)0,
outputByteArray,
out error
);
CheckErr(error, "Cl.CreateImage2D output");
//Pass the memory buffers to our kernel function
error = Cl.SetKernelArg(kernel, 0, (IntPtr)intPtrSize, inputImage2DBuffer);
error |= Cl.SetKernelArg(kernel, 1, (IntPtr)intPtrSize, outputImage2DBuffer);
CheckErr(error, "Cl.SetKernelArg");
//Create a command queue, where all of the commands for execution will be added
Cl.CommandQueue cmdQueue = Cl.CreateCommandQueue(_context, _device, (Cl.CommandQueueProperties)0, out error);
CheckErr(error, "Cl.CreateCommandQueue");
Cl.Event clevent;
//Copy input image from the host to the GPU.
IntPtr[] originPtr = new IntPtr[] { (IntPtr)0, (IntPtr)0, (IntPtr)0 }; //x, y, z
IntPtr[] regionPtr = new IntPtr[] { (IntPtr)inputImgWidth, (IntPtr)inputImgHeight, (IntPtr)1 }; //x, y, z
IntPtr[] workGroupSizePtr = new IntPtr[] { (IntPtr)inputImgWidth, (IntPtr)inputImgHeight, (IntPtr)1 };
error = Cl.EnqueueWriteImage(cmdQueue, inputImage2DBuffer, Cl.Bool.True, originPtr, regionPtr, (IntPtr)0, (IntPtr)0, inputByteArray, 0, null, out clevent);
CheckErr(error, "Cl.EnqueueWriteImage");
//Execute our kernel (OpenCL code)
error = Cl.EnqueueNDRangeKernel(cmdQueue, kernel , 2, null, workGroupSizePtr, null, 0, null, out clevent);
//Cl.WaitForEvents(1, new Cl.Event[] { clevent } );
CheckErr(error, "Cl.EnqueueNDRangeKernel");
//Wait for completion of all calculations on the GPU.
error = Cl.Finish(cmdQueue);
CheckErr(error, "Cl.Finish");
//Read the processed image from GPU to raw RGBA data byte[] array
error = Cl.EnqueueReadImage
(
cmdQueue,
outputImage2DBuffer,
Cl.Bool.True,
originPtr,
regionPtr,
(IntPtr)0,
(IntPtr)0,
outputByteArray,
0,
null,
out clevent
);
CheckErr(error, "Cl.clEnqueueReadImage");
//Clean up memory
Cl.ReleaseKernel(kernel);
Cl.ReleaseCommandQueue(cmdQueue);
Cl.ReleaseMemObject( inputImage2DBuffer);
Cl.ReleaseMemObject(outputImage2DBuffer);
//Get a pointer to our unmanaged output byte[] array
GCHandle pinnedOutputArray = GCHandle.Alloc(outputByteArray, GCHandleType.Pinned);
IntPtr outputBmpPointer = pinnedOutputArray.AddrOfPinnedObject();
//Create a new bitmap with processed data and save it to a file.
Bitmap outputBitmap = new Bitmap(inputImgWidth, inputImgHeight, inputImgStride, PixelFormat.Format32bppArgb, outputBmpPointer);
outputImage = new Image<Gray, Single>(outputBitmap);
//outputBitmap.Save(outputImagePath, System.Drawing.Imaging.ImageFormat.Png);
pinnedOutputArray.Free();
}
}
}
I've been thinking that after the
error = Cl.EnqueueNDRangeKernel(cmdQueue, kernel , 2, null, workGroupSizePtr, null, 0, null, out clevent);
line I should add
Cl.WaitForEvents(1, new Cl.Event[] { clevent } );
Then enqueue another kernel but I'm not sure how to do it.
