not able to fill the array allocated on the gpu

Question

Please, help me. I have the following code

...
#include <accelmath.h>
#include <openacc.h>
const long int G=100000;
const unsigned int GL=100000;
const long int K=G;
const int LE=1.0f;
struct Particle
{
  float x;
  float rs;
};
Particle particles[GL];
int sort[GL];
int ind01[GL];
long int MAX_ELEMENT=1;
int POSITION1;
int POSITION0;
int LIFE=0;
bool start=true;
int mini;
int count0;
int count1;
int GL1;
int js;

#pragma acc declare device_resident(ind01,POSITION0,POSITION1,mini,GL1,js,MAX_ELEMENT,count0,count1,K)
#pragma acc declare create(LIFE,particles,sort)

 void function_device()
 {
   #pragma acc update host(LIFE,particles) async
    std::cout<<"LIFE before="<<LIFE<<std::endl;
   #pragma acc update device(LIFE,particles) async
   #pragma acc parallel num_gangs(1) vector_length(1) present(particles) async
   {
     count0=0;
     count1=0;
     if(LIFE<K)
     {
       particles[LIFE].x=5.0;
       particles[LIFE].rs=MAX_ELEMENT;
       ++MAX_ELEMENT;
       ++LIFE;
      }
    }
     #pragma acc loop vector reduction(+:count0,count1)
     for(int i=0; i<LIFE; ++i)
     {
       if(particles[i].x>=LE)
       {
         sort[i]=1;
         count1=count1+1;
       }
       else
       {
         sort[i]=0;
         count0=count0+1;
       }
     }   
   #pragma acc parallel num_gangs(1) vector_length(1)
     {
       GL1=LIFE-1;
       count0=GL1;
       count1=0;
     }
     #pragma acc loop seq
     for(int i=0; i<LIFE; ++i)
     {
       if(sort[i]==1)
       {
         ind01[count1++]=i;
       }
       else
       {
         ind01[count0--]=i;
       }
     }
    #pragma acc parallel num_gangs(1) vector_length(1)
    {
      mini=GL1-count0;
      if(count1<mini) mini=count1;
      js=0;
    }
    ...
    #pragma acc update host(LIFE) async
  }
  int main(int argc, char **argv)
  {
    acc_init(acc_device_nvidia);
    int step=1;
while(start==true || LIFE>0)
{
      std::cout<<" LIFE="<<LIFE<<std::endl;
      start=false;
      function_device();   
      std::cout<<"MAIN LOOP # "<<step<<std::endl;
      ++step;
    }
  }

After i have filled in the array allocated on the gpu:

     particles[LIFE].x=5.0;

but the output of the print:

     #pragma acc update host(LIFE,particles) async
     std::cout<<"LIFE after injector="<<LIFE<<std::endl;
     for(int i=0; i<LIFE; ++i) std::cout<<" particles: "<<particles[i]<<std::endl;

and:

    #pragma acc update host(LIFE,sort) async
     std::cout<<"LIFE after 1="<<LIFE<<" c0="<<count0<<" c1="<<count1<<std::endl;
     for(int i=0; i<LIFE; ++i) std::cout<<"sort: "<<sort[i]<<std::endl;

is LIFE=1 the arrays contain only 0, as if i did not fill the array particles with 5.0f. How to fill in the array "particles"? Do i correctly launch a serial code on the gpu:

    #pragma acc parallel num_gangs(1) vector_length(1)
    {
      mini=GL1-count0;
      if(count1<mini) mini=count1;
      js=0;
    }

If there are mistakes in the usage of openacc directives, please, show them. Why the

     #pragma acc update host(LIFE) async

directive does not work without async?

Mat Colgrove Mat Colgrove · Accepted Answer · 2017-11-20T22:52:47

I updated your code to get it to work. You had a few orphaned loops which needed "parallel" on them plus since several of the variables are globals, you need to use atomic. I did encounter a compiler issue when using variables defined in the "declare create" in an atomic so needed to move a few of the variables to a data region. I'll report this to our engineers.

Note that I also put in a terminating condition when LIFE reached G, otherwise it seemed to enter an infinite loop.

#include <iostream>
#include <accelmath.h>
#include <openacc.h>
const long int G=100000;
const unsigned int GL=100000;
const long int K=G;
const int LE=1.0f;
struct Particle
{
  float x;
  float rs;
};
Particle particles[GL];
int sort[GL];
int ind01[GL];
int MAX_ELEMENT;
int POSITION1;
int POSITION0;
int LIFE;
bool start=true;
int mini;
int GL1;
int js;
int count0;
int count1;

#pragma acc declare device_resident(ind01,POSITION0,POSITION1,mini,GL1,js)
#pragma acc declare create(particles,sort)

 void function_device()
 {
   #pragma acc parallel num_gangs(1) vector_length(1) present(particles)
   {
#pragma acc atomic write
      count0=0;
#pragma acc atomic write
      count1=0;
     if(LIFE<K)
     {
       particles[LIFE].x=5.0;
       particles[LIFE].rs=MAX_ELEMENT;
#pragma acc atomic update
       ++MAX_ELEMENT;
#pragma acc atomic update
       ++LIFE;
      }
    }
     #pragma acc parallel loop
     for(int i=0; i<LIFE; ++i)
     {
       if(particles[i].x>=LE)
       {
         sort[i]=1;
  #pragma acc atomic update
         count1=count1+1;
       }
       else
       {
         sort[i]=0;
  #pragma acc atomic update
         count0=count0+1;
       }
     }
   #pragma acc parallel num_gangs(1) vector_length(1)
     {
       GL1=LIFE-1;
       count0=GL1;
       count1=0;
     }
     #pragma acc parallel loop
     for(int i=0; i<LIFE; ++i)
     {
       int cnt;
       if(sort[i]==1)
       {
        #pragma acc atomic capture
        {
          cnt = count1++;
        }
        ind01[cnt]=i;
       }
       else
       {
        #pragma acc atomic capture
        {
          cnt = count0--;
        }
        ind01[cnt]=i;
       }
     }
    #pragma acc parallel num_gangs(1) vector_length(1)
    {
      mini=GL1-count0;
      if(count1<mini) mini=count1;
      js=0;
    }
  }
  int main(int argc, char **argv)
  {
    acc_init(acc_device_nvidia);
    int step=1;
    LIFE=0;
    MAX_ELEMENT=1;
   #pragma acc data copyin(LIFE,MAX_ELEMENT,count0,count1)
{
while(start==true || (LIFE>0 && LIFE < G ))
{
      std::cout<<" LIFE="<<LIFE<<std::endl;
      start=false;
      function_device();
   #pragma acc update self(LIFE)
      std::cout<<"MAIN LOOP # "<<step<<std::endl;
      ++step;
    }
}
  }

not able to fill the array allocated on the gpu

1 Answers