Bug#848258: mesa-opencl-icd/libclc: ! does not accept pointers

["Rebecca N. Palmer" <rebecca_palmer@zoho.com>]

Control: retitle -1 mesa-opencl-icd/libclc: ! does not accept pointers

There appear to be two separate bugs here:

-get_global_id() does not work (the clinfo problem).  This has also been 
reported as #857591 and upstream, where it has a patch: 
https://bugs.freedesktop.org/attachment.cgi?id=130136

Please use #857591 for further discussion of this issue.

-the ! operator does not accept pointer-to-struct, and possibly any 
pointers (the 3 errors, as opposed to warnings, in the blender compile 
log: 
sources.debian.net/src/blender/2.78.a%2Bdfsg0-4/intern/cycles/kernel/closure/bsdf_microfacet.h/#L269 
http://sources.debian.net/src/blender/2.78.a%2Bdfsg0-4/intern/cycles/kernel/closure/alloc.h/#L65

I attach a simple test case for that (works in beignet, my hardware 
doesn't support mesa-opencl-icd).  Please try both with and without the 
#define NULL 0 - blender has that, but possibly shouldn't.
/*
apt-get install ocl-icd-opencl-dev opencl-icd
gcc opencl_pointertostruct_test.c -lOpenCL -g
gdb ./a.out
valgrind --track-origins=yes ./a.out

expected output:
built 0 built 0 kernel created 00 returned (9, 5, 8)
1 returned (2, 4, 6)
2 returned (5, 1, 2)
*/
#include <stdio.h>
#include <CL/cl.h>
int main(int argc,char** argv)
{
  cl_platform_id platforms[4];
  cl_uint num_platforms;
  cl_device_id device;
  cl_int status, ret;
  cl_context ctx;
  cl_command_queue queue;
  cl_program program,program2;
  cl_kernel kernel;
  cl_mem buffer[3];
  cl_event kernel_finished;
  FILE *f;
  size_t n = 3,program_length;
  int i;
  cl_int test_data[3][3] = {{3, 8, 5},{2, 4, 6},{5, 1, 2}};
  const char *kernel_source="#define NULL 0\ntypedef struct Fcomplex {float real;float imag;}Fcomplex;__kernel void pointer_nullcheck(__global int *p0, __global int *p1, __global int *p2){__private Fcomplex a[3];__private Fcomplex * __private ap;int i;i=get_local_id(0);if(p2[i]>2){ap=NULL;}else{ap=&a[p2[i]];}if(!ap){p0[i]=9;}else{ap->real=p1[i];ap->imag=p2[i];p0[i]=ap->real+ap->imag;}}";
  char *build_log;
  ret = 2;
  build_log=calloc(1001,1);
  clGetPlatformIDs(4,platforms,&num_platforms);
  for(i=0;i<num_platforms;i++){
  status=clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_GPU,1,&device,NULL);
  if (status == CL_SUCCESS) {
  break;
  }
  }
  if (status != CL_SUCCESS) {
  printf("Device not found");
  exit(1);
  }
  ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &status);
  if (status == CL_SUCCESS) {
    queue = clCreateCommandQueue(ctx, device, 0, &status);
    if (status == CL_SUCCESS) {
      program = clCreateProgramWithSource(ctx, 1, &kernel_source, NULL, &status);
      if (status == CL_SUCCESS) {
        //status = clBuildProgram(program, 1, &device, "", NULL, NULL);
        status = clCompileProgram(program, 1, &device, "", 0, NULL, NULL, NULL, NULL);
        clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
        printf("built %i %s",status,build_log);
        program2 = clLinkProgram(ctx, 1, &device, "", 1,&program, NULL, NULL, NULL);
        clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
        printf("built %i %s",status,build_log);
        if (status == CL_SUCCESS) {
          kernel = clCreateKernel(program2, "pointer_nullcheck", &status);
          printf("kernel created %i",status);
          if (status == CL_SUCCESS) {
            for(i=0;i<3;i++){
            buffer[i] = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR, 4*n, test_data[i], &status);
            }
            if (status == CL_SUCCESS) {
            for(i=0;i<3;i++){
              status = clSetKernelArg(kernel, i, sizeof(cl_mem), &buffer[i]);
              }
              if (status == CL_SUCCESS) {
                status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &n, &n, 0, NULL, &kernel_finished);
                if (status == CL_SUCCESS) {
                for(i=0;i<3;i++){
                  status = clEnqueueReadBuffer(queue, buffer[i], CL_TRUE, 0, n*4, test_data[i], 1, &kernel_finished, NULL);
                  }
                  if (status == CL_SUCCESS) {
                  for(i=0;i<3;i++){
                    printf("%i returned (%i, %i, %i)\n", i,test_data[i][0], test_data[i][1], test_data[i][2]);
                    }
                  }
                }
              }
            }
            for(i=0;i<3;i++){
            clReleaseMemObject(buffer[i]);
            }
          }
          clReleaseKernel(kernel);
        }
      }
      clReleaseProgram(program);
    }
    clReleaseCommandQueue(queue);
  }
  clReleaseContext(ctx);
}