Pekka,
You are of course right (and I am an idiot!).
I used to launch VexCL kernels with a single thread in a workgroup for CPU
devices, and this particular kernel used this information. Some time ago I
moved away from single-thread-on-CPU policy, but this kernel was never
updated.
The bug is fixed now:
https://github.com/ddemidov/vexcl/commit/b6983c4a33d96bf9fce45961c960af2e8e702320
.
POCL now passes all of VexCL tests (except for Boost.compute ones).
Thanks for the help, and sorry for the trouble,
Denis.
On Wed, May 22, 2013 at 9:14 PM, Pekka Jääskeläinen <
[email protected]> wrote:
> Denis,
>
> Am I getting too tired or doesn't this kernel have a race condition or
> a synchronization error?
>
> kernel void convolve(
> ulong n,
> char has_left,
> char has_right,
> int lhalo, int rhalo,
> global const real *xloc,
> global const real *xrem,
> global real *y,
> real alpha, real beta,
> local real *X
> )
> {
> int l_id = get_local_id(0);
> int block_size = get_local_size(0);
> long g_id = get_global_id(0);
> for(int i = 0, j = g_id - lhalo; i < 1 + lhalo + rhalo; i++, j++)
> X[i] = read_x(j, n, has_left, has_right, lhalo, rhalo, xloc, xrem);
> // ^---- write to shared space X by multiple work-items in parallel
> if (g_id < n) {
> real sum = stencil_oper(X + lhalo);
> if (alpha)
> y[g_id] = alpha * y[g_id] + beta * sum;
> else
> y[g_id] = beta * sum;
> }
> }
>
>
>
> The reason why it now fails is that a new pocl kernel compiler optimization
> inserts an implicit barrier to the end of the for loop to enforce better
> parallelization possibilities and then it creates a new order for writes
> to X[i]:
>
> for(int i = 0, j = g_id - lhalo; i < 1 + lhalo + rhalo; i++, j++)
> for_all_work_items {
> X[i] = read_x(j, n, has_left, has_right, lhalo, rhalo, xloc, xrem);
> }
>
> Perhaps you need a barrier before the if? Still, it seems strange to
> write to the same indices of X by all work items?
>
>
> On 05/22/2013 07:04 PM, Pekka Jääskeläinen wrote:
>
>> OK. I just decided to install a Debian Sid VM with Boost 1.53.
>> I can reproduce the issue now without Boost crashing.
>>
>> On 05/22/2013 06:25 PM, Denis Demidov wrote:
>>
>>> Hi Pekka,
>>>
>>> I think this is more a cmake problem. Does
>>>
>>> cmake -DBOOST_ROOT=/path/to/custom/**boost ..
>>>
>>> work for you? If not, you can try to:
>>>
>>> cmake -DBOOST_ROOT=/path/to/custom/**boost -DBoost_NO_SYSTEM_PATHS=1 ..
>>>
>>> For me, the first variant works, and the second does not, but somebody
>>> told me
>>> that the second version worked for them.
>>>
>>> Cheers,
>>> Denis
>>>
>>>
>>> On Wed, May 22, 2013 at 5:47 PM, Pekka Jääskeläinen <
>>> [email protected]
>>> <mailto:pekka.jaaskelainen@**tut.fi <[email protected]>>> wrote:
>>>
>>> Hi Denis,
>>>
>>>
>>> On 05/21/2013 05:23 PM, Denis Demidov wrote:
>>>
>>> There is a problem in Boost v1.49 that leads to the crash. Boost v1.52
>>> and
>>> above is known to work.
>>>
>>>
>>> Is there a way to force VexCL build to look Boost in a
>>> non-standard directory? It insists finding my Debian's too old
>>> Boost from the default /usr/lib and I'd like to use a newer version
>>> built from the source. I can text replace the paths
>>> in CMakeCache.txt but it's not very handy.
>>>
>>>
>>>
>>> --
>>> Pekka
>>>
>>>
>>>
>>>
>>> --
>>> С уважением,
>>> Денис Демидов
>>> с.н.с., отдел высокопроизводительных вычислений
>>> отделение вычислительной математики
>>> НИИММ им. Н.Г.Чеботарева
>>> Казанский федеральный университет.
>>> тел. +7(843)233-77-07
>>>
>>
>>
>>
>
> --
> Pekka
>
--
С уважением,
Денис Демидов
с.н.с., отдел высокопроизводительных вычислений
отделение вычислительной математики
НИИММ им. Н.Г.Чеботарева
Казанский федеральный университет.
тел. +7(843)233-77-07
------------------------------------------------------------------------------
Try New Relic Now & We'll Send You this Cool Shirt
New Relic is the only SaaS-based application performance monitoring service
that delivers powerful full stack analytics. Optimize and monitor your
browser, app, & servers with just a few lines of code. Try New Relic
and get this awesome Nerd Life shirt! http://p.sf.net/sfu/newrelic_d2d_may
_______________________________________________
pocl-devel mailing list
[email protected]
https://lists.sourceforge.net/lists/listinfo/pocl-devel
