[Beignet] Beignet not working on Dell Precision M3800

[Yichao Yu]

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.

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