0
$\begingroup$

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 am writing an image filter as below

float4 Convolution(__read_only image2d_t srcImg, int2 point, float * kern)
{
    const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
    int maskSize = 1;
    float4 sum = (float4)(0.0f,0.0f,0.0f,0.0f);
    for (int i = -maskSize; i <= maskSize; i++)
    {
        for(int j = -maskSize; j <= maskSize; j++) 
        {
            int2 delta = (int2)(i+maskSize,j+maskSize); 
            int2 pos = (int2)(i,j);
            sum += kern[(delta.y*3) + delta.x] * convert_float4(read_imageui(srcImg, smp, point + pos));
        }
    }
    return sum;
}


__kernel void imagingTest(__read_only  image2d_t srcImg, __write_only image2d_t dstImg)
{
    float k = 30.0L;
    float delta_t = 0.14285714285714285714285714285714L; // 1/7

    float hN[9];
    hN[0] = 0; hN[1] = 1; hN[2] = 0;
    hN[3] = 0; hN[4] =-1; hN[5] = 0; 
    hN[6] = 0; hN[7] = 0; hN[8] = 0;

    float hS[9];
    hS[0] = 0; hS[1] = 0; hS[2] = 0;
    hS[3] = 0; hS[4] =-1; hS[5] = 0; 
    hS[6] = 0; hS[7] = 1; hS[8] = 0;

    float hE[9];
    hE[0] = 0; hE[1] = 0; hE[2] = 0;
    hE[3] = 0; hE[4] =-1; hE[5] = 1; 
    hE[6] = 0; hE[7] = 0; hE[8] = 0;

    float hW[9];
    hW[0] = 0; hW[1] = 0; hW[2] = 0;
    hW[3] = 1; hW[4] =-1; hW[5] = 0; 
    hW[6] = 0; hW[7] = 0; hW[8] = 0;

    float hNE[9];
    hNE[0] = 0; hNE[1] = 0; hNE[2] = 1;
    hNE[3] = 0; hNE[4] =-1; hNE[5] = 0; 
    hNE[6] = 0; hNE[7] = 0; hNE[8] = 0;

    float hSE[9];
    hSE[0] = 0; hSE[1] = 0; hSE[2] = 0;
    hSE[3] = 0; hSE[4] =-1; hSE[5] = 0; 
    hSE[6] = 0; hSE[7] = 0; hSE[8] = 1;

    float hSW[9];
    hSW[0] = 0; hSW[1] = 0; hSW[2] = 0;
    hSW[3] = 0; hSW[4] =-1; hSW[5] = 0; 
    hSW[6] = 1; hSW[7] = 0; hSW[8] = 0;

    float hNW[9];
    hNW[0] = 1; hNW[1] = 0; hNW[2] = 0;
    hNW[3] = 0; hNW[4] =-1; hNW[5] = 0; 
    hNW[6] = 0; hNW[7] = 0; hNW[8] = 0;

    const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;

    int2 coord = (int2)(get_global_id(0), get_global_id(1));

    uint4 bgra = read_imageui(srcImg, smp, coord); 

    float4 nablaN = Convolution(srcImg, coord, hN);
    float4 nablaS = Convolution(srcImg, coord, hS);
    float4 nablaE = Convolution(srcImg, coord, hE);
    float4 nablaW = Convolution(srcImg, coord, hW);

    float4 nablaNE = Convolution(srcImg, coord, hNE);
    float4 nablaNW = Convolution(srcImg, coord, hNW);
    float4 nablaSE = Convolution(srcImg, coord, hSE);
    float4 nablaSW = Convolution(srcImg, coord, hSW);

    float4 cN  = exp(-(nablaN /k) * (nablaN /k));
    float4 cS  = exp(-(nablaS /k) * (nablaS /k));
    float4 cW  = exp(-(nablaW /k) * (nablaW /k));
    float4 cE  = exp(-(nablaE /k) * (nablaE /k));
    float4 cNE = exp(-(nablaNE/k) * (nablaNE/k));
    float4 cSE = exp(-(nablaSE/k) * (nablaSE/k));
    float4 cSW = exp(-(nablaSW/k) * (nablaSW/k));
    float4 cNW = exp(-(nablaNW/k) * (nablaNW/k));

    float4 sum = 0.5 * (nablaNE * cNE) + (nablaSE * cSE) + (nablaSW * cSW) + (nablaNW * cNW);
    sum += (nablaN * cN) + (nablaS * cS) + (nablaW * cW) + (nablaE * cE);
    sum *= delta_t; 

    bgra.x = bgra.y = bgra.z = convert_int(sum.x);

    bgra.w = 255;
    write_imageui(dstImg, coord, bgra);
}

This performs one pass of anisotropic diffusion, I would like to be able to apply this process multiple times. How do I do this?

EDIT

Here's the C# code

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

            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();
        }
    }
}
$\endgroup$
  • $\begingroup$ I'm using C# .Net Not sure if this is something that I need to do on the C# side of things $\endgroup$ – sav Aug 15 '13 at 2:41
0
$\begingroup$

OpenCL uses barriers

You need to store the results of the first "pass" in a variable then call

barrier(CLK_LOCAL_MEM_FENCE);

Once all the threads have reached the barrier, the next section of code can be executed.

This is to enforce data dependencies.

$\endgroup$

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.