[libc-commits] [libc] 3926feb - [libc] Add basic wrappers for the backend address spaces

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2023-08-01T07:41:10-05:00
New Revision: 3926feb84ebe8ffd38bb27161a15b6bb3c8fc0aa

URL: https://github.com/llvm/llvm-project/commit/3926feb84ebe8ffd38bb27161a15b6bb3c8fc0aa
DIFF: https://github.com/llvm/llvm-project/commit/3926feb84ebe8ffd38bb27161a15b6bb3c8fc0aa.diff

LOG: [libc] Add basic wrappers for the backend address spaces

The GPU makes use of different address spaces. We generally work with
global memory, thread private memory, and thread shared memory. This
patch simply adds a few preliminary wrappers to map these concepts to
the numerical values the backend uses. Obviously casts between these
will need to be checked by the user.

Reviewed By: arsenm

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

Added: 
    

Modified: 
    libc/src/__support/GPU/amdgpu/utils.h
    libc/src/__support/GPU/generic/utils.h
    libc/src/__support/GPU/nvptx/utils.h

Removed: 
    


################################################################################
diff  --git a/libc/src/__support/GPU/amdgpu/utils.h b/libc/src/__support/GPU/amdgpu/utils.h
index 78e3866bebc9c6..b8e193a793c6c3 100644
--- a/libc/src/__support/GPU/amdgpu/utils.h
+++ b/libc/src/__support/GPU/amdgpu/utils.h
@@ -20,6 +20,12 @@ namespace gpu {
 /// The number of threads that execute in lock-step in a lane.
 constexpr const uint64_t LANE_SIZE = __AMDGCN_WAVEFRONT_SIZE;
 
+/// Type aliases to the address spaces used by the AMDGPU backend.
+template <typename T> using Private = [[clang::opencl_private]] T;
+template <typename T> using Constant = [[clang::opencl_constant]] T;
+template <typename T> using Local = [[clang::opencl_local]] T;
+template <typename T> using Global = [[clang::opencl_global]] T;
+
 /// Returns the number of workgroups in the 'x' dimension of the grid.
 LIBC_INLINE uint32_t get_num_blocks_x() {
   return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();

diff  --git a/libc/src/__support/GPU/generic/utils.h b/libc/src/__support/GPU/generic/utils.h
index 71cc79654d37ea..ba985afa696ee7 100644
--- a/libc/src/__support/GPU/generic/utils.h
+++ b/libc/src/__support/GPU/generic/utils.h
@@ -18,6 +18,11 @@ namespace gpu {
 
 constexpr const uint64_t LANE_SIZE = 1;
 
+template <typename T> using Private = T;
+template <typename T> using Constant = T;
+template <typename T> using Shared = T;
+template <typename T> using Global = T;
+
 LIBC_INLINE uint32_t get_num_blocks_x() { return 1; }
 
 LIBC_INLINE uint32_t get_num_blocks_y() { return 1; }

diff  --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h
index a419e2be09b6c8..5a921ed93a272c 100644
--- a/libc/src/__support/GPU/nvptx/utils.h
+++ b/libc/src/__support/GPU/nvptx/utils.h
@@ -19,6 +19,12 @@ namespace gpu {
 /// The number of threads that execute in lock-step in a warp.
 constexpr const uint64_t LANE_SIZE = 32;
 
+/// Type aliases to the address spaces used by the NVPTX backend.
+template <typename T> using Private = [[clang::opencl_private]] T;
+template <typename T> using Constant = [[clang::opencl_constant]] T;
+template <typename T> using Local = [[clang::opencl_local]] T;
+template <typename T> using Global = [[clang::opencl_global]] T;
+
 /// Returns the number of CUDA blocks in the 'x' dimension.
 LIBC_INLINE uint32_t get_num_blocks_x() {
   return __nvvm_read_ptx_sreg_nctaid_x();