[Beignet] Beignet not working on Dell Precision M3800

[Yichao Yu]

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/backend
>> /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_algorithm
>>>> .py#L45 [4]
>>>> https://github.com/pyopencl/pyopencl/blob/master/test/test_algorithm
>>>> .py#L66 [5]
>>>> https://github.com/pyopencl/pyopencl/blob/master/test/test_algorithm
>>>> .py#L97
>>>>>
>>>>
>>>> Yours,
>>>> Yichao Yu
>>>>
>>>>>
>>>>>>>>>>Thanks for point out it, I have sent a patch to correct it.
>>>>>
>>>>
>>>> Seems fixed. THX. =)