[libc] Add basic wrappers for the backend address spaces

[jhuber6]

Summary

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.

[arsenm]

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

25–28

I don't really condone using the numbered address spaces to access these, they don't quite work the same way

[jhuber6]

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

25–28

They directly map LLVM ones to the ones documented in the backend. Since this is regular C++ we don't have any other option.

[arsenm]

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

25–28

Doesn't declaring as the opencl langAses work? I think those have attributes

[jhuber6]

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

25–28

That doesn't work for NVPTX it seems, https://godbolt.org/z/7515T83nE.

[arsenm]

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

25–28

This looks correct to me when you account for the stack-is-flat hack PTX uses

[jhuber6]

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

25–28

Ah, you're right it seems it understands constant and friends.

[jhuber6]

Changing to use opencl attrs

[arsenm]

It's not great but it's more sound than the numbers. I think the other builtin openmp and hip headers should do the same for the ocml interfacing

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

 #include <stdint.h>
 
 namespace __llvm_libc {
 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();
 }
 
 /// Returns the number of workgroups in the 'y' dimension of the grid.
 LIBC_INLINE uint32_t get_num_blocks_y() {
   return __builtin_amdgcn_grid_size_y() / __builtin_amdgcn_workgroup_size_y();

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

 
 #include <stdint.h>
 
 namespace __llvm_libc {
 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; }
 
 LIBC_INLINE uint32_t get_num_blocks_z() { return 1; }
 
 LIBC_INLINE uint64_t get_num_blocks() { return 1; }
 

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

 #include <stdint.h>
 
 namespace __llvm_libc {
 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();
 }
 
 /// Returns the number of CUDA blocks in the 'y' dimension.
 LIBC_INLINE uint32_t get_num_blocks_y() {
   return __nvvm_read_ptx_sreg_nctaid_y();