Thanks for the test. And as the patchset works fine, I will push it firstly. And you may consider to merge your local work based on this patchset to reduce the self-test overhead.
Thanks, Zhigang Gong. On Mon, May 18, 2015 at 06:01:06AM +0000, Luo, Xionghu wrote: > I've tested this patchset on my Haswell, it works now, while some > improvements could be made based on this: First, kernel version could be > checked to avoid unnecessary cl_self_test after the release of 4.2; Second, > the cl_self_test result could be written to file to avoid repeated testing; > Finally, should we ask the user to input a 'y/N' to continue since the > warning is not obvious enough when OCL_IGNOR_SELF_TEST=1? > > Luo Xionghu > Best Regards > > -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Rebecca N. Palmer > Sent: Sunday, May 17, 2015 2:00 AM > To: [email protected] > Subject: Re: [Beignet] [PATCH 1/2] Add a sanity test in clGetDeviceIDs + > [PATCH 2/2] Docs: update/clarify Haswell issues > > Sorry, both of those should have been > > Signed-off-by: Rebecca Palmer <[email protected]> > > 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](h > > ttps://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=83f8739b6fc > > + 4893fac60145326052ccb5cf653dc)) > > + 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 > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
