[Openmp-commits] [openmp] r374919 - [libomptarget][nfc] Make interface.h target independent

Jon Chesterfield via Openmp-commits openmp-commits at lists.llvm.org
Tue Oct 15 10:15:26 PDT 2019


Author: jonchesterfield
Date: Tue Oct 15 10:15:26 2019
New Revision: 374919

URL: http://llvm.org/viewvc/llvm-project?rev=374919&view=rev
Log:
[libomptarget][nfc] Make interface.h target independent

Summary:
[libomptarget][nfc] Make interface.h target independent

Move interface.h under a top level include directory.
Remove #includes to avoid the interface depending on the implementation.

Reviewers: ABataev, jdoerfert, grokos, ronlieb, RaviNarayanaswamy

Reviewed By: jdoerfert

Subscribers: mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D68615

Added:
    openmp/trunk/libomptarget/deviceRTLs/interface.h
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
Removed:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h
Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/option.h
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h

Added: openmp/trunk/libomptarget/deviceRTLs/interface.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/interface.h?rev=374919&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/interface.h (added)
+++ openmp/trunk/libomptarget/deviceRTLs/interface.h Tue Oct 15 10:15:26 2019
@@ -0,0 +1,539 @@
+//===------- interface.h - OpenMP interface definitions ---------- 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 all the definitions that are relevant to
+//  the interface. The first section contains the interface as
+//  declared by OpenMP.  The second section includes the compiler
+//  specific interfaces.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _INTERFACES_H_
+#define _INTERFACES_H_
+
+#include <stdint.h>
+
+#ifdef __CUDACC__
+#include "nvptx/src/nvptx_interface.h"
+#endif
+
+////////////////////////////////////////////////////////////////////////////////
+// OpenMP interface
+////////////////////////////////////////////////////////////////////////////////
+
+typedef uint32_t omp_lock_t;      /* arbitrary type of the right length */
+typedef uint64_t omp_nest_lock_t; /* arbitrary type of the right length */
+
+typedef enum omp_sched_t {
+  omp_sched_static = 1,  /* chunkSize >0 */
+  omp_sched_dynamic = 2, /* chunkSize >0 */
+  omp_sched_guided = 3,  /* chunkSize >0 */
+  omp_sched_auto = 4,    /* no chunkSize */
+} omp_sched_t;
+
+typedef enum omp_proc_bind_t {
+  omp_proc_bind_false = 0,
+  omp_proc_bind_true = 1,
+  omp_proc_bind_master = 2,
+  omp_proc_bind_close = 3,
+  omp_proc_bind_spread = 4
+} omp_proc_bind_t;
+
+EXTERN double omp_get_wtick(void);
+EXTERN double omp_get_wtime(void);
+
+EXTERN void omp_set_num_threads(int num);
+EXTERN int omp_get_num_threads(void);
+EXTERN int omp_get_max_threads(void);
+EXTERN int omp_get_thread_limit(void);
+EXTERN int omp_get_thread_num(void);
+EXTERN int omp_get_num_procs(void);
+EXTERN int omp_in_parallel(void);
+EXTERN int omp_in_final(void);
+EXTERN void omp_set_dynamic(int flag);
+EXTERN int omp_get_dynamic(void);
+EXTERN void omp_set_nested(int flag);
+EXTERN int omp_get_nested(void);
+EXTERN void omp_set_max_active_levels(int level);
+EXTERN int omp_get_max_active_levels(void);
+EXTERN int omp_get_level(void);
+EXTERN int omp_get_active_level(void);
+EXTERN int omp_get_ancestor_thread_num(int level);
+EXTERN int omp_get_team_size(int level);
+
+EXTERN void omp_init_lock(omp_lock_t *lock);
+EXTERN void omp_init_nest_lock(omp_nest_lock_t *lock);
+EXTERN void omp_destroy_lock(omp_lock_t *lock);
+EXTERN void omp_destroy_nest_lock(omp_nest_lock_t *lock);
+EXTERN void omp_set_lock(omp_lock_t *lock);
+EXTERN void omp_set_nest_lock(omp_nest_lock_t *lock);
+EXTERN void omp_unset_lock(omp_lock_t *lock);
+EXTERN void omp_unset_nest_lock(omp_nest_lock_t *lock);
+EXTERN int omp_test_lock(omp_lock_t *lock);
+EXTERN int omp_test_nest_lock(omp_nest_lock_t *lock);
+
+EXTERN void omp_get_schedule(omp_sched_t *kind, int *modifier);
+EXTERN void omp_set_schedule(omp_sched_t kind, int modifier);
+EXTERN omp_proc_bind_t omp_get_proc_bind(void);
+EXTERN int omp_get_cancellation(void);
+EXTERN void omp_set_default_device(int deviceId);
+EXTERN int omp_get_default_device(void);
+EXTERN int omp_get_num_devices(void);
+EXTERN int omp_get_num_teams(void);
+EXTERN int omp_get_team_num(void);
+EXTERN int omp_is_initial_device(void);
+EXTERN int omp_get_initial_device(void);
+EXTERN int omp_get_max_task_priority(void);
+
+////////////////////////////////////////////////////////////////////////////////
+// file below is swiped from kmpc host interface
+////////////////////////////////////////////////////////////////////////////////
+
+////////////////////////////////////////////////////////////////////////////////
+// kmp specifc types
+////////////////////////////////////////////////////////////////////////////////
+
+typedef enum kmp_sched_t {
+  kmp_sched_static_chunk = 33,
+  kmp_sched_static_nochunk = 34,
+  kmp_sched_dynamic = 35,
+  kmp_sched_guided = 36,
+  kmp_sched_runtime = 37,
+  kmp_sched_auto = 38,
+
+  kmp_sched_static_balanced_chunk = 45,
+
+  kmp_sched_static_ordered = 65,
+  kmp_sched_static_nochunk_ordered = 66,
+  kmp_sched_dynamic_ordered = 67,
+  kmp_sched_guided_ordered = 68,
+  kmp_sched_runtime_ordered = 69,
+  kmp_sched_auto_ordered = 70,
+
+  kmp_sched_distr_static_chunk = 91,
+  kmp_sched_distr_static_nochunk = 92,
+  kmp_sched_distr_static_chunk_sched_static_chunkone = 93,
+
+  kmp_sched_default = kmp_sched_static_nochunk,
+  kmp_sched_unordered_first = kmp_sched_static_chunk,
+  kmp_sched_unordered_last = kmp_sched_auto,
+  kmp_sched_ordered_first = kmp_sched_static_ordered,
+  kmp_sched_ordered_last = kmp_sched_auto_ordered,
+  kmp_sched_distribute_first = kmp_sched_distr_static_chunk,
+  kmp_sched_distribute_last =
+      kmp_sched_distr_static_chunk_sched_static_chunkone,
+
+  /* Support for OpenMP 4.5 monotonic and nonmonotonic schedule modifiers.
+   * Since we need to distinguish the three possible cases (no modifier,
+   * monotonic modifier, nonmonotonic modifier), we need separate bits for
+   * each modifier. The absence of monotonic does not imply nonmonotonic,
+   * especially since 4.5 says that the behaviour of the "no modifier" case
+   * is implementation defined in 4.5, but will become "nonmonotonic" in 5.0.
+   *
+   * Since we're passing a full 32 bit value, we can use a couple of high
+   * bits for these flags; out of paranoia we avoid the sign bit.
+   *
+   * These modifiers can be or-ed into non-static schedules by the compiler
+   * to pass the additional information. They will be stripped early in the
+   * processing in __kmp_dispatch_init when setting up schedules, so
+   * most of the code won't ever see schedules with these bits set.
+   */
+  kmp_sched_modifier_monotonic = (1 << 29),
+  /**< Set if the monotonic schedule modifier was present */
+  kmp_sched_modifier_nonmonotonic = (1 << 30),
+/**< Set if the nonmonotonic schedule modifier was present */
+
+#define SCHEDULE_WITHOUT_MODIFIERS(s)                                          \
+  (enum kmp_sched_t)(                                                          \
+      (s) & ~(kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic))
+#define SCHEDULE_HAS_MONOTONIC(s) (((s)&kmp_sched_modifier_monotonic) != 0)
+#define SCHEDULE_HAS_NONMONOTONIC(s)                                           \
+  (((s)&kmp_sched_modifier_nonmonotonic) != 0)
+#define SCHEDULE_HAS_NO_MODIFIERS(s)                                           \
+  (((s) & (kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic)) == \
+   0)
+
+} kmp_sched_t;
+
+/*!
+ * Enum for accesseing the reserved_2 field of the ident_t struct below.
+ */
+enum {
+  /*! Bit set to 1 when in SPMD mode. */
+  KMP_IDENT_SPMD_MODE = 0x01,
+  /*! Bit set to 1 when a simplified runtime is used. */
+  KMP_IDENT_SIMPLE_RT_MODE = 0x02,
+};
+
+/*!
+ * The ident structure that describes a source location.
+ * The struct is identical to the one in the kmp.h file.
+ * We maintain the same data structure for compatibility.
+ */
+typedef int kmp_int32;
+typedef struct ident {
+  kmp_int32 reserved_1; /**<  might be used in Fortran; see above  */
+  kmp_int32 flags; /**<  also f.flags; KMP_IDENT_xxx flags; KMP_IDENT_KMPC
+                      identifies this union member  */
+  kmp_int32 reserved_2; /**<  not really used in Fortran any more; see above */
+  kmp_int32 reserved_3; /**<  source[4] in Fortran, do not use for C++  */
+  char const *psource; /**<  String describing the source location.
+                       The string is composed of semi-colon separated fields
+                       which describe the source file, the function and a pair
+                       of line numbers that delimit the construct. */
+} ident_t;
+
+// parallel defs
+typedef ident_t kmp_Ident;
+typedef void (*kmp_ParFctPtr)(int32_t *global_tid, int32_t *bound_tid, ...);
+typedef void (*kmp_ReductFctPtr)(void *lhsData, void *rhsData);
+typedef void (*kmp_InterWarpCopyFctPtr)(void *src, int32_t warp_num);
+typedef void (*kmp_ShuffleReductFctPtr)(void *rhsData, int16_t lane_id,
+                                        int16_t lane_offset,
+                                        int16_t shortCircuit);
+typedef void (*kmp_CopyToScratchpadFctPtr)(void *reduceData, void *scratchpad,
+                                           int32_t index, int32_t width);
+typedef void (*kmp_LoadReduceFctPtr)(void *reduceData, void *scratchpad,
+                                     int32_t index, int32_t width,
+                                     int32_t reduce);
+typedef void (*kmp_ListGlobalFctPtr)(void *buffer, int idx, void *reduce_data);
+
+// task defs
+typedef struct kmp_TaskDescr kmp_TaskDescr;
+typedef int32_t (*kmp_TaskFctPtr)(int32_t global_tid, kmp_TaskDescr *taskDescr);
+typedef struct kmp_TaskDescr {
+  void *sharedPointerTable;   // ptr to a table of shared var ptrs
+  kmp_TaskFctPtr sub;         // task subroutine
+  int32_t partId;             // unused
+  kmp_TaskFctPtr destructors; // destructor of c++ first private
+} kmp_TaskDescr;
+
+// sync defs
+typedef int32_t kmp_CriticalName[8];
+
+////////////////////////////////////////////////////////////////////////////////
+// external interface
+////////////////////////////////////////////////////////////////////////////////
+
+// parallel
+EXTERN int32_t __kmpc_global_thread_num(kmp_Ident *loc);
+EXTERN void __kmpc_push_num_threads(kmp_Ident *loc, int32_t global_tid,
+                                    int32_t num_threads);
+// simd
+EXTERN void __kmpc_push_simd_limit(kmp_Ident *loc, int32_t global_tid,
+                                   int32_t simd_limit);
+// aee ... not supported
+// EXTERN void __kmpc_fork_call(kmp_Ident *loc, int32_t argc, kmp_ParFctPtr
+// microtask, ...);
+EXTERN void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid);
+EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
+                                           uint32_t global_tid);
+EXTERN uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid);
+
+// proc bind
+EXTERN void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t global_tid,
+                                  int proc_bind);
+EXTERN int omp_get_num_places(void);
+EXTERN int omp_get_place_num_procs(int place_num);
+EXTERN void omp_get_place_proc_ids(int place_num, int *ids);
+EXTERN int omp_get_place_num(void);
+EXTERN int omp_get_partition_num_places(void);
+EXTERN void omp_get_partition_place_nums(int *place_nums);
+
+// for static (no chunk or chunk)
+EXTERN void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
+                                     int32_t sched, int32_t *plastiter,
+                                     int32_t *plower, int32_t *pupper,
+                                     int32_t *pstride, int32_t incr,
+                                     int32_t chunk);
+EXTERN void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
+                                      int32_t sched, int32_t *plastiter,
+                                      uint32_t *plower, uint32_t *pupper,
+                                      int32_t *pstride, int32_t incr,
+                                      int32_t chunk);
+EXTERN void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
+                                     int32_t sched, int32_t *plastiter,
+                                     int64_t *plower, int64_t *pupper,
+                                     int64_t *pstride, int64_t incr,
+                                     int64_t chunk);
+EXTERN void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
+                                      int32_t sched, int32_t *plastiter1,
+                                      uint64_t *plower, uint64_t *pupper,
+                                      int64_t *pstride, int64_t incr,
+                                      int64_t chunk);
+EXTERN
+void __kmpc_for_static_init_4_simple_spmd(kmp_Ident *loc, int32_t global_tid,
+                                          int32_t sched, int32_t *plastiter,
+                                          int32_t *plower, int32_t *pupper,
+                                          int32_t *pstride, int32_t incr,
+                                          int32_t chunk);
+EXTERN
+void __kmpc_for_static_init_4u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
+                                           int32_t sched, int32_t *plastiter,
+                                           uint32_t *plower, uint32_t *pupper,
+                                           int32_t *pstride, int32_t incr,
+                                           int32_t chunk);
+EXTERN
+void __kmpc_for_static_init_8_simple_spmd(kmp_Ident *loc, int32_t global_tid,
+                                          int32_t sched, int32_t *plastiter,
+                                          int64_t *plower, int64_t *pupper,
+                                          int64_t *pstride, int64_t incr,
+                                          int64_t chunk);
+EXTERN
+void __kmpc_for_static_init_8u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
+                                           int32_t sched, int32_t *plastiter1,
+                                           uint64_t *plower, uint64_t *pupper,
+                                           int64_t *pstride, int64_t incr,
+                                           int64_t chunk);
+EXTERN
+void __kmpc_for_static_init_4_simple_generic(kmp_Ident *loc,
+                                             int32_t global_tid, int32_t sched,
+                                             int32_t *plastiter,
+                                             int32_t *plower, int32_t *pupper,
+                                             int32_t *pstride, int32_t incr,
+                                             int32_t chunk);
+EXTERN
+void __kmpc_for_static_init_4u_simple_generic(
+    kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter,
+    uint32_t *plower, uint32_t *pupper, int32_t *pstride, int32_t incr,
+    int32_t chunk);
+EXTERN
+void __kmpc_for_static_init_8_simple_generic(kmp_Ident *loc,
+                                             int32_t global_tid, int32_t sched,
+                                             int32_t *plastiter,
+                                             int64_t *plower, int64_t *pupper,
+                                             int64_t *pstride, int64_t incr,
+                                             int64_t chunk);
+EXTERN
+void __kmpc_for_static_init_8u_simple_generic(
+    kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter1,
+    uint64_t *plower, uint64_t *pupper, int64_t *pstride, int64_t incr,
+    int64_t chunk);
+
+EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid);
+
+// for dynamic
+EXTERN void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t global_tid,
+                                   int32_t sched, int32_t lower, int32_t upper,
+                                   int32_t incr, int32_t chunk);
+EXTERN void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t global_tid,
+                                    int32_t sched, uint32_t lower,
+                                    uint32_t upper, int32_t incr,
+                                    int32_t chunk);
+EXTERN void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t global_tid,
+                                   int32_t sched, int64_t lower, int64_t upper,
+                                   int64_t incr, int64_t chunk);
+EXTERN void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t global_tid,
+                                    int32_t sched, uint64_t lower,
+                                    uint64_t upper, int64_t incr,
+                                    int64_t chunk);
+
+EXTERN int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t global_tid,
+                                  int32_t *plastiter, int32_t *plower,
+                                  int32_t *pupper, int32_t *pstride);
+EXTERN int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t global_tid,
+                                   int32_t *plastiter, uint32_t *plower,
+                                   uint32_t *pupper, int32_t *pstride);
+EXTERN int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t global_tid,
+                                  int32_t *plastiter, int64_t *plower,
+                                  int64_t *pupper, int64_t *pstride);
+EXTERN int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t global_tid,
+                                   int32_t *plastiter, uint64_t *plower,
+                                   uint64_t *pupper, int64_t *pstride);
+
+EXTERN void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t global_tid);
+
+// Support for reducing conditional lastprivate variables
+EXTERN void __kmpc_reduce_conditional_lastprivate(kmp_Ident *loc,
+                                                  int32_t global_tid,
+                                                  int32_t varNum, void *array);
+
+// reduction
+EXTERN void __kmpc_nvptx_end_reduce(int32_t global_tid);
+EXTERN void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
+EXTERN __attribute__((deprecated)) int32_t __kmpc_nvptx_parallel_reduce_nowait(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
+EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
+    kmp_Ident *loc, int32_t global_tid, int32_t num_vars, size_t reduce_size,
+    void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
+    kmp_InterWarpCopyFctPtr cpyFct);
+EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_spmd(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
+EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_generic(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
+EXTERN int32_t __kmpc_nvptx_simd_reduce_nowait(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
+EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
+    kmp_Ident *loc, int32_t global_tid, void *global_buffer,
+    int32_t num_of_records, void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
+    kmp_InterWarpCopyFctPtr cpyFct, kmp_ListGlobalFctPtr lgcpyFct,
+    kmp_ListGlobalFctPtr lgredFct, kmp_ListGlobalFctPtr glcpyFct,
+    kmp_ListGlobalFctPtr glredFct);
+EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
+    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
+EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple_spmd(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
+    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
+EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple_generic(
+    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
+    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
+    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
+EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple(kmp_Ident *loc,
+                                                       int32_t global_tid,
+                                                       kmp_CriticalName *crit);
+EXTERN void __kmpc_nvptx_teams_end_reduce_nowait_simple(kmp_Ident *loc,
+                                                        int32_t global_tid,
+                                                        kmp_CriticalName *crit);
+EXTERN int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size);
+EXTERN int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size);
+
+// sync barrier
+EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid);
+EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid);
+EXTERN void __kmpc_barrier_simple_generic(kmp_Ident *loc_ref, int32_t tid);
+EXTERN int32_t __kmpc_cancel_barrier(kmp_Ident *loc, int32_t global_tid);
+
+// single
+EXTERN int32_t __kmpc_single(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid);
+
+// sync
+EXTERN int32_t __kmpc_master(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_end_master(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_ordered(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_end_ordered(kmp_Ident *loc, int32_t global_tid);
+EXTERN void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
+                            kmp_CriticalName *crit);
+EXTERN void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
+                                kmp_CriticalName *crit);
+EXTERN void __kmpc_flush(kmp_Ident *loc);
+
+// vote
+EXTERN __kmpc_impl_lanemask_t __kmpc_warp_active_thread_mask();
+// syncwarp
+EXTERN void __kmpc_syncwarp(__kmpc_impl_lanemask_t);
+
+// tasks
+EXTERN kmp_TaskDescr *__kmpc_omp_task_alloc(kmp_Ident *loc,
+                                            uint32_t global_tid, int32_t flag,
+                                            size_t sizeOfTaskInclPrivate,
+                                            size_t sizeOfSharedTable,
+                                            kmp_TaskFctPtr sub);
+EXTERN int32_t __kmpc_omp_task(kmp_Ident *loc, uint32_t global_tid,
+                               kmp_TaskDescr *newLegacyTaskDescr);
+EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Ident *loc, uint32_t global_tid,
+                                         kmp_TaskDescr *newLegacyTaskDescr,
+                                         int32_t depNum, void *depList,
+                                         int32_t noAliasDepNum,
+                                         void *noAliasDepList);
+EXTERN void __kmpc_omp_task_begin_if0(kmp_Ident *loc, uint32_t global_tid,
+                                      kmp_TaskDescr *newLegacyTaskDescr);
+EXTERN void __kmpc_omp_task_complete_if0(kmp_Ident *loc, uint32_t global_tid,
+                                         kmp_TaskDescr *newLegacyTaskDescr);
+EXTERN void __kmpc_omp_wait_deps(kmp_Ident *loc, uint32_t global_tid,
+                                 int32_t depNum, void *depList,
+                                 int32_t noAliasDepNum, void *noAliasDepList);
+EXTERN void __kmpc_taskgroup(kmp_Ident *loc, uint32_t global_tid);
+EXTERN void __kmpc_end_taskgroup(kmp_Ident *loc, uint32_t global_tid);
+EXTERN int32_t __kmpc_omp_taskyield(kmp_Ident *loc, uint32_t global_tid,
+                                    int end_part);
+EXTERN int32_t __kmpc_omp_taskwait(kmp_Ident *loc, uint32_t global_tid);
+EXTERN void __kmpc_taskloop(kmp_Ident *loc, uint32_t global_tid,
+                            kmp_TaskDescr *newKmpTaskDescr, int if_val,
+                            uint64_t *lb, uint64_t *ub, int64_t st, int nogroup,
+                            int32_t sched, uint64_t grainsize, void *task_dup);
+
+// cancel
+EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
+                                        int32_t cancelVal);
+EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
+                             int32_t cancelVal);
+
+// non standard
+EXTERN void __kmpc_kernel_init_params(void *ReductionScratchpadPtr);
+EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
+EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized);
+EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
+                                    int16_t RequiresDataSharing);
+EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit();
+EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime);
+EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
+                                           int16_t IsOMPRuntimeInitialized);
+EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
+                                   int16_t IsOMPRuntimeInitialized);
+EXTERN void __kmpc_kernel_end_parallel();
+EXTERN bool __kmpc_kernel_convergent_parallel(void *buffer,
+                                              __kmpc_impl_lanemask_t Mask,
+                                              bool *IsFinal,
+                                              int32_t *LaneSource);
+EXTERN void __kmpc_kernel_end_convergent_parallel(void *buffer);
+EXTERN bool __kmpc_kernel_convergent_simd(void *buffer,
+                                          __kmpc_impl_lanemask_t Mask,
+                                          bool *IsFinal, int32_t *LaneSource,
+                                          int32_t *LaneId, int32_t *NumLanes);
+EXTERN void __kmpc_kernel_end_convergent_simd(void *buffer);
+
+
+EXTERN void __kmpc_data_sharing_init_stack();
+EXTERN void __kmpc_data_sharing_init_stack_spmd();
+EXTERN void *__kmpc_data_sharing_coalesced_push_stack(size_t size,
+    int16_t UseSharedMemory);
+EXTERN void *__kmpc_data_sharing_push_stack(size_t size, int16_t UseSharedMemory);
+EXTERN void __kmpc_data_sharing_pop_stack(void *a);
+EXTERN void __kmpc_begin_sharing_variables(void ***GlobalArgs, size_t nArgs);
+EXTERN void __kmpc_end_sharing_variables();
+EXTERN void __kmpc_get_shared_variables(void ***GlobalArgs);
+
+// The slot used for data sharing by the master and worker threads. We use a
+// complete (default size version and an incomplete one so that we allow sizes
+// greater than the default).
+struct __kmpc_data_sharing_slot {
+  __kmpc_data_sharing_slot *Next;
+  __kmpc_data_sharing_slot *Prev;
+  void *PrevSlotStackPtr;
+  void *DataEnd;
+  char Data[];
+};
+EXTERN void
+__kmpc_initialize_data_sharing_environment(__kmpc_data_sharing_slot *RootS,
+                                           size_t InitialDataSize);
+EXTERN void *__kmpc_data_sharing_environment_begin(
+    __kmpc_data_sharing_slot **SavedSharedSlot, void **SavedSharedStack,
+    void **SavedSharedFrame, __kmpc_impl_lanemask_t *SavedActiveThreads,
+    size_t SharingDataSize, size_t SharingDefaultDataSize,
+    int16_t IsOMPRuntimeInitialized);
+EXTERN void __kmpc_data_sharing_environment_end(
+    __kmpc_data_sharing_slot **SavedSharedSlot, void **SavedSharedStack,
+    void **SavedSharedFrame, __kmpc_impl_lanemask_t *SavedActiveThreads,
+    int32_t IsEntryPoint);
+
+EXTERN void *
+__kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
+                                          int16_t IsOMPRuntimeInitialized);
+
+// SPMD execution mode interrogation function.
+EXTERN int8_t __kmpc_is_spmd_exec_mode();
+
+EXTERN void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode,
+                                          const void *buf, size_t size,
+                                          int16_t is_shared, const void **res);
+
+EXTERN void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
+                                              int16_t is_shared);
+
+#endif

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt?rev=374919&r1=374918&r2=374919&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt Tue Oct 15 10:15:26 2019
@@ -35,6 +35,10 @@ if(CUDA_HOST_COMPILER MATCHES clang)
   set(CUDA_HOST_COMPILER "${LIBOMPTARGET_NVPTX_ALTERNATE_GCC_HOST_COMPILER}" CACHE FILEPATH "" FORCE)
 endif()
 
