From 3926feb84ebe8ffd38bb27161a15b6bb3c8fc0aa Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Mon, 31 Jul 2023 14:32:44 -0500 Subject: [PATCH] [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 --- libc/src/__support/GPU/amdgpu/utils.h | 6 ++++++ libc/src/__support/GPU/generic/utils.h | 5 +++++ libc/src/__support/GPU/nvptx/utils.h | 6 ++++++ 3 files changed, 17 insertions(+) diff --git a/libc/src/__support/GPU/amdgpu/utils.h b/libc/src/__support/GPU/amdgpu/utils.h index 78e3866bebc9..b8e193a793c6 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 using Private = [[clang::opencl_private]] T; +template using Constant = [[clang::opencl_constant]] T; +template using Local = [[clang::opencl_local]] T; +template 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 71cc79654d37..ba985afa696e 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 using Private = T; +template using Constant = T; +template using Shared = T; +template 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 a419e2be09b6..5a921ed93a27 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 using Private = [[clang::opencl_private]] T; +template using Constant = [[clang::opencl_constant]] T; +template using Local = [[clang::opencl_local]] T; +template 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();