Histogram

common_kernel.h

#define COEFFICIENT 64

img_hist.c

#include <stdio.h>
#include <stdint.h>

#include <opencv2/imgproc/imgproc_c.h>
#include <opencv2/highgui/highgui_c.h>

#ifdef __APPLE__
# include <OpenCL/opencl.h>
#else
#include <CL/opencl.h>
#endif

#include "../common/time_utils.h"
#include "../common/ocl_utils.h"

#include "common_kernel.h"

void print_histogram(int * histogram, const int num_bins)
{
    printf("Histogram values:\n");
    for (int i = 0; i < num_bins; ++i)
    {
        printf("%d | ", histogram[i]);
    }
    printf("\n");
}

void usage(const char *prog_name)
{
    printf("Usage:\n\t");
    printf("%s <num_bins> <input image>\n", prog_name);
}

void calculate_histogram(CvMat * image, int * histogram, const int num_bins)
{
    cl_int error;

// --- Size of picture
    cl_int aantal_pixels = image->rows * image->cols;
// --- argument needed for algorithm in kernel
    cl_float bin_size = 255.f / (cl_float)num_bins;

// --- Sizeof result buffer (in bytes)
    int host_result_size = (sizeof(cl_int) * num_bins);
    
// --- Define work sizes for both kernels
    size_t global_work_sizes[2];
    // TODO: fix int deling
    global_work_sizes[0] = aantal_pixels / COEFFICIENT;
    global_work_sizes[1] = num_bins;

// --- Init kernel 1
    // Create buffers on GPU (for kernel 1)
    cl_mem image_input = clCreateBuffer(g_context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, aantal_pixels, image->data.ptr, &error);
    ocl_err(error);
    cl_mem bin_output = clCreateBuffer(g_context, CL_MEM_READ_WRITE, sizeof(cl_int) * num_bins * global_work_sizes[0], NULL, &error);
    ocl_err(error);

    // Create additional buffer on GPU (for kernel 2)
    cl_mem cl_histogram = clCreateBuffer(g_context, CL_MEM_READ_ONLY, host_result_size, NULL, &error);
    ocl_err(error);
    
    // Compile kernel
    cl_kernel kernel = clCreateKernel(g_program, "calc_bin", &error);
    
    // Set arguments
    error = clSetKernelArg(kernel, 0, sizeof(cl_int), &bin_size);
    ocl_err(error);

    error = clSetKernelArg(kernel, 1, sizeof(cl_int), &num_bins);
    ocl_err(error);

    error = clSetKernelArg(kernel, 2, sizeof(cl_mem), &image_input);
    ocl_err(error);

    error = clSetKernelArg(kernel, 3, sizeof(cl_mem), &bin_output);
    ocl_err(error);
// ---
    
// --- Start kernel 1
    ocl_err(clEnqueueNDRangeKernel(g_command_queue, kernel, 1, NULL, global_work_sizes, NULL, 0, NULL, NULL));

// --- Init kernel 2
    
    

    // Compile kernel
    cl_kernel kernel2 = clCreateKernel(g_program, "parse_bin", &error);

    // Set arguments
    error = clSetKernelArg(kernel2, 0, sizeof(cl_int), &num_bins);
    ocl_err(error);

    error = clSetKernelArg(kernel2, 1, sizeof(cl_int), global_work_sizes);
    ocl_err(error);

    error = clSetKernelArg(kernel2, 2, sizeof(cl_mem), &bin_output);
    ocl_err(error);

    error = clSetKernelArg(kernel2, 3, sizeof(cl_mem), &cl_histogram);
    ocl_err(error);
// ---
    
// --- Finish kernel 1
    ocl_err(clFinish(g_command_queue));
// ---

// --- Start kernel 2
    ocl_err(clEnqueueNDRangeKernel(g_command_queue, kernel2, 1, NULL, global_work_sizes + 1, NULL, 0, NULL, NULL));

    ocl_err(clFinish(g_command_queue));
// --- Copy result in histogram
    ocl_err(clEnqueueReadBuffer(g_command_queue, cl_histogram, CL_TRUE, 0, host_result_size, histogram, 0, NULL, NULL));
}

int main(int argc, char *argv[])
{
    if (argc != 3)
    {
        usage(argv[0]);
        exit(-1);
    }
    const int num_bins = atoi(argv[1]);
    if (num_bins <= 0)
    {
        fprintf(stderr, "Invalid bin number.\n");
        usage(argv[0]);
        exit(-1);
    }
    const char * input_file = argv[2];

    printf("(*) Log:\n");

    cl_platform_id platform = ocl_select_platform();
    cl_device_id device = ocl_select_device(platform);
    init_ocl(device);
    create_program("kernel.cl", "");

    CvMat * input_image = cvLoadImageM(input_file, CV_LOAD_IMAGE_GRAYSCALE);
    int histogram[num_bins];
    memset(histogram, '\0', num_bins * sizeof(int));
    time_measure_start("calculate_histogram");
    calculate_histogram(input_image, histogram, num_bins);
    time_measure_stop_and_print("calculate_histogram");

    printf("(*) Output:\n");
    print_histogram(histogram, num_bins);

    return 0;
}

kernel.cl

#include "common_kernel.h"

__kernel void calc_bin(__const float bin_size, __const int num_bins, __global unsigned char *image_input, __global int *result){
    int work_index = get_global_id(0);

    int offsetinresult = work_index * num_bins;
    int i;
    int help = offsetinresult;

    for(i = 0; i < num_bins; i++, help++) result[help] = 0;

    for(i = 0, help = (work_index * COEFFICIENT); i < COEFFICIENT; i++, help++){
        result[offsetinresult + ( (int) (( (float) image_input[help]) / bin_size) )]++;
    }
    return;
}

__kernel void parse_bin(__const int num_bins, __const int workgroupsKernel, __global int *result, __global int *host_result){
	int work_index = get_global_id(0);
	int i = 0;
	int offset;

	host_result[work_index] = 0;

	for(; i < workgroupsKernel; i++){
		offset = i * num_bins;
		host_result[work_index] += result[offset + work_index];
	}

	return;
}