c++ - 在平行部分循环中使用 `__pgi_gangidx()`

Question

__pgi_gangidx() 编译器扩展函数应该返回执行该函数的组的数字 id(参见 here)。但是，我一直无法弄清楚如何在并行部分循环中使用它。

在下面的代码中，我尝试了几种可能性，只有其中一种产生了所需的答案。不幸的是，这一次按顺序运行我的并行部分循环。

变量 int place 是对多个全局数组的更复杂的特定于帮派的引用的替代，因此不容易删除。

代码可以编译为:

pgc++ -fast -acc -ta=tesla,cc60 -Minfo=accel test.cpp

代码:

#include <iostream>
#include "openacc.h"

void ResetIds(int *const ids, int size){
  //Ensure everything is zeroed
  for(int i=0;i<size;i++)
    ids[i] = 0;  
}

void ShowVector(int line, int *const ids, int size){
  std::cout<<"Line "<<line<<": ";
  for(int i=0;i<size;i++)
    std::cout<<ids[i]<<" ";
  std::cout<<std::endl<<std::endl;
}

int main(){
  int gangs  = 10;
  int gwidth = 10;
  int size   = gangs*gwidth;
  int *ids   = new int[50*size];

  //Works!
  //Gives: 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14
  ResetIds(ids, size);
  #pragma acc parallel num_gangs(gangs) copy(ids[0:size])
  {
    int place = __pgi_gangidx();
    #pragma acc loop seq
    for(int i=0;i<10;i++)
      ids[place*gwidth+i] = 14;
  }  
  ShowVector(__LINE__, ids, size);

  //Gives: 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
  ResetIds(ids, size);
  #pragma acc parallel num_gangs(gangs) copy(ids[0:size])
  {
    int place = __pgi_gangidx()*gwidth;
    #pragma acc loop 
    for(int i=0;i<10;i++)
      ids[place+i] = 14;
  } 
  ShowVector(__LINE__, ids, size);

  //Gives: 14 14 14 14 14 14 14 14 14 14 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
  ResetIds(ids, size);
  #pragma acc parallel num_gangs(gangs) copy(ids[0:size])
  {
    int place = __pgi_gangidx();
    #pragma acc loop
    for(int i=0;i<10;i++)
      ids[place*gwidth+i] = 14;
  }  
  ShowVector(__LINE__, ids, size);

  //Gives: 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14 0 0 0 0 0 0 0 0 0 0 14
  ResetIds(ids, size);
  #pragma acc parallel num_gangs(gangs) copy(ids[0:size])
  {
    int place = __pgi_gangidx();
    #pragma acc loop worker
    for(int i=0;i<10;i++)
      ids[place*gwidth+i] = 14;
  }  
  ShowVector(__LINE__, ids, size);

  return 0;
}

Answer 1

抱歉回复晚了。我去参加一个 session ，所以错过了你的帖子。

“__pgi_gangidx()”工作正常。这里的问题是对“loop”指令的作用有误解。

在第一种情况下，您已将“seq”添加到循环指令中以强制循环按顺序运行。由于它在并行区域中，因此循环以“gang-redundant”模式运行，所有 gang 都将执行循环。换句话说，循环实际执行了 100 次。 (请注意，我认为这是一种不好的做法，因为如果在没有 OpenACC 的情况下按顺序运行代码会得到不同的结果)

在其他情况下，您不指定“seq”。因此，循环将在所有帮派和 vector 中并行化，并且只会运行 10 次。由于 vector 长度为​​ 128，但循环次数为 10，因此实际上只有一个团伙会更新数组，并且只会更新 10 个元素。