[libomptarget][nfc] Explicitly static function scope shared variables

[JonChesterfield]

Summary

[libomptarget][nfc] Explicitly static function scope shared variables

__shared__ in CUDA implies static in function scope. See e.g. D.2.1.1
in CUDA_C_Programming_Guide.pdf,
http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/

This is surprising for non-cuda developers, see e.g. D73239 where I thought
local variables would be thread local.

Tested by IR diff of libomptarget.bc (no change), running in tree tests,
and binary diff of the nvcc static archives (no significant change).

[jdoerfert]

LGTM.

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

static int32_t nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
   // In non-generic mode all workers participate in the teams reduction.
   // In generic mode only the team master participates in the teams
   // reduction because the workers are waiting for parallel work.
   uint32_t NumThreads =
       isSPMDExecutionMode ? GetNumberOfOmpThreads(/*isSPMDExecutionMode=*/true)
                           : /*Master thread only*/ 1;
   uint32_t TeamId = GetBlockIdInKernel();
   uint32_t NumTeams = GetNumberOfBlocksInKernel();
-  SHARED volatile bool IsLastTeam;
+  static SHARED volatile bool IsLastTeam;
 
   // Team masters of all teams write to the scratchpad.
   if (ThreadId == 0) {
     unsigned int *timestamp = GetTeamsReductionTimestamp();
     char *scratchpad = GetTeamsReductionScratchpad();
 
     scratchFct(reduce_data, scratchpad, TeamId, NumTeams);
     __kmpc_impl_threadfence();
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   // In non-generic mode all workers participate in the teams reduction.
   // In generic mode only the team master participates in the teams
   // reduction because the workers are waiting for parallel work.
   uint32_t NumThreads =
       checkSPMDMode(loc) ? GetNumberOfOmpThreads(/*isSPMDExecutionMode=*/true)
                          : /*Master thread only*/ 1;
   uint32_t TeamId = GetBlockIdInKernel();
   uint32_t NumTeams = GetNumberOfBlocksInKernel();
-  SHARED unsigned Bound;
-  SHARED unsigned ChunkTeamCount;
+  static SHARED unsigned Bound;
+  static SHARED unsigned ChunkTeamCount;
 
   // Block progress for teams greater than the current upper
   // limit. We always only allow a number of teams less or equal
   // to the number of slots in the buffer.
   bool IsMaster = isMaster(loc, ThreadId);
   while (IsMaster) {
     // Atomic read
     Bound = __kmpc_atomic_add((uint32_t *)&IterCnt, 0u);