diff --git a/libc/src/__support/CMakeLists.txt b/libc/src/__support/CMakeLists.txt
--- a/libc/src/__support/CMakeLists.txt
+++ b/libc/src/__support/CMakeLists.txt
@@ -212,6 +212,7 @@
 add_subdirectory(FPUtil)
 add_subdirectory(OSUtil)
 add_subdirectory(StringUtil)
+add_subdirectory(GPU)
 add_subdirectory(RPC)
 
 # Thread support is used by other "File". So, we add the "threads"
diff --git a/libc/src/__support/GPU/CMakeLists.txt b/libc/src/__support/GPU/CMakeLists.txt
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/CMakeLists.txt
@@ -0,0 +1,16 @@
+if(NOT LIBC_TARGET_ARCHITECTURE_IS_GPU)
+  return()
+endif()
+
+foreach(target nvptx amdgpu generic)
+  add_subdirectory(${target})
+  list(APPEND target_gpu_utils libc.src.__support.GPU.${target}.${target}_utils)
+endforeach()
+
+add_header_library(
+  utils
+  HDRS
+    utils.h
+  DEPENDS
+    ${target_gpu_utils}
+)
diff --git a/libc/src/__support/GPU/amdgpu/CMakeLists.txt b/libc/src/__support/GPU/amdgpu/CMakeLists.txt
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/amdgpu/CMakeLists.txt
@@ -0,0 +1,7 @@
+add_header_library(
+  amdgpu_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)
diff --git a/libc/src/__support/GPU/amdgpu/utils.h b/libc/src/__support/GPU/amdgpu/utils.h
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/amdgpu/utils.h
@@ -0,0 +1,24 @@
+//===-------------- AMDGPU implementation of GPU utils ----------*- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_SUPPORT_GPU_AMDGPU_IO_H
+#define LLVM_LIBC_SRC_SUPPORT_GPU_AMDGPU_IO_H
+
+#include "src/__support/common.h"
+
+#include <stdint.h>
+
+namespace __llvm_libc {
+
+LIBC_INLINE uint32_t get_block_id_x() {
+  return __builtin_amdgcn_workgroup_id_x();
+}
+
+} // namespace __llvm_libc
+
+#endif
diff --git a/libc/src/__support/GPU/generic/CMakeLists.txt b/libc/src/__support/GPU/generic/CMakeLists.txt
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/generic/CMakeLists.txt
@@ -0,0 +1,7 @@
+add_header_library(
+  generic_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)
diff --git a/libc/src/__support/GPU/generic/utils.h b/libc/src/__support/GPU/generic/utils.h
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/generic/utils.h
@@ -0,0 +1,22 @@
+//===-------------- Generic implementation of GPU utils ---------*- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_SUPPORT_GPU_GENERIC_IO_H
+#define LLVM_LIBC_SRC_SUPPORT_GPU_GENERIC_IO_H
+
+#include "src/__support/common.h"
+
+#include <stdint.h>
+
+namespace __llvm_libc {
+
+LIBC_INLINE uint32_t get_block_id_x() { return 0; }
+
+} // namespace __llvm_libc
+
+#endif
diff --git a/libc/src/__support/GPU/nvptx/CMakeLists.txt b/libc/src/__support/GPU/nvptx/CMakeLists.txt
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/nvptx/CMakeLists.txt
@@ -0,0 +1,7 @@
+add_header_library(
+  nvptx_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)
diff --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/nvptx/utils.h
@@ -0,0 +1,22 @@
+//===-------------- NVPTX implementation of GPU utils -----------*- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_SUPPORT_GPU_NVPTX_IO_H
+#define LLVM_LIBC_SRC_SUPPORT_GPU_NVPTX_IO_H
+
+#include "src/__support/common.h"
+
+#include <stdint.h>
+
+namespace __llvm_libc {
+
+LIBC_INLINE uint32_t get_block_id_x() { return __nvvm_read_ptx_sreg_ctaid_x(); }
+
+} // namespace __llvm_libc
+
+#endif
diff --git a/libc/src/__support/GPU/utils.h b/libc/src/__support/GPU/utils.h
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/GPU/utils.h
@@ -0,0 +1,22 @@
+//===---------------- Implementation of GPU utils ---------------*- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_SUPPORT_GPU_UTIL_H
+#define LLVM_LIBC_SRC_SUPPORT_GPU_UTIL_H
+
+#include "src/__support/macros/properties/architectures.h"
+
+#if defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+#include "amdgpu/utils.h"
+#elif defined(LIBC_TARGET_ARCH_IS_NVPTX)
+#include "nvptx/utils.h"
+#else
+#include "generic/utils.h"
+#endif
+
+#endif // LLVM_LIBC_SRC_SUPPORT_OSUTIL_IO_H
diff --git a/libc/src/__support/RPC/CMakeLists.txt b/libc/src/__support/RPC/CMakeLists.txt
--- a/libc/src/__support/RPC/CMakeLists.txt
+++ b/libc/src/__support/RPC/CMakeLists.txt
@@ -6,6 +6,7 @@
   DEPENDS
     libc.src.__support.common
     libc.src.__support.CPP.atomic
+    libc.src.__support.GPU.utils
 )
 
 add_object_library(
diff --git a/libc/src/__support/RPC/rpc.h b/libc/src/__support/RPC/rpc.h
--- a/libc/src/__support/RPC/rpc.h
+++ b/libc/src/__support/RPC/rpc.h
@@ -20,6 +20,7 @@
 
 #include "rpc_util.h"
 #include "src/__support/CPP/atomic.h"
+#include "src/__support/GPU/utils.h"
 
 #include <stdint.h>