[nfc][libomptarget] Reorganise support header

[JonChesterfield]

Summary

[nfc][libomptarget] Reorganise support header

All functions defined in support implementation are now declared in support.h
Reordered functions in support implementation to match the sequence in support.h
Added include guards to support.h
Added #include interface to support.h to provide kmp_Ident declaration
Move supporti.h to support.cu and s/INLINE/EXTERN/g
Add remaining includes to support.cu

A minor side effect is to change the name mangling of the support functions to
extern "C". If this matters another macro along the lines of INLINE/EXTERN
can be added - perhaps DEVICE as that's the obvious implementation.

[jdoerfert]

There seem to be unrelated moves but other than that this looks good to me. Consider not moving stuff if it's not necessary and also using a different header guard name.

LGTM.

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

14

The device RTL is not consistent here (__XXX and _XXX are used) and it is different from LLVM. I'd vote for LLVM style, thus OMPTARGET_SUPPORT_H

[JonChesterfield]

• Review comments, retrieve inline annotation and mangled symbols

[JonChesterfield]

Using EXTERN means dropping the __always_inline__ annotation on the functions. That was unintentional, changed back to INLINE in the implementation.

Also introduced DEVICE as an annotation for the non-inline header functions. We don't currently restrict the symbols exposed from the library and names like PadBytes are generic enough that I can imagine them colliding with application symbols. The mangled names are known to not collide, to the extent that no-one has reported such a collision.

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

14

Yeah, I'd noticed the variation. I like your suggestion, will change it.

[JonChesterfield]

• revert code motion

[JonChesterfield]

Reverted code motion, which decreases the diff size considerably.

[jdoerfert]

LGTM. Thx!

[JonChesterfield]

Discussion at https://reviews.llvm.org/rG764c8420e4b8fc11a9fa94d00f4ee617aa754cb2

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

set(cuda_src_files
       src/cancel.cu
       src/critical.cu
       src/data_sharing.cu
       src/libcall.cu
       src/loop.cu
       src/omptarget-nvptx.cu
       src/parallel.cu
       src/reduction.cu
+      src/support.cu
       src/sync.cu
       src/task.cu
   )
 
   set(omp_data_objects src/omp_data.cu)
 
   # Get the compute capability the user requested or use SM_35 by default.
   # SM_35 is what clang uses by default.

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

 getMyTopTaskDescriptor(bool isSPMDExecutionMode);
 INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
 
 ////////////////////////////////////////////////////////////////////////////////
 // inlined implementation
 ////////////////////////////////////////////////////////////////////////////////
 
 #include "omptarget-nvptxi.h"
-#include "supporti.h"
 
 #endif

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

 //===--------- support.h - NVPTX OpenMP support functions -------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper to some functions natively supported by the GPU.
 //
 //===----------------------------------------------------------------------===//
 
+#ifndef _OMPTARGET_SUPPORT_H_
+#define _OMPTARGET_SUPPORT_H_
+
+#include "interface.h"
 #include "target_impl.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 enum ExecutionMode {
   Generic = 0x00u,
   Spmd = 0x01u,
   ModeMask = 0x01u,
 };
 
 enum RuntimeMode {
   RuntimeInitialized = 0x00u,
   RuntimeUninitialized = 0x02u,
   RuntimeMask = 0x02u,
 };
 
-INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
-INLINE bool isGenericMode();
-INLINE bool isSPMDMode();
-INLINE bool isRuntimeUninitialized();
-INLINE bool isRuntimeInitialized();
+EXTERN void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
+EXTERN bool isGenericMode();
+EXTERN bool isSPMDMode();
+EXTERN bool isRuntimeUninitialized();
+EXTERN bool isRuntimeInitialized();
+
+////////////////////////////////////////////////////////////////////////////////
+// Execution Modes based on location parameter fields
+////////////////////////////////////////////////////////////////////////////////
+
+EXTERN bool checkSPMDMode(kmp_Ident *loc);
+
+EXTERN bool checkGenericMode(kmp_Ident *loc);
+
+EXTERN bool checkRuntimeUninitialized(kmp_Ident *loc);
+
+EXTERN bool checkRuntimeInitialized(kmp_Ident *loc);
 
 ////////////////////////////////////////////////////////////////////////////////
 // get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 // get low level ids of resources
-INLINE int GetThreadIdInBlock();
-INLINE int GetBlockIdInKernel();
-INLINE int GetNumberOfBlocksInKernel();
-INLINE int GetNumberOfThreadsInBlock();
-INLINE unsigned GetWarpId();
-INLINE unsigned GetLaneId();
+EXTERN int GetThreadIdInBlock();
+EXTERN int GetBlockIdInKernel();
+EXTERN int GetNumberOfBlocksInKernel();
+EXTERN int GetNumberOfThreadsInBlock();
+EXTERN unsigned GetWarpId();
+EXTERN unsigned GetLaneId();
 
 // get global ids to locate tread/team info (constant regardless of OMP)
