Session 17: Introduction to Heterogeneous Programming Concepts

//
// saxpy.c
//

#include <fcntl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <unistd.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <OpenCL/opencl.h>

// Use a static data size for simplicity

#define DATA_SIZE (1024)
#define SEPARATOR       ("----------------------------------------------------------------------\n")

// Simple kernel for computing the saxpy operation with real single precision
// vectors

const char *KernelSource = "\n" \
"__kernel void saxpy(                                                   \n" \
"   const unsigned int n,                                               \n" \
"   const float a,                                                      \n" \
"   __global float* x,                                                  \n" \
"   __global float* y)                                                  \n" \
"{                                                                      \n" \
"   int i = get_global_id(0);                                           \n" \
"   if(i < n)                                                           \n" \
"       y[i] = a * x[i] + y[i];                                         \n" \
"}                                                                      \n" \
"\n";

// Prototype for device reporting function

cl_int report_and_mark_devices(cl_device_id *, cl_uint, int *, int *, int *);

int main(int argc, char** argv)
{
  cl_int err;                         // error code returned from api calls

  float x[DATA_SIZE];                 // original data vector x given to device
  float y[DATA_SIZE];                 // original data vector y given to device
  float a;
  float results[DATA_SIZE];           // results returned from device
  unsigned int correct;               // number of correct results returned
  cl_device_id device_id;             // compute device id
  int gpu = 1;

  if (argc == 1) {
    printf("%s -trying to use CL_DEVICE_GPU\n",argv[0]);
  } else if (argc == 2) {
    gpu = atoi(argv[1]);
    if (gpu != 0 && gpu != 1) {
      printf("Usage: %s 0 -use CL_DEVICE_CPU\n",argv[0]);
      printf("Usage: %s 1 -use CL_DEVICE_GPU or CL_DEVICE_ACCELERATOR\n",argv[0]);
      return EXIT_FAILURE;
    }
    printf("%s -trying to use %s\n",argv[0], gpu ? "CL_DEVICE_GPU or CL_DEVICE_ACCELERATOR" :
        "CL_DEVICE_CPU");
  }

  // Fill our data set with random float values

  int i = 0;
  a = 2.0;
  unsigned int count = DATA_SIZE;
  for(i = 0; i < count; i++) {
    x[i] = rand() / (float)RAND_MAX;
    y[i] = rand() / (float)RAND_MAX;
  }

  // Trying to identify one platform:

  cl_platform_id platform;
  cl_uint num_platforms;
  err = clGetPlatformIDs(1,&platform,&num_platforms);

  if (err != CL_SUCCESS) {
    printf("Error: Failed to get a platform id!\n");
    return EXIT_FAILURE;
  }

  // Found one platform. Query specific information about the found platform using the function
  //
  // Trying to query platform specific information...

  size_t returned_size = 0;
  cl_char platform_name[1024] = {0}, platform_prof[1024] = {0}, platform_vers[1024] = {0}, platform_exts[1024] = {0};
  err  = clGetPlatformInfo(platform, CL_PLATFORM_NAME,       sizeof(platform_name), platform_name, &returned_size);
  err |= clGetPlatformInfo(platform, CL_PLATFORM_VERSION,    sizeof(platform_vers), platform_vers, &returned_size);
  err |= clGetPlatformInfo(platform, CL_PLATFORM_PROFILE,    sizeof(platform_prof), platform_prof, &returned_size);
  err |= clGetPlatformInfo(platform, CL_PLATFORM_EXTENSIONS, sizeof(platform_exts), platform_exts, &returned_size);

  if (err != CL_SUCCESS) {
    printf("Error: Failed to get platform infor!\n");
    return EXIT_FAILURE;
  }

  printf("\nPlatform information\n");
  printf(SEPARATOR);
  printf("Platform name:       %s\n", (char *)platform_name);
  printf("Platform version:    %s\n", (char *)platform_vers);
  printf("Platform profile:    %s\n", (char *)platform_prof);
  printf("Platform extensions: %s\n", ((char)platform_exts[0] != '\0') ? (char *)platform_exts : "NONE");

  // Getting a list of available compute devices on our platform by function
  //
  // Get all available devices (up to 4)

  cl_uint num_devices;
  cl_device_id devices[4];
  err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 4, devices, &num_devices);

  if (err != CL_SUCCESS) {
    printf("Failed to collect device list on this platform!\n");
    return EXIT_FAILURE;
  }

  printf(SEPARATOR);
  printf("\nFound %d compute devices!:\n",num_devices);

  // Collect and report device information, return indices of devices of type CL_DEVICE_TYPE_CPU,
  // CL_DEVICE_TYPE_GPU and CL_DEVICE_TYPE_ACCELERATOR in array devices

  int a_cpu=-1, a_gpu=-1, an_accelerator=-1;
  err = report_and_mark_devices(devices,num_devices,&a_cpu,&a_gpu,&an_accelerator);

  if (err != CL_SUCCESS) {
    printf("Failed to report information about the devices on this platform!\n");
    return EXIT_FAILURE;
  }

  // Checking for availability of the required device

  if (gpu == 0) {     // No accelerator or gpu, just cpu
    if (a_cpu == -1) {
      printf("No cpus available, weird...\n");
      return EXIT_FAILURE;
    }
    device_id = devices[a_cpu];
    printf("There is a cpu, using it\n");
  }
  else {              // Trying to find a gpu, or if that fails, an accelerator
    if (a_gpu != -1) { // There is a gpu in our platform
      device_id = devices[a_gpu];
      printf("Found a gpu, using it\n");
    } else if (an_accelerator != -1) {
      device_id = devices[an_accelerator];
      printf("No gpu but found an accelerator, using it\n");
    } else {
      printf("No cpu, no gpu, nor an accelerator... where am I running???\n");
      return EXIT_FAILURE;
    }
  }

  // We have a compute device of required type! Next, create a compute context on it.

  printf("\n");
  printf(SEPARATOR);
  printf("\nCreating a compute context for the required device\n");

  cl_context context;                 // compute context
  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &err);

  if (!context) {
    printf("Error: Failed to create a compute context!\n");
    return EXIT_FAILURE;
  }

  // Create a command-queue on a specific device. The context must be a valid OpenCL context.

  printf("\n");
  printf(SEPARATOR);
  printf("\nCreating a command queue\n");

  cl_command_queue commands;
  commands = clCreateCommandQueue(context, device_id, 0, &err);

  if (!commands) {
    printf("Error: Failed to create a command queue!\n");
    return EXIT_FAILURE;
  }

  // Create the compute program object for our context and load the source code from the source buffer

  printf("\n");
  printf(SEPARATOR);
  printf("\nCreating the compute program from source\n");
  printf("%s\n", KernelSource);

  cl_program program;                 // compute program
  program = clCreateProgramWithSource(context, 1, (const char **) &KernelSource, NULL, &err);

  if (!program) {
    printf("Error: Failed to create compute program!\n");
    return EXIT_FAILURE;
  }

  // Build the program executable

  printf(SEPARATOR);
  printf("\nCompiling the program executable\n");

  err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

  if (err != CL_SUCCESS) {
    size_t len;
    char buffer[2048];
    printf("Error: Failed to build program executable!\n");

    // See page 98...
    clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
    printf("%s\n", buffer);
    exit(1);
  }

  // Create the compute kernel object in the program we wish to run

  printf("\n");
  printf(SEPARATOR);
  printf("\nCreating the compute kernel from program executable\n");

  cl_kernel kernel;                   // compute kernel
  kernel = clCreateKernel(program, "saxpy", &err);

  if (!kernel || err != CL_SUCCESS) {
    printf("Error: Failed to create compute kernel!\n");
    exit(1);
  }

  // Create buffer objects for the input and input/output arrays in device memory for our calculation

  printf("\n");
  printf(SEPARATOR);
  printf("\nCreating the input and input/output arrays in device memory\n");

  cl_mem dx;                       // device memory used for the input array x
  cl_mem dy;                       // device memory used for the input/output array y
  dx = clCreateBuffer(context, CL_MEM_READ_ONLY,  sizeof(float) * count, NULL, NULL);
  dy = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * count, NULL, NULL);

  if (!dx || !dy) {
    printf("Error: Failed to allocate device memory!\n");
    exit(1);
  }

  // Write our data set into the input buffer object in device memory

  printf("\n");
  printf(SEPARATOR);
  printf("\nWriting the data into the input and input/output arrays in device memory\n");

  err  = clEnqueueWriteBuffer(commands, dx, CL_TRUE, 0, sizeof(float) * count, x, 0, NULL, NULL);
  err |= clEnqueueWriteBuffer(commands, dy, CL_TRUE, 0, sizeof(float) * count, y, 0, NULL, NULL);

  if (err != CL_SUCCESS) {
    printf("Error: Failed to write to source array!\n");
    exit(1);
  }

  // Setting the arguments to our compute kernel in order to execute it.

  printf("\n");
  printf(SEPARATOR);
  printf("\nSetting the kernel arguments\n");

  err  = 0;
  err  = clSetKernelArg(kernel, 0, sizeof(unsigned int), &count);
  err |= clSetKernelArg(kernel, 1, sizeof(float),        &a);
  err |= clSetKernelArg(kernel, 2, sizeof(cl_mem),       &dx);
  err |= clSetKernelArg(kernel, 3, sizeof(cl_mem),       &dy);

  if (err != CL_SUCCESS) {
    printf("Error: Failed to set kernel arguments! %d\n", err);
    exit(1);
  }

  // Get the maximum kernel work group size for executing the kernel on the device

  size_t local;                       // local domain size for our calculation
  err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);

  if (err != CL_SUCCESS) {
    printf("Error: Failed to retrieve kernel work group info! %d\n", err);
    exit(1);
  }

  // Execute the kernel over the entire range of our logically 1d configuration
  // using the maximum kernel work group size

  size_t global;                      // global domain size for our calculation
  global = count;
  printf("\n");
  printf(SEPARATOR);
  printf("\nExecuting the kernel over 1d range %d.\n", (int)global);
  printf("Maximum work group size for this kernel is %d on this device\n", (int)local);

  err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);

  if (err) {
    printf("Error: Failed to execute kernel!\n");
    return EXIT_FAILURE;
  }

  // Wait for the command commands to get serviced before reading back results

  clFinish(commands);

  // Read back the results from the device to verify the output. In order to fully mimic
  // the saxpy operation, the result should be read into the array y, we read it to array results
  // in order to verify the computations.

  printf("\n");
  printf(SEPARATOR);
  printf("\nReading back the results from the device\n");

  err = clEnqueueReadBuffer( commands, dy, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL );

  if (err != CL_SUCCESS) {
    printf("Error: Failed to read output array! %d\n", err);
    exit(1);
  }

  // Validate our results

  correct = 0;
  for(i = 0; i < count; i++) {
    if(results[i] == a*x[i] + y[i])
      correct++;
  }

  // Print a brief summary detailing the results

  printf("\n");
  printf(SEPARATOR);
  printf("\nComputed '%d/%d' correct values!\n", correct, count);

  // Shutdown and cleanup

  clReleaseMemObject(dx);
  clReleaseMemObject(dy);
  clReleaseProgram(program);
  clReleaseKernel(kernel);
  clReleaseCommandQueue(commands);
  clReleaseContext(context);

  return 0;
}

