[clang] Invert condition refactored in llvm#160935 (llvm#161583)

The PR llvm#160935 incorrectly replaced `llvm::sys::fs::getUniqueID()` with
`llvm::vfs::FileSystem::exists()` in a condition. That's incorrect,
since the first function returns `std::error_code` that evaluates to
`true` when there is an error (file doesn't exist), while the new code
does the opposite. This PR fixes that issue by inverting the
conditional.

Co-authored-by: ronlieb <ron.lieberman@amd.com>

Author: jansvoboda11

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -1539,7 +1539,7 @@ static llvm::TargetRegionEntryInfo getEntryInfoFromPresumedLoc(
     SourceManager &SM = CGM.getContext().getSourceManager();
     PresumedLoc PLoc = SM.getPresumedLoc(BeginLoc);
 
-    if (CGM.getFileSystem()->exists(PLoc.getFilename()))
+    if (!CGM.getFileSystem()->exists(PLoc.getFilename()))
       PLoc = SM.getPresumedLoc(BeginLoc, /*UseLineDirectives=*/false);
 
     return std::pair<std::string, uint64_t>(PLoc.getFilename(), PLoc.getLine());

clang/test/OpenMP/amdgcn_save_temps.c

@@ -0,0 +1,23 @@
+
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -E -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -save-temps=cwd %s -o %t-openmp-amdgcn-amd-amdhsa-gfx90a.i
+// RUN: %clang_cc1 -fopenmp  -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -save-temps=cwd -emit-llvm-bc %s -o %t-x86_64-unknown-unknown.bc
+// RUN: %clang_cc1 -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -save-temps=cwd -emit-llvm -fopenmp-is-target-device -x cpp-output %t-openmp-amdgcn-amd-amdhsa-gfx90a.i -fopenmp-host-ir-file-path %t-x86_64-unknown-unknown.bc -o - | FileCheck %s
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+#define N 1000
+
+int test_amdgcn_save_temps() {
+  int arr[N];
+#pragma omp target
+  for (int i = 0; i < N; i++) {
+    arr[i] = 1;
+  }
+  return arr[0];
+}
+#endif
+
+// CHECK: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_amdgcn_save_temps