[beignet] 01/06: Do a self-test instead of unconditionally disabling __local on Haswell
Rebecca Palmer
rnpalmer-guest at moszumanska.debian.org
Tue May 19 10:12:46 UTC 2015
This is an automated email from the git hooks/post-receive script.
rnpalmer-guest pushed a commit to branch master
in repository beignet.
commit 61f18d3510489b8d9c14582df8b508d044856e5b
Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
Date: Tue May 19 08:46:39 2015 +0100
Do a self-test instead of unconditionally disabling __local on Haswell
This allows it to work if a fixed Linux is installed
---
debian/changelog | 2 +-
debian/patches/haswell.patch | 33 ------------
debian/patches/self-test.patch | 116 +++++++++++++++++++++++++++++++++++++++++
debian/patches/series | 2 +-
4 files changed, 118 insertions(+), 35 deletions(-)
diff --git a/debian/changelog b/debian/changelog
index 2a701c6..d6e73c6 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -11,7 +11,7 @@ beignet (1.0.3-1) UNRELEASED; urgency=medium
* Drop Link-against-terminfo.patch, llvm-config --system-libs and
related dependencies, as they are not needed when dynamically
linking libllvm.
- * Make __local on Haswell an error instead of silently doing nothing,
+ * Disable devices where __local silently does nothing,
and document workarounds.
* Make libx11-dev dependency explicit.
diff --git a/debian/patches/haswell.patch b/debian/patches/haswell.patch
deleted file mode 100644
index 5e2ea0d..0000000
--- a/debian/patches/haswell.patch
+++ /dev/null
@@ -1,33 +0,0 @@
-Description: Better handle known Haswell bugs
-
-Print a helpful error message instead of silently doing nothing
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-Forwarded: http://lists.freedesktop.org/archives/beignet/2015-April/005589.html
-
---- beignet-1.0.3.orig/src/cl_command_queue_gen7.c
-+++ beignet-1.0.3/src/cl_command_queue_gen7.c
-@@ -343,6 +343,10 @@ cl_command_queue_ND_range_gen7(cl_comman
- /* Curbe step 1: fill the constant urb buffer data shared by all threads */
- if (ker->curbe) {
- kernel.slm_sz = cl_curbe_fill(ker, work_dim, global_wk_off, global_wk_sz, local_wk_sz, thread_n);
-+ if (kernel.slm_sz > 0 && cl_driver_get_ver(ctx->drv) == 75){
-+ fprintf(stderr, "Beignet: Shared local memory does not work on Haswell, see /usr/share/doc/beignet-opencl-icd/README.Debian\n");
-+ return CL_OUT_OF_RESOURCES;
-+ }
- if (kernel.slm_sz > ker->program->ctx->device->local_mem_size) {
- fprintf(stderr, "Beignet: Out of shared local memory %d.\n", kernel.slm_sz);
- return CL_OUT_OF_RESOURCES;
---- beignet-1.0.3.orig/src/intel/intel_batchbuffer.c
-+++ beignet-1.0.3/src/intel/intel_batchbuffer.c
-@@ -135,6 +135,10 @@ intel_batchbuffer_flush(intel_batchbuffe
- }
- if (drm_intel_gem_bo_context_exec(batch->buffer, batch->intel->ctx, used, flag) < 0) {
- fprintf(stderr, "drm_intel_gem_bo_context_exec() failed: %s\n", strerror(errno));
-+ if (errno == EINVAL && IS_GEN75(batch->intel->device_id)) {
-+ fprintf(stderr, "This is a known bug on Haswell systems, see /usr/share/doc/beignet-opencl-icd/README.Debian\n"
-+ "'sudo echo 0 > /sys/module/i915/parameters/enable_cmd_parser' usually helps\n");
-+ }
- err = -1;
- }
-
diff --git a/debian/patches/self-test.patch b/debian/patches/self-test.patch
new file mode 100644
index 0000000..b6ce93f
--- /dev/null
+++ b/debian/patches/self-test.patch
@@ -0,0 +1,116 @@
+Description: Test device in clGetDeviceIDs
+
+Run a small kernel to check that the device works (including the
+__local memory space, known to be problematic on Haswell), and
+hide non-working devices
+
+Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
+Forwarded: http://cgit.freedesktop.org/beignet/commit/?id=dfcc554d3bbc3bcc0810b3da5ee26b4345e6b4f7
+
+--- beignet-1.0.3.orig/src/cl_device_id.c
++++ beignet-1.0.3/src/cl_device_id.c
+@@ -458,6 +458,72 @@ 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"
++ "This is a known bug on Haswell systems, see /usr/share/doc/beignet-opencl-icd/README.Debian\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\nSee /usr/share/doc/beignet-opencl-icd/README.Debian\n", status);
++ }
++ return ret;
++}
++
+ LOCAL cl_int
+ cl_get_device_ids(cl_platform_id platform,
+ cl_device_type device_type,
+@@ -469,6 +535,20 @@ cl_get_device_ids(cl_platform_id plat
+
+ /* 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;
+--- beignet-1.0.3.orig/utests/setenv.sh.in
++++ beignet-1.0.3/utests/setenv.sh.in
+@@ -6,6 +6,8 @@ export OCL_PCH_PATH=@LOCAL_OCL_PCH_OBJEC
+ 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
+ obj-`dpkg-architecture -qDEB_HOST_MULTIARCH`/utests/utest_run -a
+ export OCL_STRICT_CONFORMANCE=0
+ obj-`dpkg-architecture -qDEB_HOST_MULTIARCH`/utests/utest_run -a
diff --git a/debian/patches/series b/debian/patches/series
index b368e53..9275b5b 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -7,4 +7,4 @@ shared-llvm.patch
builtin_pow-fix-spurious-failure.patch
tgamma-accuracy.patch
python3.patch
-haswell.patch
+self-test.patch
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/beignet.git
More information about the Pkg-opencl-commits
mailing list