https://gcc.gnu.org/bugzilla/show_bug.cgi?id=120216

--- Comment #3 from Benjamin Schulz <schulz.benjamin at googlemail dot com> ---
I now noticed that nvidia's hmm is denoted as 

#pragma omp requires unified_address

I do not know what unified_shared memory then is. Perhaps thats really reserved
for onboard gpu's then. 


The following code works with gcc now. Interestingly, it fails with clang...
It appears to work. But I guess I have to make checks whether it really runs on
the gpu...

That way, well, it appears, that one can now use stl vectors in a target region
without manually copying anything as this is done by the gpu driver...


Interestingly, that very same code fails with clang.


#include "omp.h"
#include <math.h>
#include <vector>
#include <numeric>
#include <stdio.h>
using namespace std;

#pragma omp requires unified_address

int main()
{

 std::vector<double> v(10,1) ;
double*d1=v.data();

// A first attempt with mapping the data field of an stl vector.
#pragma omp target enter data map(tofrom: d1[0:10])
#pragma omp target teams distribute parallel for
    for(size_t i=1;i<10;i++)
    {
        printf("%f\n",d1[i]);
    }
#pragma omp target exit data map(delete: d1[0:10])

//Now we change the data on the host.
#pragma omp parallel for
for (size_t i=1;i<10;i++)
    v[i]=20;

int b=omp_is_initial_device();
if (b==1) printf("is initial device called from a host region. This statement
should appear\n");

//no mapping with unified_address necessary, so commented out.
//#pragma omp target enter data map(tofrom: d1[0:10])

//we create a target region, check whether we are really on a target,
//we call a function of stl vector on the device, (the function pointers should
point to the same function on host and device),
//and we look whether the values changed at host by a loop appear changed in
the device.

#pragma omp target teams distribute
    for(size_t i=1;i<10;i++)
    {
    int b=omp_is_initial_device();
    if (b==1) printf("from a target region: is initial device. This statemend
should not appear.");
    //since the host and target pointer coincide, the  data changed by the host
should appear and the stl function should be accessible on the target.
    printf("%f\n",v.at(i));
    }
}

Reply via email to