[clang] Use the VFS to get the OpenMP entry info

[jansvoboda11]

This PR uses the VFS to get the OpenMP entry info instead of going straight to the real file system. This matches the behavior of other input files of the compiler.

[llvmbot]

@llvm/pr-subscribers-clang-codegen

Author: Jan Svoboda (jansvoboda11)

Changes

This PR uses the VFS to get the OpenMP entry info instead of going straight to the real file system. This matches the behavior of other input files of the compiler.

---

Full diff: https://github.com/llvm/llvm-project/pull/160935.diff

1 Files Affected:

• (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+1-3)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 0136f69172aa5..75bde3f72c4c2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1542,10 +1542,8 @@ static llvm::TargetRegionEntryInfo getEntryInfoFromPresumedLoc(
     SourceManager &SM = CGM.getContext().getSourceManager();
     PresumedLoc PLoc = SM.getPresumedLoc(BeginLoc);
 
-    llvm::sys::fs::UniqueID ID;
-    if (llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID)) {
+    if (CGM.getFileSystem()->exists(PLoc.getFilename()))
       PLoc = SM.getPresumedLoc(BeginLoc, /*UseLineDirectives=*/false);
-    }
 
     return std::pair<std::string, uint64_t>(PLoc.getFilename(), PLoc.getLine());
   };

[llvmbot]

@llvm/pr-subscribers-clang

Author: Jan Svoboda (jansvoboda11)

Changes

This PR uses the VFS to get the OpenMP entry info instead of going straight to the real file system. This matches the behavior of other input files of the compiler.

---

Full diff: https://github.com/llvm/llvm-project/pull/160935.diff

1 Files Affected:

• (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+1-3)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 0136f69172aa5..75bde3f72c4c2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1542,10 +1542,8 @@ static llvm::TargetRegionEntryInfo getEntryInfoFromPresumedLoc(
     SourceManager &SM = CGM.getContext().getSourceManager();
     PresumedLoc PLoc = SM.getPresumedLoc(BeginLoc);
 
-    llvm::sys::fs::UniqueID ID;
-    if (llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID)) {
+    if (CGM.getFileSystem()->exists(PLoc.getFilename()))
       PLoc = SM.getPresumedLoc(BeginLoc, /*UseLineDirectives=*/false);
-    }
 
     return std::pair<std::string, uint64_t>(PLoc.getFilename(), PLoc.getLine());
   };

[carlobertolli]

@jansvoboda11 this change breaks -save-temps builds.
Reproducer:

#include <stdio.h>
#include <omp.h>

int main()
{
  int N = 10;

  int a[N];
  int b[N];

  int i;

  for (i=0; i<N; i++)
    a[i]=0;

  for (i=0; i<N; i++)
    b[i]=i;

#pragma omp target parallel for
  {
    for (int j = 0; j< N; j++)
      a[j]=b[j];
  }

  int rc = 0;
  for (i=0; i<N; i++)
    if (a[i] != b[i] ) {
      rc++;
      printf ("Wrong varlue: a[%d]=%d\n", i, a[i]);
    }

  if (!rc)
    printf("Success\n");

  return rc;
}

Build without -save-temps and make llvm print out list of target regions it identified:
clang -mllvm -amdgpu-dump-hsa-metadata -O2 -fopenmp --offload-arch=gfx90a -D__OFFLOAD_ARCH_gfx90a__ veccopy.c

Excerpt from resulting diagnostic shows that there are multiple entry points:
.name: __omp_offloading_811_2ea3964_main_l19
.private_segment_fixed_size: 0
.sgpr_count: 14
.sgpr_spill_count: 0
.symbol: __omp_offloading_811_2ea3964_main_l19.kd

Build with -save-temps and make llvm print out list of target regions it identified:
clang -mllvm -amdgpu-dump-hsa-metadata -O2 -save-temps -fopenmp --offload-arch=gfx90a -D__OFFLOAD_ARCH_gfx90a__ veccopy.c

Excerpt from resulting diagnostic shows that there are no kernels AMDGPU HSA Metadata

amdhsa.kernels: []
amdhsa.target: amdgcn-amd-amdhsa--gfx90a
amdhsa.version:

• 1
• 2
  ...

Please fix?

Debugging showed me that this line
if (CGM.getFileSystem()->exists(PLoc.getFilename())) {
is now succeeding (it was not before and I presume that was incorrect, hence your fix).
When using -save-temps, the following statement
PLoc = SM.getPresumedLoc(BeginLoc, /*UseLineDirectives=*/false);
is now identifying the preprocessed <src_file>.i as the actual PLoc filename and not the input source file.

When not using -save-temps, the preprocessed <src_file>.i does not exist in the build folder and the source file remains the same.

For some reason, this has an effect on clang's ability to identify target regions when building for the host, which seems fragile at best.

[jansvoboda11]

Thanks for bringing this to my attention, @carlobertolli. I'm happy to take a look, but I'm unable to reproduce on macOS. Could you please create a self-contained LIT test independent of the host platform?

[carlobertolli]

here we go:
#161472

[carlobertolli]

My guess as to why this stopped working is because the function in OMPIRBuilder.cpp
OpenMPIRBuilder::getTargetEntryUniqueInfo(FileIdentifierInfoCallbackTy CallBack,
vfs::FileSystem &VFS,
StringRef ParentName) {

is referring to the source_file.c instead of the preprocessed source_file.i, and when we scan for target regions, we cannot find a match.