Sorry, both of those should have been Signed-off-by: Rebecca Palmer <rebecca_pal...@zoho.com>
As usual, I can only test on Ivy Bridge, so someone should probably check that they actually catch the no-__local-on-Haswell bug. On 16/05/15 18:48, 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 > > > > Reflect recent beignet and Linux changes. > > diff --git a/docs/Beignet.mdwn b/docs/Beignet.mdwn > index ec528b5..c0650bb 100644 > --- a/docs/Beignet.mdwn > +++ b/docs/Beignet.mdwn > @@ -142,7 +142,7 @@ Supported Targets > > * 3rd Generation Intel Core Processors > * Intel “Bay Trail” platforms with Intel HD Graphics > - * 4th Generation Intel Core Processors, need kernel patch currently, see > the "Known Issues" section. > + * 4th Generation Intel Core Processors "Haswell", need kernel patch > currently, see the "Known Issues" section. > * 5th Generation Intel Core Processors "Broadwell". > > Known Issues > @@ -163,22 +163,34 @@ Known Issues > But this command is a little bit dangerous, as if your kernel really hang, > then the gpu will lock up > forever until a reboot. > > -* Almost all unit tests fail. > - There is a known issue in some versions of linux kernel which enable > register whitelist feature > - but miss some necessary registers which are required for beignet. For > non-HSW platforms, the > - problematic version are around 3.15 and 3.16 which have commit f0a346b... > but haven't commit > - c9224f... If it is the case, you can apply c9224f... manually and rebuild > the kernel or just > - disable the parse command by invoke the following command (use Ubuntu as > an example): > +* "Beignet: self-test failed" and almost all unit tests fail. > + Linux 3.15 and 3.16 (commits > [f0a346b](https://git.kernel.org/cgit/linux/kernel/git/torvalds/linux.git/commit/?id=f0a346bdafaf6fc4a51df9ddf1548fd888f860d8) > + to > [c9224fa](https://git.kernel.org/cgit/linux/kernel/git/torvalds/linux.git/commit/?id=c9224faa59c3071ecfa2d4b24592f4eb61e57069)) > + enable the register whitelist by default but miss some registers needed > + for Beignet. > + > + This can be fixed by upgrading Linux, or by disabling the whitelist: > > `# echo 0 > /sys/module/i915/parameters/enable_cmd_parser` > > - For HSW platforms, this issue exists in all linux kernel version after > 3.15. We always need > - to execute the above command. > - > -* Some unit test cases, maybe 20 to 30, fail on 4th Generation (HSW) > platform. > - _The 4th Generation Intel Core Processors's support requires some Linux > kernel > - modification_. You need to apply the patch at: > - > [https://01.org/zh/beignet/downloads/linux-kernel-patch-hsw-support](https://01.org/zh/beignet/downloads/linux-kernel-patch-hsw-support) > + On Haswell hardware, Beignet 1.0.1 to 1.0.3 also required the > + above workaround on later Linux versions, but this _should not_ be > + required in current (after > [83f8739](http://cgit.freedesktop.org/beignet/commit/?id=83f8739b6fc4893fac60145326052ccb5cf653dc)) > + git master. > + > +* "Beignet: self-test failed" and 15-30 unit tests fail on 4th Generation > (Haswell) hardware. > + On Haswell, shared local memory (\_\_local) does not work at all on > + Linux <= 4.0, and requires the i915.enable_ppgtt=2 [boot > parameter](https://wiki.ubuntu.com/Kernel/KernelBootParameters) > + on Linux 4.1. > + > + This will be fixed in Linux 4.2; older versions can be fixed with > + [this > patch](https://01.org/zh/beignet/downloads/linux-kernel-patch-hsw-support). > + > + If you do not need \_\_local, you can override the self-test with > + > + `export OCL_IGNORE_SELF_TEST=1` > + > + but using \_\_local after this may silently give wrong results. > > * Precision issue. > Currently Gen does not provide native support of high precision math > functions > _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet