gem5  v22.0.0.2
hsa_queue_entry.hh
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2017-2018 Advanced Micro Devices, Inc.
3  * All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  *
8  * 1. Redistributions of source code must retain the above copyright notice,
9  * this list of conditions and the following disclaimer.
10  *
11  * 2. Redistributions in binary form must reproduce the above copyright notice,
12  * this list of conditions and the following disclaimer in the documentation
13  * and/or other materials provided with the distribution.
14  *
15  * 3. Neither the name of the copyright holder nor the names of its
16  * contributors may be used to endorse or promote products derived from this
17  * software without specific prior written permission.
18  *
19  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
20  * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
21  * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
22  * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
23  * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
24  * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
25  * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
26  * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
27  * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
28  * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
29  * POSSIBILITY OF SUCH DAMAGE.
30  */
31 
41 #ifndef __GPU_COMPUTE_HSA_QUEUE_ENTRY__
42 #define __GPU_COMPUTE_HSA_QUEUE_ENTRY__
43 
44 #include <bitset>
45 #include <cstdint>
46 #include <cstring>
47 #include <iostream>
48 #include <vector>
49 
50 #include "base/intmath.hh"
51 #include "base/types.hh"
52 #include "dev/hsa/hsa_packet.hh"
53 #include "dev/hsa/hsa_queue.hh"
55 
56 namespace gem5
57 {
58 
60 {
61  public:
62  HSAQueueEntry(std::string kernel_name, uint32_t queue_id,
63  int dispatch_id, void *disp_pkt, AMDKernelCode *akc,
64  Addr host_pkt_addr, Addr code_addr)
65  : kernName(kernel_name),
66  _wgSize{{(int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_x,
67  (int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_y,
68  (int)((_hsa_dispatch_packet_t*)disp_pkt)->workgroup_size_z}},
69  _gridSize{{(int)((_hsa_dispatch_packet_t*)disp_pkt)->grid_size_x,
70  (int)((_hsa_dispatch_packet_t*)disp_pkt)->grid_size_y,
71  (int)((_hsa_dispatch_packet_t*)disp_pkt)->grid_size_z}},
72  numVgprs(akc->workitem_vgpr_count),
73  numSgprs(akc->wavefront_sgpr_count),
74  _queueId(queue_id), _dispatchId(dispatch_id), dispPkt(disp_pkt),
75  _hostDispPktAddr(host_pkt_addr),
76  _completionSignal(((_hsa_dispatch_packet_t*)disp_pkt)
77  ->completion_signal),
78  codeAddress(code_addr),
79  kernargAddress(((_hsa_dispatch_packet_t*)disp_pkt)->kernarg_address),
81  _ldsSize((int)((_hsa_dispatch_packet_t*)disp_pkt)->
82  group_segment_size),
83  _privMemPerItem((int)((_hsa_dispatch_packet_t*)disp_pkt)->
84  private_segment_size),
85  _contextId(0), _wgId{{ 0, 0, 0 }},
87  _globalWgId(0), dispatchComplete(false)
88 
89  {
90  // Precompiled BLIT kernels actually violate the spec a bit
91  // and don't set many of the required akc fields. For these kernels,
92  // we need to rip register usage from the resource registers.
93  //
94  // We can't get an exact number of registers from the resource
95  // registers because they round, but we can get an upper bound on it
96  if (!numVgprs)
97  numVgprs = (akc->granulated_workitem_vgpr_count + 1) * 4;
98 
99  // TODO: Granularity changes for GFX9!
100  if (!numSgprs)
101  numSgprs = (akc->granulated_wavefront_sgpr_count + 1) * 8;
102 
103  initialVgprState.reset();
104  initialSgprState.reset();
105 
106  for (int i = 0; i < MAX_DIM; ++i) {
108  _numWgTotal *= _numWg[i];
109  }
110 
111  parseKernelCode(akc);
112  }
113 
114  const std::string&
115  kernelName() const
116  {
117  return kernName;
118  }
119 
120  int
121  wgSize(int dim) const
122  {
123  assert(dim < MAX_DIM);
124  return _wgSize[dim];
125  }
126 
127  int
128  gridSize(int dim) const
129  {
130  assert(dim < MAX_DIM);
131  return _gridSize[dim];
132  }
133 
134  int
136  {
137  return numVgprs;
138  }
139 
140  int
142  {
143  return numSgprs;
144  }
145 
146  uint32_t
147  queueId() const
148  {
149  return _queueId;
150  }
151 
152  int
153  dispatchId() const
154  {
155  return _dispatchId;
156  }
157 
158  void*
160  {
161  return dispPkt;
162  }
163 
164  Addr
166  {
167  return _hostDispPktAddr;
168  }
169 
170  Addr
172  {
173  return _completionSignal;
174  }
175 
176  Addr
177  codeAddr() const
178  {
179  return codeAddress;
180  }
181 
182  Addr
183  kernargAddr() const
184  {
185  return kernargAddress;
186  }
187 
188  int
189  ldsSize() const
190  {
191  return _ldsSize;
192  }
193 
194  int privMemPerItem() const { return _privMemPerItem; }
195 
196  int
197  contextId() const
198  {
199  return _contextId;
200  }
201 
202  bool
203  dispComplete() const
204  {
205  return dispatchComplete;
206  }
207 
208  int
209  wgId(int dim) const
210  {
211  assert(dim < MAX_DIM);
212  return _wgId[dim];
213  }
214 
215  void
216  wgId(int dim, int val)
217  {
218  assert(dim < MAX_DIM);
219  _wgId[dim] = val;
220  }
221 
222  int
223  globalWgId() const
224  {
225  return _globalWgId;
226  }
227 
228  void
230  {
231  _globalWgId = val;
232  }
233 
234  int
235  numWg(int dim) const
236  {
237  assert(dim < MAX_DIM);
238  return _numWg[dim];
239  }
240 
241  void
243  {
244  ++_numWgCompleted;
245  }
246 
247  int
249  {
250  return _numWgCompleted;
251  }
252 
253  int
254  numWgTotal() const
255  {
256  return _numWgTotal;
257  }
258 
259  void
261  {
262  ++_wgId[0];
263  ++_globalWgId;
264 
265  if (wgId(0) * wgSize(0) >= gridSize(0)) {
266  _wgId[0] = 0;
267  ++_wgId[1];
268 
269  if (wgId(1) * wgSize(1) >= gridSize(1)) {
270  _wgId[1] = 0;
271  ++_wgId[2];
272 
273  if (wgId(2) * wgSize(2) >= gridSize(2)) {
274  dispatchComplete = true;
275  }
276  }
277  }
278  }
279 
280  int
282  {
283  return numWgArrivedAtBarrier;
284  }
285 
286  bool vgprBitEnabled(int bit) const
287  {
288  return initialVgprState.test(bit);
289  }
290 
291  bool sgprBitEnabled(int bit) const
292  {
293  return initialSgprState.test(bit);
294  }
295 
301 
308 
309  // the maximum number of dimensions for a grid or workgroup
310  const static int MAX_DIM = 3;
311 
312  /* getter */
313  int
315  return _outstandingInvs;
316  }
317 
323  bool
325  {
326  return (_outstandingInvs != -1);
327  }
328 
334  void
336  {
338  assert(_outstandingInvs >= 0);
339  }
340 
344  void
346  {
347  _outstandingInvs = 0;
348  }
349 
353  bool
354  isInvDone() const
355  {
356  assert(_outstandingInvs >= 0);
357  return (_outstandingInvs == 0);
358  }
359 
360  int
362  {
363  return _outstandingWbs;
364  }
365 
371  void
373  {
374  _outstandingWbs += val;
375  assert(_outstandingWbs >= 0);
376  }
377 
378  private:
379  void
381  {
388  akc->enable_sgpr_queue_ptr);
413 
418  initialVgprState.set(WorkitemIdX, true);
421  }
422 
423  // name of the kernel associated with the AQL entry
424  std::string kernName;
425  // workgroup Size (3 dimensions)
426  std::array<int, MAX_DIM> _wgSize;
427  // grid Size (3 dimensions)
428  std::array<int, MAX_DIM> _gridSize;
429  // total number of VGPRs per work-item
430  int numVgprs;
431  // total number of SGPRs per wavefront
432  int numSgprs;
433  // id of AQL queue in which this entry is placed
434  uint32_t _queueId;
436  // raw AQL packet pointer
437  void *dispPkt;
438  // host-side addr of the dispatch packet
440  // pointer to bool
442  // base address of the raw machine code
444  // base address of the kernel args
463  int _ldsSize;
466  std::array<int, MAX_DIM> _wgId;
467  std::array<int, MAX_DIM> _numWg;
470  // The number of completed work groups
474 
475  std::bitset<NumVectorInitFields> initialVgprState;
476  std::bitset<NumScalarInitFields> initialSgprState;
477 };
478 
479 } // namespace gem5
480 
481 #endif // __GPU_COMPUTE_HSA_QUEUE_ENTRY__
gem5::_amd_queue_t
Definition: hsa_queue.hh:64
gem5::HSAQueueEntry::numWg
int numWg(int dim) const
Definition: hsa_queue_entry.hh:235
gem5::HSAQueueEntry::_hostDispPktAddr
Addr _hostDispPktAddr
Definition: hsa_queue_entry.hh:439
gem5::FlatScratchInit
@ FlatScratchInit
Definition: kernel_code.hh:61
gem5::HSAQueueEntry::parseKernelCode
void parseKernelCode(AMDKernelCode *akc)
Definition: hsa_queue_entry.hh:380
gem5::HSAQueueEntry::HSAQueueEntry
HSAQueueEntry(std::string kernel_name, uint32_t queue_id, int dispatch_id, void *disp_pkt, AMDKernelCode *akc, Addr host_pkt_addr, Addr code_addr)
Definition: hsa_queue_entry.hh:62
gem5::AMDKernelCode::enable_vgpr_workitem_id
uint32_t enable_vgpr_workitem_id
Definition: kernel_code.hh:132
gem5::WorkitemIdY
@ WorkitemIdY
Definition: kernel_code.hh:77
kernel_code.hh
gem5::HSAQueueEntry::_numWgTotal
int _numWgTotal
Definition: hsa_queue_entry.hh:468
gem5::HSAQueueEntry::privMemPerItem
int privMemPerItem() const
Definition: hsa_queue_entry.hh:194
gem5::HSAQueueEntry::_contextId
int _contextId
Definition: hsa_queue_entry.hh:465
gem5::HSAQueueEntry::_completionSignal
Addr _completionSignal
Definition: hsa_queue_entry.hh:441
gem5::AMDKernelCode::enable_sgpr_workgroup_id_x
uint32_t enable_sgpr_workgroup_id_x
Definition: kernel_code.hh:128
gem5::HSAQueueEntry::completionSignal
Addr completionSignal() const
Definition: hsa_queue_entry.hh:171
gem5::PrivateSegBuf
@ PrivateSegBuf
Definition: kernel_code.hh:56
gem5::AMDKernelCode::enable_sgpr_dispatch_id
uint32_t enable_sgpr_dispatch_id
Definition: kernel_code.hh:152
gem5::HSAQueueEntry::_numWgCompleted
int _numWgCompleted
Definition: hsa_queue_entry.hh:471
gem5::AMDKernelCode::enable_sgpr_grid_workgroup_count_z
uint32_t enable_sgpr_grid_workgroup_count_z
Definition: kernel_code.hh:157
gem5::HSAQueueEntry::_dispatchId
int _dispatchId
Definition: hsa_queue_entry.hh:435
gem5::HSAQueueEntry::numWgTotal
int numWgTotal() const
Definition: hsa_queue_entry.hh:254
gem5::HSAQueueEntry::contextId
int contextId() const
Definition: hsa_queue_entry.hh:197
gem5::HSAQueueEntry::kernargAddress
Addr kernargAddress
Definition: hsa_queue_entry.hh:445
gem5::AMDKernelCode::enable_sgpr_dispatch_ptr
uint32_t enable_sgpr_dispatch_ptr
Definition: kernel_code.hh:149
gem5::AMDKernelCode::enable_sgpr_private_segment_wave_byte_offset
uint32_t enable_sgpr_private_segment_wave_byte_offset
Definition: kernel_code.hh:125
gem5::HSAQueueEntry::_outstandingInvs
int _outstandingInvs
Number of outstanding invs for the kernel.
Definition: hsa_queue_entry.hh:454
gem5::HSAQueueEntry::_privMemPerItem
int _privMemPerItem
Definition: hsa_queue_entry.hh:464
gem5::HSAQueueEntry
Definition: hsa_queue_entry.hh:59
gem5::AMDKernelCode::enable_sgpr_workgroup_id_z
uint32_t enable_sgpr_workgroup_id_z
Definition: kernel_code.hh:130
gem5::HSAQueueEntry::queueId
uint32_t queueId() const
Definition: hsa_queue_entry.hh:147
gem5::HSAQueueEntry::wgId
void wgId(int dim, int val)
Definition: hsa_queue_entry.hh:216
gem5::X86ISA::val
Bitfield< 63 > val
Definition: misc.hh:769
gem5::WorkgroupInfo
@ WorkgroupInfo
Definition: kernel_code.hh:69
gem5::AMDKernelCode::enable_sgpr_private_segment_buffer
uint32_t enable_sgpr_private_segment_buffer
Definition: kernel_code.hh:148
gem5::HSAQueueEntry::wgId
int wgId(int dim) const
Definition: hsa_queue_entry.hh:209
gem5::HSAQueueEntry::markWgDispatch
void markWgDispatch()
Definition: hsa_queue_entry.hh:260
gem5::PrivSegWaveByteOffset
@ PrivSegWaveByteOffset
Definition: kernel_code.hh:70
gem5::HSAQueueEntry::kernName
std::string kernName
Definition: hsa_queue_entry.hh:424
gem5::ArmISA::i
Bitfield< 7 > i
Definition: misc_types.hh:67
gem5::HSAQueueEntry::kernargAddr
Addr kernargAddr() const
Definition: hsa_queue_entry.hh:183
gem5::HSAQueueEntry::numScalarRegs
int numScalarRegs() const
Definition: hsa_queue_entry.hh:141
hsa_queue.hh
gem5::AMDKernelCode::enable_sgpr_private_segment_size
uint32_t enable_sgpr_private_segment_size
Definition: kernel_code.hh:154
gem5::HSAQueueEntry::dispatchComplete
bool dispatchComplete
Definition: hsa_queue_entry.hh:473
gem5::HSAQueueEntry::numWgAtBarrier
int numWgAtBarrier() const
Definition: hsa_queue_entry.hh:281
gem5::HSAQueueEntry::amdQueue
_amd_queue_t amdQueue
Keep a copy of the AMD HSA queue because we need info from some of its fields to initialize register ...
Definition: hsa_queue_entry.hh:307
gem5::WorkgroupIdY
@ WorkgroupIdY
Definition: kernel_code.hh:67
gem5::WorkgroupIdZ
@ WorkgroupIdZ
Definition: kernel_code.hh:68
gem5::HSAQueueEntry::numWgArrivedAtBarrier
int numWgArrivedAtBarrier
Definition: hsa_queue_entry.hh:469
gem5::HSAQueueEntry::isInvStarted
bool isInvStarted()
Whether invalidate has started or finished -1 is the initial value indicating inv has not started for...
Definition: hsa_queue_entry.hh:324
gem5::PrivateSegSize
@ PrivateSegSize
Definition: kernel_code.hh:62
gem5::HSAQueueEntry::markInvDone
void markInvDone()
Forcefully change the state to be inv done.
Definition: hsa_queue_entry.hh:345
gem5::HSAQueueEntry::_queueId
uint32_t _queueId
Definition: hsa_queue_entry.hh:434
gem5::AMDKernelCode::enable_sgpr_workgroup_info
uint32_t enable_sgpr_workgroup_info
Definition: kernel_code.hh:131
gem5::HSAQueueEntry::dispPktPtr
void * dispPktPtr()
Definition: hsa_queue_entry.hh:159
gem5::HSAQueueEntry::MAX_DIM
const static int MAX_DIM
Definition: hsa_queue_entry.hh:310
gem5::WorkitemIdX
@ WorkitemIdX
Definition: kernel_code.hh:76
gem5::HSAQueueEntry::notifyWgCompleted
void notifyWgCompleted()
Definition: hsa_queue_entry.hh:242
gem5::HSAQueueEntry::numVectorRegs
int numVectorRegs() const
Definition: hsa_queue_entry.hh:135
gem5::HSAQueueEntry::codeAddress
Addr codeAddress
Definition: hsa_queue_entry.hh:443
gem5::AMDKernelCode::enable_sgpr_flat_scratch_init
uint32_t enable_sgpr_flat_scratch_init
Definition: kernel_code.hh:153
gem5::HSAQueueEntry::vgprBitEnabled
bool vgprBitEnabled(int bit) const
Definition: hsa_queue_entry.hh:286
gem5::KernargSegPtr
@ KernargSegPtr
Definition: kernel_code.hh:59
gem5::DispatchId
@ DispatchId
Definition: kernel_code.hh:60
gem5::HSAQueueEntry::_numWg
std::array< int, MAX_DIM > _numWg
Definition: hsa_queue_entry.hh:467
gem5::AMDKernelCode::enable_sgpr_kernarg_segment_ptr
uint32_t enable_sgpr_kernarg_segment_ptr
Definition: kernel_code.hh:151
gem5::HSAQueueEntry::outstandingInvs
int outstandingInvs()
Definition: hsa_queue_entry.hh:314
gem5::HSAQueueEntry::isInvDone
bool isInvDone() const
Is invalidate done?
Definition: hsa_queue_entry.hh:354
hsa_packet.hh
gem5::HSAQueueEntry::hostDispPktAddr
Addr hostDispPktAddr() const
Definition: hsa_queue_entry.hh:165
gem5::HSAQueueEntry::dispPkt
void * dispPkt
Definition: hsa_queue_entry.hh:437
gem5::HSAQueueEntry::updateOutstandingInvs
void updateOutstandingInvs(int val)
update the number of pending invalidate requests
Definition: hsa_queue_entry.hh:335
gem5::HSAQueueEntry::wgSize
int wgSize(int dim) const
Definition: hsa_queue_entry.hh:121
gem5::HSAQueueEntry::_globalWgId
int _globalWgId
Definition: hsa_queue_entry.hh:472
gem5::AMDKernelCode::enable_sgpr_queue_ptr
uint32_t enable_sgpr_queue_ptr
Definition: kernel_code.hh:150
gem5::HSAQueueEntry::codeAddr
Addr codeAddr() const
Definition: hsa_queue_entry.hh:177
gem5::_hsa_dispatch_packet_t
Definition: hsa_packet.hh:53
gem5::HSAQueueEntry::dispComplete
bool dispComplete() const
Definition: hsa_queue_entry.hh:203
gem5::HSAQueueEntry::_gridSize
std::array< int, MAX_DIM > _gridSize
Definition: hsa_queue_entry.hh:428
gem5::HSAQueueEntry::globalWgId
int globalWgId() const
Definition: hsa_queue_entry.hh:223
gem5::HSAQueueEntry::gridSize
int gridSize(int dim) const
Definition: hsa_queue_entry.hh:128
gem5::GridWorkgroupCountX
@ GridWorkgroupCountX
Definition: kernel_code.hh:63
gem5::HSAQueueEntry::hostAMDQueueAddr
Addr hostAMDQueueAddr
Host-side addr of the amd_queue_t on which this task was queued.
Definition: hsa_queue_entry.hh:300
gem5::HSAQueueEntry::initialSgprState
std::bitset< NumScalarInitFields > initialSgprState
Definition: hsa_queue_entry.hh:476
gem5::WorkitemIdZ
@ WorkitemIdZ
Definition: kernel_code.hh:78
gem5::HSAQueueEntry::numVgprs
int numVgprs
Definition: hsa_queue_entry.hh:430
gem5::GridWorkgroupCountZ
@ GridWorkgroupCountZ
Definition: kernel_code.hh:65
gem5::HSAQueueEntry::numWgCompleted
int numWgCompleted() const
Definition: hsa_queue_entry.hh:248
gem5::HSAQueueEntry::initialVgprState
std::bitset< NumVectorInitFields > initialVgprState
Definition: hsa_queue_entry.hh:475
gem5::Addr
uint64_t Addr
Address type This will probably be moved somewhere else in the near future.
Definition: types.hh:147
gem5::HSAQueueEntry::globalWgId
void globalWgId(int val)
Definition: hsa_queue_entry.hh:229
gem5::divCeil
static constexpr T divCeil(const T &a, const U &b)
Definition: intmath.hh:110
gem5::HSAQueueEntry::kernelName
const std::string & kernelName() const
Definition: hsa_queue_entry.hh:115
gem5::HSAQueueEntry::outstandingWbs
int outstandingWbs() const
Definition: hsa_queue_entry.hh:361
gem5::AMDKernelCode::enable_sgpr_workgroup_id_y
uint32_t enable_sgpr_workgroup_id_y
Definition: kernel_code.hh:129
gem5::HSAQueueEntry::_ldsSize
int _ldsSize
Definition: hsa_queue_entry.hh:463
gem5::QueuePtr
@ QueuePtr
Definition: kernel_code.hh:58
gem5::WorkgroupIdX
@ WorkgroupIdX
Definition: kernel_code.hh:66
types.hh
gem5::HSAQueueEntry::updateOutstandingWbs
void updateOutstandingWbs(int val)
Update the number of pending writeback requests.
Definition: hsa_queue_entry.hh:372
gem5::GridWorkgroupCountY
@ GridWorkgroupCountY
Definition: kernel_code.hh:64
gem5::DispatchPtr
@ DispatchPtr
Definition: kernel_code.hh:57
gem5::HSAQueueEntry::sgprBitEnabled
bool sgprBitEnabled(int bit) const
Definition: hsa_queue_entry.hh:291
gem5::HSAQueueEntry::_outstandingWbs
int _outstandingWbs
Number of outstanding wbs for the kernel values: 0: 1)initial value, flush has not started for the ke...
Definition: hsa_queue_entry.hh:462
gem5::HSAQueueEntry::_wgSize
std::array< int, MAX_DIM > _wgSize
Definition: hsa_queue_entry.hh:426
gem5::AMDKernelCode::enable_sgpr_grid_workgroup_count_y
uint32_t enable_sgpr_grid_workgroup_count_y
Definition: kernel_code.hh:156
gem5::HSAQueueEntry::numSgprs
int numSgprs
Definition: hsa_queue_entry.hh:432
gem5::HSAQueueEntry::ldsSize
int ldsSize() const
Definition: hsa_queue_entry.hh:189
intmath.hh
gem5
Reference material can be found at the JEDEC website: UFS standard http://www.jedec....
Definition: gpu_translation_state.hh:37
gem5::HSAQueueEntry::dispatchId
int dispatchId() const
Definition: hsa_queue_entry.hh:153
gem5::HSAQueueEntry::_wgId
std::array< int, MAX_DIM > _wgId
Definition: hsa_queue_entry.hh:466
gem5::AMDKernelCode::enable_sgpr_grid_workgroup_count_x
uint32_t enable_sgpr_grid_workgroup_count_x
Definition: kernel_code.hh:155
gem5::AMDKernelCode
Definition: kernel_code.hh:82

Generated on Thu Jul 28 2022 13:32:33 for gem5 by doxygen 1.8.17