[Beignet] [PATCH 1/2] Add a sanity test in clGetDeviceIDs + [PATCH 2/2] Docs: update/clarify Haswell issues
Luo, Xionghu
xionghu.luo at intel.com
Sun May 17 23:27:40 PDT 2015
Zhigang just notified me that no need to require user input since beignet is only a library for UI apps.
Luo Xionghu
Best Regards
-----Original Message-----
From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of Luo, Xionghu
Sent: Monday, May 18, 2015 2:01 PM
To: 'Rebecca N. Palmer'
Cc: 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
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