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


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());

            if (devicesList.Count <= 0)
                Console.WriteLine("No devices found.");

            _device = devicesList[0];

            if (Cl.GetDeviceInfo(_device, Cl.DeviceInfo.ImageSupport, out error).CastTo<Cl.Bool>() == Cl.Bool.False)
                Console.WriteLine("No image support.");

            _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);

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

                //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),

            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
                Cl.MemFlags.CopyHostPtr | Cl.MemFlags.ReadOnly,
                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
                Cl.MemFlags.CopyHostPtr | Cl.MemFlags.WriteOnly,
                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
                out clevent

            CheckErr(error, "Cl.clEnqueueReadImage");

            //Clean up memory

            Cl.ReleaseMemObject( inputImage2DBuffer);

            //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);


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.

  • 1
    \$\begingroup\$ What OpenCL .NET library are you using? \$\endgroup\$ – S.Richmond Aug 15 '13 at 6:35
  • \$\begingroup\$ There are more than one? \$\endgroup\$ – sav Aug 15 '13 at 6:35
  • \$\begingroup\$ Yeah there are a few: openclnet.codeplex.com and opentk.com \$\endgroup\$ – S.Richmond Aug 15 '13 at 6:37
  • \$\begingroup\$ I'm using openclnet.codeplex.com \$\endgroup\$ – sav Aug 15 '13 at 6:38
  • \$\begingroup\$ Although I'm not particular about which one I use if another is easier \$\endgroup\$ – sav Aug 15 '13 at 6:38

I've figured it out. (I'm using cloo now)

I needed to do this in the opencl code.

I put in a barrier.


So for example

__kernel void AnisoDiff2D
    __global double* img1,
    __global double* img2
    const int x = get_global_id(0);
    const int y = get_global_id(1);
    int index   = CoordinateToIndex(  x, y  );

    //write some stuff to img2
    //img2[index] = ...

    //put in a barrier, this ensures all the processing is done before proceding

    //Set img1 based on some calculation from img2
    //img1[index] = f(img2[index]);

    //Repeat this as many times as you need

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Not the answer you're looking for? Browse other questions tagged or ask your own question.