[BUG] nvc 26.5: OpenMP reduction on volatile long returns zero at -O2/-O3, works at -O0

Accelerated Computing HPC Compilers nvc, nvc++ and nvfortran
1

Hello,

I encountered a suspected optimisation-dependent issue with an OpenMP sum reduction on a volatile long. With NVIDIA nvc 26.5-0, the volatile counter remains zero at -O2 and -O3, while an otherwise identical non-volatile counter produces the expected result.

Both counters work at -O0. GCC 12.2.0 produces the expected results at all three optimisation levels.

Minimal reproducer:

Save as volatile_reduction.c:

int main(void)
{
volatile long v = 0;
long p = 0;

#pragma omp parallel for reduction(+:v)
for (long i = 0; i < 1000; i++)
    v++;

#pragma omp parallel for reduction(+:p)
for (long i = 0; i < 1000; i++)
    p++;

printf("volatile: %ld  plain: %ld  (both should be 1000)\n",
       v, p);

return (v != 1000 || p != 1000);

}

Compile and Run:

nvc -O3 -mp volatile_reduction.c -o reduction_nvc
OMP_DYNAMIC=FALSE OMP_NUM_THREADS=1 ./reduction_nvc
OMP_DYNAMIC=FALSE OMP_NUM_THREADS=4 ./reduction_nvc

Actual output for both thread counts:

volatile: 0 plain: 1000 (both should be 1000)

Both runs return exit status 1.

For Comparison:

gcc -O3 -fopenmp volatile_reduction.c -o reduction_gcc
OMP_DYNAMIC=FALSE OMP_NUM_THREADS=1 ./reduction_gcc
OMP_DYNAMIC=FALSE OMP_NUM_THREADS=4 ./reduction_gcc

Both GCC runs print the expected output and return exit status 0:

volatile: 1000 plain: 1000 (both should be 1000)

Optimisation-level comparison:

I compiled the same source separately at each optimisation level. Each row below was tested with both one and four threads, with identical results for both thread counts.

Compiler Optimisation Volatile counter Plain counter
nvc 26.5-0 -O0 1000 1000
nvc 26.5-0 -O2 0 1000
nvc 26.5-0 -O3 0 1000
GCC 12.2.0 -O0 1000 1000
GCC 12.2.0 -O2 1000 1000
GCC 12.2.0 -O3 1000 1000

Environment:

  • System: NCI Gadi
  • OS: Rocky Linux 8.10 (Green Obsidian)
  • Architecture: x86_64
  • CPU: Intel Xeon Platinum 8268 @ 2.90 GHz
  • NVIDIA module: nvidia-hpc-sdk/26.5
  • GCC module: gcc/12.2.0
  • CPU OpenMP execution, without GPU offloading
  • OMP_DYNAMIC=FALSE

The original teaching example uses volatile to make unsynchronised accesses observable before introducing a reduction. The reproducer above includes the reduction, so I would expect both counters to produce 1000.

Is this a known optimisation issue involving volatile reduction variables? Please let me know if additional compiler diagnostics would help.

I also reproduced this with 25.1,5,7,9

Thank you.

2

Hi ecktons and welcome!

I was able to reproduce the issue here and filed a report, TPR #39038.

Though I’m wondering if using a volatile variable in a reduction is a good idea. Per the standard, OpenMP reduction variables are meant to be ordinary list items of arithmetic type. While the standard doesn’t explicitly forbid volatile, it doesn’t define what it means there, so I think you’re relying on unspecified behavior.

In a reduction, each thread works on a private copy and the copies are combined at the end. Those private copies aren’t shared, so volatile has no useful job to do. It only makes the code slower and the semantics murkier. To get the reduction to work properly, the compiler effectively would not have the private copies inherit “volatile”.

-Mat

3

Thanks Mat,

In the context I was exploring the volatile variable was only defined as such to discouraging the compiler from keeping this counter in a register, to observe the effects of race conditions in OpenMP with and without the reduction. I agree that volatile isn’t needed for the corrected reduction. Obviously this is a very specific case and wouldn’t be something seen often however I thought it was something interesting to report.

Thanks again for looking into it!

Regards,

Sam