[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. ```c #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. ```console $ 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.