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

[Tom Stellard]

On Thu, Jan 09, 2014 at 12:49:51PM +0000, Dorrington, Albert wrote:
> 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

I'm surprised that this works, since the r600 GPU does not support OpenCL
(Note that R600 is the name of the target and also one of the individual
GPUs supported by the compiler).  The  argument of -mcpu= needs to be
GPU you are compiling the code for.  So if you have a redwood GPU you
would need to pass -mcpu=redwood.

However, the main issue here is that clover does not support
clCreateProgramWithBinary() yet.  If you are interested in implementing
this, I can give you some pointers.  Just send an email to the list or
ping me on irc (nick: tstellar on #radeon @ irc.freednode.net).

-Tom

> 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
> 

> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/mesa-dev