1

I am trying to implement a blur-function I wrote in c#. I want it to run in an openCL Kernel

The function is as follows:

     private static int cSharpBlur(double[,] blur, int width, int height, int[,] imageToInt, int[,]   outputImage, int i, int j)
    {
        int maskSize = 1;
        double sum = 0.0f;
        // Collect neighbor values and multiply with gaussian
        
        for (int a = -maskSize; a < maskSize + 1; a++)
        {
            for (int b = -maskSize; b < maskSize + 1; b++)
            {
                sum += blur[a+1, b+1] * imageToInt[Clamp(i+a,0,width-1), Clamp(j+b,0,height-1)];
                
            }
        }
        byte[] values = BitConverter.GetBytes(imageToInt[i,j]);
        int alpha = values[3];
        int alphasum = alpha * (int)sum;
        values[3] = (byte)alphasum;
        int newValue = BitConverter.ToInt32(values,0);
        return newValue;
    }

Now I obviously don't have .GetBytes and BitConverter.ToInt32 in openCL. Neither do I have a 2 dimensional array.

I solved this via __kernel void gaussianBlur(__global int* imageToInt, int width, int height, __global double* blurBuffer, int blurBufferSize, __global int* outputBuffer){ int col = get_global_id(0); int row = get_global_id(1);

double sum = 0.0f;
for (int a = -1; a < 2; a++)
{
    for (int b = -1; b < 2; b++)
    {
        sum += blurBuffer[a+b+2] * imageToInt[col+width*row];
    }
}
outputBuffer[col + width * row] = sum;

What's missing is the entire getBytes and ToInt stuff.

How can I do that in openCL?

Thanks in advance and have a great weekend!

1 Answer 1

2

For getBytes, use as_uchar4:

uchar4 values = as_uchar4(imageToInt[col+width*row]);

and for ToInt32, use as_int:

int newValue = as_int(values);

So your kernel should look something like this:

__kernel void gaussianBlur(__global int* imageToInt, int width, int height, __global double* blurBuffer, int blurBufferSize, __global int* outputBuffer) {
    int col = get_global_id(0); int row = get_global_id(1);
    int maskSize = 1;
    double sum = 0.0f;
    for(int a = -maskSize; a < maskSize + 1; a++) { // Collect neighbor values and multiply with gaussian
        for(int b = -maskSize; b < maskSize + 1; b++) {
            sum += blurBuffer[a+1+(b+1)*2*maskSize] * imageToInt[col+width*row];
        }
    }
    uchar4 values = as_uchar4(imageToInt[col+width*row]);
    int alpha = (int)values.s3;
    int alphasum = alpha * (int)sum;
    values.s3 = (uchar)alphasum;
    int newValue = as_int(values);
    outputBuffer[col+width*row] = newValue;
}

These as_... functions come built-in with OpenCL C, and enable you to reinterpret the bits that make up a number, as long as the total number of bits remains the same. In your case, an int is made up of 4 bytes, just like the uchar4 vector data type. With uchar4, you can address the individual bytes with .s1, .s1, .s2, .s3, or with .x, .y, .z, .w.

The as_... functions can also be used to get the individual bits of a float number for example:

float x = 1.0f;
int bits = as_int(x);

In plain C, you can also manually write such a function:

uint as_int(const float x) {
    return *(int*)&x;
}

OpenCL C comes packed with math functionality, more than most other languages. All of the built-in functions are listed in this super helpful reference card.

Sign up to request clarification or add additional context in comments.

1 Comment

Thanks a lot! It worked out really well :)

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Start asking to get answers

Find the answer to your question by asking.

Ask question

Explore related questions

See similar questions with these tags.