[libomptarget] Replace INLINE with DEVICE to work around nvcc deleting forceinline symbols

[JonChesterfield]

Summary

[libomptarget] Replace INLINE with DEVICE to work around nvcc deleting forceinline symbols

There is a problem with some installs of nvcc deleting the forceinline symbols
from support.cu. Adding the people who reported this as reviewers. This patch is
an educated guess at a fix, posting it without a local repro for the failure. See D69859 for an alternative.

[ABataev]

What about the inlining, will it affect it?

[JonChesterfield]

In D69857#1734286, @ABataev wrote:

What about the inlining, will it affect it?

I've given up on guessing what nvcc will do. Changing the INLINE macro to __inline__ (and away from the possibly buggy __always_inline__) is a nicer fix if it works.

Working theories:

• nvcc requires declaration and definition to have identical attributes
• nvcc __always_inline__ is broken
• nvcc changes the semantics of code based on #include hierarchy

I'm still trying to reproduce the missing symbols locally so don't have much scope for experimentation.

^ for always_inline, please read forceinline...

[ABataev]

In D69857#1734288, @JonChesterfield wrote:

In D69857#1734286, @ABataev wrote:

What about the inlining, will it affect it?

I've given up on guessing what nvcc will do. Changing the INLINE macro to __inline__ (and away from the possibly buggy __always_inline__) is a nicer fix if it works.

Working theories:

• nvcc requires declaration and definition to have identical attributes
• nvcc __always_inline__ is broken
• nvcc changes the semantics of code based on #include hierarchy

I'm still trying to reproduce the missing symbols locally so don't have much scope for experimentation.

Try to replace with __inline__, I will try this patch locally with my version of Cuda. Also, I don't think it is nvcc, I think we get something bad with clang/LLVM inlining, when we try to link against .bc library.

[JonChesterfield]

Try to replace with __inline__, I will try this patch locally with my version of Cuda. Also, I don't think it is nvcc, I think we get something bad with clang/LLVM inlining, when we try to link against .bc library.

Thanks. D69859 for the forceinline to inline change. That's appealing in that we used to use inline.

[ABataev]

This one ends up with the multiply defined symbols instead :)

[JonChesterfield]

In D69857#1734541, @ABataev wrote:

This one ends up with the multiply defined symbols instead :)

How? The functions are only defined in support.cu and it's only compiled once. Where are the duplicates from?

[ABataev]

In D69857#1734561, @JonChesterfield wrote:

In D69857#1734541, @ABataev wrote:

This one ends up with the multiply defined symbols instead :)

How? The functions are only defined in support.cu and it's only compiled once. Where are the duplicates from?

No idea, llvm-link reports multiply defined functions.

[JonChesterfield]

In D69857#1734569, @ABataev wrote:

In D69857#1734561, @JonChesterfield wrote:

In D69857#1734541, @ABataev wrote:

This one ends up with the multiply defined symbols instead :)

How? The functions are only defined in support.cu and it's only compiled once. Where are the duplicates from?

No idea, llvm-link reports multiply defined functions.

Why is llvm-link involved? I thought nvcc generated object files (ptxas, possibly sass).

I believe there are two compilation paths in the nvptx cmake. One includes everything into unity.cu and passes that to nvcc, which builds a single object file. The other builds each file individually with clang and passes the result to llvm-link. The most likely explanation at this point is that I've totally misunderstood what the cmake file is supposed to be doing, and thus failed to modify it appropriately.

[ABataev]

In D69857#1734592, @JonChesterfield wrote:

In D69857#1734569, @ABataev wrote:

In D69857#1734561, @JonChesterfield wrote:

In D69857#1734541, @ABataev wrote:

This one ends up with the multiply defined symbols instead :)

How? The functions are only defined in support.cu and it's only compiled once. Where are the duplicates from?

No idea, llvm-link reports multiply defined functions.

Why is llvm-link involved? I thought nvcc generated object files (ptxas, possibly sass).

I believe there are two compilation paths in the nvptx cmake. One includes everything into unity.cu and passes that to nvcc, which builds a single object file. The other builds each file individually with clang and passes the result to llvm-link. The most likely explanation at this point is that I've totally misunderstood what the cmake files is supposed to be doing, and thus failed to modify it appropriately.

.bc files are linked with llvm-link. the library is compiled by both, nvcc (to generate .a library) and clang + llvm tools (to generate .bc library). .bc is used for generating of the most optimal code with inlined runtime library.

[JonChesterfield]

I think we should apply this patch anyway. It fixes the nvcc build which we have multiple reports of breakage from.

I can't see a reason why clang would conjure duplicate symbols, and have failed to reproduce anything like that, so am hopeful that it's a quirk of your current build setup.

Fundamentally I believe that nvcc would fail to handle inline functions correctly, as it fails to do a bunch of things correctly, and I believe that clang will build this correctly from a clean directory.

Otherwise, please provide enough details of your build setup that I can reproduce the failure.

[ABataev]

In D69857#1734819, @JonChesterfield wrote:

I think we should apply this patch anyway. It fixes the nvcc build which we have multiple reports of breakage from.

I can't see a reason why clang would conjure duplicate symbols, and have failed to reproduce anything like that, so am hopeful that it's a quirk of your current build setup.

