[libc] Add a support library for GPU utilities

The GPU has many features that can only be accessed through builtin or
intrinsic functions. Furthermore, these functions are unique for each
GPU target. This patch outlines an interface to create a common `libc`
interface to access these. Currently I only implement a function for the
CUDA equivalent of `blockIdx.x`. More will be added in the future.

Reviewed By: sivachandra

Differential Revision: https://reviews.llvm.org/D148635

libc/src/__support/CMakeLists.txt

@@ -212,6 +212,7 @@ add_header_library(
 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"

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}
+)

libc/src/__support/GPU/amdgpu/CMakeLists.txt

@@ -0,0 +1,7 @@
+add_header_library(
+  amdgpu_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

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

libc/src/__support/GPU/generic/CMakeLists.txt

@@ -0,0 +1,7 @@
+add_header_library(
+  generic_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

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

libc/src/__support/GPU/nvptx/CMakeLists.txt

@@ -0,0 +1,7 @@
+add_header_library(
+  nvptx_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

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

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

libc/src/__support/RPC/CMakeLists.txt

@@ -6,6 +6,7 @@ add_header_library(
   DEPENDS
     libc.src.__support.common
     libc.src.__support.CPP.atomic
+    libc.src.__support.GPU.utils
 )
 
 add_object_library(

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>