<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=us-ascii">
<meta name="Generator" content="Microsoft Word 14 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
        {font-family:Calibri;
        panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
        {font-family:"Lucida Console";
        panose-1:2 11 6 9 4 5 4 2 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
        {margin:0in;
        margin-bottom:.0001pt;
        font-size:11.0pt;
        font-family:"Calibri","sans-serif";}
a:link, span.MsoHyperlink
        {mso-style-priority:99;
        color:blue;
        text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
        {mso-style-priority:99;
        color:purple;
        text-decoration:underline;}
span.EmailStyle17
        {mso-style-type:personal-compose;
        font-family:"Calibri","sans-serif";
        color:windowtext;}
.MsoChpDefault
        {mso-style-type:export-only;
        font-family:"Calibri","sans-serif";}
@page WordSection1
        {size:8.5in 11.0in;
        margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
        {page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal">I am not sure if this is the appropriate list on which to ask this question, if not hopefully someone can suggest an alternative.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">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.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Unfortunately, while clover is loading the binary, I end up getting a segmentation fault:<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal" style="margin-left:.5in;text-autospace:none"><span style="font-size:9.0pt;font-family:"Lucida Console"">Program received signal SIGSEGV, Segmentation fault.<o:p></o:p></span></p>
<p class="MsoNormal" style="margin-left:.5in;text-autospace:none"><span style="font-size:9.0pt;font-family:"Lucida Console"">proc (v=..., is=...) at core/module.cpp:50<o:p></o:p></span></p>
<p class="MsoNormal" style="margin-left:.5in;text-autospace:none"><span style="font-size:9.0pt;font-family:"Lucida Console"">50            T x;<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">I have pasted the source code I am using below, for both the kernel and the host code.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">I am compiling both with the following commands:<o:p></o:p></p>
<p class="MsoNormal">clang -target r600-unknown-unknown -x cl -S -emit-llvm -mcpu=r600 kernel.cl -o kernel.clbin
<o:p></o:p></p>
<p class="MsoNormal">clang -g -L/usr/local/lib -lOpenCL offline_host.c -o offline_host<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">I have LLVM/Clang 3.4RC3 installed and Mesa 10.0.1.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">If anyone has suggestions, or can point me to the appropriate mailing list or documentation,  I’d appreciate it.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Thanks!<o:p></o:p></p>
<p class="MsoNormal">-Al<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Source code for “kernel.cl”<o:p></o:p></p>
<p class="MsoNormal">====================<o:p></o:p></p>
<p class="MsoNormal">__kernel void vecAdd(__global float* a)<o:p></o:p></p>
<p class="MsoNormal">{<o:p></o:p></p>
<p class="MsoNormal">  int gid = get_global_id(0);<o:p></o:p></p>
<p class="MsoNormal">  a[gid] += a[gid];<o:p></o:p></p>
<p class="MsoNormal">}<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Source code for “offline_host.c”<o:p></o:p></p>
<p class="MsoNormal">==========================<o:p></o:p></p>
<p class="MsoNormal">#include <stdio.h><o:p></o:p></p>
<p class="MsoNormal">#include <stdlib.h><o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">#ifdef __APPLE__<o:p></o:p></p>
<p class="MsoNormal">#include <OpenCL/opencl.h><o:p></o:p></p>
<p class="MsoNormal">#else<o:p></o:p></p>
<p class="MsoNormal">#include <CL/cl.h><o:p></o:p></p>
<p class="MsoNormal">#endif<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">#define MEM_SIZE (128)<o:p></o:p></p>
<p class="MsoNormal">#define MAX_BINARY_SIZE (0x100000)<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">int main()<o:p></o:p></p>
<p class="MsoNormal">{<o:p></o:p></p>
<p class="MsoNormal">  cl_platform_id platform_id = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_device_id device_id = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_context context = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_command_queue command_queue = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_mem memobj = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_program program = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_kernel kernel = NULL;<o:p></o:p></p>
<p class="MsoNormal">  cl_uint ret_num_devices;<o:p></o:p></p>
<p class="MsoNormal">  cl_uint ret_num_platforms;<o:p></o:p></p>
<p class="MsoNormal">  cl_int ret;<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  float mem[MEM_SIZE];<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  FILE *fp;<o:p></o:p></p>
<p class="MsoNormal">  char fileName[] = "kernel.clbin";<o:p></o:p></p>
<p class="MsoNormal">  size_t binary_size;<o:p></o:p></p>
<p class="MsoNormal">  char *binary_buf;<o:p></o:p></p>
<p class="MsoNormal">  cl_int binary_status;<o:p></o:p></p>
<p class="MsoNormal">  cl_int i;<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Load kernel binary */<o:p></o:p></p>
<p class="MsoNormal">  fp = fopen(fileName, "r");<o:p></o:p></p>
<p class="MsoNormal">  if (!fp) {<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  }<o:p></o:p></p>
<p class="MsoNormal">  binary_buf = (char *)malloc(MAX_BINARY_SIZE);<o:p></o:p></p>
<p class="MsoNormal">  binary_size = fread(binary_buf, 1, MAX_BINARY_SIZE, fp);<o:p></o:p></p>
<p class="MsoNormal">  fclose(fp);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Initialize input data */<o:p></o:p></p>
<p class="MsoNormal">  for (i = 0; i < MEM_SIZE; i++) {<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  }<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Get platform/device information */<o:p></o:p></p>
<p class="MsoNormal">  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);<o:p></o:p></p>
<p class="MsoNormal">  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, &ret_num_devices);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Create OpenCL context*/<o:p></o:p></p>
<p class="MsoNormal">  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Create command queue */<o:p></o:p></p>
<p class="MsoNormal">  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Create memory buffer */<o:p></o:p></p>
<p class="MsoNormal">  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &ret);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Transfer data over to the memory buffer */<o:p></o:p></p>
<p class="MsoNormal">  ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Create kernel program from the kernel binary */<o:p></o:p></p>
<p class="MsoNormal">  program = clCreateProgramWithBinary(context, 1, &device_id, (const size_t *)&binary_size,<o:p></o:p></p>
<p class="MsoNormal">      (const unsigned char **)&binary_buf, &binary_status, &ret);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Create OpenCL kernel */<o:p></o:p></p>
<p class="MsoNormal">  kernel = clCreateKernel(program, "vecAdd", &ret);<o:p></o:p></p>
<p class="MsoNormal">  printf("err:%d\n", ret);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Set OpenCL kernel arguments */<o:p></o:p></p>
<p class="MsoNormal">  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  size_t global_work_size[3] = {MEM_SIZE, 0, 0};<o:p></o:p></p>
<p class="MsoNormal">  size_t local_work_size[3] = {MEM_SIZE, 0, 0};<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Execute OpenCL kernel */<o:p></o:p></p>
<p class="MsoNormal">  ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Copy result from the memory buffer */<o:p></o:p></p>
<p class="MsoNormal">  ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, MEM_SIZE * sizeof(float), mem, 0, NULL, NULL);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Display results */<o:p></o:p></p>
<p class="MsoNormal">  for (i=0; i < MEM_SIZE; i++) {<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  }<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  /* Finalization */<o:p></o:p></p>
<p class="MsoNormal">  ret = clFlush(command_queue);<o:p></o:p></p>
<p class="MsoNormal">  ret = clFinish(command_queue);<o:p></o:p></p>
<p class="MsoNormal">  ret = clReleaseKernel(kernel);<o:p></o:p></p>
<p class="MsoNormal">  ret = clReleaseProgram(program);<o:p></o:p></p>
<p class="MsoNormal">  ret = clReleaseMemObject(memobj);<o:p></o:p></p>
<p class="MsoNormal">  ret = clReleaseCommandQueue(command_queue);<o:p></o:p></p>
<p class="MsoNormal">  ret = clReleaseContext(context);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  free(binary_buf);<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">  return 0;<o:p></o:p></p>
<p class="MsoNormal">}<o:p></o:p></p>
<p class="MsoNormal"><i><span style="font-size:10.0pt"><o:p> </o:p></span></i></p>
</div>
</body>
</html>