example.c

Go to the documentation of this file.

/* -*- coding: latin-1 -*- */
/** \file examples/C/example.c
 * \brief
 * Simple C example to show how run a kernel with assert*() and PRINT*() macros
 * and test them.
 *
 * Piece of assertOpenCL
 * --- GPLv3 --- Copyright (C) 2018 Olivier Pirson
 * --- http://www.opimedia.be/
 * --- September 19, 2018
 */

// \cond
#include <assert.h>
#include <stdint.h>
#include <stdlib.h>
#include <stdbool.h>
#include <stdio.h>
#include <string.h>

#include <CL/cl.h>
// \endcond



/* ******
 * Type *
 ********/
typedef struct {
  int32_t code;
  const char* const name;
} error_name_t;



/* ***********
 * Constants *
 *************/

/** \brief
 * List of all error codes and names extracted from /usr/include/CL/cl.h
 */
const error_name_t errors_map[61] = {
#include "opencl_errors_map.h"
};


/** \brief
 * Number of elements in errors_map.
 */
const unsigned int errors_map_size = 61;



/* ************
 * Prototypes *
 **************/

/** \brief
 * Return the directory part of path in a (new allocated) string.
 */
char*
dirname_alloc(const char* path);


/** \brief
 * Return the error name corresponding to the error code.
 */
const char*
error_name(cl_int code);


/** \brief
 * Read the file and return its content to a (new allocated) string.
 * If failed then print a error message and exit.
 */
char*
file_to_alloc_string(const char* filename);


/** \brief
 * Return the given id device of the given platform OpenCL,
 * or exit if doesn't exists.
 *
 * @param platform_i
 * @param device_i
 *
 * @return OpenCL device id
 */
cl_device_id
get_device_id(unsigned int platform_i, unsigned int device_i);


/** \brief
 * Return a (new allocated) string corresponding to device info parameter.
 */
char*
get_device_info_alloc_string(cl_device_id device_id, cl_device_info param_name);


/** \brief
 * If code != 0
 * then print an error message corresponding to the error code.
 */
void
print_error(cl_int code, const char* message);


/** \brief
 * Run the kernel ../kernel/example.cl.

 * If debug
 * then run the kernel in debug mode,
 * else run the kernel with the macro NDEBUG defined.
 */
void
run_example(unsigned int nb_work_group, unsigned int nb_work_items_by_work_group,
            cl_device_id device_id, bool debug, const char* path);



/* ***********
 * Functions *
 *************/

char*
dirname_alloc(const char* path) {
  const char sep = '/';
  const size_t size = strlen(path);

  if ((size == 0) || (path[size - 1] == sep)) {  // path is empty or a directory name
    char* part = (char*)malloc(sizeof(char) * (size + 1));

    strcpy(part, path);

    return part;
  }
  else {                                         // path has the structure of a filename
    bool found_sep = false;
    size_t j = size;

    while (j > 0) {  // search the last sep
      --j;
      if (path[j] == sep) {
        found_sep = true;

        break;
      }
    }

    char* part = (char*)malloc(sizeof(char) * (j + 1));

    if (found_sep) {
      strncpy(part, path, j + 1);
      part[j + 1] = '\0';
    }
    else {
      part[0] = '\0';
    }

    return part;
  }
}


const char*
error_name(cl_int code) {
  for (unsigned int i = 0; i < errors_map_size; ++i) {
    if (errors_map[i].code == code) {
      return errors_map[i].name;
    }
  }

  return "unknow";
}


char*
file_to_alloc_string(const char* filename) {
  const char* const first_prefix = "#line 1 \"";
  const char* const first_suffix = "\"\n";
  const size_t first_size = strlen(first_prefix) + strlen(filename) + strlen(first_suffix);

  FILE* const fin = fopen(filename, "rb");

  fseek(fin, 0l, SEEK_END);

  const size_t file_size = ftell(fin);
  const size_t size = first_size + file_size;
  char* const s = malloc(size + 1);

  sprintf(s, "%s%s%s", first_prefix, filename, first_suffix);

  rewind(fin);
  if (fread((void*)(s + first_size), 1, size, fin) != file_size) {
    fprintf(stderr, "! Impossible to read file \"%s\"\n", filename);

    exit(EXIT_FAILURE);
  }
  fclose(fin);

  s[size] = '\0';

  return s;
}