+get_filename_component(devicertl_base_directory
+  ${CMAKE_CURRENT_SOURCE_DIR}
+  DIRECTORY)
+
 if(LIBOMPTARGET_DEP_CUDA_FOUND)
   libomptarget_say("Building CUDA offloading device RTL.")
 
@@ -83,7 +87,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
   # yet supported by the CUDA toolchain on the device.
   set(BUILD_SHARED_LIBS OFF)
   set(CUDA_SEPARABLE_COMPILATION ON)
-
+  list(APPEND CUDA_NVCC_FLAGS -I${devicertl_base_directory})
   cuda_add_library(omptarget-nvptx STATIC ${cuda_src_files} ${omp_data_objects}
       OPTIONS ${CUDA_ARCH} ${CUDA_DEBUG})
 
@@ -117,7 +121,8 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
     libomptarget_say("Building CUDA LLVM bitcode offloading device RTL.")
 
     # Set flags for LLVM Bitcode compilation.
-    set(bc_flags ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER_FLAGS})
+    set(bc_flags ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER_FLAGS}
+                 -I${devicertl_base_directory})
     if(${LIBOMPTARGET_NVPTX_DEBUG})
       set(bc_flags ${bc_flags} -DOMPTARGET_NVPTX_DEBUG=-1)
     else()

