[libc-commits] [libc] 814dfb0 - [libc] Add a support library for GPU utilities
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Wed Apr 19 06:02:08 PDT 2023
Author: Joseph Huber
Date: 2023-04-19T08:01:56-05:00
New Revision: 814dfb016aad7ceae2e3fda19659e0bb20f10464
URL: https://github.com/llvm/llvm-project/commit/814dfb016aad7ceae2e3fda19659e0bb20f10464
DIFF: https://github.com/llvm/llvm-project/commit/814dfb016aad7ceae2e3fda19659e0bb20f10464.diff
LOG: [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
Added:
libc/src/__support/GPU/CMakeLists.txt
libc/src/__support/GPU/amdgpu/CMakeLists.txt
libc/src/__support/GPU/amdgpu/utils.h
libc/src/__support/GPU/generic/CMakeLists.txt
libc/src/__support/GPU/generic/utils.h
libc/src/__support/GPU/nvptx/CMakeLists.txt
libc/src/__support/GPU/nvptx/utils.h
libc/src/__support/GPU/utils.h
Modified:
libc/src/__support/CMakeLists.txt
libc/src/__support/RPC/CMakeLists.txt
libc/src/__support/RPC/rpc.h
Removed:
################################################################################
diff --git a/libc/src/__support/CMakeLists.txt b/libc/src/__support/CMakeLists.txt
index 29d5b980ad3e8..7f60e5df42185 100644
--- 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
index 0000000000000..5a899215f4b6e
--- /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
index 0000000000000..f2b98fc03b218
--- /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
index 0000000000000..be90cb3edbc7a
--- /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
index 0000000000000..68ba7d1ec80e9
--- /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
index 0000000000000..d54551f953712
--- /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
index 0000000000000..0d3f8c7933c86
--- /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
index 0000000000000..fa361cdbdf5c5
--- /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
index 0000000000000..f3277f42a32d2
--- /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 c1a971b54b365..9c578884c6a76 100644
--- 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 43660fd8e1c9c..196a62daa970c 100644
--- 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>
More information about the libc-commits
mailing list