Add a sanity test in clGetDeviceIDs

Run a small __local-using kernel in clGetDeviceIDs; if this returns
the wrong result, return CL_DEVICE_NOT_FOUND.

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.

Signed-off-by: Rebecca Palmer <rebecca_palmer@zoho.com>
Reviewed-by: "Luo, Xionghu" <xionghu.luo@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

2 files changed, 84 insertions, 0 deletions

diff --git a/src/cl_device_id.c b/src/cl_device_id.c
index 0fd4a690..0c1c52f5 100644
--- a/src/cl_device_id.c
+++ b/src/cl_device_id.c
@@ -436,6 +436,74 @@ brw_gt3_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,
@@ -447,6 +515,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 ac06b105..67e3bf1f 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