[Beignet] [PATCH 1/2] Add a sanity test in clGetDeviceIDs + [PATCH 2/2] Docs: update/clarify Haswell issues
Zhigang Gong
zhigang.gong at linux.intel.com
Sun May 17 22:53:27 PDT 2015
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:beignet-bounces at lists.freedesktop.org] On Behalf Of Rebecca N. Palmer
> Sent: Sunday, May 17, 2015 2:00 AM
> To: beignet at lists.freedesktop.org
> 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 <rebecca_palmer at 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](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
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list