cl_int report_and_mark_devices(cl_device_id *devices, cl_uint num_devices, int *a_cpu, int *a_gpu, int *an_accelerator)
{
  int i, type_name_index = 0;
  cl_int err = 0;
  size_t returned_size;
  size_t max_workgroup_size = 0;
  cl_uint max_compute_units = 0, vec_width_char = 0, vec_width_short = 0;
  cl_uint vec_width_int = 0, vec_width_long = 0, vec_width_float = 0, vec_width_double = 0;
  char vendor_name[1024] = {0}, device_name[1024] = {0}, device_version[1024] = {0};
  cl_ulong global_mem_size;
  cl_device_type device_type;
  char type_names[3][27]={"CL_DEVICE_TYPE_CPU        " , "CL_DEVICE_TYPE_GPU        " , "CL_DEVICE_TYPE_ACCELERATOR"};

  for (i=0;i<num_devices;i++) {
    err = clGetDeviceInfo(devices[i], CL_DEVICE_TYPE,                          sizeof(device_type),    &device_type,        &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_VENDOR,                        sizeof(vendor_name),    vendor_name,         &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_NAME,                          sizeof(device_name),    device_name,         &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_VERSION,                       sizeof(device_version), device_version,      &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE,           sizeof(size_t),         &max_workgroup_size, &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_MAX_COMPUTE_UNITS,             sizeof(cl_uint),        &max_compute_units,  &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_GLOBAL_MEM_SIZE,               sizeof(cl_ulong),       &global_mem_size,    &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR,   sizeof(cl_uint),        &vec_width_char,     &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT,  sizeof(cl_uint),        &vec_width_short,    &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT,    sizeof(cl_uint),        &vec_width_int,      &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG,   sizeof(cl_uint),        &vec_width_long,     &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT,  sizeof(cl_uint),        &vec_width_float,    &returned_size);
    err|= clGetDeviceInfo(devices[i], CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, sizeof(cl_uint),        &vec_width_double,   &returned_size);

    if (err != CL_SUCCESS) {
      printf("Error: Failed to retrieve device info!\n");
      return EXIT_FAILURE;
    }

    if (device_type == CL_DEVICE_TYPE_CPU) {
      *a_cpu = i;
      type_name_index = 0;
    }

    if (device_type == CL_DEVICE_TYPE_GPU) {
      *a_gpu = i;
      type_name_index = 1;
    }

    if (device_type == CL_DEVICE_TYPE_ACCELERATOR) {
      *an_accelerator = i;
      type_name_index = 2;
    }

    printf("\nDevice information:\n");
    printf(SEPARATOR);
    printf("Type:               %s\n", type_names[type_name_index]);
    printf("Vendor:             %s\n", vendor_name);
    printf("Device:             %s\n", device_name);
    printf("Version:            %s\n", device_version);
    printf("Max workgroup size: %d\n", (int)max_workgroup_size);
    printf("Max compute units:  %d\n", (int)max_compute_units);
    printf("Global mem size:    %ld\n", (long)global_mem_size);
    printf(SEPARATOR);

    printf("\nPreferred vector widths by type:\n");

    printf(SEPARATOR);
    printf("Vector char:  %d\n",   (int)vec_width_char);
    printf("Vector short: %d\n",   (int)vec_width_short);
    printf("Vector int:   %d\n",   (int)vec_width_int);
    printf("Vector long:  %d\n",   (int)vec_width_long);
    printf("Vector float: %d\n",   (int)vec_width_float);
    printf("Vector dble:  %d\n",   (int)vec_width_double);
    printf(SEPARATOR);
    printf("\n");
  }
  return err;
}