[Beignet] Beignet not working on Dell Precision M3800

[Yichao Yu]

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