diff --git a/libc/src/__support/GPU/amdgpu/utils.h b/libc/src/__support/GPU/amdgpu/utils.h index 78e3866bebc9c..b8e193a793c6c 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 71cc79654d37e..ba985afa696ee 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 a419e2be09b6c..5a921ed93a272 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();