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

   1 //===-------------- AMDGPU implementation of GPU utils ----------*- C++ -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8
   9 #ifndef LLVM_LIBC_SRC___SUPPORT_GPU_AMDGPU_IO_H
  10 #define LLVM_LIBC_SRC___SUPPORT_GPU_AMDGPU_IO_H
  11
  12 #include "src/__support/common.h"
  13 #include "src/__support/macros/config.h"
  14
  15 #include <stdint.h>
  16
  17 namespace LIBC_NAMESPACE_DECL {
  18 namespace gpu {
  19
  20 /// Type aliases to the address spaces used by the AMDGPU backend.
  21 template <typename T> using Private = [[clang::opencl_private]] T;
  22 template <typename T> using Constant = [[clang::opencl_constant]] T;
  23 template <typename T> using Local = [[clang::opencl_local]] T;
  24 template <typename T> using Global = [[clang::opencl_global]] T;
  25
  26 /// Returns the number of workgroups in the 'x' dimension of the grid.
  27 LIBC_INLINE uint32_t get_num_blocks_x() {
  28   return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
  29 }
  30
  31 /// Returns the number of workgroups in the 'y' dimension of the grid.
  32 LIBC_INLINE uint32_t get_num_blocks_y() {
  33   return __builtin_amdgcn_grid_size_y() / __builtin_amdgcn_workgroup_size_y();
  34 }
  35
  36 /// Returns the number of workgroups in the 'z' dimension of the grid.
  37 LIBC_INLINE uint32_t get_num_blocks_z() {
  38   return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
  39 }
  40
  41 /// Returns the total number of workgruops in the grid.
  42 LIBC_INLINE uint64_t get_num_blocks() {
  43   return get_num_blocks_x() * get_num_blocks_y() * get_num_blocks_z();
  44 }
  45
  46 /// Returns the 'x' dimension of the current AMD workgroup's id.
  47 LIBC_INLINE uint32_t get_block_id_x() {
  48   return __builtin_amdgcn_workgroup_id_x();
  49 }
  50
  51 /// Returns the 'y' dimension of the current AMD workgroup's id.
  52 LIBC_INLINE uint32_t get_block_id_y() {
  53   return __builtin_amdgcn_workgroup_id_y();
  54 }
  55
  56 /// Returns the 'z' dimension of the current AMD workgroup's id.
  57 LIBC_INLINE uint32_t get_block_id_z() {
  58   return __builtin_amdgcn_workgroup_id_z();
  59 }
  60
  61 /// Returns the absolute id of the AMD workgroup.
  62 LIBC_INLINE uint64_t get_block_id() {
  63   return get_block_id_x() + get_num_blocks_x() * get_block_id_y() +
  64          get_num_blocks_x() * get_num_blocks_y() * get_block_id_z();
  65 }
  66
  67 /// Returns the number of workitems in the 'x' dimension.
  68 LIBC_INLINE uint32_t get_num_threads_x() {
  69   return __builtin_amdgcn_workgroup_size_x();
  70 }
  71
  72 /// Returns the number of workitems in the 'y' dimension.
  73 LIBC_INLINE uint32_t get_num_threads_y() {
  74   return __builtin_amdgcn_workgroup_size_y();
  75 }
  76
  77 /// Returns the number of workitems in the 'z' dimension.
  78 LIBC_INLINE uint32_t get_num_threads_z() {
  79   return __builtin_amdgcn_workgroup_size_z();
  80 }
  81
  82 /// Returns the total number of workitems in the workgroup.
  83 LIBC_INLINE uint64_t get_num_threads() {
  84   return get_num_threads_x() * get_num_threads_y() * get_num_threads_z();
  85 }
  86
  87 /// Returns the 'x' dimension id of the workitem in the current AMD workgroup.
  88 LIBC_INLINE uint32_t get_thread_id_x() {
  89   return __builtin_amdgcn_workitem_id_x();
  90 }
  91
  92 /// Returns the 'y' dimension id of the workitem in the current AMD workgroup.
  93 LIBC_INLINE uint32_t get_thread_id_y() {
  94   return __builtin_amdgcn_workitem_id_y();
  95 }
  96
  97 /// Returns the 'z' dimension id of the workitem in the current AMD workgroup.
  98 LIBC_INLINE uint32_t get_thread_id_z() {
  99   return __builtin_amdgcn_workitem_id_z();
 100 }
 101
 102 /// Returns the absolute id of the thread in the current AMD workgroup.
 103 LIBC_INLINE uint64_t get_thread_id() {
 104   return get_thread_id_x() + get_num_threads_x() * get_thread_id_y() +
 105          get_num_threads_x() * get_num_threads_y() * get_thread_id_z();
 106 }
 107
 108 /// Returns the size of an AMD wavefront, either 32 or 64 depending on hardware
 109 /// and compilation options.
 110 LIBC_INLINE uint32_t get_lane_size() {
 111   return __builtin_amdgcn_wavefrontsize();
 112 }
 113
 114 /// Returns the id of the thread inside of an AMD wavefront executing together.
 115 [[clang::convergent]] LIBC_INLINE uint32_t get_lane_id() {
 116   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
 117 }
 118
 119 /// Returns the bit-mask of active threads in the current wavefront.
 120 [[clang::convergent]] LIBC_INLINE uint64_t get_lane_mask() {
 121   return __builtin_amdgcn_read_exec();
 122 }
 123
 124 /// Copies the value from the first active thread in the wavefront to the rest.
 125 [[clang::convergent]] LIBC_INLINE uint32_t broadcast_value(uint64_t,
 126                                                            uint32_t x) {
 127   return __builtin_amdgcn_readfirstlane(x);
 128 }
 129
 130 /// Returns a bitmask of threads in the current lane for which \p x is true.
 131 [[clang::convergent]] LIBC_INLINE uint64_t ballot(uint64_t lane_mask, bool x) {
 132   // the lane_mask & gives the nvptx semantics when lane_mask is a subset of
 133   // the active threads
 134   return lane_mask & __builtin_amdgcn_ballot_w64(x);
 135 }
 136
 137 /// Waits for all the threads in the block to converge and issues a fence.
 138 [[clang::convergent]] LIBC_INLINE void sync_threads() {
 139   __builtin_amdgcn_s_barrier();
 140   __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
 141 }
 142
 143 /// Waits for all pending memory operations to complete in program order.
 144 [[clang::convergent]] LIBC_INLINE void memory_fence() {
 145   __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "");
 146 }
 147
 148 /// Wait for all threads in the wavefront to converge, this is a noop on AMDGPU.
 149 [[clang::convergent]] LIBC_INLINE void sync_lane(uint64_t) {
 150   __builtin_amdgcn_wave_barrier();
 151 }
 152
 153 /// Shuffles the the lanes inside the wavefront according to the given index.
 154 [[clang::convergent]] LIBC_INLINE uint32_t shuffle(uint64_t, uint32_t idx,
 155                                                    uint32_t x) {
 156   return __builtin_amdgcn_ds_bpermute(idx << 2, x);
 157 }
 158
 159 /// Returns the current value of the GPU's processor clock.
 160 /// NOTE: The RDNA3 and RDNA2 architectures use a 20-bit cycle counter.
 161 LIBC_INLINE uint64_t processor_clock() { return __builtin_readcyclecounter(); }
 162
 163 /// Returns a fixed-frequency timestamp. The actual frequency is dependent on
 164 /// the card and can only be queried via the driver.
 165 LIBC_INLINE uint64_t fixed_frequency_clock() {
 166   return __builtin_readsteadycounter();
 167 }
 168
 169 /// Terminates execution of the associated wavefront.
 170 [[noreturn]] LIBC_INLINE void end_program() { __builtin_amdgcn_endpgm(); }
 171
 172 /// Returns a unique identifier for the process cluster the current wavefront is
 173 /// executing on. Here we use the identifier for the compute unit (CU) and
 174 /// shader engine.
 175 /// FIXME: Currently unimplemented on AMDGPU until we have a simpler interface
 176 /// than the one at
 177 /// https://github.com/ROCm/clr/blob/develop/hipamd/include/hip/amd_detail/amd_device_functions.h#L899
 178 LIBC_INLINE uint32_t get_cluster_id() { return 0; }
 179
 180 } // namespace gpu
 181 } // namespace LIBC_NAMESPACE_DECL
 182
 183 #endif