Removed: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h?rev=374918&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h (removed)
@@ -1,538 +0,0 @@
-//===------- interface.h - NVPTX OpenMP interface definitions ---- 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 debug macros to be used in the application.
-//
-//  This file contains all the definitions that are relevant to
-//  the interface. The first section contains the interface as
-//  declared by OpenMP.  The second section includes the compiler
-//  specific interfaces.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef _INTERFACES_H_
-#define _INTERFACES_H_
-
-#include "option.h"
-#include "target_impl.h"
-
-////////////////////////////////////////////////////////////////////////////////
-// OpenMP interface
-////////////////////////////////////////////////////////////////////////////////
-
-typedef uint32_t omp_lock_t;      /* arbitrary type of the right length */
-typedef uint64_t omp_nest_lock_t; /* arbitrary type of the right length */
-
-typedef enum omp_sched_t {
-  omp_sched_static = 1,  /* chunkSize >0 */
-  omp_sched_dynamic = 2, /* chunkSize >0 */
-  omp_sched_guided = 3,  /* chunkSize >0 */
-  omp_sched_auto = 4,    /* no chunkSize */
-} omp_sched_t;
-
-typedef enum omp_proc_bind_t {
-  omp_proc_bind_false = 0,
-  omp_proc_bind_true = 1,
-  omp_proc_bind_master = 2,
-  omp_proc_bind_close = 3,
-  omp_proc_bind_spread = 4
-} omp_proc_bind_t;
-
-EXTERN double omp_get_wtick(void);
-EXTERN double omp_get_wtime(void);
-
-EXTERN void omp_set_num_threads(int num);
-EXTERN int omp_get_num_threads(void);
-EXTERN int omp_get_max_threads(void);
-EXTERN int omp_get_thread_limit(void);
-EXTERN int omp_get_thread_num(void);
-EXTERN int omp_get_num_procs(void);
-EXTERN int omp_in_parallel(void);
-EXTERN int omp_in_final(void);
-EXTERN void omp_set_dynamic(int flag);
-EXTERN int omp_get_dynamic(void);
-EXTERN void omp_set_nested(int flag);
-EXTERN int omp_get_nested(void);
-EXTERN void omp_set_max_active_levels(int level);
-EXTERN int omp_get_max_active_levels(void);
-EXTERN int omp_get_level(void);
-EXTERN int omp_get_active_level(void);
-EXTERN int omp_get_ancestor_thread_num(int level);
-EXTERN int omp_get_team_size(int level);
-
-EXTERN void omp_init_lock(omp_lock_t *lock);
-EXTERN void omp_init_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_destroy_lock(omp_lock_t *lock);
-EXTERN void omp_destroy_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_set_lock(omp_lock_t *lock);
-EXTERN void omp_set_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_unset_lock(omp_lock_t *lock);
-EXTERN void omp_unset_nest_lock(omp_nest_lock_t *lock);
-EXTERN int omp_test_lock(omp_lock_t *lock);
-EXTERN int omp_test_nest_lock(omp_nest_lock_t *lock);
-
-EXTERN void omp_get_schedule(omp_sched_t *kind, int *modifier);
-EXTERN void omp_set_schedule(omp_sched_t kind, int modifier);
-EXTERN omp_proc_bind_t omp_get_proc_bind(void);
-EXTERN int omp_get_cancellation(void);
-EXTERN void omp_set_default_device(int deviceId);
-EXTERN int omp_get_default_device(void);
-EXTERN int omp_get_num_devices(void);
-EXTERN int omp_get_num_teams(void);
-EXTERN int omp_get_team_num(void);
-EXTERN int omp_is_initial_device(void);
-EXTERN int omp_get_initial_device(void);
-EXTERN int omp_get_max_task_priority(void);
-
-////////////////////////////////////////////////////////////////////////////////
-// file below is swiped from kmpc host interface
-////////////////////////////////////////////////////////////////////////////////
-
-////////////////////////////////////////////////////////////////////////////////
-// kmp specifc types
-////////////////////////////////////////////////////////////////////////////////
-
-typedef enum kmp_sched_t {
-  kmp_sched_static_chunk = 33,
-  kmp_sched_static_nochunk = 34,
-  kmp_sched_dynamic = 35,
-  kmp_sched_guided = 36,
-  kmp_sched_runtime = 37,
-  kmp_sched_auto = 38,
-
-  kmp_sched_static_balanced_chunk = 45,
-
-  kmp_sched_static_ordered = 65,
-  kmp_sched_static_nochunk_ordered = 66,
-  kmp_sched_dynamic_ordered = 67,
-  kmp_sched_guided_ordered = 68,
-  kmp_sched_runtime_ordered = 69,
-  kmp_sched_auto_ordered = 70,
-
-  kmp_sched_distr_static_chunk = 91,
-  kmp_sched_distr_static_nochunk = 92,
-  kmp_sched_distr_static_chunk_sched_static_chunkone = 93,
-
-  kmp_sched_default = kmp_sched_static_nochunk,
-  kmp_sched_unordered_first = kmp_sched_static_chunk,
-  kmp_sched_unordered_last = kmp_sched_auto,
-  kmp_sched_ordered_first = kmp_sched_static_ordered,
-  kmp_sched_ordered_last = kmp_sched_auto_ordered,
-  kmp_sched_distribute_first = kmp_sched_distr_static_chunk,
-  kmp_sched_distribute_last =
-      kmp_sched_distr_static_chunk_sched_static_chunkone,
-
-  /* Support for OpenMP 4.5 monotonic and nonmonotonic schedule modifiers.
-   * Since we need to distinguish the three possible cases (no modifier,
-   * monotonic modifier, nonmonotonic modifier), we need separate bits for
-   * each modifier. The absence of monotonic does not imply nonmonotonic,
-   * especially since 4.5 says that the behaviour of the "no modifier" case
-   * is implementation defined in 4.5, but will become "nonmonotonic" in 5.0.
-   *
-   * Since we're passing a full 32 bit value, we can use a couple of high
-   * bits for these flags; out of paranoia we avoid the sign bit.
-   *
-   * These modifiers can be or-ed into non-static schedules by the compiler
-   * to pass the additional information. They will be stripped early in the
-   * processing in __kmp_dispatch_init when setting up schedules, so
-   * most of the code won't ever see schedules with these bits set.
-   */
-  kmp_sched_modifier_monotonic = (1 << 29),
-  /**< Set if the monotonic schedule modifier was present */
-  kmp_sched_modifier_nonmonotonic = (1 << 30),
-/**< Set if the nonmonotonic schedule modifier was present */
-
-#define SCHEDULE_WITHOUT_MODIFIERS(s)                                          \
-  (enum kmp_sched_t)(                                                          \
-      (s) & ~(kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic))
-#define SCHEDULE_HAS_MONOTONIC(s) (((s)&kmp_sched_modifier_monotonic) != 0)
-#define SCHEDULE_HAS_NONMONOTONIC(s)                                           \
-  (((s)&kmp_sched_modifier_nonmonotonic) != 0)
-#define SCHEDULE_HAS_NO_MODIFIERS(s)                                           \
-  (((s) & (kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic)) == \
-   0)
-
-} kmp_sched_t;
-
-/*!
- * Enum for accesseing the reserved_2 field of the ident_t struct below.
- */
-enum {
-  /*! Bit set to 1 when in SPMD mode. */
-  KMP_IDENT_SPMD_MODE = 0x01,
-  /*! Bit set to 1 when a simplified runtime is used. */
-  KMP_IDENT_SIMPLE_RT_MODE = 0x02,
-};
-
-/*!
- * The ident structure that describes a source location.
- * The struct is identical to the one in the kmp.h file.
- * We maintain the same data structure for compatibility.
- */
-typedef int kmp_int32;
-typedef struct ident {
-  kmp_int32 reserved_1; /**<  might be used in Fortran; see above  */
-  kmp_int32 flags; /**<  also f.flags; KMP_IDENT_xxx flags; KMP_IDENT_KMPC
-                      identifies this union member  */
-  kmp_int32 reserved_2; /**<  not really used in Fortran any more; see above */
-  kmp_int32 reserved_3; /**<  source[4] in Fortran, do not use for C++  */
-  char const *psource; /**<  String describing the source location.
-                       The string is composed of semi-colon separated fields
-                       which describe the source file, the function and a pair
-                       of line numbers that delimit the construct. */
-} ident_t;
-
-// parallel defs
-typedef ident_t kmp_Ident;
-typedef void (*kmp_ParFctPtr)(int32_t *global_tid, int32_t *bound_tid, ...);
-typedef void (*kmp_ReductFctPtr)(void *lhsData, void *rhsData);
-typedef void (*kmp_InterWarpCopyFctPtr)(void *src, int32_t warp_num);
-typedef void (*kmp_ShuffleReductFctPtr)(void *rhsData, int16_t lane_id,
-                                        int16_t lane_offset,
-                                        int16_t shortCircuit);
-typedef void (*kmp_CopyToScratchpadFctPtr)(void *reduceData, void *scratchpad,
-                                           int32_t index, int32_t width);
-typedef void (*kmp_LoadReduceFctPtr)(void *reduceData, void *scratchpad,
-                                     int32_t index, int32_t width,
-                                     int32_t reduce);
-typedef void (*kmp_ListGlobalFctPtr)(void *buffer, int idx, void *reduce_data);
-
-// task defs
-typedef struct kmp_TaskDescr kmp_TaskDescr;
-typedef int32_t (*kmp_TaskFctPtr)(int32_t global_tid, kmp_TaskDescr *taskDescr);
-typedef struct kmp_TaskDescr {
-  void *sharedPointerTable;   // ptr to a table of shared var ptrs
-  kmp_TaskFctPtr sub;         // task subroutine
-  int32_t partId;             // unused
-  kmp_TaskFctPtr destructors; // destructor of c++ first private
-} kmp_TaskDescr;
-
-// sync defs
-typedef int32_t kmp_CriticalName[8];
-
-////////////////////////////////////////////////////////////////////////////////
-// external interface
-////////////////////////////////////////////////////////////////////////////////
-
-// parallel
-EXTERN int32_t __kmpc_global_thread_num(kmp_Ident *loc);
-EXTERN void __kmpc_push_num_threads(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t num_threads);
-// simd
-EXTERN void __kmpc_push_simd_limit(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t simd_limit);
-// aee ... not supported
-// EXTERN void __kmpc_fork_call(kmp_Ident *loc, int32_t argc, kmp_ParFctPtr
-// microtask, ...);
-EXTERN void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
-                                           uint32_t global_tid);
-EXTERN uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid);
-
-// proc bind
-EXTERN void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t global_tid,
-                                  int proc_bind);
-EXTERN int omp_get_num_places(void);
-EXTERN int omp_get_place_num_procs(int place_num);
-EXTERN void omp_get_place_proc_ids(int place_num, int *ids);
-EXTERN int omp_get_place_num(void);
-EXTERN int omp_get_partition_num_places(void);
-EXTERN void omp_get_partition_place_nums(int *place_nums);
-
-// for static (no chunk or chunk)
-EXTERN void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
-                                     int32_t sched, int32_t *plastiter,
-                                     int32_t *plower, int32_t *pupper,
-                                     int32_t *pstride, int32_t incr,
-                                     int32_t chunk);
-EXTERN void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
-                                      int32_t sched, int32_t *plastiter,
-                                      uint32_t *plower, uint32_t *pupper,
-                                      int32_t *pstride, int32_t incr,
-                                      int32_t chunk);
-EXTERN void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
-                                     int32_t sched, int32_t *plastiter,
-                                     int64_t *plower, int64_t *pupper,
-                                     int64_t *pstride, int64_t incr,
-                                     int64_t chunk);
-EXTERN void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
-                                      int32_t sched, int32_t *plastiter1,
-                                      uint64_t *plower, uint64_t *pupper,
-                                      int64_t *pstride, int64_t incr,
-                                      int64_t chunk);
-EXTERN
-void __kmpc_for_static_init_4_simple_spmd(kmp_Ident *loc, int32_t global_tid,
-                                          int32_t sched, int32_t *plastiter,
-                                          int32_t *plower, int32_t *pupper,
-                                          int32_t *pstride, int32_t incr,
-                                          int32_t chunk);
-EXTERN
-void __kmpc_for_static_init_4u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
-                                           int32_t sched, int32_t *plastiter,
-                                           uint32_t *plower, uint32_t *pupper,
-                                           int32_t *pstride, int32_t incr,
-                                           int32_t chunk);
-EXTERN
-void __kmpc_for_static_init_8_simple_spmd(kmp_Ident *loc, int32_t global_tid,
-                                          int32_t sched, int32_t *plastiter,
-                                          int64_t *plower, int64_t *pupper,
-                                          int64_t *pstride, int64_t incr,
-                                          int64_t chunk);
-EXTERN
-void __kmpc_for_static_init_8u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
-                                           int32_t sched, int32_t *plastiter1,
-                                           uint64_t *plower, uint64_t *pupper,
-                                           int64_t *pstride, int64_t incr,
-                                           int64_t chunk);
-EXTERN
-void __kmpc_for_static_init_4_simple_generic(kmp_Ident *loc,
-                                             int32_t global_tid, int32_t sched,
-                                             int32_t *plastiter,
-                                             int32_t *plower, int32_t *pupper,
-                                             int32_t *pstride, int32_t incr,
-                                             int32_t chunk);
-EXTERN
-void __kmpc_for_static_init_4u_simple_generic(
-    kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter,
-    uint32_t *plower, uint32_t *pupper, int32_t *pstride, int32_t incr,
-    int32_t chunk);
-EXTERN
-void __kmpc_for_static_init_8_simple_generic(kmp_Ident *loc,
-                                             int32_t global_tid, int32_t sched,
-                                             int32_t *plastiter,
-                                             int64_t *plower, int64_t *pupper,
-                                             int64_t *pstride, int64_t incr,
-                                             int64_t chunk);
-EXTERN
-void __kmpc_for_static_init_8u_simple_generic(
-    kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter1,
-    uint64_t *plower, uint64_t *pupper, int64_t *pstride, int64_t incr,
-    int64_t chunk);
-
-EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid);
-
-// for dynamic
-EXTERN void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t sched, int32_t lower, int32_t upper,
-                                   int32_t incr, int32_t chunk);
-EXTERN void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t sched, uint32_t lower,
-                                    uint32_t upper, int32_t incr,
-                                    int32_t chunk);
-EXTERN void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t sched, int64_t lower, int64_t upper,
-                                   int64_t incr, int64_t chunk);
-EXTERN void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t sched, uint64_t lower,
-                                    uint64_t upper, int64_t incr,
-                                    int64_t chunk);
-
-EXTERN int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t global_tid,
-                                  int32_t *plastiter, int32_t *plower,
-                                  int32_t *pupper, int32_t *pstride);
-EXTERN int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t *plastiter, uint32_t *plower,
-                                   uint32_t *pupper, int32_t *pstride);
-EXTERN int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t global_tid,
-                                  int32_t *plastiter, int64_t *plower,
-                                  int64_t *pupper, int64_t *pstride);
-EXTERN int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t *plastiter, uint64_t *plower,
-                                   uint64_t *pupper, int64_t *pstride);
-
-EXTERN void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t global_tid);
-
-// Support for reducing conditional lastprivate variables
-EXTERN void __kmpc_reduce_conditional_lastprivate(kmp_Ident *loc,
-                                                  int32_t global_tid,
-                                                  int32_t varNum, void *array);
-
-// reduction
-EXTERN void __kmpc_nvptx_end_reduce(int32_t global_tid);
-EXTERN void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
-EXTERN __attribute__((deprecated)) int32_t __kmpc_nvptx_parallel_reduce_nowait(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
-    kmp_Ident *loc, int32_t global_tid, int32_t num_vars, size_t reduce_size,
-    void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
-    kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_spmd(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_simple_generic(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_simd_reduce_nowait(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
-    kmp_Ident *loc, int32_t global_tid, void *global_buffer,
-    int32_t num_of_records, void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
-    kmp_InterWarpCopyFctPtr cpyFct, kmp_ListGlobalFctPtr lgcpyFct,
-    kmp_ListGlobalFctPtr lgredFct, kmp_ListGlobalFctPtr glcpyFct,
-    kmp_ListGlobalFctPtr glredFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
-    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple_spmd(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
-    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple_generic(
-    int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
-    kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct,
-    kmp_CopyToScratchpadFctPtr sratchFct, kmp_LoadReduceFctPtr ldFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple(kmp_Ident *loc,
-                                                       int32_t global_tid,
-                                                       kmp_CriticalName *crit);
-EXTERN void __kmpc_nvptx_teams_end_reduce_nowait_simple(kmp_Ident *loc,
-                                                        int32_t global_tid,
-                                                        kmp_CriticalName *crit);
-EXTERN int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size);
-EXTERN int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size);
-
-// sync barrier
-EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid);
-EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid);
-EXTERN void __kmpc_barrier_simple_generic(kmp_Ident *loc_ref, int32_t tid);
-EXTERN int32_t __kmpc_cancel_barrier(kmp_Ident *loc, int32_t global_tid);
-
-// single
-EXTERN int32_t __kmpc_single(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid);
-
-// sync
-EXTERN int32_t __kmpc_master(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_master(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_ordered(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_ordered(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
-                            kmp_CriticalName *crit);
-EXTERN void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
-                                kmp_CriticalName *crit);
-EXTERN void __kmpc_flush(kmp_Ident *loc);
-
-// vote
-EXTERN __kmpc_impl_lanemask_t __kmpc_warp_active_thread_mask();
-// syncwarp
-EXTERN void __kmpc_syncwarp(__kmpc_impl_lanemask_t);
-
-// tasks
-EXTERN kmp_TaskDescr *__kmpc_omp_task_alloc(kmp_Ident *loc,
-                                            uint32_t global_tid, int32_t flag,
-                                            size_t sizeOfTaskInclPrivate,
-                                            size_t sizeOfSharedTable,
-                                            kmp_TaskFctPtr sub);
-EXTERN int32_t __kmpc_omp_task(kmp_Ident *loc, uint32_t global_tid,
-                               kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Ident *loc, uint32_t global_tid,
-                                         kmp_TaskDescr *newLegacyTaskDescr,
-                                         int32_t depNum, void *depList,
-                                         int32_t noAliasDepNum,
-                                         void *noAliasDepList);
-EXTERN void __kmpc_omp_task_begin_if0(kmp_Ident *loc, uint32_t global_tid,
-                                      kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN void __kmpc_omp_task_complete_if0(kmp_Ident *loc, uint32_t global_tid,
-                                         kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN void __kmpc_omp_wait_deps(kmp_Ident *loc, uint32_t global_tid,
-                                 int32_t depNum, void *depList,
-                                 int32_t noAliasDepNum, void *noAliasDepList);
-EXTERN void __kmpc_taskgroup(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_end_taskgroup(kmp_Ident *loc, uint32_t global_tid);
-EXTERN int32_t __kmpc_omp_taskyield(kmp_Ident *loc, uint32_t global_tid,
-                                    int end_part);
-EXTERN int32_t __kmpc_omp_taskwait(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_taskloop(kmp_Ident *loc, uint32_t global_tid,
-                            kmp_TaskDescr *newKmpTaskDescr, int if_val,
-                            uint64_t *lb, uint64_t *ub, int64_t st, int nogroup,
-                            int32_t sched, uint64_t grainsize, void *task_dup);
-
-// cancel
-EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
-                                        int32_t cancelVal);
-EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
-                             int32_t cancelVal);
-
-// non standard
-EXTERN void __kmpc_kernel_init_params(void *ReductionScratchpadPtr);
-EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
-EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized);
-EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
-                                    int16_t RequiresDataSharing);
-EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit();
-EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime);
-EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn,
-                                           int16_t IsOMPRuntimeInitialized);
-EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
-                                   int16_t IsOMPRuntimeInitialized);
-EXTERN void __kmpc_kernel_end_parallel();
-EXTERN bool __kmpc_kernel_convergent_parallel(void *buffer,
-                                              __kmpc_impl_lanemask_t Mask,
-                                              bool *IsFinal,
-                                              int32_t *LaneSource);
-EXTERN void __kmpc_kernel_end_convergent_parallel(void *buffer);
-EXTERN bool __kmpc_kernel_convergent_simd(void *buffer,
-                                          __kmpc_impl_lanemask_t Mask,
-                                          bool *IsFinal, int32_t *LaneSource,
-                                          int32_t *LaneId, int32_t *NumLanes);
-EXTERN void __kmpc_kernel_end_convergent_simd(void *buffer);
-
-
-EXTERN void __kmpc_data_sharing_init_stack();
-EXTERN void __kmpc_data_sharing_init_stack_spmd();
-EXTERN void *__kmpc_data_sharing_coalesced_push_stack(size_t size,
-    int16_t UseSharedMemory);
-EXTERN void *__kmpc_data_sharing_push_stack(size_t size, int16_t UseSharedMemory);
-EXTERN void __kmpc_data_sharing_pop_stack(void *a);
-EXTERN void __kmpc_begin_sharing_variables(void ***GlobalArgs, size_t nArgs);
-EXTERN void __kmpc_end_sharing_variables();
-EXTERN void __kmpc_get_shared_variables(void ***GlobalArgs);
-
-// The slot used for data sharing by the master and worker threads. We use a
-// complete (default size version and an incomplete one so that we allow sizes
-// greater than the default).
-struct __kmpc_data_sharing_slot {
-  __kmpc_data_sharing_slot *Next;
-  __kmpc_data_sharing_slot *Prev;
-  void *PrevSlotStackPtr;
-  void *DataEnd;
-  char Data[];
-};
-EXTERN void
-__kmpc_initialize_data_sharing_environment(__kmpc_data_sharing_slot *RootS,
-                                           size_t InitialDataSize);
-EXTERN void *__kmpc_data_sharing_environment_begin(
-    __kmpc_data_sharing_slot **SavedSharedSlot, void **SavedSharedStack,
-    void **SavedSharedFrame, __kmpc_impl_lanemask_t *SavedActiveThreads,
-    size_t SharingDataSize, size_t SharingDefaultDataSize,
-    int16_t IsOMPRuntimeInitialized);
-EXTERN void __kmpc_data_sharing_environment_end(
-    __kmpc_data_sharing_slot **SavedSharedSlot, void **SavedSharedStack,
-    void **SavedSharedFrame, __kmpc_impl_lanemask_t *SavedActiveThreads,
-    int32_t IsEntryPoint);
-
-EXTERN void *
-__kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
-                                          int16_t IsOMPRuntimeInitialized);
-
-// SPMD execution mode interrogation function.
-EXTERN int8_t __kmpc_is_spmd_exec_mode();
-
-EXTERN void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode,
-                                          const void *buf, size_t size,
-                                          int16_t is_shared, const void **res);
-
-EXTERN void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
-                                              int16_t is_shared);
-
-#endif

Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h?rev=374919&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h Tue Oct 15 10:15:26 2019
@@ -0,0 +1,17 @@
+//===--- nvptx_interface.h - OpenMP interface definitions -------- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _NVPTX_INTERFACE_H_
+#define _NVPTX_INTERFACE_H_
+
+#include <stdint.h>
+
+#define EXTERN extern "C" __device__
+typedef uint32_t __kmpc_impl_lanemask_t;
+
+#endif

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/option.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/option.h?rev=374919&r1=374918&r2=374919&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/option.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/option.h Tue Oct 15 10:15:26 2019
@@ -12,6 +12,8 @@
 #ifndef _OPTION_H_
 #define _OPTION_H_
 
+#include "interface.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
@@ -54,7 +56,6 @@
 // misc options (by def everythig here is device)
 ////////////////////////////////////////////////////////////////////////////////
 
-#define EXTERN extern "C" __device__
 #define INLINE __forceinline__ __device__
 #define NOINLINE __noinline__ __device__
 #ifndef TRUE

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h?rev=374919&r1=374918&r2=374919&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/target_impl.h Tue Oct 15 10:15:26 2019
@@ -26,7 +26,6 @@ INLINE uint64_t __kmpc_impl_pack(uint32_
   return val;
 }
 
-typedef uint32_t __kmpc_impl_lanemask_t;
 static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
     UINT32_C(0xffffffff);
 




More information about the Openmp-commits mailing list