cl_device_id
get_device_id(unsigned int platform_i, unsigned int device_i) {
  // Get number of platforms
  cl_uint nb;
  cl_int error = clGetPlatformIDs(0, NULL, &nb);

  print_error(error, "clGetPlatformIDs(..., &nb)");
  if (platform_i >= nb) {
      exit(EXIT_FAILURE);
  }

  // Get all platform ids
  cl_platform_id* platform_ids;

  platform_ids = (cl_platform_id*)malloc(sizeof(cl_platform_id) * nb);
  error = clGetPlatformIDs(nb, platform_ids, NULL);
  print_error(error, "clGetPlatformIDs(..., NULL)");

  // The wanted platform
  const cl_platform_id platform_id = platform_ids[platform_i];

  free(platform_ids);

  // Get number of devices for the wanted platform
  error = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, 0, NULL, &nb);
  print_error(error, "clGetDeviceIDs(..., &nb)");
  if (device_i >= nb) {
    exit(EXIT_FAILURE);
  }

  // Get all device ids for the wanted platform
  cl_device_id* device_ids;

  device_ids = (cl_device_id*)malloc(sizeof(cl_device_id) * nb);
  error = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, nb, device_ids, NULL);
  print_error(error, "clGetDeviceIDs(..., NULL)");

  // The wanted device
  const cl_device_id device_id = device_ids[device_i];

  free(device_ids);

  return device_id;
}


char*
get_device_info_alloc_string(cl_device_id device_id, cl_device_info param_name) {
  char* result;

  // Get length of result
  size_t size;
  cl_int error = clGetDeviceInfo(device_id, param_name, 0, NULL, &size);

  print_error(error, "clGetDeviceInfo(..., &size)");

  // Get the wanted string
  result = (char*)malloc(sizeof(char) * (size + 1));
  error = clGetDeviceInfo(device_id, param_name, size, result, NULL);
  print_error(error, "clGetDeviceInfo(..., NULL)");

  return result;
}


void
print_error(cl_int code, const char* message) {
  if (code != 0) {
    if (message == NULL) {
      fprintf(stderr, "! OpenCL error %i %s\n", code, error_name(code));
    }
    else {
      fprintf(stderr, "! OpenCL error %i %s %s\n", code, error_name(code), message);
    }
    fflush(stderr);
  }
}


void
run_example(unsigned int nb_work_group, unsigned int nb_work_items_by_work_group,
            cl_device_id device_id, bool debug, const char* path) {
  // Host buffer
  uint32_t h_outs[2];
  const size_t h_outs_byte_size = sizeof(h_outs[0])*2;

  uint64_t h_asserts[2] = {0, 0};
  const size_t h_asserts_byte_size = sizeof(h_asserts[0])*2;

  float h_assert_float[1] = {0};
  const size_t h_assert_float_byte_size = sizeof(h_assert_float[0]);

  assert(sizeof(float) == 4);


  // OpenCL context
  cl_int error;
  const cl_context context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &error);

  print_error(error, "clCreateContext");


  // OpenCL kernel
  const char* const prefix = "-I";
  const char* const suffix = "../../OpenCL/";
  const char* const suffix_ndebug = " -DNDEBUG";
  const size_t options_max_size =
    strlen(prefix) + strlen(path) + strlen(suffix) + strlen(suffix_ndebug);
  char* const options = malloc(sizeof(char) * (options_max_size + 1));

  sprintf(options, "%s%s%s", prefix, path, suffix);

  if (debug) {
    fputs("OpenCL in DEBUG mode!\n", stderr);
    fflush(stderr);
  }
  else {        // transmits NDEBUG macro to kernel
    strcat(options, suffix_ndebug);
  }

  const char* const suffix_filename = "../kernel/example.cl";
  char* const kernel_filename = malloc(sizeof(char)
                                       * (options_max_size + strlen(suffix_filename) + 1));

  sprintf(kernel_filename, "%s%s", path, suffix_filename);

  char* const kernel_src = file_to_alloc_string(kernel_filename);

  const cl_program program = clCreateProgramWithSource(context, 1, (const char**)&kernel_src, NULL, &error);

  print_error(error, "clCreateProgramWithSource");
  free(kernel_src);

  error = clBuildProgram(program, 1, &device_id, options, NULL, NULL);
  print_error(error, "clBuildProgram");
  free(options);

  const cl_kernel kernel = clCreateKernel(program, "example", &error);

  print_error(error, "clCreateKernel");


  // OpenCL queue
#if __OPENCL_C_VERSION__ < 200  // CL_VERSION_2_0
  cl_command_queue queue = clCreateCommandQueue(context, device_id, CL_QUEUE_PROFILING_ENABLE, &error);
#else
  const cl_queue_properties properties[] = {CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0};
  cl_command_queue queue = clCreateCommandQueueWithProperties(context, device_id, properties, &error);
