[Openmp-commits] [PATCH] D69424: [NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: ABataev, jdoerfert, grokos.
Herald added a project: OpenMP.
Herald added a subscriber: openmp-commits.

[NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h

Strictly there is one remaining difference wrt amdgcn - parallelLevel is
volatile qualified on amdgcn and not on nvptx. Determining whether this is
correct - and how to represent the different semantics of 'volatile' under
various conditions - is beyond the scope of this code motion patch.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D69424

Files:
  openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
  openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h


Index: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -12,10 +12,30 @@
 #ifndef _TARGET_IMPL_H_
 #define _TARGET_IMPL_H_
 
+#include <cuda.h>
 #include <stdint.h>
 
 #include "option.h"
 
+// Data sharing related quantities, need to match what is used in the compiler.
+enum DATA_SHARING_SIZES {
+  // The maximum number of workers in a kernel.
+  DS_Max_Worker_Threads = 992,
+  // The size reserved for data in a shared memory slot.
+  DS_Slot_Size = 256,
+  // The slot size that should be reserved for a working warp.
+  DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
+  // The maximum number of warps in use
+  DS_Max_Warp_Number = 32,
+  // The size of the preallocated shared memory buffer per team
+  DS_Shared_Memory_Size = 128,
+};
+
+/// Device environment data
+struct omptarget_device_environmentTy {
+  int32_t debug_level;
+};
+
 INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "l"(val));
 }
Index: openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -15,16 +15,13 @@
 #define __OMPTARGET_NVPTX_H
 
 // std includes
-#include <stdint.h>
-#include <stdlib.h>
-
 #include <inttypes.h>
-
-// cuda includes
-#include <cuda.h>
 #include <math.h>
+#include <stdint.h>
+#include <stdlib.h>
 
 // local includes
+#include "target_impl.h"
 #include "debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "option.h"    // choices we have
@@ -86,20 +83,6 @@
 extern __device__ __shared__ omptarget_nvptx_SharedArgs
     omptarget_nvptx_globalArgs;
 
-// Data sharing related quantities, need to match what is used in the compiler.
-enum DATA_SHARING_SIZES {
-  // The maximum number of workers in a kernel.
-  DS_Max_Worker_Threads = 992,
-  // The size reserved for data in a shared memory slot.
-  DS_Slot_Size = 256,
-  // The slot size that should be reserved for a working warp.
-  DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
-  // The maximum number of warps in use
-  DS_Max_Warp_Number = 32,
-  // The size of the preallocated shared memory buffer per team
-  DS_Shared_Memory_Size = 128,
-};
-
 // Data structure to keep in shared memory that traces the current slot, stack,
 // and frame pointer as well as the active threads that didn't exit the current
 // environment.
@@ -343,11 +326,6 @@
   uint64_t cnt;
 };
 
-/// Device envrionment data
-struct omptarget_device_environmentTy {
-  int32_t debug_level;
-};
-
 /// Memory manager for statically allocated memory.
 class omptarget_nvptx_SimpleMemoryManager {
 private:


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D69424.226379.patch
Type: text/x-patch
Size: 2986 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20191025/8918c4d1/attachment-0001.bin>