diff --git a/openmp/libomptarget/deviceRTLs/common/device_environment.h b/openmp/libomptarget/deviceRTLs/common/device_environment.h
--- a/openmp/libomptarget/deviceRTLs/common/device_environment.h
+++ b/openmp/libomptarget/deviceRTLs/common/device_environment.h
@@ -13,12 +13,12 @@
 #ifndef _OMPTARGET_DEVICE_ENVIRONMENT_H_
 #define _OMPTARGET_DEVICE_ENVIRONMENT_H_
 
-#include "target_impl.h"
+#include "common/target.h"
 
 struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
-extern DEVICE omptarget_device_environmentTy omptarget_device_environment;
+__LEAGUE_VAR(extern omptarget_device_environmentTy, omptarget_device_environment)
 
 #endif
diff --git a/openmp/libomptarget/deviceRTLs/common/src/cancel.cpp b/openmp/libomptarget/deviceRTLs/common/src/cancel.cpp
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/common/src/cancel.cpp
@@ -0,0 +1,33 @@
+//===------ cancel.cpp - NVPTX OpenMP cancel interface ------------ c++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Interface to be used in the implementation of OpenMP cancel.
+//
+//===----------------------------------------------------------------------===//
+
+//#include "common/debug.h"
+#include "common/target.h"
+
+struct kmp_Ident;
+
+__DEVICE_SCOPE_BEGIN()
+
+int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
+                                 int32_t cancelVal) {
+  //PRINT(LD_IO, "call kmpc_cancellationpoint(cancel val %d)\n", (int)cancelVal);
+  // disabled
+  return 1;
+}
+
+int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid, int32_t cancelVal) {
+  //PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
+  // disabled
+  return 0;
+}
+
+__DEVICE_SCOPE_END()
diff --git a/openmp/libomptarget/deviceRTLs/common/src/cancel.cu b/openmp/libomptarget/deviceRTLs/common/src/cancel.cu
deleted file mode 100644
--- a/openmp/libomptarget/deviceRTLs/common/src/cancel.cu
+++ /dev/null
@@ -1,28 +0,0 @@
-//===------ cancel.cu - NVPTX OpenMP cancel interface ------------ CUDA -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// Interface to be used in the implementation of OpenMP cancel.
-//
-//===----------------------------------------------------------------------===//
-
-#include "interface.h"
-#include "common/debug.h"
-
-EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
-                                        int32_t cancelVal) {
-  PRINT(LD_IO, "call kmpc_cancellationpoint(cancel val %d)\n", (int)cancelVal);
-  // disabled
-  return 0;
-}
-
-EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
-                             int32_t cancelVal) {
-  PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
-  // disabled
-  return 0;
-}
diff --git a/openmp/libomptarget/deviceRTLs/common/src/critical.cu b/openmp/libomptarget/deviceRTLs/common/src/critical.cpp
rename from openmp/libomptarget/deviceRTLs/common/src/critical.cu
rename to openmp/libomptarget/deviceRTLs/common/src/critical.cpp
--- a/openmp/libomptarget/deviceRTLs/common/src/critical.cu
+++ b/openmp/libomptarget/deviceRTLs/common/src/critical.cpp
@@ -1,4 +1,4 @@
-//===------ critical.cu - NVPTX OpenMP critical ------------------ CUDA -*-===//
+//===------ critical.cpp - NVPTX OpenMP critical ------------------ c++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -10,19 +10,24 @@
 //
 //===----------------------------------------------------------------------===//
 
+//#include "common/debug.h"
+#include "nvptx_interface.h"
 #include "interface.h"
-#include "common/debug.h"
 
-EXTERN
+#include "common/target.h"
+
+__DEVICE_SCOPE_BEGIN()
+
 void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
                      kmp_CriticalName *lck) {
-  PRINT0(LD_IO, "call to kmpc_critical()\n");
+  //PRINT0(LD_IO, "call to kmpc_critical()\n");
   omp_set_lock((omp_lock_t *)lck);
 }
 
-EXTERN
 void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
                          kmp_CriticalName *lck) {
-  PRINT0(LD_IO, "call to kmpc_end_critical()\n");
+  //PRINT0(LD_IO, "call to kmpc_end_critical()\n");
   omp_unset_lock((omp_lock_t *)lck);
 }
+
+__DEVICE_SCOPE_END()
diff --git a/openmp/libomptarget/deviceRTLs/common/support.h b/openmp/libomptarget/deviceRTLs/common/support.h
--- a/openmp/libomptarget/deviceRTLs/common/support.h
+++ b/openmp/libomptarget/deviceRTLs/common/support.h
@@ -13,8 +13,8 @@
 #ifndef OMPTARGET_SUPPORT_H
 #define OMPTARGET_SUPPORT_H
 