#endif

  print_error(error, "clCreateCommandQueue");


  // OpenCL buffers
  cl_mem d_outs = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, h_outs_byte_size, h_outs, &error);

  print_error(error, "d_outs = clCreateBuffer");

  cl_mem d_asserts;
  cl_mem d_assert_float;

  if (debug) {
    d_asserts = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                               h_asserts_byte_size, h_asserts, &error);
    print_error(error, "d_asserts = clCreateBuffer");

    d_assert_float = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                                    h_assert_float_byte_size, h_assert_float, &error);
    print_error(error, "d_assert_float = clCreateBuffer");
  }


  // Params: just two parameters for the example
  const cl_uint n = 666;

  error = clSetKernelArg(kernel, 0, sizeof(n), &n);
  print_error(error, "clSetKernelArg(..., 0, ...)");

  error = clSetKernelArg(kernel, 1, sizeof(d_outs), &d_outs);
  print_error(error, "clSetKernelArg(..., 1, ...)");

  if (debug) {  // extra parameters to receive assertion information
    const unsigned int nb_args = 2;

    error = clSetKernelArg(kernel, nb_args, sizeof(d_asserts), &d_asserts);
    print_error(error, "clSetKernelArg(..., d_asserts)");

    error = clSetKernelArg(kernel, nb_args + 1, sizeof(d_assert_float), &d_assert_float);
    print_error(error, "clSetKernelArg(..., d_assert_float)");
  }


  // Run
  const size_t global_size = nb_work_items_by_work_group * nb_work_group;
  const size_t local_size = nb_work_items_by_work_group;

  puts("===== run kernel =====");
  fflush(stdout);

  error = clEnqueueNDRangeKernel(queue,
                                 kernel,
                                 1,
                                 NULL,
                                 &global_size,
                                 &local_size,
                                 0,
                                 NULL,
                                 NULL);
  print_error(error, "clEnqueueNDRangeKernel");

  clFinish(queue);
  clFlush(queue);
  fflush(stdout);
  fflush(stderr);
  puts("===== end kernel =====");
  fflush(stdout);


  // Results
  error = clEnqueueReadBuffer(queue, d_outs, CL_TRUE, 0,
                              h_outs_byte_size, h_outs, 0, NULL, NULL);
  print_error(error, "clEnqueueReadBuffer(..., d_outs, ...)");

  if (debug) {
    error = clEnqueueReadBuffer(queue, d_asserts, CL_TRUE, 0,
                                h_asserts_byte_size, h_asserts, 0, NULL, NULL);
    print_error(error, "clEnqueueReadBuffer(..., d_asserts, ...)");

    error = clEnqueueReadBuffer(queue, d_assert_float, CL_TRUE, 0,
                                h_assert_float_byte_size, h_assert_float, 0, NULL, NULL);
    print_error(error, "clEnqueueReadBuffer(..., d_assert_float, ...)");

    const uint64_t line = h_asserts[0];

    if (line != 0) {  // there had (at least) an assertion
      const uint64_t uint64_value = h_asserts[1];
      const int64_t sint64_value = (int64_t)h_asserts[1];
      const float float_value = h_assert_float[0];

      fprintf(stderr, "%s:%lu\tAssertion failed | Maybe\t%lu\t%li | Maybe\t%f\n",
              kernel_filename, line, uint64_value, sint64_value, float_value);
      /*
        Maybe incoherent assert information because the parallel execution of work items.
        But each element of these information concern assertion(s) that failed.
      */
      fflush(stderr);
    }
  }

  printf("Results: (%u, %u)\n", h_outs[0], h_outs[1]);

  free(kernel_filename);


  // Free OpenCL resources
  clReleaseMemObject(d_outs);
  if (debug) {
    clReleaseMemObject(d_asserts);
    clReleaseMemObject(d_assert_float);
  }

  clReleaseCommandQueue(queue);
  clReleaseProgram(program);
  clReleaseKernel(kernel);
  clReleaseContext(context);
}



/* ******
 * Main *
 ********/

/** \brief
 * Get the optional parameter --device platform:device
 * and run the kernel ../kernel/example.cl
 */
int
main(int argc, const char* const argv[]) {
  bool debug =
#ifdef NDEBUG
    false;
#else
    true;
#endif
  signed int device_i = 0;
  signed int platform_i = 0;

  // Read parameters
  {
    int i = 1;

    while (i < argc) {
      const char* arg = argv[i];

      if ((strcmp(arg, "--debug") == 0) || (strcmp(arg, "--ndebug") == 0)) {
        debug = (strcmp(arg, "--debug") == 0);
      }
      else if (strcmp(arg, "--device") == 0) {
        ++i;
        if (i >= argc) {
          return EXIT_FAILURE;
        }

        const char* both_i = argv[i];
        const size_t len = strlen(both_i);
        bool only_platform = true;
        size_t j = 0;

        while (j < len) {
          if (both_i[j] == ':') {
            only_platform = false;

            break;
          }
          ++j;
        }

        if (only_platform) {
          platform_i = atoi(both_i);
        }
        else {
          platform_i = atoi(both_i);
          device_i = atoi(both_i + j + 1);
        }

        if ((platform_i < 0) || (device_i < 0)) {
          return EXIT_FAILURE;
        }
      }

      ++i;
    }
  }

  // Get wanted device
  const cl_device_id device_id = get_device_id(platform_i, device_i);

  // Get device name and print
  {
    char* device_name = get_device_info_alloc_string(device_id, CL_DEVICE_NAME);

    printf("Device %u:%u %s\n", (unsigned int)platform_i, (unsigned int)device_i, device_name);
    free(device_name);
    fflush(stdout);
  }

  // Get current path
  char* path = dirname_alloc(argv[0]);

  // Run
  run_example(3, 4, device_id, debug, path);

  free(path);

  return EXIT_SUCCESS;
}