[Mesa-dev] OpenCL Clang/Clover Offline Compilation issue

[Dorrington, Albert]

I am not sure if this is the appropriate list on which to ask this question, if not hopefully someone can suggest an alternative.

Under Linux, I am attempting to perform an offline compile of an OpenCL kernel example using Clang, and then load that binary using the clCreateProgramWithBinary() function.

Unfortunately, while clover is loading the binary, I end up getting a segmentation fault:

Program received signal SIGSEGV, Segmentation fault.
proc (v=..., is=...) at core/module.cpp:50
50            T x;

I have pasted the source code I am using below, for both the kernel and the host code.

I am compiling with the following commands:
clang -target r600-unknown-unknown -x cl -S -emit-llvm -mcpu=r600 kernel.cl -o kernel.clbin
clang -g -L/usr/local/lib -lOpenCL offline_host.c -o offline_host

I have LLVM/Clang 3.4RC3 installed and Mesa 10.0.1.

If anyone has suggestions, or can point me to the appropriate mailing list or documentation,  I'd appreciate it.

Thanks!
-Al


Source code for "kernel.cl"
====================
__kernel void vecAdd(__global float* a)
{
  int gid = get_global_id(0);
  a[gid] += a[gid];
}


Source code for "offline_host.c"
==========================
#include <stdio.h>
#include <stdlib.h>

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

#define MEM_SIZE (128)
#define MAX_BINARY_SIZE (0x100000)

int main()
{
  cl_platform_id platform_id = NULL;
  cl_device_id device_id = NULL;
  cl_context context = NULL;
  cl_command_queue command_queue = NULL;
  cl_mem memobj = NULL;
  cl_program program = NULL;
  cl_kernel kernel = NULL;
  cl_uint ret_num_devices;
  cl_uint ret_num_platforms;
  cl_int ret;

  float mem[MEM_SIZE];

  FILE *fp;
  char fileName[] = "kernel.clbin";
  size_t binary_size;
  char *binary_buf;
  cl_int binary_status;
  cl_int i;

  /* Load kernel binary */
  fp = fopen(fileName, "r");
  if (!fp) {
    fprintf(stderr, "Failed to load kernel.\n");
  exit(1);
  }
  binary_buf = (char *)malloc(MAX_BINARY_SIZE);
  binary_size = fread(binary_buf, 1, MAX_BINARY_SIZE, fp);
  fclose(fp);

  /* Initialize input data */
  for (i = 0; i < MEM_SIZE; i++) {
    mem[i] = i;
  }

  /* Get platform/device information */
  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);

  /* Create OpenCL context*/
  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

  /* Create command queue */
  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

  /* Create memory buffer */
  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &ret);

  /* Transfer data over to the memory buffer */
  ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);

  /* Create kernel program from the kernel binary */
  program = clCreateProgramWithBinary(context, 1, &device_id, (const size_t *)&binary_size,
      (const unsigned char **)&binary_buf, &binary_status, &ret);

  /* Create OpenCL kernel */
  kernel = clCreateKernel(program, "vecAdd", &ret);
  printf("err:%d\n", ret);

  /* Set OpenCL kernel arguments */
  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

  size_t global_work_size[3] = {MEM_SIZE, 0, 0};
  size_t local_work_size[3] = {MEM_SIZE, 0, 0};

  /* Execute OpenCL kernel */
  ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);

  /* Copy result from the memory buffer */
  ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);

  /* Display results */
  for (i=0; i < MEM_SIZE; i++) {
    printf("mem[%d] : $f\n", i, mem[i]);
  }

  /* Finalization */
  ret = clFlush(command_queue);
  ret = clFinish(command_queue);
  ret = clReleaseKernel(kernel);
  ret = clReleaseProgram(program);
  ret = clReleaseMemObject(memobj);
  ret = clReleaseCommandQueue(command_queue);
  ret = clReleaseContext(context);

  free(binary_buf);

  return 0;
}


Al Dorrington
Software Engineer Sr
Lockheed Martin, Mission Systems and Training
albert.dorrington at lmco.com<mailto:albert.dorrington at lmco.com> / 607-751-4859

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/mesa-dev/attachments/20140109/2fa37dda/attachment.html>