Linux/AM5728: OpenCL image processing support

Part Number: AM5728

Tool/software: Linux

I am beginning to investigate OpenCL on the AM5728. I have OpenCL working and can do basic things like vector addition. I am now interested in image processing with OpenCL. I tried to implement the same "blur" function found in this blog post:  "https://anteru.net/blog/2012/11/05/2022/index.html".

I am using TI Processor SDK with Kernel 4.4.41.

When OpenCL tries to compile the kernel, it fails with CL_BUILD_PROGRAM_FAILURE ( -11 ) because of the functions "read_imagef" and "write_imagef".

Does TI OpenCL on the AM5728 not support image functions?

Error message:

Defined and initialized filter 

 undefined    first referenced     
  symbol          in file          
 ---------    ----------------     
 read_imagef  /tmp/opencltJqDR7.obj
 write_imagef /tmp/opencltJqDR7.obj

error #10234-D: unresolved symbols remain
error #10010: errors encountered during linking; "/tmp/opencltJqDR7.out" not
   built

>> Compilation failure
OpenCL call failed with error -11

Here is the code:

/**
 * OCLControl::blurImage()
 * Purpose: Blue Image with OpenCL
 *
 * @param input is a char* representing the image to process
 * @param width is stride of input image
 * @param height is height of input image
 * @param depth is bit depth of input image
 *
 */
void OCLControl::blurImage(char *input, uint width, uint height, uint depth) {

    // Error tracking
    int error = -1;

    // Simple Gaussian blur filter
    float filter [] = {
        1, 2, 1,
        2, 4, 2,
        1, 2, 1
    };

    // Normalize the filter
    for (int i = 0; i < 9; ++i) {
        filter [i] /= 16.0f;
    }

#ifdef OCL_DEBUG
    qDebug() << "Defined and initialized filter";
#endif

    cl_program program = CreateProgram(QString("ocl.cl"));
    CheckError(clBuildProgram(program, 1, &m_DeviceId, "-D FILTER_SIZE=1", nullptr, nullptr));

#ifdef OCL_DEBUG
    qDebug() << "Created program";
#endif

    // www.khronos.org/.../clCreateKernel.html
    cl_kernel kernel = clCreateKernel(program, "Filter", &error);
    CheckError(error);

#ifdef OCL_DEBUG
    qDebug() << "Created kernel";
#endif

    // Create image from DMABUF
    const auto image = LoadImageFromCharArray(input, width, height, depth);

    // www.khronos.org/.../clCreateImage2D.html
    static const cl_image_format format = { CL_RGBA, CL_UNORM_INT8 };
    cl_mem inputImage = clCreateImage2D(m_Context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &format,
                                        image.width, image.height, 0,
                                        // This is a bug in the spec
                                        const_cast<char*> (image.pixel.data ()),
                                        &error);
    CheckError(error);

#ifdef OCL_DEBUG
    qDebug() << "Created 2D input Image";
#endif

    // Create output image
    cl_mem outputImage = clCreateImage2D(m_Context, CL_MEM_WRITE_ONLY, &format,
                                         image.width, image.height, 0,
                                         nullptr, &error);
    CheckError (error);

#ifdef OCL_DEBUG
    qDebug() << "Created 2D output Image";
#endif

    // Create a buffer for the filter weights
    // www.khronos.org/.../clCreateBuffer.html
    cl_mem filterWeightsBuffer = clCreateBuffer(m_Context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof (float) * 9, filter, &error);
    CheckError (error);

    // Setup the kernel arguments
    clSetKernelArg (kernel, 0, sizeof (cl_mem), &inputImage);
    clSetKernelArg (kernel, 1, sizeof (cl_mem), &filterWeightsBuffer);
    clSetKernelArg (kernel, 2, sizeof (cl_mem), &outputImage);

#ifdef OCL_DEBUG
    qDebug() << "Set Kernel Arguments";
#endif

    // Run the processing
    // www.khronos.org/.../clEnqueueNDRangeKernel.html
    std::size_t offset [3] = { 0 };
    std::size_t size [3] = { image.width, image.height, 1 };
    CheckError(clEnqueueNDRangeKernel(m_CommandQueue, kernel, 2, offset, size, nullptr, 0, nullptr, nullptr));

    // Prepare the result image, set to black
    Image result = image;
    std::fill(result.pixel.begin (), result.pixel.end (), 0);

#ifdef OCL_DEBUG
    qDebug() << "Submitted image for processing";
#endif

    // Get the result back to the host
    std::size_t origin [3] = { 0 };
    std::size_t region [3] = { result.width, result.height, 1 };
    clEnqueueReadImage (m_CommandQueue, outputImage, CL_TRUE,
                        origin, region, 0, 0,
                        result.pixel.data(), 0, nullptr, nullptr);

#ifdef OCL_DEBUG
    qDebug() << "Read back image";
#endif

    // Save for debugging
    SaveImage(RGBAtoRGB(result), "output.ppm");

#ifdef OCL_DEBUG
    qDebug() << "Saved Image as file to output";
#endif

    // Cleanup
    clReleaseMemObject(outputImage);
    clReleaseMemObject(filterWeightsBuffer);
    clReleaseMemObject(inputImage);
    clReleaseKernel (kernel);
    clReleaseProgram (program);

}