[Beignet] Beignet not working on Dell Precision M3800

[Yang, Rong R]

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.
-----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. =)