[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

diff --git a/libc/src/__support/CMakeLists.txt b/libc/src/__support/CMakeLists.txt
index 29d5b980ad3e8f385518e996adb4fbb55519af84..7f60e5df421850951dd1997d89dbea44e6b2a3a6 100644 (file)
--- a/libc/src/__support/CMakeLists.txt
+++ b/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"

diff --git a/libc/src/__support/GPU/CMakeLists.txt b/libc/src/__support/GPU/CMakeLists.txt
new file mode 100644 (file)
index 0000000..5a89921
--- /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 (file)
index 0000000..f2b98fc
--- /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 (file)
index 0000000..be90cb3
--- /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 (file)
index 0000000..68ba7d1
--- /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 (file)
index 0000000..d54551f
--- /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 (file)
index 0000000..0d3f8c7
--- /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 (file)
index 0000000..fa361cd
--- /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 (file)
index 0000000..f3277f4
--- /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
index c1a971b54b3652bcf8a30c395514f60d2b0cef17..9c578884c6a760805874bc77785841bc6bee6fb4 100644 (file)
--- a/libc/src/__support/RPC/CMakeLists.txt
+++ b/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(

diff --git a/libc/src/__support/RPC/rpc.h b/libc/src/__support/RPC/rpc.h
index 43660fd8e1c9ca288e62845b30d581b2142d10be..196a62daa970cd33e7961bee01b00a3a61320485 100644 (file)
--- 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>