[libomptarget][devicertl] Wrap source in declare target pragmas

[JonChesterfield]

Summary

[libomptarget][devicertl] Wrap source in declare target pragmas

Factored out of D93135 / D94745. C++ and cuda ignore unknown pragmas
so this is a NFC for the current implementation language. Removes noise
from patches for building deviceRTL as openmp.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip

30

One nit. There is a blank line between declare target and the include, so I’d also insert a blank line between the last line of code and the end declare target.

[JonChesterfield]

• insert whitespace requested at review

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip

30

Sure, may as well.

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip

 //===-- amdgcn_locks.hip - AMDGCN OpenMP GPU lock implementation -- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // A 'thread' maps onto a lane of the wavefront. This means a per-thread lock
 // cannot be implemented - if one thread gets the lock, it can't continue on to
 // the next instruction in order to do anything as the other threads are waiting
 // to take the lock.
 // These functions will be implemented to provide the documented semantics for
 // a SIMD => wavefront mapping once that is implemented.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/debug.h"
 
 static DEVICE void warn() {
   PRINT0(LD_ALL, "Locks are not supported in this thread mapping model");
 }
 
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_set_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_unset_lock(omp_lock_t *) { warn(); }
 DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock) { warn(); }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip

 //===-------- amdgcn_smid.hip - AMDGCN smid implementation -------- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "target_impl.h"
 
 // Partially derived fom hcc_detail/device_functions.h
 
 // HW_ID Register bit structure
 // WAVE_ID     3:0     Wave buffer slot number. 0-9.
 // SIMD_ID     5:4     SIMD which the wave is assigned to within the CU.

 
 DEVICE uint32_t __kmpc_impl_smid() {
   uint32_t cu_id = __builtin_amdgcn_s_getreg(
       ENCODE_HWREG(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID));
   uint32_t se_id = __builtin_amdgcn_s_getreg(
       ENCODE_HWREG(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID));
   return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

 //===------- target_impl.hip - AMDGCN OpenMP GPU implementation --- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "target_impl.h"
 
 // Implementations initially derived from hcc
 
 // Initialized with a 64-bit mask with bits set in positions less than the
 // thread's lane number in the warp
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {

 
 EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads() {
   return GetNumberOfThreadsInBlock();
 }
 
 // Stub implementations
 DEVICE void *__kmpc_impl_malloc(size_t) { return nullptr; }
 DEVICE void __kmpc_impl_free(void *) {}
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/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.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "interface.h"
 #include "common/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;
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/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
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "interface.h"
 #include "common/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);
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/data_sharing.cu

 //===----- data_sharing.cu - OpenMP GPU data sharing ------------- 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 data sharing environments
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
+
 #include "common/omptarget.h"
 #include "target_impl.h"
 
 // Return true if this is the master thread.
 INLINE static bool IsMasterThread(bool isSPMDExecutionMode) {
   return !isSPMDExecutionMode && GetMasterThreadID() == GetThreadIdInBlock();
 }
 
if (isSPMDExecutionMode) {
     return;
   }
   __kmpc_impl_threadfence();
   ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   omptarget_nvptx_simpleMemoryManager.Release();
 }
 
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/libcall.cu

 //===------------ libcall.cu - OpenMP GPU user calls ------------- 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 implements the OpenMP runtime functions that can be
 // invoked by the user in an OpenMP region
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "common/target_atomic.h"
 #include "target_impl.h"
 
 EXTERN double omp_get_wtick(void) {
   double rc = __kmpc_impl_get_wtick();
   PRINT(LD_IO, "omp_get_wtick() returns %g\n", rc);
EXTERN void omp_unset_lock(omp_lock_t *lock) {
   PRINT0(LD_IO, "call omp_unset_lock()\n");
 }
 
 EXTERN int omp_test_lock(omp_lock_t *lock) {
   int rc = __kmpc_impl_test_lock(lock);
   PRINT(LD_IO, "call omp_test_lock() return %d\n", rc);
   return rc;
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/loop.cu

 //===------------ loop.cu - NVPTX OpenMP loop constructs --------- 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 the KMPC interface
 // for the loop construct plus other worksharing constructs that use the same
 // interface as loops.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "target_impl.h"
 #include "common/target_atomic.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////
 // template class that encapsulate all the helper functions
void __kmpc_for_static_init_8u_simple_generic(
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/false);
 }
 
 EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid) {
   PRINT0(LD_IO, "call kmpc_for_static_fini\n");
 }
+
+#pragma omp end declare target

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

 //===------------ omp_data.cu - 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.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "common/device_environment.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // global device environment
 ////////////////////////////////////////////////////////////////////////////////
 

 // Scratchpad for teams reduction.
 ////////////////////////////////////////////////////////////////////////////////
 DEVICE SHARED void *ReductionScratchpadPtr;
 
 ////////////////////////////////////////////////////////////////////////////////
 // Data sharing related variables.
 ////////////////////////////////////////////////////////////////////////////////
 DEVICE SHARED omptarget_nvptx_SharedArgs omptarget_nvptx_globalArgs;
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/omptarget.cu

 //===--- omptarget.cu - OpenMP GPU initialization ---------------- 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 initialization code for the GPU
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "target_impl.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime) {
   }
 }
 
 // Return true if the current target region is executed in SPMD mode.
 EXTERN int8_t __kmpc_is_spmd_exec_mode() {
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_is_spmd_exec_mode\n");
   return isSPMDMode();
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/parallel.cu

 //    in the parallel loop is that for each barrier in the parallel
 //    region, these non-included threads will cycle through the
 //    syncthread A. Thus they must preserve their current threadId that
 //    is larger than thread in team.
 //
 //    To make a long story short...
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "target_impl.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // support for parallel that goes parallel (1 static level only)
 ////////////////////////////////////////////////////////////////////////////////
 
EXTERN void __kmpc_push_num_teams(kmp_Ident *loc, int32_t tid,
   ASSERT0(LT_FUSSY, 0,
           "should never have anything with new teams on device");
 }
 
 EXTERN void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t tid,
                                   int proc_bind) {
   PRINT(LD_IO, "call kmpc_push_proc_bind %d\n", (int)proc_bind);
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

 //===---- reduction.cu - GPU OpenMP reduction implementation ----- 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 reduction with KMPC interface.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "common/target_atomic.h"
 #include "target_impl.h"
 
 EXTERN
 void __kmpc_nvptx_end_reduce(int32_t global_tid) {}
 
if (IsMaster && ChunkTeamCount == num_of_records - 1) {
     // Allow SIZE number of teams to proceed writing their
     // intermediate results to the global buffer.
     __kmpc_atomic_add((uint32_t *)&IterCnt, uint32_t(num_of_records));
   }
 
   return 0;
 }
 