Fundamentally I believe that nvcc would fail to handle inline functions correctly, as it fails to do a bunch of things correctly, and I believe that clang will build this correctly from a clean directory.

Otherwise, please provide enough details of your build setup that I can reproduce the failure.

Bad decision. I would suggest to revert your pr3vious patches and restore intial code layout and then try step by step fixing it.

[JonChesterfield]

In D69857#1734826, @ABataev wrote:

Bad decision. I would suggest to revert your pr3vious patches and restore intial code layout and then try step by step fixing it.

Well, we can't leave trunk broken. So sure, feel free to revert them.

However I cannot debug errors which only occur on other developers machines when using closed source compilers. So I can't fix it later either.

[ABataev]

In D69857#1734842, @JonChesterfield wrote:

In D69857#1734826, @ABataev wrote:

Bad decision. I would suggest to revert your pr3vious patches and restore intial code layout and then try step by step fixing it.

Well, we can't leave trunk broken. So sure, feel free to revert them.

It takes too much time, better you to revert them as yoh're much more familiar with the changes you landed.

However I cannot debug errors which only occur on other developers machines when using closed source compilers. So I can't fix it later either.

I can try them locally and can say if they break something or not before yoj will try to commit them.

[JonChesterfield]

In D69857#1734877, @ABataev wrote:

It takes too much time, better you to revert them as yoh're much more familiar with the changes you landed.

I can try them locally and can say if they break something or not before yoj will try to commit them.

Patch at D69885.

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

 #include "support.h"
 #include "debug.h"
 #include "omptarget-nvptx.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
+DEVICE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
   execution_param = EMode;
   execution_param |= RMode;
 }
 
-INLINE bool isGenericMode() { return (execution_param & ModeMask) == Generic; }
+DEVICE bool isGenericMode() { return (execution_param & ModeMask) == Generic; }
 
-INLINE bool isSPMDMode() { return (execution_param & ModeMask) == Spmd; }
+DEVICE bool isSPMDMode() { return (execution_param & ModeMask) == Spmd; }
 
-INLINE bool isRuntimeUninitialized() {
+DEVICE bool isRuntimeUninitialized() {
   return (execution_param & RuntimeMask) == RuntimeUninitialized;
 }
 
-INLINE bool isRuntimeInitialized() {
+DEVICE bool isRuntimeInitialized() {
   return (execution_param & RuntimeMask) == RuntimeInitialized;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Modes based on location parameter fields
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE bool checkSPMDMode(kmp_Ident *loc) {
+DEVICE 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) {
+DEVICE bool checkGenericMode(kmp_Ident *loc) {
   return !checkSPMDMode(loc);
 }
 
-INLINE bool checkRuntimeUninitialized(kmp_Ident *loc) {
+DEVICE 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) {
+DEVICE 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; }
+DEVICE int GetThreadIdInBlock() { return threadIdx.x; }
 
-INLINE int GetBlockIdInKernel() { return blockIdx.x; }
+DEVICE int GetBlockIdInKernel() { return blockIdx.x; }
 
-INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; }
+DEVICE int GetNumberOfBlocksInKernel() { return gridDim.x; }
 
-INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+DEVICE int GetNumberOfThreadsInBlock() { return blockDim.x; }
 
-INLINE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 
-INLINE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
+DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (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() {
+DEVICE int GetMasterThreadID() {
   return (GetNumberOfThreadsInBlock() - 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(); }
+DEVICE 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) {
+DEVICE 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;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // OpenMP Thread Support Layer
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE int GetOmpThreadId(int threadId, bool isSPMDExecutionMode) {
+DEVICE 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) {
+DEVICE 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;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Team id linked to OpenMP
 
-INLINE int GetOmpTeamId() {
+DEVICE int GetOmpTeamId() {
   // omp_team_num
   return GetBlockIdInKernel(); // assume 1 block per team
 }
 
-INLINE int GetNumberOfOmpTeams() {
+DEVICE int GetNumberOfOmpTeams() {
   // omp_num_teams
   return GetNumberOfBlocksInKernel(); // assume 1 block per team
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Masters
 
-INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
+DEVICE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Parallel level
 
-INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
+DEVICE 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) {
+DEVICE 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) {
+DEVICE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
   if (!isSPMDExecutionMode)
     return GetNumberOfWorkersInTeam();
   return GetNumberOfThreadsInBlock();
 }
 
-INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode) {
+DEVICE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode) {
   return GetNumberOfProcsInDevice(isSPMDExecutionMode);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE unsigned long PadBytes(unsigned long size,
+DEVICE 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
+DEVICE 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) {
+DEVICE 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;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE unsigned int *GetTeamsReductionTimestamp() {
+DEVICE unsigned int *GetTeamsReductionTimestamp() {
   return static_cast<unsigned int *>(ReductionScratchpadPtr);
 }
 
-INLINE char *GetTeamsReductionScratchpad() {
+DEVICE char *GetTeamsReductionScratchpad() {
   return static_cast<char *>(ReductionScratchpadPtr) + 256;
 }
 
-INLINE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr) {
+DEVICE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr) {
   ReductionScratchpadPtr = ScratchpadPtr;
 }