[OpenMP] OpenMPOpt miscompiles sprintf function from the GPU implementation

[jhuber6]

The OpenMPOpt pass currently miscompiles trivial uses of the sprintf function. This was recently enabled with the AMDGPU support of varargs. The following example prints an empty string with openmp-opt enabled, and 1 with it disabled.

#include <stdio.h>

int main() {
  char str[256];
#pragma omp target
  {
    snprintf(str, 255, "%d\n", 1);
  }
  fputs(str, stdout);
}

This can be compiled with the GPU libc using the following, with the disabled version printing 1 as expected.

$ clang sprintf.c -fopenmp --offload-arch=gfx1030 -O1
$ clang sprintf.c -fopenmp --offload-arch=gfx1030 -O1 -mllvm -openmp-opt-disable

The output before the offending openmp-opt pass is available here https://godbolt.org/z/ecEjEKvv3.

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Joseph Huber (jhuber6)

The OpenMPOpt pass currently miscompiles trivial uses of the `sprintf` function. This was recently enabled with the AMDGPU support of varargs. The following example prints an empty string with `openmp-opt` enabled, and `1` with it disabled.

#include <stdio.h>

int main() {
  char str[256];
#pragma omp target
  {
    snprintf(str, 255, "%d\n", 1);
  }
  fputs(str, stdout);
}

This can be compiled with the GPU libc using the following, with the disabled version printing 1 as expected.

$ clang sprintf.c -fopenmp --offload-arch=gfx1030 -O1
$ clang sprintf.c -fopenmp --offload-arch=gfx1030 -O1 -mllvm -openmp-opt-disable

The output before the offending openmp-opt pass is available here https://godbolt.org/z/ecEjEKvv3.