[Openmp-commits] [openmp] 74bb5ee - [nfc][libomptarget] Move smid() into target_impl

[via Openmp-commits]

Author: Jon Chesterfield
Date: 2019-10-30T13:39:15Z
New Revision: 74bb5ee67413db7e3e3351b7fde65db7e4568e02

URL: https://github.com/llvm/llvm-project/commit/74bb5ee67413db7e3e3351b7fde65db7e4568e02
DIFF: https://github.com/llvm/llvm-project/commit/74bb5ee67413db7e3e3351b7fde65db7e4568e02.diff

LOG: [nfc][libomptarget] Move smid() into target_impl

Summary: [nfc][libomptarget] Move smid() into target_impl

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: openmp-commits

Tags: #openmp

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

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
index c84c0554424a..0481bdf854fb 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
@@ -25,12 +25,6 @@ extern __device__
 // init entry points
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE static unsigned smid() {
-  unsigned id;
-  asm("mov.u32 %0, %%smid;" : "=r"(id));
-  return id;
-}
-
 EXTERN void __kmpc_kernel_init_params(void *Ptr) {
   PRINT(LD_IO, "call to __kmpc_kernel_init_params with version %f\n",
         OMPTARGET_NVPTX_VERSION);
@@ -53,7 +47,7 @@ EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime) {
   PRINT0(LD_IO, "call to __kmpc_kernel_init for master\n");
 
   // Get a state object from the queue.
-  int slot = smid() % MAX_SM;
+  int slot = __kmpc_impl_smid() % MAX_SM;
   usedSlotIdx = slot;
   omptarget_nvptx_threadPrivateContext =
       omptarget_nvptx_device_State[slot].Dequeue();
@@ -98,7 +92,7 @@ EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime,
                                                   : RuntimeUninitialized);
   int threadId = GetThreadIdInBlock();
   if (threadId == 0) {
-    usedSlotIdx = smid() % MAX_SM;
+    usedSlotIdx = __kmpc_impl_smid() % MAX_SM;
     parallelLevel[0] =
         1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
   } else if (GetLaneId() == 0) {

diff  --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
index bd212ca18cf6..bbce9f1c5119 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -91,6 +91,12 @@ INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
   return res;
 }
 
+INLINE uint32_t __kmpc_impl_smid() {
+  uint32_t id;
+  asm("mov.u32 %0, %%smid;" : "=r"(id));
+  return id;
+}
+
 INLINE uint32_t __kmpc_impl_ffs(uint32_t x) { return __ffs(x); }
 
 INLINE uint32_t __kmpc_impl_popc(uint32_t x) { return __popc(x); }