diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
@@ -132,8 +132,9 @@
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
-  printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),
-         (int)(threadIdx.x & 0x1F), parameters...);
+  printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
+         (int)(__kmpc_impl_get_warp_id()),
+         (int)(__kmpc_impl_get_thread_id_in_warp()), parameters...);
 }
 
 #endif
@@ -145,8 +146,8 @@
                            Arguments... parameters) {
   if (!cond)
     printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
-           (int)(threadIdx.x / WARPSIZE), (int)(threadIdx.x & 0x1F),
-           parameters...);
+           (int)(__kmpc_impl_get_warp_id()),
+           (int)(__kmpc_impl_get_thread_id_in_warp()), parameters...);
   assert(cond);
 }
 
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -147,4 +147,10 @@
 #endif // CUDA_VERSION
 }
 
+INLINE uint32_t __kmpc_impl_get_warp_id() { return threadIdx.x / WARPSIZE; }
+
+INLINE uint32_t __kmpc_impl_get_thread_id_in_warp() {
+  return threadIdx.x % WARPSIZE;
+}
+
 #endif