Basic Image Kernel Trying to get my first image kernel to run

I got the image2D creation steps to work by choosing RGBA. It calls the kernel, but it doesn’t seem to be doing anything.

I create the output image using one of the following:
// outputImage = clCreateImage2D(context, CL_MEM_WRITE_ONLY, &imageFormat, pOutImage->width, pOutImage->height, 0, NULL, &err);
outputImage = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, &imageFormat, pOutImage->width, pOutImage->height, pOutImage->stride, pOutImage->image, &err);
if (err != CL_SUCCESS) {
printf(“Error creating output image rc=%d\n”,err);
rc=11;
goto cleanup;
}

I enqueue the kernel with the following code:
// since I am not trying to share stuff between work items in the work group,
// we can just pass in NULL and let OpenCL divide stuff up

        globalSizes[0]=pInImage->width;
        globalSizes[1]=pInImage->height;
        
        err = clEnqueueNDRangeKernel(commands, kernel, 2, NULL, globalSizes, NULL, 0, NULL, NULL);
        if (err != CL_SUCCESS) {
            printf("Error Enqueueing Kernel rc=%d\n",err);
            rc=14;
            goto cleanup;
        }

I use the following Kernel, which should set the output image to a solid color (not black).
const sampler_t kernelSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void kernelFunc(__read_only image2d_t inImage, __write_only image2d_t outImage,
int width, int height, int stride, int bpp)
{

size_t id = get_global_id(0);
int x,y;
int2 pos;
uint4 pixel;

x = id % width;
y = id / width;
if (y >= height)
    y=0;
pos.x = x;
pos.y = y;

/* pixel = read_imageui(inImage, kernelSampler, pos); */
pixel.x = 128;
pixel.y = 128;
pixel.z = 64;
pixel.w = 255;
write_imageui(outImage, pos, pixel);

return;
}

If I don’t use CL_MEM_*_PTR when creating the output image, then I get a black image.

If I do use CL_MEM*_PTR when creating the output image, then I get the original contents of the output image (from host memory).

Any suggestions?

Dean

I got the image2D creation steps to work by choosing RGBA. It calls the kernel, but it doesn’t seem to be doing anything.

I create the output image using one of the following:
// outputImage = clCreateImage2D(context, CL_MEM_WRITE_ONLY, &imageFormat, pOutImage->width, pOutImage->height, 0, NULL, &err);
outputImage = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, &imageFormat, pOutImage->width, pOutImage->height, pOutImage->stride, pOutImage->image, &err);
if (err != CL_SUCCESS) {
printf(“Error creating output image rc=%d\n”,err);
rc=11;
goto cleanup;
}

I enqueue the kernel with the following code:
// since I am not trying to share stuff between work items in the work group,
// we can just pass in NULL and let OpenCL divide stuff up

        globalSizes[0]=pInImage->width;
        globalSizes[1]=pInImage->height;
        
        err = clEnqueueNDRangeKernel(commands, kernel, 2, NULL, globalSizes, NULL, 0, NULL, NULL);
        if (err != CL_SUCCESS) {
            printf("Error Enqueueing Kernel rc=%d\n",err);
            rc=14;
            goto cleanup;
        }

I use the following Kernel, which should set the output image to a solid color (not black).
const sampler_t kernelSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void kernelFunc(__read_only image2d_t inImage, __write_only image2d_t outImage,
int width, int height, int stride, int bpp)
{

size_t id = get_global_id(0);
int x,y;
int2 pos;
uint4 pixel;

x = id % width;
y = id / width;
if (y >= height)
    y=0;
pos.x = x;
pos.y = y;

/* pixel = read_imageui(inImage, kernelSampler, pos); */
pixel.x = 128;
pixel.y = 128;
pixel.z = 64;
pixel.w = 255;
write_imageui(outImage, pos, pixel);

return;
}

If I don’t use CL_MEM_*_PTR when creating the output image, then I get a black image.

If I do use CL_MEM*_PTR when creating the output image, then I get the original contents of the output image (from host memory).

Any suggestions?

Dean

Why don’t you use a 2D range kernel properly to avoid manually computation of x and y coordinates? There is no need to check image borders as the coordinates will be clamped thanks to the CLK_ADDRESS_CLAMP sampler.

use a kernel code like
int2 pos;
pos.x = get_global_id(0);
pos.y = get_global_id(1);

write_imageui(outImage, pos, pixel);

Why don’t you use a 2D range kernel properly to avoid manually computation of x and y coordinates? There is no need to check image borders as the coordinates will be clamped thanks to the CLK_ADDRESS_CLAMP sampler.

use a kernel code like
int2 pos;
pos.x = get_global_id(0);
pos.y = get_global_id(1);

write_imageui(outImage, pos, pixel);

You hit the nail on the head. It makes sense after you think about it. You told OpenCL how many dimensions and how big each dimension is.

It is now RUNNING!

Thanks!!!

You hit the nail on the head. It makes sense after you think about it. You told OpenCL how many dimensions and how big each dimension is.

It is now RUNNING!

Thanks!!!