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.
get_num_blocks_x()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.
get_num_blocks_y()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.
get_num_blocks_z()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.
get_num_blocks()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.
get_block_id_x()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.
get_block_id_y()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.
get_block_id_z()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.
get_block_id()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.
get_num_threads_x()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.
get_num_threads_y()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.
get_num_threads_z()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.
get_num_threads()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.
get_thread_id_x()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.
get_thread_id_y()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.
get_thread_id_z()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.
get_thread_id()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.
get_lane_size()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.
get_lane_id()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.
get_lane_mask()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.
broadcast_value(uint64_t,uint32_t x)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.
ballot(uint64_t lane_mask,bool x)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.
sync_threads()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.
memory_fence()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.
sync_lane(uint64_t)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.
shuffle(uint64_t,uint32_t idx,uint32_t x)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.
processor_clock()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.
fixed_frequency_clock()165 LIBC_INLINE uint64_t fixed_frequency_clock() {
166 return __builtin_readsteadycounter();
167 }
168
169 /// Terminates execution of the associated wavefront.
end_program()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
get_cluster_id()178 LIBC_INLINE uint32_t get_cluster_id() { return 0; }
179
180 } // namespace gpu
181 } // namespace LIBC_NAMESPACE_DECL
182
183 #endif
184