[Beignet] [PATCH 1/2] Add a sanity test in clGetDeviceIDs

[Zhigang Gong]

Thanks for the patches. The whole patch set LGTM. I will push them latter
Actually xionghu is doing very similar thing and will submit patch to record the self-test
result some where in the system such as $HOME/.config/beignet/beignetrc with the kernel
tag. Then we can avoid duplciate testing on the same system with the same kernel.

On Sat, May 16, 2015 at 06:48:37PM +0100, Rebecca N. Palmer wrote:
> Run a small __local-using kernel in clGetDeviceIDs; if this returns
> the wrong result, return CL_DEVICE_NOT_FOUND.
> ---
> 
> > just check kernel version is not
> > an ideal method for those unofficial kernels with back porting patches. Then we have the
> > following open questions in my mind:
> > 
> >   How do we check whether the i915 KMD support secure batch buffer execution if the batch
> >   buffer pass the cmd parser check under full-ppgtt mode in UMD?
> > 
> >   How do we check whether the i915 KMD support secure batch buffer execution with aliasing
> >   ppgtt after the merging of the patch "drm/i915: Arm cmd parser with aliasing ppgtt only" in UMD?
> 
> As far as I can see, there's no way to tell in advance (except
> unreliably with a global version check) whether __local-using batches
> will be accepted...so the easiest solution is probably to just try
> running one and see what result we get.
> 
> diff --git a/src/cl_device_id.c b/src/cl_device_id.c
> index 6aa6b3b..218b7a5 100644
> --- a/src/cl_device_id.c
> +++ b/src/cl_device_id.c
> @@ -545,6 +545,74 @@ skl_gt4_break:
>    return ret;
>  }
>  
> +/* Runs a small kernel to check that the device works; returns
> + * 0 for success, 1 for silently wrong result, 2 for error */
> +LOCAL cl_int
> +cl_self_test(cl_device_id device)
> +{
> +  cl_int status, ret;
> +  cl_context ctx;
> +  cl_command_queue queue;
> +  cl_program program;
> +  cl_kernel kernel;
> +  cl_mem buffer;
> +  cl_event kernel_finished;
> +  size_t n = 3;
> +  cl_int test_data[3] = {3, 7, 5};
> +  const char* kernel_source = "__kernel void self_test(__global int *buf) {"
> +  "  __local int tmp[3];"
> +  "  tmp[get_local_id(0)] = buf[get_local_id(0)];"
> +  "  barrier(CLK_LOCAL_MEM_FENCE);"
> +  "  buf[get_global_id(0)] = tmp[2 - get_local_id(0)] + buf[get_global_id(0)];"
> +  "}"; // using __local to catch the "no SLM on Haswell" problem
> +  ret = 2;
> +  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);
> +        if (status == CL_SUCCESS) {
> +          kernel = clCreateKernel(program, "self_test", &status);
> +          if (status == CL_SUCCESS) {
> +            buffer = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR, n*4, test_data, &status);
> +            if (status == CL_SUCCESS) {
> +              status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
> +              if (status == CL_SUCCESS) {
> +                status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &n, &n, 0, NULL, &kernel_finished);
> +                if (status == CL_SUCCESS) {
> +                  status = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, n*4, test_data, 1, &kernel_finished, NULL);
> +                  if (status == CL_SUCCESS) {
> +                    if (test_data[0] == 8 && test_data[1] == 14 && test_data[2] == 8){
> +                      ret = 0;
> +                    } else {
> +                      ret = 1;
> +                      printf("Beignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (%i, %i, %i)\n"
> +                      "See README.md or http://www.freedesktop.org/wiki/Software/Beignet/\n",
> +                      test_data[0], test_data[1], test_data[2]);
> +                    }
> +                  }
> +                }
> +              }
> +            }
> +            clReleaseMemObject(buffer);
> +          }
> +          clReleaseKernel(kernel);
> +        }
> +      }
> +      clReleaseProgram(program);
> +    }
> +    clReleaseCommandQueue(queue);
> +  }
> +  clReleaseContext(ctx);
> +  if (ret == 2) {
> +    printf("Beignet: self-test failed: error %i\n"
> +    "See README.md or http://www.freedesktop.org/wiki/Software/Beignet/\n", status);
> +  }
> +  return ret;
> +}
> +
>  LOCAL cl_int
>  cl_get_device_ids(cl_platform_id    platform,
>                    cl_device_type    device_type,
> @@ -556,6 +624,20 @@ cl_get_device_ids(cl_platform_id    platform,
>  
>    /* Do we have a usable device? */
>    device = cl_get_gt_device();
> +  if (device && cl_self_test(device)) {
> +    int disable_self_test = 0;
> +    // can't use BVAR (backend/src/sys/cvar.hpp) here as it's C++
> +    const char *env = getenv("OCL_IGNORE_SELF_TEST");
> +    if (env != NULL) {
> +      sscanf(env, "%i", &disable_self_test);
> +    }
> +    if (disable_self_test) {
> +      printf("Beignet: Warning - overriding self-test failure\n");
> +    } else {
> +      printf("Beignet: disabling non-working device\n");
> +      device = 0;
> +    }
> +  }
>    if (!device) {
>      if (num_devices)
>        *num_devices = 0;
> diff --git a/utests/setenv.sh.in b/utests/setenv.sh.in
> index ac06b10..67e3bf1 100644
> --- a/utests/setenv.sh.in
> +++ b/utests/setenv.sh.in
> @@ -6,3 +6,5 @@ export OCL_PCH_PATH=@LOCAL_OCL_PCH_OBJECT@
>  export OCL_KERNEL_PATH=@CMAKE_CURRENT_SOURCE_DIR@/../kernels
>  export OCL_GBE_PATH=@LOCAL_GBE_OBJECT_DIR@
>  export OCL_INTERP_PATH=@LOCAL_INTERP_OBJECT_DIR@
> +#disable self-test so we can get something more precise than "doesn't work"
> +export OCL_IGNORE_SELF_TEST=1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet