У меня есть следующий простой фрагмент кода, состоящий из 4 файлов:
//Data.h:
#ifndef DATA_H
#define DATA_H
constexpr int N=10000000;
namespace data{
float particles[N];
float outputArray[N];
}
#endif
//Process.h:
#ifndef PROCESS_H
#define PROCESS_H
template <typename ProcessImpl>
class Process{
public:
using Base_t = ProcessImpl;
Process(ProcessImpl arg):fProcessImpl(arg){}
void Get1(int N, float * outputArray) const;
void Get2(int N) const;
private:
ProcessImpl fProcessImpl;
};
template <class ProcessImpl>
void Process<ProcessImpl>::Get1(int N, float * outputArray) const
{
#pragma acc parallel loop gang vector present(outputArray)
for(int ind=0; ind < N; ++ind){outputCSArray[ind]=fProcessImpl.Get1(ind);}
}
template <class ProcessImpl>
void Process<ProcessImpl>::Get2(int N) const
{
#pragma acc parallel loop gang vector
for (int ind = 0u; ind < N; ++ind){fProcessImpl.Get2(ind);}
}
#endif
//ProcessImpl.h:
#ifndef PROCESSIMPL_H
#define PROCESSIMPL_H
#include "Data.h"
using namespace data;
class ProcessImpl
{
public:
inline float Get1(int ind, float * outputArray) const;
inline void Get2(int ind) const;
};
float ProcessImpl::Get1(int ind, float * outputArray) const
{
outputArray[ind]=particles[ind];
return particles[ind+1];
}
void ProcessImpl::Get2(int ind) const
{
particles[ind]=2*particles[ind];
}
#endif
//main.cpp:
#include <iostream>
#include "Data.h"
#include "Process.h"
#include "ProcessImpl.h"
#include <accelmath.h>
#include <openacc.h>
using namespace data;
using Process_t = Process<ProcessImpl>;
Process_t process = Process_t(typename Process_t::Base_t());
int main(int argc, char **argv)
{
#pragma acc data create(particles,outputArray)
{
#pragma acc parallel loop gang vector present(particles)
for(int i=0; i<N; ++i) particles[i]=static_cast<float>(i);
#pragma acc update host(particles)
for(int i=0; i<100; ++i) std::cout<<particles[i]<<" ";
std::cout<<std::endl;
process.Get2(N);
#pragma acc update host(particles)
for(int i=0; i<100; ++i) std::cout<<particles[i]<<" ";
std::cout<<std::endl;
}
return 0;
}
Он корректно работает на CPU с компилятором PGI 19.4. Но моя задача - запустить код на GPU. Я использую PGI 19.4 + OpenAcc. Используйте простой файл CMakeLists.txt и строку компиляции (GPU Nvidia GeForce 650 Ti, вычислительные возможности 3.0):
cmake . -DCMAKE_C_COMPILER=pgcc -DCMAKE_CXX_COMPILER=pgc++
-DCMAKE_C_FLAGS="-acc -Minfo=acc -mcmodel=medium -ta=tesla:cc30"
-DCMAKE_CXX_FLAGS="-acc -Minfo=acc -mcmodel=medium -ta=tesla:cc30"
Сбой компиляции:
> Scanning dependencies of target Test
[ 50%] Building CXX object CMakeFiles/Test.dir/main.cpp.o
main:
16, Generating create(_ZN4data11outputArrayE[:])
Generating present(_ZN4data9particlesE[:])
Generating create(_ZN4data9particlesE[:])
Generating Tesla code
18, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
18, Generating update self(_ZN4data9particlesE[:])
23, Generating update self(_ZN4data9particlesE[:])
PGCC-W-0155-External and Static variables are not supported in acc routine -
_ZN4data9particlesE (/home/70-gaa/source/13OpenAccTest/main.cpp: 19)
ProcessImpl::Get2(int) const:
4, include "ProcessImpl.h"
18, Generating implicit acc routine seq
Process<ProcessImpl>::Get2(int) const:
3, include "Process.h"
25, Generating Tesla code
27, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
25, Generating implicit copyin(this[:])
PGCC/x86-64 Linux 19.4-0: compilation completed with warnings
[100%] Linking CXX executable Test
nvlink error : Undefined reference to '_ZNK11ProcessImpl4Get2Ei' in
'CMakeFiles/Test.dir/main.cpp.o'
pgacclnk: child process exit status 2: /opt/pgi/linux86-64-llvm/19.4/bin/pgnvd
CMakeFiles/Test.dir/build.make:83: recipe for target 'Test' failed
make[2]: *** [Test] Error 2
CMakeFiles/Makefile2:72: recipe for target 'CMakeFiles/Test.dir/all' failed
make[1]: *** [CMakeFiles/Test.dir/all] Error 2
Makefile:83: recipe for target 'all' failed
make: *** [all] Error 2
Используя pggdecode, обнаружил, что "_ZNK11ProcessImpl4Get2Ei" - это искаженное имя для ProcessImpl :: Get2 (int) const. Я удалил ключевое слово inline из ProcessImpl.h и попытался добавить copyin (process) в #pragma acc data create (частицы, outputArray) в main (), но это не помогло. Работайте в Fedora 23 с gcc 5.3.1.
В полном коде я избегаю проблемы множественного определения массивов частиц и outputArray, включая их в один файл .cpp, потому что OpenAcc не позволяет использовать extern ключевое слово. Может быть, это и нехорошо (если умеете лучше, посоветуйте, пожалуйста), но работает.
Вопросы:
Как правильно передать версию массивов частиц и outputArray для GPU в Get1 () и Get2 () в < em> Process.h и заставьте Get1 () и Get2 () в ProcessImpl.h работать с массивами, размещенными на GPU. ? А как скомпилировать этот код?
Как OpenAcc позволяет напрямую обращаться к глобально распределенным массивам, скопированным в GPU в коде в вычислительных регионах OpenAcc, без передачи на них указателя в качестве параметра вызывающей функции?
Спасибо.
