Issue Compiling using Clang14.x with OpenMP Offloading

[Ramas31]

Hello all,

Any OpenMP offloading program with schedule(auto) clause does not compile successfully with the latest LLVM-CLANG(14.x) version. However replacing the same with schedule(static) works without any errors. It did not show any such issues with previous versions.

You could try compiling the below code

int main()                                                                      
{                                                                               
#pragma omp target teams distribute parallel for schedule(auto)                 
  for(int i=0; i<1000; i++)                                                     
  {}                                                                            
  return 0;                                                                     
}

I compile it using the below command

clang -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target -march=sm_70 -o test test.c

Below is the complete error observed

ptxas /tmp/test-f7e4d8.s, line 266; error : Module-scoped variables in .local state space are not allowed with ABI
ptxas fatal : Ptx assembly aborted due to errors
clang-14: **error: ptxas command failed with exit code 255 (use -v to see invocation)**
clang version 14.0.1 (https://github.com/llvm/llvm-project.git 0e27d08cdeb338766a477fba071b3df7a06ea6e2)
Target: x86_64-unknown-linux-gnu
Thread model: posix
clang-14: **note:** diagnostic msg: Error generating preprocessed source(s).

Additionally, using O2/O3 optimization level during compilation works for this sample code but not always. Some programs still give the ptxas error during compilation even after using O2/O3.

Issue was discussed on LLVM- discourse https://discourse.llvm.org/t/issue-compiling-using-clang14-x-with-openmp-offloading/62445

Please solve this issue and let us know

Thank You Rama

[llvmbot]

@llvm/issue-subscribers-openmp

[shiltian]

We generally don't support dynamic scheduling on device yet because of this, a thread local global variable (sounds weird).