1 //===------- Mapping.cpp - OpenMP device runtime mapping helpers -- 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 //
10 //===----------------------------------------------------------------------===//
11 
12 #include "Mapping.h"
13 #include "State.h"
14 #include "Types.h"
15 #include "Utils.h"
16 
17 #pragma omp declare target
18 
19 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
20 
21 using namespace _OMP;
22 
23 namespace _OMP {
24 namespace impl {
25 
26 /// AMDGCN Implementation
27 ///
28 ///{
29 #pragma omp begin declare variant match(device = {arch(amdgcn)})
30 
31 constexpr const llvm::omp::GV &getGridValue() {
32   return llvm::omp::AMDGPUGridValues;
33 }
34 
35 uint32_t getGridDim(uint32_t n, uint16_t d) {
36   uint32_t q = n / d;
37   return q + (n > q * d);
38 }
39 
40 uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,
41                          uint16_t group_size) {
42   uint32_t r = grid_size - group_id * group_size;
43   return (r < group_size) ? r : group_size;
44 }
45 
46 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
47 
48 LaneMaskTy lanemaskLT() {
49   uint32_t Lane = mapping::getThreadIdInWarp();
50   int64_t Ballot = mapping::activemask();
51   uint64_t Mask = ((uint64_t)1 << Lane) - (uint64_t)1;
52   return Mask & Ballot;
53 }
54 
55 LaneMaskTy lanemaskGT() {
56   uint32_t Lane = mapping::getThreadIdInWarp();
57   if (Lane == (mapping::getWarpSize() - 1))
58     return 0;
59   int64_t Ballot = mapping::activemask();
60   uint64_t Mask = (~((uint64_t)0)) << (Lane + 1);
61   return Mask & Ballot;
62 }
63 
64 uint32_t getThreadIdInWarp() {
65   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
66 }
67 
68 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
69 
70 uint32_t getBlockSize() {
71   // TODO: verify this logic for generic mode.
72   return getWorkgroupDim(__builtin_amdgcn_workgroup_id_x(),
73                          __builtin_amdgcn_grid_size_x(),
74                          __builtin_amdgcn_workgroup_size_x());
75 }
76 
77 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }
78 
79 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }
80 
81 uint32_t getNumberOfBlocks() {
82   return getGridDim(__builtin_amdgcn_grid_size_x(),
83                     __builtin_amdgcn_workgroup_size_x());
84 }
85 
86 uint32_t getNumberOfProcessorElements() {
87   // TODO
88   return mapping::getBlockSize();
89 }
90 
91 uint32_t getWarpId() {
92   return mapping::getThreadIdInBlock() / mapping::getWarpSize();
93 }
94 
95 uint32_t getNumberOfWarpsInBlock() {
96   return mapping::getBlockSize() / mapping::getWarpSize();
97 }
98 
99 #pragma omp end declare variant
100 ///}
101 
102 /// NVPTX Implementation
103 ///
104 ///{
105 #pragma omp begin declare variant match(                                       \
106     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
107 
108 constexpr const llvm::omp::GV &getGridValue() {
109   return llvm::omp::NVPTXGridValues;
110 }
111 
112 LaneMaskTy activemask() {
113   unsigned int Mask;
114   asm("activemask.b32 %0;" : "=r"(Mask));
115   return Mask;
116 }
117 
118 LaneMaskTy lanemaskLT() {
119   __kmpc_impl_lanemask_t Res;
120   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(Res));
121   return Res;
122 }
123 
124 LaneMaskTy lanemaskGT() {
125   __kmpc_impl_lanemask_t Res;
126   asm("mov.u32 %0, %%lanemask_gt;" : "=r"(Res));
127   return Res;
128 }
129 
130 uint32_t getThreadIdInWarp() {
131   return mapping::getThreadIdInBlock() & (mapping::getWarpSize() - 1);
132 }
133 
134 uint32_t getThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
135 
136 uint32_t getBlockSize() {
137   return __nvvm_read_ptx_sreg_ntid_x() -
138          (!mapping::isSPMDMode() * mapping::getWarpSize());
139 }
140 
141 uint32_t getKernelSize() { return __nvvm_read_ptx_sreg_nctaid_x(); }
142 
143 uint32_t getBlockId() { return __nvvm_read_ptx_sreg_ctaid_x(); }
144 
145 uint32_t getNumberOfBlocks() { return __nvvm_read_ptx_sreg_nctaid_x(); }
146 
147 uint32_t getNumberOfProcessorElements() {
148   return __nvvm_read_ptx_sreg_ntid_x();
149 }
150 
151 uint32_t getWarpId() {
152   return mapping::getThreadIdInBlock() / mapping::getWarpSize();
153 }
154 
155 uint32_t getNumberOfWarpsInBlock() {
156   return (mapping::getBlockSize() + mapping::getWarpSize() - 1) /
157          mapping::getWarpSize();
158 }
159 
160 #pragma omp end declare variant
161 ///}
162 
163 uint32_t getWarpSize() { return getGridValue().GV_Warp_Size; }
164 
165 } // namespace impl
166 } // namespace _OMP
167 
168 bool mapping::isMainThreadInGenericMode() {
169   if (mapping::isSPMDMode() || icv::Level)
170     return false;
171 
172   // Check if this is the last warp in the block.
173   uint32_t MainTId = (mapping::getNumberOfProcessorElements() - 1) &
174                      ~(mapping::getWarpSize() - 1);
175   return mapping::getThreadIdInBlock() == MainTId;
176 }
177 
178 bool mapping::isLeaderInWarp() {
179   __kmpc_impl_lanemask_t Active = mapping::activemask();
180   __kmpc_impl_lanemask_t LaneMaskLT = mapping::lanemaskLT();
181   return utils::popc(Active & LaneMaskLT) == 0;
182 }
183 
184 LaneMaskTy mapping::activemask() { return impl::activemask(); }
185 
186 LaneMaskTy mapping::lanemaskLT() { return impl::lanemaskLT(); }
187 
188 LaneMaskTy mapping::lanemaskGT() { return impl::lanemaskGT(); }
189 
190 uint32_t mapping::getThreadIdInWarp() { return impl::getThreadIdInWarp(); }
191 
192 uint32_t mapping::getThreadIdInBlock() { return impl::getThreadIdInBlock(); }
193 
194 uint32_t mapping::getBlockSize() { return impl::getBlockSize(); }
195 
196 uint32_t mapping::getKernelSize() { return impl::getKernelSize(); }
197 
198 uint32_t mapping::getBlockId() { return impl::getBlockId(); }
199 
200 uint32_t mapping::getNumberOfBlocks() { return impl::getNumberOfBlocks(); }
201 
202 uint32_t mapping::getNumberOfProcessorElements() {
203   return impl::getNumberOfProcessorElements();
204 }
205 
206 uint32_t mapping::getWarpId() { return impl::getWarpId(); }
207 
208 uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }
209 
210 uint32_t mapping::getNumberOfWarpsInBlock() {
211   return impl::getNumberOfWarpsInBlock();
212 }
213 
214 /// Execution mode
215 ///
216 ///{
217 static int SHARED(IsSPMDMode);
218 
219 void mapping::init(bool IsSPMD) {
220   if (!mapping::getThreadIdInBlock())
221     IsSPMDMode = IsSPMD;
222 }
223 
224 bool mapping::isSPMDMode() { return IsSPMDMode; }
225 
226 bool mapping::isGenericMode() { return !isSPMDMode(); }
227 ///}
228 
229 #pragma omp end declare target
230