[OMPX] Change `thread_dim` to `block_dim` and the original `block_dim` to `grid_dim`

[tianshilei1992]

Summary

There is no threadDim in CUDA. Instead, it is blockDim. Then the current
blockDim is gridDim in CUDA.

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

__attribute__((noinline)) uint32_t __kmpc_get_warp_size() {
   return impl::getWarpSize();
 }
 }
 
 #define _TGT_KERNEL_LANGUAGE(NAME, MAPPER_NAME)                                \
   extern "C" int ompx_##NAME(int Dim) { return mapping::MAPPER_NAME(Dim); }
 
 _TGT_KERNEL_LANGUAGE(thread_id, getThreadIdInBlock)
-_TGT_KERNEL_LANGUAGE(thread_dim, getNumberOfThreadsInBlock)
 _TGT_KERNEL_LANGUAGE(block_id, getBlockIdInKernel)
-_TGT_KERNEL_LANGUAGE(block_dim, getNumberOfBlocksInKernel)
+_TGT_KERNEL_LANGUAGE(block_dim, getNumberOfThreadsInBlock)
+_TGT_KERNEL_LANGUAGE(grid_dim, getNumberOfBlocksInKernel)
 
 #pragma omp end declare target

openmp/libomptarget/test/api/ompx_3d.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 #include <omp.h>
 #include <ompx.h>
 #include <stdio.h>
 
 void foo(int device) {
   int tid = 0, bid = 0, bdim = 0;
 #pragma omp target teams distribute parallel for map(from                      \
                                                      : tid, bid, bdim)         \
     device(device) thread_limit(2) num_teams(5)
   for (int i = 0; i < 1000; ++i) {
     if (i == 42) {
-      tid = ompx_thread_dim_x();
+      tid = ompx_block_dim_x();
       bid = ompx_block_id_x();
-      bdim = ompx_block_dim_x();
+      bdim = ompx_grid_dim_x();
     }
   }
   // CHECK: tid: 2, bid: 1, bdim: 5
   // CHECK: tid: 2, bid: 0, bdim: 1
   printf("tid: %i, bid: %i, bdim: %i\n", tid, bid, bdim);
 }
 
 int isGPU() { return 0; }

openmp/libomptarget/test/api/ompx_3d.cpp

 // RUN: %libomptarget-compilexx-run-and-check-generic
 
 #include <omp.h>
 #include <ompx.h>
 #include <stdio.h>
 
 void foo(int device) {
   int tid = 0, bid = 0, bdim = 0;
 #pragma omp target teams distribute parallel for map(from                      \
                                                      : tid, bid, bdim)         \
     device(device) thread_limit(2) num_teams(5)
   for (int i = 0; i < 1000; ++i) {
     if (i == 42) {
-      tid = ompx::thread_dim_x();
+      tid = ompx::block_dim_x();
       bid = ompx::block_id_x();
-      bdim = ompx::block_dim_x();
+      bdim = ompx::grid_dim_x();
     }
   }
   // CHECK: tid: 2, bid: 1, bdim: 5
   // CHECK: tid: 2, bid: 0, bdim: 1
   printf("tid: %i, bid: %i, bdim: %i\n", tid, bid, bdim);
 }
 
 int isGPU() { return 0; }

openmp/libomptarget/test/offloading/ompx_coords.c

 
 #pragma omp target data map(tofrom : X [0:N])
 #pragma omp target teams num_teams(NT) thread_limit(TL)
   {
 #pragma omp parallel
     {
       int tid = ompx_thread_id_x();
       int bid = ompx_block_id_x();
-      int tdim = ompx_thread_dim_x();
+      int tdim = ompx_block_dim_x();
       int gid = tid + bid * tdim;
       if (gid < N) {
         X[gid].tid = tid;
         X[gid].bid = bid;
         X[gid].tdim = tdim;
       };
     }
   }

openmp/libomptarget/test/offloading/ompx_saxpy_mixed.c

 
 #pragma omp target data map(to : X [0:N]) map(Y [0:N])
 #pragma omp target teams num_teams(NT) thread_limit(TL)
   {
 #pragma omp parallel
     {
       int tid = ompx_thread_id_x();
       int bid = ompx_block_id_x();
-      int tdim = ompx_thread_dim_x();
+      int tdim = ompx_block_dim_x();
       int gid = tid + bid * tdim;
       if (gid < N)
         Y[gid] = a * X[gid] + Y[gid];
     }
   }
 
   float maxError = 0.0f;
   for (int i = 0; i < N; i++) {

openmp/runtime/src/include/ompx.h.var

 /// ompx_{thread,block}_{id,dim}
 ///{
 #pragma omp begin declare variant match(device = {kind(cpu)})
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(NAME, VALUE)                     \
   static inline int ompx_##NAME(int Dim) { return VALUE; }
 
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(thread_id,
                                       omp_get_ancestor_thread_num(Dim + 1))
-_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(thread_dim, omp_get_team_size(Dim + 1))
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(block_dim, omp_get_team_size(Dim + 1))
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(block_id, 0)
-_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(block_dim, 1)
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(grid_dim, 1)
 #undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C
 ///}
 
 /// ompx_{sync_block}_{,divergent}
 ///{
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C(RETTY, NAME, ARGS, BODY)         \
   static inline RETTY ompx_##NAME(ARGS) { BODY; }
 

 ///{
 #define _TGT_KERNEL_LANGUAGE_DECL_GRID_C(NAME)                                 \
   int ompx_##NAME(int Dim);                                                    \
   static inline int ompx_##NAME##_x() { return ompx_##NAME(ompx_dim_x); }      \
   static inline int ompx_##NAME##_y() { return ompx_##NAME(ompx_dim_y); }      \
   static inline int ompx_##NAME##_z() { return ompx_##NAME(ompx_dim_z); }
 
 _TGT_KERNEL_LANGUAGE_DECL_GRID_C(thread_id)
-_TGT_KERNEL_LANGUAGE_DECL_GRID_C(thread_dim)
-_TGT_KERNEL_LANGUAGE_DECL_GRID_C(block_id)
 _TGT_KERNEL_LANGUAGE_DECL_GRID_C(block_dim)
+_TGT_KERNEL_LANGUAGE_DECL_GRID_C(block_id)
+_TGT_KERNEL_LANGUAGE_DECL_GRID_C(grid_dim)
 #undef _TGT_KERNEL_LANGUAGE_DECL_GRID_C
 ///}
 
 #ifdef __cplusplus
 }
 #endif
 
 #ifdef __cplusplus

 ///{
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(NAME)                          \
   static inline int NAME(int Dim) noexcept { return ompx_##NAME(Dim); }        \
   static inline int NAME##_x() noexcept { return NAME(ompx_dim_x); }           \
   static inline int NAME##_y() noexcept { return NAME(ompx_dim_y); }           \
   static inline int NAME##_z() noexcept { return NAME(ompx_dim_z); }
 
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(thread_id)
-_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(thread_dim)
-_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(block_id)
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(block_dim)
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(block_id)
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(grid_dim)
 #undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX
 ///}
 
 /// ompx_{sync_block}_{,divergent}
 ///{
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_CXX(RETTY, NAME, ARGS, CALL_ARGS)  \
   static inline RETTY NAME(ARGS) {               \
     return ompx_##NAME(CALL_ARGS);                                             \