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);
}