Error using atomics in OpenMP offload region

[olupton]

Hello,

I have been working on porting a code that uses the Eigen library to use OpenMP target offload with NVC++ 22.1.
Trying to use Eigen in offloaded regions sometimes produces the following error:

NVC++-S-0000-Internal compiler error. accel.c::acc_remove_atomics: ILI_ALT is missing for atomicrmw opcode     402  (stripped_tmp.cpp: 402)
NVC++/x86-64 Linux 22.1-0: compilation completed with severe errors

which I have managed to reproduce with the small example:

#include <atomic>
int main() {
  #pragma omp target
  {
    std::atomic<int> x{};
    x -= 1;
  }
}

or with no #includes

void sub() {
  int i{}, j{};
  __atomic_sub_fetch(&i, j, int());
}
int main() {
  #pragma omp target
  {
    sub();
  }
}

compiled with nvc++ -mp=gpu test.cpp.

On our system nvc++ is configured to use headers from GCC 11.2.0. Please let me know if any other information would be helpful.

In the meantime, we have a workaround where we compile the relevant Eigen functions as CUDA code and then call those CUDA-defined functions from OpenMP offload regions, but this is a bit of a frustrating kludge.

[MatColgrove]

Thanks olupton,

Using std::atomic in OpenMP offload regions is expected to work, hence I’ve filed a problem report, TPR #31305, and sent it to engineering for investigation.

-Mat

[MatColgrove]

Hi Olli,

TPR #31305 has been fixed in our 22.5 release.

-Mat

% cat atomic.cpp
#include <atomic>
int main() {
#pragma omp target
{
std::atomic<int> x{};
x -= 1;
}
}
% nvc++ -mp=gpu atomic.cpp -V22.3
NVC++-S-0000-Internal compiler error. accel.c::acc_remove_atomics: ILI_ALT is missing for atomicrmw opcode     343  (atomic.cpp: 343)
NVC++/x86-64 Linux 22.3-0: compilation completed with severe errors
% nvc++ -mp=gpu atomic.cpp -V22.5
% cat atomic_sub.cpp
void sub() {
int i{}, j{};
__atomic_sub_fetch(&i, j, int());
}
int main() {
#pragma omp target
{
sub();
}
}
% nvc++ -mp=gpu atomic_sub.cpp -V22.3
NVC++-S-0000-Internal compiler error. accel.c::acc_remove_atomics: ILI_ALT is missing for atomicrmw opcode       4  (atomic_sub.cpp: 4)
NVC++/x86-64 Linux 22.3-0: compilation completed with severe errors
% nvc++ -mp=gpu atomic_sub.cpp -V22.5
%