[libc-commits] [PATCH] D156731: [libc] Add basic wrappers for the backend address spaces

[Joseph Huber via Phabricator via libc-commits]

jhuber6 updated this revision to Diff 545800.
jhuber6 added a comment.

Changing to use opencl attrs


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D156731/new/

https://reviews.llvm.org/D156731

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


Index: libc/src/__support/GPU/nvptx/utils.h
===================================================================
--- libc/src/__support/GPU/nvptx/utils.h
+++ libc/src/__support/GPU/nvptx/utils.h
@@ -19,6 +19,12 @@
 /// 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();
Index: libc/src/__support/GPU/generic/utils.h
===================================================================
--- libc/src/__support/GPU/generic/utils.h
+++ libc/src/__support/GPU/generic/utils.h
@@ -18,6 +18,11 @@
 
 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; }
Index: libc/src/__support/GPU/amdgpu/utils.h
===================================================================
--- libc/src/__support/GPU/amdgpu/utils.h
+++ libc/src/__support/GPU/amdgpu/utils.h
@@ -20,6 +20,12 @@
 /// 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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D156731.545800.patch
Type: text/x-patch
Size: 2228 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/libc-commits/attachments/20230731/930f6270/attachment-0001.bin>