-INLINE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
-INLINE int GetMasterThreadID();
-INLINE int GetNumberOfWorkersInTeam();
+EXTERN int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
+EXTERN int GetMasterThreadID();
+EXTERN int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
-INLINE int GetOmpThreadId(int threadId,
+EXTERN int GetOmpThreadId(int threadId,
                           bool isSPMDExecutionMode);    // omp_thread_num
-INLINE int GetOmpTeamId();                              // omp_team_num
+EXTERN int GetOmpTeamId();                              // omp_team_num
 
 // get OpenMP number of threads and team
-INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
-INLINE int GetNumberOfOmpTeams();                           // omp_num_teams
+EXTERN int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
+EXTERN int GetNumberOfOmpTeams();                           // omp_num_teams
 
 // get OpenMP number of procs
-INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
-INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
+EXTERN int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
+EXTERN int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
 
 // masters
-INLINE int IsTeamMaster(int ompThreadId);
+EXTERN int IsTeamMaster(int ompThreadId);
 
 // Parallel level
-INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
-INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+EXTERN void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+EXTERN void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
 
 ////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
 // safe alloc and free
-INLINE void *SafeMalloc(size_t size, const char *msg); // check if success
-INLINE void *SafeFree(void *ptr, const char *msg);
+EXTERN void *SafeMalloc(size_t size, const char *msg); // check if success
+EXTERN void *SafeFree(void *ptr, const char *msg);
 // pad to a alignment (power of 2 only)
-INLINE unsigned long PadBytes(unsigned long size, unsigned long alignment);
+EXTERN unsigned long PadBytes(unsigned long size, unsigned long alignment);
 #define ADD_BYTES(_addr, _bytes)                                               \
   ((void *)((char *)((void *)(_addr)) + (_bytes)))
 #define SUB_BYTES(_addr, _bytes)                                               \
   ((void *)((char *)((void *)(_addr)) - (_bytes)))
 
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
-INLINE unsigned int *GetTeamsReductionTimestamp();
-INLINE char *GetTeamsReductionScratchpad();
-INLINE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
+EXTERN unsigned int *GetTeamsReductionTimestamp();
+EXTERN char *GetTeamsReductionScratchpad();
+EXTERN void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
+
+#endif

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

-//===--------- supporti.h - NVPTX OpenMP support functions ------- CUDA -*-===//
+//===--------- support.cu - NVPTX OpenMP support functions ------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper implementation to some functions natively supported by the GPU.
 //
 //===----------------------------------------------------------------------===//
 
+#include "support.h"
+#include "debug.h"
+#include "omptarget-nvptx.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
-#include "target_impl.h"
-
-INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
+EXTERN void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
   execution_param = EMode;
   execution_param |= RMode;
 }
 
-INLINE bool isGenericMode() { return (execution_param & ModeMask) == Generic; }
+EXTERN bool isGenericMode() { return (execution_param & ModeMask) == Generic; }
 
-INLINE bool isSPMDMode() { return (execution_param & ModeMask) == Spmd; }
+EXTERN bool isSPMDMode() { return (execution_param & ModeMask) == Spmd; }
 