-#include "interface.h"
 #include "target_impl.h"
+#include "interface.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
diff --git a/openmp/libomptarget/deviceRTLs/common/target.h b/openmp/libomptarget/deviceRTLs/common/target.h
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/common/target.h
@@ -0,0 +1,43 @@
+//===---- target.h - OpenMP defines and helpers for target code --- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines and helpers for target code.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OMPTARGET_TARGET_H
+#define OMPTARGET_TARGET_H
+
+#include <cstdint>
+
+#define __p(STR) _Pragma(STR)
+#define __p2(STR) __p(#STR)
+
+#define __DEVICE_SCOPE_BEGIN()                                                 \
+  extern "C" {                                                                 \
+  __p("omp declare target")
+
+#define __DEVICE_SCOPE_END()                                                   \
+  __p("omp end declare target")                                                \
+  } /* extern "C" */
+
+#define __CONSTEXPR static constexpr __attribute__((nothrow, always_inline))
+
+#define __LEAGUE_VAR(TYPE, NAME)                                               \
+  TYPE NAME [[clang::loader_uninitialized]];                                   \
+  __p2(omp declare variant to(NAME))
+
+#define __TEAM_VAR(TYPE, NAME)                                                 \
+  TYPE NAME[[ clang::loader_uninitialized, omp::allocator::access:pteam ]];    \
+  __p2(omp declare variant to(NAME))
+
+#define __THREAD_VAR(TYPE, NAME)                                               \
+  TYPE NAME[[ clang::loader_uninitialized, omp::allocator::access:thread ]];   \
+  __p2(omp declare variant to(NAME))
+
+#endif
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
--- a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
@@ -50,9 +50,12 @@
   # propagating host flags.
   set(CUDA_PROPAGATE_HOST_FLAGS OFF)
 
+  set(cxx_src_files
+      ${devicertl_common_directory}/src/cancel.cpp
+      ${devicertl_common_directory}/src/critical.cpp
+  )
+
   set(cuda_src_files
-      ${devicertl_common_directory}/src/cancel.cu
-      ${devicertl_common_directory}/src/critical.cu
       ${devicertl_common_directory}/src/data_sharing.cu
       ${devicertl_common_directory}/src/libcall.cu
       ${devicertl_common_directory}/src/loop.cu
@@ -93,6 +96,8 @@
   # yet supported by the CUDA toolchain on the device.
   set(BUILD_SHARED_LIBS OFF)
   set(CUDA_SEPARABLE_COMPILATION ON)
+  list(APPEND CXX_FLAGS -I${devicertl_base_directory}
+                        -I${devicertl_nvptx_directory}/src)
   list(APPEND CUDA_NVCC_FLAGS -I${devicertl_base_directory}
                               -I${devicertl_nvptx_directory}/src)
   cuda_add_library(omptarget-nvptx STATIC ${cuda_src_files} ${omp_data_objects}
@@ -152,8 +157,28 @@
     foreach(sm ${nvptx_sm_list})
       set(cuda_arch --cuda-gpu-arch=sm_${sm})
 
-      # Compile CUDA files to bitcode.
+      # Compile C++ files to bitcode.
       set(bc_files "")
+      foreach(src ${cxx_src_files})
+        get_filename_component(infile ${src} ABSOLUTE)
+        get_filename_component(outfile ${src} NAME)
+
+        add_custom_command(OUTPUT ${outfile}-sm_${sm}.bc
+          COMMAND ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER} ${bc_flags} ${cuda_arch}
+            -c ${infile} -o ${outfile}-sm_${sm}.bc -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda
+            -Xclang -fopenmp-is-device
+            -Xclang -aux-triple -Xclang x86_64-unknown-linux
+          DEPENDS ${infile}
+          IMPLICIT_DEPENDS CXX ${infile}
+          COMMENT "Building LLVM bitcode ${outfile}-sm_${sm}.bc"
+          VERBATIM
+        )
+        set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${outfile}-sm_${sm}.bc)
+
+        list(APPEND bc_files ${outfile}-sm_${sm}.bc)
+      endforeach()
+
+      # Compile CUDA files to bitcode.
       foreach(src ${cuda_src_files})
         get_filename_component(infile ${src} ABSOLUTE)
         get_filename_component(outfile ${src} NAME)
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h b/openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
@@ -11,7 +11,12 @@
 
 #include <stdint.h>
 
+#ifndef _OPENMP
 #define EXTERN extern "C" __device__
+#else
+#define EXTERN extern "C"
+#endif
+
 typedef uint32_t __kmpc_impl_lanemask_t;
 typedef uint32_t omp_lock_t; /* arbitrary type of the right length */