[nfc][libomptarget] Decrease coupling between files

[JonChesterfield]

Summary

[nfc][libomptarget] Decrease coupling between files

debug.h used the symbol omptarget_device_environment so implicitly required
an include of omptarget-nvptx.h to compile. Similarly interface.h uses size_t.

Moving this declaration to a new header means cancel, critical can now build
without omptarget-nvptx.h. After this change, debug.h, cancel.cu, critical.cu
could move under a common source directory.

[ABataev]

LG

openmp/libomptarget/deviceRTLs/interface.h

 //  declared by OpenMP.  The second section includes the compiler
 //  specific interfaces.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _INTERFACES_H_
 #define _INTERFACES_H_
 
+#include <stddef.h>
 #include <stdint.h>
 
 #ifdef __CUDACC__
 #include "nvptx/src/nvptx_interface.h"
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // OpenMP interface

openmp/libomptarget/deviceRTLs/nvptx/src/cancel.cu

 //===------ cancel.cu - NVPTX OpenMP cancel interface ------------ CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Interface to be used in the implementation of OpenMP cancel.
 //
 //===----------------------------------------------------------------------===//
 
-#include "omptarget-nvptx.h"
+#include "interface.h"
+#include "debug.h"
 
 EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
                                         int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancellationpoint(cancel val %d)\n", (int)cancelVal);
   // disabled
   return 0;
 }
 
 EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
                              int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
   // disabled
   return 0;
 }

openmp/libomptarget/deviceRTLs/nvptx/src/critical.cu

 //===------ critical.cu - NVPTX OpenMP critical ------------------ CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the implementation of critical with KMPC interface
 //
 //===----------------------------------------------------------------------===//
 
-#include <stdio.h>
-
-#include "omptarget-nvptx.h"
+#include "interface.h"
+#include "debug.h"
 
 EXTERN
 void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
                      kmp_CriticalName *lck) {
   PRINT0(LD_IO, "call to kmpc_critical()\n");
   omp_set_lock((omp_lock_t *)lck);
 }
 
 EXTERN
 void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
                          kmp_CriticalName *lck) {
   PRINT0(LD_IO, "call to kmpc_end_critical()\n");
   omp_unset_lock((omp_lock_t *)lck);
 }

openmp/libomptarget/deviceRTLs/nvptx/src/debug.h

 //   WARNING(flag, str, args)      : if warning flag is on, print the warning
 //   WON(flag)                     : return true if warning flag is on
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_NVPTX_DEBUG_H_
 #define _OMPTARGET_NVPTX_DEBUG_H_
 
+#include "device_environment.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // set desired level of debugging
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LD_SET_NONE 0ULL /* none */
 #define LD_SET_ALL -1ULL /* all */
 
 // pos 1

openmp/libomptarget/deviceRTLs/nvptx/src/device_environment.h

+//===---- device_environment.h - OpenMP GPU device environment --- CUDA -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Global device environment
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPTARGET_DEVICE_ENVIRONMENT_H_
+#define _OMPTARGET_DEVICE_ENVIRONMENT_H_
+
+#include "target_impl.h"
+
+struct omptarget_device_environmentTy {
+  int32_t debug_level;
+};
+
+extern __device__ omptarget_device_environmentTy omptarget_device_environment;
+
+#endif

openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu

 //===------------ omp_data.cu - NVPTX OpenMP GPU objects --------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the data objects used on the GPU device.
 //
 //===----------------------------------------------------------------------===//
 
 #include "omptarget-nvptx.h"
+#include "device_environment.h"
 
 ////////////////////////////////////////////////////////////////////////////////
-// global device envrionment
+// global device environment
 ////////////////////////////////////////////////////////////////////////////////
 
 __device__ omptarget_device_environmentTy omptarget_device_environment;
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data holding OpenMP state information
 ////////////////////////////////////////////////////////////////////////////////
 

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

private:
   int64_t chunk[MAX_THREADS_PER_TEAM];
   int64_t loopUpperBound[MAX_THREADS_PER_TEAM];
   // state for dispatch with dyn/guided OR static (never use both at a time)
   int64_t nextLowerBound[MAX_THREADS_PER_TEAM];
   int64_t stride[MAX_THREADS_PER_TEAM];
   uint64_t cnt;
 };
 
-/// Device envrionment data
-struct omptarget_device_environmentTy {
-  int32_t debug_level;
-};
-
 /// Memory manager for statically allocated memory.
 class omptarget_nvptx_SimpleMemoryManager {
 private:
   __align__(128) struct MemDataTy {
     volatile unsigned keys[OMP_STATE_COUNT];
   } MemData[MAX_SM];
 
   INLINE static uint32_t hash(unsigned key) {
     return key & (OMP_STATE_COUNT - 1);
   }
 
 public:
   INLINE void Release();
   INLINE const void *Acquire(const void *buf, size_t size);
 };
 
 ////////////////////////////////////////////////////////////////////////////////
-// global device envrionment
-////////////////////////////////////////////////////////////////////////////////
-
-extern __device__ omptarget_device_environmentTy omptarget_device_environment;
-
-////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__ omptarget_nvptx_SimpleMemoryManager
     omptarget_nvptx_simpleMemoryManager;
 extern __device__ __shared__ uint32_t usedMemIdx;