Clarification on using OpenACC in a shared library

OpenACC and CUDA Fortran
brentl
Posts: 256
Joined: Jul 20 2004

Re: Clarification on using OpenACC in a shared library

Post by brentl » Tue Oct 01, 2019 10:04 am

Yeah, that makes sense. The global data in declare directives is a case of what needs to get setup at program init time, which we can handle when you link with pgc++.

These are user functions you are trying to inline, correct? Not in header files...

DavidGutzwiller
Posts: 10
Joined: Jan 30 2015

Re: Clarification on using OpenACC in a shared library

Post by DavidGutzwiller » Tue Oct 01, 2019 11:04 am

Yes, these are user functions/methods. Some defined in header files, which seem to work, and some in source (*.C) files which are giving me the problems.

DavidGutzwiller
Posts: 10
Joined: Jan 30 2015

Re: Clarification on using OpenACC in a shared library

Post by DavidGutzwiller » Tue Oct 01, 2019 12:41 pm

I should also note that some of these routines are nested.

For example:
I have class A where the parallelism is exposed. Class A hosts a pointer to class B, which in turn calls a utility function implemented elsewhere. A system of create/attach methods offloads everything to the GPU and keeps the pointer tree up to date.

--in file A.C--
#pragma acc parallel loop present(B)
for (i=1, i<n; i++)
{
B->foo(i)
}

--in file B.H--
class B
{
#pragma acc routine seq
void foo(int i)
}

--in file B.C--
B::foo(int i)
{
C->foo(i)
}

--in file util.H--
#pragma acc routine seq
void utility(int i)

--in file util.C--
utility(int i)
{
<a bunch of math>
}

As currently configured, I encounter compile errors like:
ptxas fatal : Unresolved extern function '_Z9utilityRKiPKdRS1_RiS4_'

I suppose the path around this would be to first create an inline library with -Mextract, followed by compiling with -Minline=lib:<library>. Is this correct?

brentl
Posts: 256
Joined: Jul 20 2004

Re: Clarification on using OpenACC in a shared library

Post by brentl » Wed Oct 02, 2019 7:48 am

If you have "acc routines" marked appropriately, you should not get the unresolved error at link time. Inlining should be an optimization, and not required for correctness.

If you are using our compilers on OpenPOWER, or our LLVM compiler (default since 19.1) on X86, we do not support IPA yet, so -Minline is your only option for inlining across files, as you have discovered.

DavidGutzwiller
Posts: 10
Joined: Jan 30 2015

Re: Clarification on using OpenACC in a shared library

Post by DavidGutzwiller » Wed Oct 02, 2019 9:41 am

To be clear, the "ptxas fatal : Unresolved extern function" error occurs at compile time, not linking. This occurs due to the nordc flag, which is where I still am not quite clear. I have found that "nordc" is still mandatory to get any acc pragmas or API calls to work. From your earlier comments this does not seem to be correct.

I'll take another look into my build system, perhaps I am missing something.

Post Reply