-INLINE bool isRuntimeUninitialized() {
+EXTERN bool isRuntimeUninitialized() {
   return (execution_param & RuntimeMask) == RuntimeUninitialized;
 }
 
-INLINE bool isRuntimeInitialized() {
+EXTERN bool isRuntimeInitialized() {
   return (execution_param & RuntimeMask) == RuntimeInitialized;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Modes based on location parameter fields
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE bool checkSPMDMode(kmp_Ident *loc) {
+EXTERN bool checkSPMDMode(kmp_Ident *loc) {
   if (!loc)
     return isSPMDMode();
 
   // If SPMD is true then we are not in the UNDEFINED state so
   // we can return immediately.
   if (loc->reserved_2 & KMP_IDENT_SPMD_MODE)
     return true;
 
   // If not in SPMD mode and runtime required is a valid
   // combination of flags so we can return immediately.
   if (!(loc->reserved_2 & KMP_IDENT_SIMPLE_RT_MODE))
     return false;
 
   // We are in underfined state.
   return isSPMDMode();
 }
 
-INLINE bool checkGenericMode(kmp_Ident *loc) {
+EXTERN bool checkGenericMode(kmp_Ident *loc) {
   return !checkSPMDMode(loc);
 }
 
-INLINE bool checkRuntimeUninitialized(kmp_Ident *loc) {
+EXTERN bool checkRuntimeUninitialized(kmp_Ident *loc) {
   if (!loc)
     return isRuntimeUninitialized();
 
   // If runtime is required then we know we can't be
   // in the undefined mode. We can return immediately.
   if (!(loc->reserved_2 & KMP_IDENT_SIMPLE_RT_MODE))
     return false;
 
   // If runtime is required then we need to check is in
   // SPMD mode or not. If not in SPMD mode then we end
   // up in the UNDEFINED state that marks the orphaned
   // functions.
   if (loc->reserved_2 & KMP_IDENT_SPMD_MODE)
     return true;
 
   // Check if we are in an UNDEFINED state. Undefined is denoted by
   // non-SPMD + noRuntimeRequired which is a combination that
   // cannot actually happen. Undefined states is used to mark orphaned
   // functions.
   return isRuntimeUninitialized();
 }
 
-INLINE bool checkRuntimeInitialized(kmp_Ident *loc) {
+EXTERN bool checkRuntimeInitialized(kmp_Ident *loc) {
   return !checkRuntimeUninitialized(loc);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // support: get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // Calls to the NVPTX layer  (assuming 1D layout)
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE int GetThreadIdInBlock() { return threadIdx.x; }
+EXTERN int GetThreadIdInBlock() { return threadIdx.x; }
 
-INLINE int GetBlockIdInKernel() { return blockIdx.x; }
+EXTERN int GetBlockIdInKernel() { return blockIdx.x; }
 
-INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; }
+EXTERN int GetNumberOfBlocksInKernel() { return gridDim.x; }
 
-INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+EXTERN int GetNumberOfThreadsInBlock() { return blockDim.x; }
 
-INLINE unsigned GetWarpId() { return threadIdx.x / WARPSIZE; }
+EXTERN unsigned GetWarpId() { return threadIdx.x / WARPSIZE; }
 
-INLINE unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); }
+EXTERN unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // Calls to the Generic Scheme Implementation Layer (assuming 1D layout)
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-// The master thread id is the first thread (lane) of the last warp.
-// Thread id is 0 indexed.
-// E.g: If NumThreads is 33, master id is 32.
-//      If NumThreads is 64, master id is 32.
-//      If NumThreads is 97, master id is 96.
-//      If NumThreads is 1024, master id is 992.
-//
-// Called in Generic Execution Mode only.
-INLINE int GetMasterThreadID() { return (blockDim.x - 1) & ~(WARPSIZE - 1); }
-
-// The last warp is reserved for the master; other warps are workers.
-// Called in Generic Execution Mode only.
-INLINE int GetNumberOfWorkersInTeam() { return GetMasterThreadID(); }
-
 ////////////////////////////////////////////////////////////////////////////////
 // get thread id in team
 
 // This function may be called in a parallel region by the workers
 // or a serial region by the master.  If the master (whose CUDA thread
 // id is GetMasterThreadID()) calls this routine, we return 0 because
 // it is a shadow for the first worker.
-INLINE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode) {
+EXTERN int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode) {
   // Implemented using control flow (predication) instead of with a modulo
   // operation.
   int tid = GetThreadIdInBlock();
   if (!isSPMDExecutionMode && tid >= GetMasterThreadID())
     return 0;
   else
     return tid;
 }
 
+// The master thread id is the first thread (lane) of the last warp.
+// Thread id is 0 indexed.
+// E.g: If NumThreads is 33, master id is 32.
+//      If NumThreads is 64, master id is 32.
+//      If NumThreads is 97, master id is 96.
+//      If NumThreads is 1024, master id is 992.
+//
+// Called in Generic Execution Mode only.
+EXTERN int GetMasterThreadID() { return (blockDim.x - 1) & ~(WARPSIZE - 1); }
+
+// The last warp is reserved for the master; other warps are workers.
+// Called in Generic Execution Mode only.
+EXTERN int GetNumberOfWorkersInTeam() { return GetMasterThreadID(); }
+
 ////////////////////////////////////////////////////////////////////////////////
 //
 // OpenMP Thread Support Layer
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) {
+EXTERN int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) {
   // omp_thread_num
   int rc;
   if ((parallelLevel[GetWarpId()] & (OMP_ACTIVE_PARALLEL_LEVEL - 1)) > 1) {
     rc = 0;
   } else if (isSPMDExecutionMode) {
     rc = GetThreadIdInBlock();
   } else {
     omptarget_nvptx_TaskDescr *currTaskDescr =
         omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId);
     ASSERT0(LT_FUSSY, currTaskDescr, "expected a top task descr");
     rc = currTaskDescr->ThreadId();
   }
   return rc;
 }
 
-INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode) {
+EXTERN int GetOmpTeamId() {
+  // omp_team_num
+  return GetBlockIdInKernel(); // assume 1 block per team
+}
+
+EXTERN int GetNumberOfOmpThreads(bool isSPMDExecutionMode) {
   // omp_num_threads
   int rc;
   int Level = parallelLevel[GetWarpId()];
   if (Level != OMP_ACTIVE_PARALLEL_LEVEL + 1) {
     rc = 1;
   } else if (isSPMDExecutionMode) {
     rc = GetNumberOfThreadsInBlock();
   } else {
     rc = threadsInTeam;
   }
 
   return rc;
 }
 
+EXTERN int GetNumberOfOmpTeams() {
+  // omp_num_teams
+  return GetNumberOfBlocksInKernel(); // assume 1 block per team
+}
+
 ////////////////////////////////////////////////////////////////////////////////
-// Team id linked to OpenMP
+// get OpenMP number of procs
 
-INLINE int GetOmpTeamId() {
-  // omp_team_num
-  return GetBlockIdInKernel(); // assume 1 block per team
+EXTERN int GetNumberOfProcsInTeam(bool isSPMDExecutionMode) {
+  return GetNumberOfProcsInDevice(isSPMDExecutionMode);
 }
 
-INLINE int GetNumberOfOmpTeams() {
-  // omp_num_teams
-  return GetNumberOfBlocksInKernel(); // assume 1 block per team
+// Get the number of processors in the device.
+EXTERN int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
+  if (!isSPMDExecutionMode)
+    return GetNumberOfWorkersInTeam();
+  return GetNumberOfThreadsInBlock();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Masters
 
-INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
+EXTERN int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Parallel level
 
-INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
+EXTERN void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
   __kmpc_impl_syncwarp(Mask);
   __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt();
   unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] +=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
   __kmpc_impl_syncwarp(Mask);
 }
 
-INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
+EXTERN void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
   __kmpc_impl_syncwarp(Mask);
   __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt();
   unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] -=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
   __kmpc_impl_syncwarp(Mask);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
-// get OpenMP number of procs
-
-// Get the number of processors in the device.
-INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
-  if (!isSPMDExecutionMode)
-    return GetNumberOfWorkersInTeam();
-  return GetNumberOfThreadsInBlock();
-}
-
-INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode) {
-  return GetNumberOfProcsInDevice(isSPMDExecutionMode);
-}
-
-////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE unsigned long PadBytes(unsigned long size,
-                              unsigned long alignment) // must be a power of 2
-{
-  // compute the necessary padding to satisfy alignment constraint
-  ASSERT(LT_FUSSY, (alignment & (alignment - 1)) == 0,
-         "alignment %lu is not a power of 2\n", alignment);
-  return (~(unsigned long)size + 1) & (alignment - 1);
-}
-
-INLINE void *SafeMalloc(size_t size, const char *msg) // check if success
+EXTERN void *SafeMalloc(size_t size, const char *msg) // check if success
 {
   void *ptr = malloc(size);
   PRINT(LD_MEM, "malloc data of size %llu for %s: 0x%llx\n",
         (unsigned long long)size, msg, (unsigned long long)ptr);
   return ptr;
 }
 
