OpenACCArray swap function

Question

while trying to create an object oriented OpenACC implementation I stumbled upon this question.

From there I took the code provided by @mat-colgrove at the GTC15 (code available at http://www.pgroup.com/lit/samples/gtc15_S5233.tar).

Since I am interested how to use objects to manage data on with OpenACC I posted another question. I was quite impressed by the ease of the OpenACCArray::swap function, so I created a small example to test it (see gist).

First I tried to just swap and hope that it is sufficient to swap the pointers on the host, but this ends in a fatal memory error. (presumably because the size and capacity members are not updated on the device)
A safer approach, that I assumed to work is to update the host, swap arrays and update device. This runs but creates wrong results.

I am compiling for nvidia accelerators.

score 1 · Accepted Answer · answered Aug 25 '16 at 21:48

Looks like this is my fault since I didn't test the swap routine.

The problem here is while the code is swapping the data on the host, the device copy of the objects still point to the old array. The fix is to re-attach (i.e. set the object's device pointers to the correct arrays) the lists.

    void swap(OpenACCArray<type>& x)
    {
        type* tmp_list = list;
        int tmp_size = _size;
        int tmp_capacity = _capacity;
        list = x.list;
        _size = x._size;
        _capacity = x._capacity;
        x.list = tmp_list;
        x._size = tmp_size;
        x._capacity = tmp_capacity;
#ifdef _OPENACC
#pragma acc update device(_size,_capacity,x._size,x._capacity)
        acc_attach((void**)&list);
        acc_attach((void**)&x.list);
#endif

    }

"acc_attach" is a PGI extension that hopefully will be adopted in the OpenACC 3.0 standard.

Thanks for trying things out and let me know if you encounter other issues. - Mat

OpenACCArray swap function

1 Answers1