OpenCL

[acminor]

Using clinfo, I discovered that the RX5700XT supports OpenCL 2.1.  So I decided to try out support for SPIR-V IL loading. However, I have run into some issues. Can anyone help point out my error.

simple_add.cl

__kernel void add(__global float *a, __global float *b, __global float *output) {
  int id = get_global_id(0);

  output[id] = a[id] + b[id];
}

program.c

#define CL_TARGET_OPENCL_VERSION 210
#include <CL/cl.h>

#include <stdbool.h>
#include <string.h>
#include <stdio.h>

int read_source(const char *file_name, char **source_code, size_t *length) {
  FILE *file = fopen(file_name, "rb");
  fseek(file, -1, SEEK_END);
  long int offset = ftell(file);

  rewind(file);

  *source_code = malloc((offset+1)*sizeof(char));
  fread(*source_code, sizeof(char), offset, file);
  (*source_code)[offset] = '\0';

  *length = offset;

  fclose(file);

  return 0;
}

typedef struct {
  cl_context context;
  cl_command_queue command_queue;
  cl_program program;
  cl_kernel kernel;
  cl_int err;
  cl_uint num_of_platforms;
  cl_platform_id platform_id;
  cl_platform_id *platform_ids;
  cl_device_id device_id;
  cl_uint num_of_devices;
} OCL_Meta;

int main() {
  OCL_Meta meta = {0};

  clGetPlatformIDs(0, NULL, &meta.num_of_platforms);
  meta.platform_ids = malloc(meta.num_of_platforms * sizeof(cl_platform_id));
  if (clGetPlatformIDs(meta.num_of_platforms, meta.platform_ids, NULL) != CL_SUCCESS) {
    printf("Unable to get any platform_ids\n");
    return 1;
  }

  bool found_desired_platform = false;
  const char* desired_platform = "AMD Accelerated Parallel Processing";
  //const char* desired_platform = "Portable Computing Language";
  for (cl_uint i = 0; i < meta.num_of_platforms; i++) {
    cl_platform_id platform = meta.platform_ids[i];

    size_t string_len = 0;
    clGetPlatformInfo(platform, CL_PLATFORM_NAME, 0, NULL, &string_len);

    char *string = malloc(string_len + 1);
    string[string_len] = '\0';
    clGetPlatformInfo(platform, CL_PLATFORM_NAME, string_len+1, string, NULL);

    if (strcmp(desired_platform, string) == 0) {
      found_desired_platform = true;
      meta.platform_id = platform;
      free(string);
      break;
    } else {
      free(string);
    }
  }

  if (!found_desired_platform) {
    printf("%s device platform not found, exiting\n", desired_platform);
    return 1;
  } else {
    printf("%s device platform found and selected\n", desired_platform);
  }

  // NOTE assumes only one device per platform
  if (clGetDeviceIDs(meta.platform_id, CL_DEVICE_TYPE_ALL, 1, &meta.device_id, &meta.num_of_devices)
      != CL_SUCCESS) {
    printf("Unable to get device_id\n");
    return 1;
  }

  cl_context_properties properties[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties) meta.platform_id, 0 };
  meta.context = clCreateContext(properties, 1, &meta.device_id, NULL, NULL, &meta.err);
  meta.command_queue = clCreateCommandQueueWithProperties(meta.context, meta.device_id, NULL, &meta.err);

  char *kernel_il = NULL;
  size_t kernel_il_len = 0;
  read_source("simple_add.spv", &kernel_il, &kernel_il_len);
  meta.program = clCreateProgramWithIL(meta.context, kernel_il, kernel_il_len, &meta.err);

  if ((meta.err = clBuildProgram(meta.program, 0, NULL, NULL, NULL, NULL)) != CL_SUCCESS) {
    printf("Err: %i\n", meta.err);
    return 1;
  }

  meta.kernel = clCreateKernel(meta.program, "simple_add", &meta.err);

  const size_t work_size = 16;
  float a[work_size];
  float b[work_size];
  for(int i = 0; i < work_size; i++) {
    a[i] = i+1;
    b[i] = work_size - i;
  }

  cl_mem a_buf = clCreateBuffer(meta.context, CL_MEM_READ_ONLY,
                                sizeof(float)*work_size, NULL, NULL);
  cl_mem b_buf = clCreateBuffer(meta.context, CL_MEM_READ_ONLY,
                                sizeof(float)*work_size, NULL, NULL);
  cl_mem out_buf = clCreateBuffer(meta.context, CL_MEM_READ_ONLY,
                                  sizeof(float)*work_size, NULL, NULL);

  clEnqueueWriteBuffer(meta.command_queue, a_buf, CL_TRUE, 0,
                       sizeof(float)*work_size, a, 0, NULL, NULL);
  clEnqueueWriteBuffer(meta.command_queue, b_buf, CL_TRUE, 0,
                       sizeof(float)*work_size, b, 0, NULL, NULL);

  clSetKernelArg(meta.kernel, 0, sizeof(cl_mem), &a_buf);
  clSetKernelArg(meta.kernel, 1, sizeof(cl_mem), &b_buf);
  clSetKernelArg(meta.kernel, 2, sizeof(cl_mem), &out_buf);

  meta.err = clEnqueueNDRangeKernel(meta.command_queue, meta.kernel, 1, NULL,
                                    &work_size, NULL, 0, NULL, NULL);

  printf("err: %i\n", meta.err);

  clFlush(meta.command_queue);
  clFinish(meta.command_queue);

  float c[work_size];
  clEnqueueReadBuffer(meta.command_queue, out_buf, CL_TRUE, 0,
                      sizeof(float)*work_size, c, 0, NULL, NULL);

  printf("a: ");
  for (int i = 0; i < work_size; i++) {
    printf("%.2f ", a[i]);
  }
  printf("\n");

  printf("b: ");
  for (int i = 0; i < work_size; i++) {
    printf("%.2f ", b[i]);
  }
  printf("\n");

  printf("a+b: ");
  for (int i = 0; i < work_size; i++) {
    printf("%.2f ", c[i]);
  }
  printf("\n");

  return 0;
}

Build Commands

clang -cc1 -triple spir64-unknown-unknown -finclude-default-header -x cl -std=cl2.0 simple_add.cl -O0 -emit-llvm-bc -o simple_add.bc

llvm-spirv simple_add.bc -o simple_add.spv

gcc program.c -lOpenC

I keep gettting error code -48 when running the kernel with clEnqueqeNDRangeKernel which according to https://streamhpc.com/blog/2013-04-28/opencl-error-codes/ means that the kernel is invalid.

Is there any compiling options or program functions that I need to call or call differently to make this work. I tried searching the internet but found no solutions.

[dipak]

As I know, currently we don't have SPIR-V support in OpenCL.

Thanks.