-INLINE void *SafeFree(void *ptr, const char *msg) {
+EXTERN void *SafeFree(void *ptr, const char *msg) {
   PRINT(LD_MEM, "free data ptr 0x%llx for %s\n", (unsigned long long)ptr, msg);
   free(ptr);
   return NULL;
 }
 
+EXTERN unsigned long PadBytes(unsigned long size,
+                              unsigned long alignment) // must be a power of 2
+{
+  // compute the necessary padding to satisfy alignment constraint
+  ASSERT(LT_FUSSY, (alignment & (alignment - 1)) == 0,
+         "alignment %lu is not a power of 2\n", alignment);
+  return (~(unsigned long)size + 1) & (alignment - 1);
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE unsigned int *GetTeamsReductionTimestamp() {
+EXTERN unsigned int *GetTeamsReductionTimestamp() {
   return static_cast<unsigned int *>(ReductionScratchpadPtr);
 }
 
-INLINE char *GetTeamsReductionScratchpad() {
+EXTERN char *GetTeamsReductionScratchpad() {
   return static_cast<char *>(ReductionScratchpadPtr) + 256;
 }
 
-INLINE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr) {
+EXTERN void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr) {
   ReductionScratchpadPtr = ScratchpadPtr;
 }

openmp/libomptarget/deviceRTLs/nvptx/unity.cu

 #include "src/critical.cu"
 #include "src/data_sharing.cu"
 #include "src/libcall.cu"
 #include "src/loop.cu"
 #include "src/omp_data.cu"
 #include "src/omptarget-nvptx.cu"
 #include "src/parallel.cu"
 #include "src/reduction.cu"
+#include "src/support.cu"
 #include "src/sync.cu"
 #include "src/task.cu"