diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip b/openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip
@@ -0,0 +1,21 @@
+//===----- amdgcn_threadfence.hip - AMDGCN fence implementation --- HIP -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "target_impl.h"
+
+EXTERN void __kmpc_impl_threadfence() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+}
+
+EXTERN void __kmpc_impl_threadfence_block() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+}
+
+EXTERN void __kmpc_impl_threadfence_system() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+}
diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
--- a/openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
@@ -11,11 +11,6 @@
 
 #include "target_impl.h"
 
-// inc requires an amdgcn specific intrinsic which is not yet available
-DEVICE unsigned atomicInc(unsigned *address);
-DEVICE unsigned atomicInc(unsigned *address, unsigned max);
-DEVICE int atomicInc(int *address);
-
 namespace {
 
 template <typename T> DEVICE T atomicAdd(T *address, T val) {
@@ -38,5 +33,15 @@
   return compare;
 }
 
+// Variants of atomicInc with and without wraparound MAX value
+DEVICE uint32_t atomicInc(uint32_t *address) {
+  return __builtin_amdgcn_atomic_inc32(address, UINT32_MAX, __ATOMIC_SEQ_CST,
+                                       "");
+}
+
+DEVICE uint32_t atomicInc(uint32_t *address, uint32_t max) {
+  return __builtin_amdgcn_atomic_inc32(address, max, __ATOMIC_SEQ_CST, "");
+}
+
 } // namespace
 #endif