Nvc++ compilation fails with option -O2 and -mp or -mp=gpu for specific OpenMP-5 offload code

When attempting to compile the attached C++ code with -O2 and -mp or -mp=gpu, compilation with nvc++ fails with error message:

/hpc_sdk/Linux_aarch64/26.1/compilers/share/llvm/bin/opt: /tmp/nvc++Bx0RnruxizbaC.ll:15296:11: error: use of undefined value ‘%L_T158_5218’
br label %L_T158_5218, !dbg !2028

The code is aimed at OpenMP-5 target offload and contains #pragma omp target directives. If all of the omp target directives are removed, compilation with nvc++ is successful. The code compiles and runs correctly with the current clang++ compiler built with offload support. The same failure occurs on Linux_x86_64 systems with nvc++ … again with llvm/bin/opt.

I would really like to have success with nvc++ … so please take a look and advise if there is a user-side fix, or if there is an issue with the optimizer for nvc++.

serial.cc.txt (47.2 KB)

Hi walkup -

This definitely appears to be a compiler bug. I reported this to engineering as a TPR#38251, and I’ll get back to you when I hear anything.

Compiling this myself and looking at the error message, it actually points me to the problematic line:

scamp:~$ nvc++ test.cpp -o test -llapack -lblas -mp -O2
/proj/nv/Linux_x86_64/295247-dev/compilers/share/llvm/bin/opt: /tmp/nvc++yk0Wci7y0ff-l.ll:15849:11: error: use of undefined value '%L_T158_5367'
        br label %L_T158_5367, !dbg !2042 ; "test.cpp":1152

So that corresponds to the break here:

         if ( (l == lmax - 1) && (iter <= maxiter) ) {
            if (myrank == 0) printf("re-starting iterations ...\n");
            restart = 1;
            break;
         }

If instead of calling that break, I instead write a new if statement that keys onto whether restart = 1 - then I can get the code to compile:

         if ( (l == lmax - 1) && (iter <= maxiter) ) {
            if (myrank == 0) printf("re-starting iterations ...\n");
            restart = 1;
         }

         if ( restart == 1) {
         break ;
         }
...
scamp:~$ nvc++ test.cpp -o test -llapack -lblas -mp -O2
"test.cpp", line 99: warning: variable "wflag" was set but never used [set_but_not_used]
     int ch, fflag, aflag, bflag, wflag, rflag;
                                  ^
Remark: individual warnings can be suppressed with "--diag_suppress <warning-name>"
"test.cpp", line 99: warning: variable "rflag" was set but never used [set_but_not_used]
     int ch, fflag, aflag, bflag, wflag, rflag;
                                         ^
scamp@dev-sky5:~$

Obviously I don’t have the right input file to run your code, but I don’t see any reason that wouldn’t be logically the same think, assuming you’re breaking anytime restart is set to 1.