+#pragma omp end declare target

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

 //===--------- support.cu - GPU 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.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/support.h"
 #include "common/debug.h"
 #include "common/omptarget.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////

 DEVICE unsigned int *GetTeamsReductionTimestamp() {
   return static_cast<unsigned int *>(ReductionScratchpadPtr);
 }
 
 DEVICE char *GetTeamsReductionScratchpad() {
   return static_cast<char *>(ReductionScratchpadPtr) + 256;
 }
 
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/sync.cu

 //===------------ sync.cu - GPU OpenMP synchronizations ---------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Include all synchronization.
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 #include "target_impl.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // KMP Ordered calls
 ////////////////////////////////////////////////////////////////////////////////
 

 ////////////////////////////////////////////////////////////////////////////////
 // Syncwarp
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_syncwarp(__kmpc_impl_lanemask_t Mask) {
   PRINT0(LD_IO, "call __kmpc_syncwarp\n");
   __kmpc_impl_syncwarp(Mask);
 }
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/task.cu

 //    -  private (of size given by task_alloc call). Accessed by
 //       task+sizeof(klegacy_TaskDescr)
 //        * private data *
 //    - shared: X. Accessed by shared ptr in klegacy_TaskDescr
 //        * pointer table to shared variables *
 //    - end
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "common/omptarget.h"
 
 EXTERN kmp_TaskDescr *__kmpc_omp_task_alloc(
     kmp_Ident *loc,     // unused
     uint32_t global_tid, // unused
     int32_t flag, // unused (because in our impl, all are immediately exec
     size_t sizeOfTaskInclPrivate, size_t sizeOfSharedTable,
if (*lb > *ub)
     return;
 
   // the compiler has already stored lb and ub in the kmp_TaskDescr structure
   // as we are using a single task to execute the entire loop, we can leave
   // the initial task_t untouched
 
   __kmpc_omp_task_with_deps(loc, global_tid, newKmpTaskDescr, 0, 0, 0, 0);
 }
+
+#pragma omp end declare target

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

 //===---------- target_impl.cu - NVPTX OpenMP GPU options ------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
+#pragma omp declare target
 
 #include "target_impl.h"
 #include "common/debug.h"
 #include "common/target_atomic.h"
 
 #include <cuda.h>
 
 DEVICE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {

 }
 
 DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock) {
   return __kmpc_atomic_add(lock, 0u);
 }
 
 DEVICE void *__kmpc_impl_malloc(size_t x) { return malloc(x); }
 DEVICE void __kmpc_impl_free(void *x) { free(x); }
+
+#pragma omp end declare target