Index: openmp/libomptarget/DeviceRTL/src/Mapping.cpp
===================================================================
--- openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -372,4 +372,12 @@
 }
 }
 
+#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)
+
 #pragma omp end declare target
Index: openmp/libomptarget/DeviceRTL/src/exports
===================================================================
--- openmp/libomptarget/DeviceRTL/src/exports
+++ openmp/libomptarget/DeviceRTL/src/exports
@@ -1,4 +1,5 @@
 omp_*
+ompx_*
 *llvm_*
 __kmpc_*
 
Index: openmp/libomptarget/test/api/ompx_3d.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/api/ompx_3d.c
@@ -0,0 +1,27 @@
+// 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();
+      bid = ompx_block_id_x();
+      bdim = ompx_block_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 main() {
+  foo(omp_get_default_device());
+  foo(omp_get_initial_device());
+}
Index: openmp/libomptarget/test/api/ompx_3d.cpp
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/api/ompx_3d.cpp
@@ -0,0 +1,27 @@
+// 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();
+      bid = ompx::block_id_x();
+      bdim = ompx::block_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 main() {
+  foo(omp_get_default_device());
+  foo(omp_get_initial_device());
+}
Index: openmp/runtime/cmake/LibompExports.cmake
===================================================================
--- openmp/runtime/cmake/LibompExports.cmake
+++ openmp/runtime/cmake/LibompExports.cmake
@@ -50,6 +50,7 @@
 add_custom_command(TARGET omp POST_BUILD
   COMMAND ${CMAKE_COMMAND} -E make_directory ${LIBOMP_EXPORTS_CMN_DIR}
   COMMAND ${CMAKE_COMMAND} -E copy omp.h ${LIBOMP_EXPORTS_CMN_DIR}
+  COMMAND ${CMAKE_COMMAND} -E copy ompx.h ${LIBOMP_EXPORTS_CMN_DIR}
 )
 if(${LIBOMP_OMPT_SUPPORT})
   add_custom_command(TARGET omp POST_BUILD
Index: openmp/runtime/src/CMakeLists.txt
===================================================================
--- openmp/runtime/src/CMakeLists.txt
+++ openmp/runtime/src/CMakeLists.txt
@@ -12,6 +12,7 @@
 
 # Configure omp.h, kmp_config.h and omp-tools.h if necessary
 configure_file(${LIBOMP_INC_DIR}/omp.h.var omp.h @ONLY)
+configure_file(${LIBOMP_INC_DIR}/ompx.h.var ompx.h @ONLY)
 configure_file(kmp_config.h.cmake kmp_config.h @ONLY)
 if(${LIBOMP_OMPT_SUPPORT})
   configure_file(${LIBOMP_INC_DIR}/omp-tools.h.var omp-tools.h @ONLY)
@@ -393,6 +394,7 @@
 install(
   FILES
   ${CMAKE_CURRENT_BINARY_DIR}/omp.h
+  ${CMAKE_CURRENT_BINARY_DIR}/ompx.h
   DESTINATION ${LIBOMP_HEADERS_INSTALL_PATH}
 )
 if(${LIBOMP_OMPT_SUPPORT})