[libc] Add a support library for GPU utilities

[jhuber6]

Summary

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.

[jhuber6]

Adding a 'generic' folder. This will be a shim that tries to return something
reasonable when called from a non-GPU target. This will primarily be used for
RPC which needs to be compiled on the CPU and GPU. Also this makes my language
server happy because otherwise it doesn't define the function.

[JonChesterfield]

It seems like an abstraction over these things shows up in essentially every GPU codebase. I think they're a good candidate for clang intrinsics that translate to the corresponding platform specific incantation.

However there are currently no such intrinsics and it's hard to anticipate how well received such a proposal would be so I'm very much I'm favour of the approach you're taking here. Long term maybe the file would shrink to zero as intrinsics are created.

It seems credible that this could be a header only library. I know openmp devicertl does it's own thing based around variant dispatch but maybe it wouldn't be the worst thing to include header files from libc for part of that. It would use intrinsics if they existed I think.

[jhuber6]

So, one

In D148635#4277933, @JonChesterfield wrote:

It seems like an abstraction over these things shows up in essentially every GPU codebase. I think they're a good candidate for clang intrinsics that translate to the corresponding platform specific incantation.

However there are currently no such intrinsics and it's hard to anticipate how well received such a proposal would be so I'm very much I'm favour of the approach you're taking here. Long term maybe the file would shrink to zero as intrinsics are created.

It seems credible that this could be a header only library. I know openmp devicertl does it's own thing based around variant dispatch but maybe it wouldn't be the worst thing to include header files from libc for part of that. It would use intrinsics if they existed I think.

Yeah, separation here is done via the macros. Intrinsics would be a solution here, but it's difficult to find a common set sometimes, since not every Nvidia builtin matches an AMD one. it would actually be possible to export these via the libc interface as an extension. But providing this as a library might be less optimal than handling it in LLVM.

[jdoerfert]

I'm not too happy with the name but the concept seems like something we need (to duplicate or move) here.
Works for me if we cannot find a better name(ing scheme).

[jhuber6]

In D148635#4278410, @jdoerfert wrote:

I'm not too happy with the name but the concept seems like something we need (to duplicate or move) here.
Works for me if we cannot find a better name(ing scheme).

What would a better name be? I just went with utils since a lot of the other random libc stuff is XXXutil.

[sivachandra]

libc/src/__support/GPU/CMakeLists.txt

11

Not a strong opinion but you can simplify few things here if you use files like nvptx_utils.h etc. instead of nested files nvptx/utils.h etc.

libc/src/__support/RPC/rpc.h

23

Is this new inclusion being used anywhere or is it only to make rpc.h expose utils.h implicitly? For proper layering, we normally avoid such implicit dependencies and inclusions.

[jhuber6]

libc/src/__support/GPU/CMakeLists.txt

11

I wanted to keep a single utils.h file that you can include and get the correct definition for that target.

libc/src/__support/RPC/rpc.h

23

A future patch will require it, so I just put it here to make sure that it worked, I can remove it for this patch if needed.

libc/src/__support/CMakeLists.txt

add_header_library(
   DEPENDS
     .integer_to_string
     libc.src.__support.OSUtil.osutil
 )
 
 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"
 # before "File".
 add_subdirectory(threads)
 
 add_subdirectory(File)

libc/src/__support/GPU/CMakeLists.txt

+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

+add_header_library(
+  amdgpu_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

libc/src/__support/GPU/amdgpu/utils.h

+//===-------------- 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

+add_header_library(
+  generic_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

libc/src/__support/GPU/generic/utils.h

+//===-------------- 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

+add_header_library(
+  nvptx_utils
+  HDRS
+    utils.h
+  DEPENDS
+    libc.src.__support.common
+)

libc/src/__support/GPU/nvptx/utils.h

+//===-------------- 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

+//===---------------- 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

 add_header_library(
   rpc
   HDRS
     rpc.h
     rpc_util.h
   DEPENDS
     libc.src.__support.common
     libc.src.__support.CPP.atomic
+    libc.src.__support.GPU.utils
 )
 
 add_object_library(
   rpc_client
   SRCS
     rpc_client.cpp
   HDRS
     rpc_client.h
   DEPENDS
     .rpc
 )

libc/src/__support/RPC/rpc.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIBC_SRC_SUPPORT_RPC_RPC_H
 #define LLVM_LIBC_SRC_SUPPORT_RPC_RPC_H
 
 #include "rpc_util.h"
 #include "src/__support/CPP/atomic.h"
+#include "src/__support/GPU/utils.h"
 
 #include <stdint.h>
 
 namespace __llvm_libc {
 namespace rpc {
 
 /// A list of opcodes that we use to invoke certain actions on the server. We
 /// reserve the first 255 values for internal libc usage.