[Beignet] Beignet not working on Dell Precision M3800

Sat Jun 14 15:02:02 PDT 2014

On Sat, Jun 14, 2014 at 11:51 AM, Yichao Yu <yyc1992 at gmail.com> wrote:
> Sorry for the delay. I was busy graduating and didn't have much time
> in the past two weeks for testing.
>
> The gpu hang is still there and I haven't been able to make a c

Sorry I was wrong, it seems that the hang only happens before I
upgrade beignet. There are still a lot of failing tests but the screen
does not freeze anymore.

> version of the test program. However, I have found another problem
> with the newly merged opencl1.2 APIs when testing sth else.
>
> The c test program to trigger the issue is here[1]. When running on my
> Haswell CPU, beignet hangs in clWaitForEvents with the backtrace
>
> #0  0x00007ffff78cc9d0 in __nanosleep_nocancel () from
> /usr/lib/libc.so.6 #1  0x00007ffff78f6c94 in usleep () from
> /usr/lib/libc.so.6 #2  0x00007ffff73dfc8a in clWaitForEvents
> (num_events=1,      event_list=0x7fffffffda58)     at
> /home/yuyichao/projects/mlinux/pkg/all/beignet-git/src/beignet/src/cl_api.c:1316
> #3  0x00007ffff7bc861e in clWaitForEvents (num_events=1,
> event_list=0x7fffffffda58) at ocl_icd_loader.c:873 #4
> 0x00000000004009aa in main () at beignet-bug2.c:34
>
> It seems that the problem only happens for the event returned by
> clEnqueueBarrierWithWaitList when the wait list is not empty. I hope I
> am not using the api in the wrong way but I don't have another working
> opencl 1.2 implementation (pocl crashes on clEnqueueBarrier*...) to
> test it.........
>
> [1] https://gist.github.com/yuyichao/8b661d51c81f1c85466e
>
> On Wed, Jun 4, 2014 at 7:29 AM, Yichao Yu <yyc1992 at gmail.com> wrote:
>> On Tue, Jun 3, 2014 at 11:15 PM, Yang, Rong R <rong.r.yang at intel.com> wrote:
>>> Printf is not a built in function OpenCL 1.1, so beignet don't support it now. However, beignet are supporting it, maybe you could use it soon.
>>>
>>
>> However, even if the function is not defined, shouldn't the compiler
>> return a error (opencl error) rather than raising a exception and
>> abort?
>>
>>> Yes, the patch about 3D pipe have not push now, but You can apply by manual and try it.
>>
>> I'm afraid I don't have time to test it soon...
>>
>>>
>>> -----Original Message-----
>>> From: Yichao Yu [mailto:yyc1992 at gmail.com]
>>> Sent: Thursday, May 29, 2014 8:40 PM
>>> To: Yang, Rong R
>>> Cc: beignet at lists.freedesktop.org
>>> Subject: Re: [Beignet] Beignet not working on Dell Precision M3800
>>>
>>> On Thu, May 29, 2014 at 4:46 AM, Yang, Rong R <rong.r.yang at intel.com> wrote:
>>>> I have checked this issue, it is a beignet compiler bug, should be fix by patch "GBE: Change 64bit integer storage in register".
>>>>
>>>> For the first problem, I have sent some patch, can you try them? The patch " HSW: Restore L3 control register to disable SLM mode." fix a 3D pipe affect by Beignet bug. May be the same problem you met.
>>>
>>> I am testing using the current master
>>>
>>> c34eba71bd5a518906d6d5d3ba26e44327cab251
>>> GBE: fix one illegal instruction when replace a uniform dst.
>>>
>>> So the patch u mentioned for 3D pipe doesn't seem to be included yet.
>>>
>>> Here are what I saw,
>>> 1, `printf("%d\n", i);` works on pocl but still crashes the compiler on beignet with the same error.
>>> 2, the c example I gave works but the original python version does not... Will figure out the difference once I get more time.
>>> 3, the interference with opengl seems to be different. The same effect I mentioned last time shows up when sth is running on the GPU but recovers afterward. However, it now gives your email a funny texture by replacing some of the characters with another one...[1] (o in this
>>> case...) I also remember seeing this problem randomly sometime before but it was not as reproducible...
>>>
>>> I guess I will test again once those 3d pipe fixing patches are applied.
>>>
>>> [1] http://wstaw.org/m/2014/05/29/plasma-desktopzSP722.png
>>>
>>> Yichao Yu
>>>
>>>> -----Original Message-----
>>>> From: Yichao Yu [mailto:yyc1992 at gmail.com]
>>>> Sent: Wednesday, May 28, 2014 11:49 PM
>>>> To: Yang, Rong R
>>>> Cc: beignet at lists.freedesktop.org
>>>> Subject: Re: [Beignet] Beignet not working on Dell Precision M3800
>>>>
>>>> On Wed, May 28, 2014 at 11:45 AM, Yichao Yu <yyc1992 at gmail.com> wrote:
>>>>> On Wed, May 28, 2014 at 10:39 AM, Yichao Yu <yyc1992 at gmail.com> wrote:
>>>>>>> The second problem is that there seems to be sth wrong if I run two
>>>>>>> tests in series. More specifically, `test_elwise_kernel`[3],
>>>>>>> `test_elwise_kernel_with_option`[4] and
>>>>>>> `test_ranged_elwise_kernel`[5] can all pass if I run them
>>>>>>> individually. However, if I run them together, only the first one
>>>>>>> can pass... I will try to reproduce this in C...
>>>>>>
>>>>>> Sorry this is NOT what happened... I was not using the right
>>>>>> parameter to select the tests and there isn't any (at least no
>>>>>> evidence for it) interference between kernels.
>>>>>> The problem is rather the test_elsize_kernel_with_option and
>>>>>> test_ranged_elwise_kernel are not working..
>>>>>> Also the failing one sometimes (~2 times in 8) hang the wm for ~10s...
>>>>>> will try to make a c version....
>>>>>>
>>>>>
>>>>> And it seems that none of them is actually working, just that when
>>>>> the difference is calculated using OpenCL, it always returns 0...
>>>>>
>>>>> so here[1] is the c version. The problem seems to be related to the
>>>>> use of get_local_size and/or get_group_id in the kernel. When I was
>>>>> using a simple kernel with `int i = get_global_id(0);`, everything
>>>>> works fine.
>>>>
>>>> I haven't applied the patch for using local memory in the kernel. Does that patch affect not only local memory but also local size somehow?
>>>>
>>>>>
>>>>> [1] https://gist.github.com/yuyichao/242fd2a812088930af91
>>>>>
>>>>> P.S. I was trying to use printf in the kernel and it seems to crash
>>>>> the compiler..... Not sure if I was using it correctly but I guess it
>>>>> shouldn't crash in any case...
>>>>>
>>>>> here is the error:
>>>>> ```
>>>>> ASSERTION FAILED: it != instrinsicMap.map.end()   at file
>>>>> /home/yuyichao/projects/mlinux/pkg/all/beignet-git/src/beignet/backen
>>>>> d /src/llvm/llvm_gen_backend.cpp, function void
>>>>> gbe::GenWriter::regAllocateCallInst(llvm::CallInst&),
>>>>> line 2115 [1]    28951 trace trap (core dumped)  ./beignet-bug
>>>>> ```
>>>>>
>>>>> with the following kernel (not sure if it is valid haven't use printf
>>>>> before....),
>>>>>
>>>>> ```
>>>>> __kernel void fill_one(__global float *out, long n) {
>>>>>     int i = get_global_id(0);
>>>>>     printf("%d\n", i);
>>>>>     if (i < n) {
>>>>>         out[i] = 1;
>>>>>     }
>>>>> }
>>>>> ```
>>>>> (this kernel (without printf) works btw....)
>>>>>
>>>>> Yichao Yu
>>>>>
>>>>>>>
>>>>>>> [1] http://wstaw.org/m/2014/05/28/plasma-desktopObn722.png
>>>>>>> [2] http://wstaw.org/m/2014/05/28/plasma-desktopWbB722.png
>>>>>>> [3]
>>>>>>> https://github.com/pyopencl/pyopencl/blob/master/test/test_algorith
>>>>>>> m
>>>>>>> .py#L45 [4]
>>>>>>> https://github.com/pyopencl/pyopencl/blob/master/test/test_algorith
>>>>>>> m
>>>>>>> .py#L66 [5]
>>>>>>> https://github.com/pyopencl/pyopencl/blob/master/test/test_algorith
>>>>>>> m
>>>>>>> .py#L97
>>>>>>>>
>>>>>>>
>>>>>>> Yours,
>>>>>>> Yichao Yu
>>>>>>>
>>>>>>>>
>>>>>>>>>>>>>Thanks for point out it, I have sent a patch to correct it.
>>>>>>>>
>>>>>>>
>>>>>>> Seems fixed. THX. =)