Использование `__pgi_gangidx ()` в цикле параллельного сечения

__pgi_gangidx() Функция расширения компилятора должна возвращать числовой идентификатор группы, выполняющей функцию (см. Вот). Тем не менее, я не смог понять, как использовать его в цикле параллельного сечения.

В приведенном ниже коде я пробую несколько возможностей, только одна из которых дает желаемый ответ. Этот, к сожалению, выполняет мой параллельный цикл секций последовательно.

Переменная 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;
}