Coverage Report

Created: 2019-07-24 05:18

/Users/buildslave/jenkins/workspace/clang-stage2-coverage-R/llvm/lib/Target/AMDGPU/AMDKernelCodeT.h
Line
Count
Source
1
//===-- AMDGPUKernelCodeT.h - Print AMDGPU assembly code ---------*- 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
/// \file AMDKernelCodeT.h
9
//===----------------------------------------------------------------------===//
10
11
#ifndef AMDKERNELCODET_H
12
#define AMDKERNELCODET_H
13
14
#include "llvm/MC/SubtargetFeature.h"
15
16
#include <cstddef>
17
#include <cstdint>
18
19
#include "llvm/Support/Debug.h"
20
//---------------------------------------------------------------------------//
21
// AMD Kernel Code, and its dependencies                                     //
22
//---------------------------------------------------------------------------//
23
24
typedef uint8_t hsa_powertwo8_t;
25
typedef uint32_t hsa_ext_code_kind_t;
26
typedef uint8_t hsa_ext_brig_profile8_t;
27
typedef uint8_t hsa_ext_brig_machine_model8_t;
28
typedef uint64_t hsa_ext_control_directive_present64_t;
29
typedef uint16_t hsa_ext_exception_kind16_t;
30
typedef uint32_t hsa_ext_code_kind32_t;
31
32
typedef struct hsa_dim3_s {
33
  uint32_t x;
34
  uint32_t y;
35
  uint32_t z;
36
} hsa_dim3_t;
37
38
/// The version of the amd_*_code_t struct. Minor versions must be
39
/// backward compatible.
40
typedef uint32_t amd_code_version32_t;
41
enum amd_code_version_t {
42
  AMD_CODE_VERSION_MAJOR = 0,
43
  AMD_CODE_VERSION_MINOR = 1
44
};
45
46
// Sets val bits for specified mask in specified dst packed instance.
47
#define AMD_HSA_BITS_SET(dst, mask, val)                                       \
48
1.16k
  dst &= (~(1 << mask ## _SHIFT) & ~mask);                                     \
49
1.16k
  dst |= (((val) << mask ## _SHIFT) & mask)
50
51
// Gets bits for specified mask from specified src packed instance.
52
#define AMD_HSA_BITS_GET(src, mask)                                            \
53
  ((src & mask) >> mask ## _SHIFT)                                             \
54
55
/// The values used to define the number of bytes to use for the
56
/// swizzle element size.
57
enum amd_element_byte_size_t {
58
  AMD_ELEMENT_2_BYTES = 0,
59
  AMD_ELEMENT_4_BYTES = 1,
60
  AMD_ELEMENT_8_BYTES = 2,
61
  AMD_ELEMENT_16_BYTES = 3
62
};
63
64
/// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and
65
/// COMPUTE_PGM_RSRC2 registers.
66
typedef uint64_t amd_compute_pgm_resource_register64_t;
67
68
/// Every amd_*_code_t has the following properties, which are composed of
69
/// a number of bit fields. Every bit field has a mask (AMD_CODE_PROPERTY_*),
70
/// bit width (AMD_CODE_PROPERTY_*_WIDTH, and bit shift amount
71
/// (AMD_CODE_PROPERTY_*_SHIFT) for convenient access. Unused bits must be 0.
72
///
73
/// (Note that bit fields cannot be used as their layout is
74
/// implementation defined in the C standard and so cannot be used to
75
/// specify an ABI)
76
typedef uint32_t amd_code_property32_t;
77
enum amd_code_property_mask_t {
78
79
  /// Enable the setup of the SGPR user data registers
80
  /// (AMD_CODE_PROPERTY_ENABLE_SGPR_*), see documentation of amd_kernel_code_t
81
  /// for initial register state.
82
  ///
83
  /// The total number of SGPRuser data registers requested must not
84
  /// exceed 16. Any requests beyond 16 will be ignored.
85
  ///
86
  /// Used to set COMPUTE_PGM_RSRC2.USER_SGPR (set to total count of
87
  /// SGPR user data registers enabled up to 16).
88
89
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT = 0,
90
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH = 1,
91
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT,
92
93
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT = 1,
94
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH = 1,
95
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT,
96
97
  AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT = 2,
98
  AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH = 1,
99
  AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT,
100
101
  AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT = 3,
102
  AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH = 1,
103
  AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT,
104
105
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT = 4,
106
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH = 1,
107
  AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT,
108
109
  AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT = 5,
110
  AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH = 1,
111
  AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT,
112
113
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT = 6,
114
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH = 1,
115
  AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT,
116
117
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT = 7,
118
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH = 1,
119
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT,
120
121
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT = 8,
122
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH = 1,
123
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT,
124
125
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT = 9,
126
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH = 1,
127
  AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z = ((1 << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT,
128
129
  AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT = 10,
130
  AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_WIDTH = 1,
131
  AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32 = ((1 << AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT,
132
133
  AMD_CODE_PROPERTY_RESERVED1_SHIFT = 11,
134
  AMD_CODE_PROPERTY_RESERVED1_WIDTH = 5,
135
  AMD_CODE_PROPERTY_RESERVED1 = ((1 << AMD_CODE_PROPERTY_RESERVED1_WIDTH) - 1) << AMD_CODE_PROPERTY_RESERVED1_SHIFT,
136
137
  /// Control wave ID base counter for GDS ordered-append. Used to set
138
  /// COMPUTE_DISPATCH_INITIATOR.ORDERED_APPEND_ENBL. (Not sure if
139
  /// ORDERED_APPEND_MODE also needs to be settable)
140
  AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT = 16,
141
  AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH = 1,
142
  AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS = ((1 << AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH) - 1) << AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT,
143
144
  /// The interleave (swizzle) element size in bytes required by the
145
  /// code for private memory. This must be 2, 4, 8 or 16. This value
146
  /// is provided to the finalizer when it is invoked and is recorded
147
  /// here. The hardware will interleave the memory requests of each
148
  /// lane of a wavefront by this element size to ensure each
149
  /// work-item gets a distinct memory memory location. Therefore, the
150
  /// finalizer ensures that all load and store operations done to
151
  /// private memory do not exceed this size. For example, if the
152
  /// element size is 4 (32-bits or dword) and a 64-bit value must be
153
  /// loaded, the finalizer will generate two 32-bit loads. This
154
  /// ensures that the interleaving will get the work-item
155
  /// specific dword for both halves of the 64-bit value. If it just
156
  /// did a 64-bit load then it would get one dword which belonged to
157
  /// its own work-item, but the second dword would belong to the
158
  /// adjacent lane work-item since the interleaving is in dwords.
159
  ///
160
  /// The value used must match the value that the runtime configures
161
  /// the GPU flat scratch (SH_STATIC_MEM_CONFIG.ELEMENT_SIZE). This
162
  /// is generally DWORD.
163
  ///
164
  /// uSE VALUES FROM THE AMD_ELEMENT_BYTE_SIZE_T ENUM.
165
  AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT = 17,
166
  AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH = 2,
167
  AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE = ((1 << AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH) - 1) << AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT,
168
169
  /// Are global memory addresses 64 bits. Must match
170
  /// amd_kernel_code_t.hsail_machine_model ==
171
  /// HSA_MACHINE_LARGE. Must also match
172
  /// SH_MEM_CONFIG.PTR32 (GFX6 (SI)/GFX7 (CI)),
173
  /// SH_MEM_CONFIG.ADDRESS_MODE (GFX8 (VI)+).
174
  AMD_CODE_PROPERTY_IS_PTR64_SHIFT = 19,
175
  AMD_CODE_PROPERTY_IS_PTR64_WIDTH = 1,
176
  AMD_CODE_PROPERTY_IS_PTR64 = ((1 << AMD_CODE_PROPERTY_IS_PTR64_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_PTR64_SHIFT,
177
178
  /// Indicate if the generated ISA is using a dynamically sized call
179
  /// stack. This can happen if calls are implemented using a call
180
  /// stack and recursion, alloca or calls to indirect functions are
181
  /// present. In these cases the Finalizer cannot compute the total
182
  /// private segment size at compile time. In this case the
183
  /// workitem_private_segment_byte_size only specifies the statically
184
  /// know private segment size, and additional space must be added
185
  /// for the call stack.
186
  AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT = 20,
187
  AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH = 1,
188
  AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK = ((1 << AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT,
189
190
  /// Indicate if code generated has support for debugging.
191
  AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT = 21,
192
  AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH = 1,
193
  AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED = ((1 << AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT,
194
195
  AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_SHIFT = 22,
196
  AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_WIDTH = 1,
197
  AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED = ((1 << AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_SHIFT,
198
199
  AMD_CODE_PROPERTY_RESERVED2_SHIFT = 23,
200
  AMD_CODE_PROPERTY_RESERVED2_WIDTH = 9,
201
  AMD_CODE_PROPERTY_RESERVED2 = ((1 << AMD_CODE_PROPERTY_RESERVED2_WIDTH) - 1) << AMD_CODE_PROPERTY_RESERVED2_SHIFT
202
};
203
204
/// The hsa_ext_control_directives_t specifies the values for the HSAIL
205
/// control directives. These control how the finalizer generates code. This
206
/// struct is used both as an argument to hsaFinalizeKernel to specify values for
207
/// the control directives, and is used in HsaKernelCode to record the values of
208
/// the control directives that the finalize used when generating the code which
209
/// either came from the finalizer argument or explicit HSAIL control
210
/// directives. See the definition of the control directives in HSA Programmer's
211
/// Reference Manual which also defines how the values specified as finalizer
212
/// arguments have to agree with the control directives in the HSAIL code.
213
typedef struct hsa_ext_control_directives_s {
214
  /// This is a bit set indicating which control directives have been
215
  /// specified. If the value is 0 then there are no control directives specified
216
  /// and the rest of the fields can be ignored. The bits are accessed using the
217
  /// hsa_ext_control_directives_present_mask_t. Any control directive that is not
218
  /// enabled in this bit set must have the value of all 0s.
219
  hsa_ext_control_directive_present64_t enabled_control_directives;
220
221
  /// If enableBreakExceptions is not enabled then must be 0, otherwise must be
222
  /// non-0 and specifies the set of HSAIL exceptions that must have the BREAK
223
  /// policy enabled. If this set is not empty then the generated code may have
224
  /// lower performance than if the set is empty. If the kernel being finalized
225
  /// has any enablebreakexceptions control directives, then the values specified
226
  /// by this argument are unioned with the values in these control
227
  /// directives. If any of the functions the kernel calls have an
228
  /// enablebreakexceptions control directive, then they must be equal or a
229
  /// subset of, this union.
230
  hsa_ext_exception_kind16_t enable_break_exceptions;
231
232
  /// If enableDetectExceptions is not enabled then must be 0, otherwise must be
233
  /// non-0 and specifies the set of HSAIL exceptions that must have the DETECT
234
  /// policy enabled. If this set is not empty then the generated code may have
235
  /// lower performance than if the set is empty. However, an implementation
236
  /// should endeavour to make the performance impact small. If the kernel being
237
  /// finalized has any enabledetectexceptions control directives, then the
238
  /// values specified by this argument are unioned with the values in these
239
  /// control directives. If any of the functions the kernel calls have an
240
  /// enabledetectexceptions control directive, then they must be equal or a
241
  /// subset of, this union.
242
  hsa_ext_exception_kind16_t enable_detect_exceptions;
243
244
  /// If maxDynamicGroupSize is not enabled then must be 0, and any amount of
245
  /// dynamic group segment can be allocated for a dispatch, otherwise the value
246
  /// specifies the maximum number of bytes of dynamic group segment that can be
247
  /// allocated for a dispatch. If the kernel being finalized has any
248
  /// maxdynamicsize control directives, then the values must be the same, and
249
  /// must be the same as this argument if it is enabled. This value can be used
250
  /// by the finalizer to determine the maximum number of bytes of group memory
251
  /// used by each work-group by adding this value to the group memory required
252
  /// for all group segment variables used by the kernel and all functions it
253
  /// calls, and group memory used to implement other HSAIL features such as
254
  /// fbarriers and the detect exception operations. This can allow the finalizer
255
  /// to determine the expected number of work-groups that can be executed by a
256
  /// compute unit and allow more resources to be allocated to the work-items if
257
  /// it is known that fewer work-groups can be executed due to group memory
258
  /// limitations.
259
  uint32_t max_dynamic_group_size;
260
261
  /// If maxFlatGridSize is not enabled then must be 0, otherwise must be greater
262
  /// than 0. See HSA Programmer's Reference Manual description of
263
  /// maxflatgridsize control directive.
264
  uint32_t max_flat_grid_size;
265
266
  /// If maxFlatWorkgroupSize is not enabled then must be 0, otherwise must be
267
  /// greater than 0. See HSA Programmer's Reference Manual description of
268
  /// maxflatworkgroupsize control directive.
269
  uint32_t max_flat_workgroup_size;
270
271
  /// If requestedWorkgroupsPerCu is not enabled then must be 0, and the
272
  /// finalizer is free to generate ISA that may result in any number of
273
  /// work-groups executing on a single compute unit. Otherwise, the finalizer
274
  /// should attempt to generate ISA that will allow the specified number of
275
  /// work-groups to execute on a single compute unit. This is only a hint and
276
  /// can be ignored by the finalizer. If the kernel being finalized, or any of
277
  /// the functions it calls, has a requested control directive, then the values
278
  /// must be the same. This can be used to determine the number of resources
279
  /// that should be allocated to a single work-group and work-item. For example,
280
  /// a low value may allow more resources to be allocated, resulting in higher
281
  /// per work-item performance, as it is known there will never be more than the
282
  /// specified number of work-groups actually executing on the compute
283
  /// unit. Conversely, a high value may allocate fewer resources, resulting in
284
  /// lower per work-item performance, which is offset by the fact it allows more
285
  /// work-groups to actually execute on the compute unit.
286
  uint32_t requested_workgroups_per_cu;
287
288
  /// If not enabled then all elements for Dim3 must be 0, otherwise every
289
  /// element must be greater than 0. See HSA Programmer's Reference Manual
290
  /// description of requiredgridsize control directive.
291
  hsa_dim3_t required_grid_size;
292
293
  /// If requiredWorkgroupSize is not enabled then all elements for Dim3 must be
294
  /// 0, and the produced code can be dispatched with any legal work-group range
295
  /// consistent with the dispatch dimensions. Otherwise, the code produced must
296
  /// always be dispatched with the specified work-group range. No element of the
297
  /// specified range must be 0. It must be consistent with required_dimensions
298
  /// and max_flat_workgroup_size. If the kernel being finalized, or any of the
299
  /// functions it calls, has a requiredworkgroupsize control directive, then the
300
  /// values must be the same. Specifying a value can allow the finalizer to
301
  /// optimize work-group id operations, and if the number of work-items in the
302
  /// work-group is less than the WAVESIZE then barrier operations can be
303
  /// optimized to just a memory fence.
304
  hsa_dim3_t required_workgroup_size;
305
306
  /// If requiredDim is not enabled then must be 0 and the produced kernel code
307
  /// can be dispatched with 1, 2 or 3 dimensions. If enabled then the value is
308
  /// 1..3 and the code produced must only be dispatched with a dimension that
309
  /// matches. Other values are illegal. If the kernel being finalized, or any of
310
  /// the functions it calls, has a requireddimsize control directive, then the
311
  /// values must be the same. This can be used to optimize the code generated to
312
  /// compute the absolute and flat work-group and work-item id, and the dim
313
  /// HSAIL operations.
314
  uint8_t required_dim;
315
316
  /// Reserved. Must be 0.
317
  uint8_t reserved[75];
318
} hsa_ext_control_directives_t;
319
320
/// AMD Kernel Code Object (amd_kernel_code_t). GPU CP uses the AMD Kernel
321
/// Code Object to set up the hardware to execute the kernel dispatch.
322
///
323
/// Initial Kernel Register State.
324
///
325
/// Initial kernel register state will be set up by CP/SPI prior to the start
326
/// of execution of every wavefront. This is limited by the constraints of the
327
/// current hardware.
328
///
329
/// The order of the SGPR registers is defined, but the Finalizer can specify
330
/// which ones are actually setup in the amd_kernel_code_t object using the
331
/// enable_sgpr_* bit fields. The register numbers used for enabled registers
332
/// are dense starting at SGPR0: the first enabled register is SGPR0, the next
333
/// enabled register is SGPR1 etc.; disabled registers do not have an SGPR
334
/// number.
335
///
336
/// The initial SGPRs comprise up to 16 User SRGPs that are set up by CP and
337
/// apply to all waves of the grid. It is possible to specify more than 16 User
338
/// SGPRs using the enable_sgpr_* bit fields, in which case only the first 16
339
/// are actually initialized. These are then immediately followed by the System
340
/// SGPRs that are set up by ADC/SPI and can have different values for each wave
341
/// of the grid dispatch.
342
///
343
/// SGPR register initial state is defined as follows:
344
///
345
/// Private Segment Buffer (enable_sgpr_private_segment_buffer):
346
///   Number of User SGPR registers: 4. V# that can be used, together with
347
///   Scratch Wave Offset as an offset, to access the Private/Spill/Arg
348
///   segments using a segment address. It must be set as follows:
349
///     - Base address: of the scratch memory area used by the dispatch. It
350
///       does not include the scratch wave offset. It will be the per process
351
///       SH_HIDDEN_PRIVATE_BASE_VMID plus any offset from this dispatch (for
352
///       example there may be a per pipe offset, or per AQL Queue offset).
353
///     - Stride + data_format: Element Size * Index Stride (???)
354
///     - Cache swizzle: ???
355
///     - Swizzle enable: SH_STATIC_MEM_CONFIG.SWIZZLE_ENABLE (must be 1 for
356
///       scratch)
357
///     - Num records: Flat Scratch Work Item Size / Element Size (???)
358
///     - Dst_sel_*: ???
359
///     - Num_format: ???
360
///     - Element_size: SH_STATIC_MEM_CONFIG.ELEMENT_SIZE (will be DWORD, must
361
///       agree with amd_kernel_code_t.privateElementSize)
362
///     - Index_stride: SH_STATIC_MEM_CONFIG.INDEX_STRIDE (will be 64 as must
363
///       be number of wavefront lanes for scratch, must agree with
364
///       amd_kernel_code_t.wavefrontSize)
365
///     - Add tid enable: 1
366
///     - ATC: from SH_MEM_CONFIG.PRIVATE_ATC,
367
///     - Hash_enable: ???
368
///     - Heap: ???
369
///     - Mtype: from SH_STATIC_MEM_CONFIG.PRIVATE_MTYPE
370
///     - Type: 0 (a buffer) (???)
371
///
372
/// Dispatch Ptr (enable_sgpr_dispatch_ptr):
373
///   Number of User SGPR registers: 2. 64 bit address of AQL dispatch packet
374
///   for kernel actually executing.
375
///
376
/// Queue Ptr (enable_sgpr_queue_ptr):
377
///   Number of User SGPR registers: 2. 64 bit address of AmdQueue object for
378
///   AQL queue on which the dispatch packet was queued.
379
///
380
/// Kernarg Segment Ptr (enable_sgpr_kernarg_segment_ptr):
381
///   Number of User SGPR registers: 2. 64 bit address of Kernarg segment. This
382
///   is directly copied from the kernargPtr in the dispatch packet. Having CP
383
///   load it once avoids loading it at the beginning of every wavefront.
384
///
385
/// Dispatch Id (enable_sgpr_dispatch_id):
386
///   Number of User SGPR registers: 2. 64 bit Dispatch ID of the dispatch
387
///   packet being executed.
388
///
389
/// Flat Scratch Init (enable_sgpr_flat_scratch_init):
390
///   Number of User SGPR registers: 2. This is 2 SGPRs.
391
///
392
///   For CI/VI:
393
///     The first SGPR is a 32 bit byte offset from SH_MEM_HIDDEN_PRIVATE_BASE
394
///     to base of memory for scratch for this dispatch. This is the same offset
395
///     used in computing the Scratch Segment Buffer base address. The value of
396
///     Scratch Wave Offset must be added by the kernel code and moved to
397
///     SGPRn-4 for use as the FLAT SCRATCH BASE in flat memory instructions.
398
///
399
///     The second SGPR is 32 bit byte size of a single work-item's scratch
400
///     memory usage. This is directly loaded from the dispatch packet Private
401
///     Segment Byte Size and rounded up to a multiple of DWORD.
402
///
403
///     \todo [Does CP need to round this to >4 byte alignment?]
404
///
405
///     The kernel code must move to SGPRn-3 for use as the FLAT SCRATCH SIZE in
406
///     flat memory instructions. Having CP load it once avoids loading it at
407
///     the beginning of every wavefront.
408
///
409
///   For PI:
410
///     This is the 64 bit base address of the scratch backing memory for
411
///     allocated by CP for this dispatch.
412
///
413
/// Private Segment Size (enable_sgpr_private_segment_size):
414
///   Number of User SGPR registers: 1. The 32 bit byte size of a single
415
///   work-item's scratch memory allocation. This is the value from the dispatch
416
///   packet. Private Segment Byte Size rounded up by CP to a multiple of DWORD.
417
///
418
///   \todo [Does CP need to round this to >4 byte alignment?]
419
///
420
///   Having CP load it once avoids loading it at the beginning of every
421
///   wavefront.
422
///
423
///   \todo [This will not be used for CI/VI since it is the same value as
424
///   the second SGPR of Flat Scratch Init. However, it is need for PI which
425
///   changes meaning of Flat Scratchg Init..]
426
///
427
/// Grid Work-Group Count X (enable_sgpr_grid_workgroup_count_x):
428
///   Number of User SGPR registers: 1. 32 bit count of the number of
429
///   work-groups in the X dimension for the grid being executed. Computed from
430
///   the fields in the HsaDispatchPacket as
431
///   ((gridSize.x+workgroupSize.x-1)/workgroupSize.x).
432
///
433
/// Grid Work-Group Count Y (enable_sgpr_grid_workgroup_count_y):
434
///   Number of User SGPR registers: 1. 32 bit count of the number of
435
///   work-groups in the Y dimension for the grid being executed. Computed from
436
///   the fields in the HsaDispatchPacket as
437
///   ((gridSize.y+workgroupSize.y-1)/workgroupSize.y).
438
///
439
///   Only initialized if <16 previous SGPRs initialized.
440
///
441
/// Grid Work-Group Count Z (enable_sgpr_grid_workgroup_count_z):
442
///   Number of User SGPR registers: 1. 32 bit count of the number of
443
///   work-groups in the Z dimension for the grid being executed. Computed
444
///   from the fields in the HsaDispatchPacket as
445
///   ((gridSize.z+workgroupSize.z-1)/workgroupSize.z).
446
///
447
///   Only initialized if <16 previous SGPRs initialized.
448
///
449
/// Work-Group Id X (enable_sgpr_workgroup_id_x):
450
///   Number of System SGPR registers: 1. 32 bit work group id in X dimension
451
///   of grid for wavefront. Always present.
452
///
453
/// Work-Group Id Y (enable_sgpr_workgroup_id_y):
454
///   Number of System SGPR registers: 1. 32 bit work group id in Y dimension
455
///   of grid for wavefront.
456
///
457
/// Work-Group Id Z (enable_sgpr_workgroup_id_z):
458
///   Number of System SGPR registers: 1. 32 bit work group id in Z dimension
459
///   of grid for wavefront. If present then Work-group Id Y will also be
460
///   present
461
///
462
/// Work-Group Info (enable_sgpr_workgroup_info):
463
///   Number of System SGPR registers: 1. {first_wave, 14'b0000,
464
///   ordered_append_term[10:0], threadgroup_size_in_waves[5:0]}
465
///
466
/// Private Segment Wave Byte Offset
467
/// (enable_sgpr_private_segment_wave_byte_offset):
468
///   Number of System SGPR registers: 1. 32 bit byte offset from base of
469
///   dispatch scratch base. Must be used as an offset with Private/Spill/Arg
470
///   segment address when using Scratch Segment Buffer. It must be added to
471
///   Flat Scratch Offset if setting up FLAT SCRATCH for flat addressing.
472
///
473
///
474
/// The order of the VGPR registers is defined, but the Finalizer can specify
475
/// which ones are actually setup in the amd_kernel_code_t object using the
476
/// enableVgpr*  bit fields. The register numbers used for enabled registers
477
/// are dense starting at VGPR0: the first enabled register is VGPR0, the next
478
/// enabled register is VGPR1 etc.; disabled registers do not have an VGPR
479
/// number.
480
///
481
/// VGPR register initial state is defined as follows:
482
///
483
/// Work-Item Id X (always initialized):
484
///   Number of registers: 1. 32 bit work item id in X dimension of work-group
485
///   for wavefront lane.
486
///
487
/// Work-Item Id X (enable_vgpr_workitem_id > 0):
488
///   Number of registers: 1. 32 bit work item id in Y dimension of work-group
489
///   for wavefront lane.
490
///
491
/// Work-Item Id X (enable_vgpr_workitem_id > 0):
492
///   Number of registers: 1. 32 bit work item id in Z dimension of work-group
493
///   for wavefront lane.
494
///
495
///
496
/// The setting of registers is being done by existing GPU hardware as follows:
497
///   1) SGPRs before the Work-Group Ids are set by CP using the 16 User Data
498
///      registers.
499
///   2) Work-group Id registers X, Y, Z are set by SPI which supports any
500
///      combination including none.
501
///   3) Scratch Wave Offset is also set by SPI which is why its value cannot
502
///      be added into the value Flat Scratch Offset which would avoid the
503
///      Finalizer generated prolog having to do the add.
504
///   4) The VGPRs are set by SPI which only supports specifying either (X),
505
///      (X, Y) or (X, Y, Z).
506
///
507
/// Flat Scratch Dispatch Offset and Flat Scratch Size are adjacent SGRRs so
508
/// they can be moved as a 64 bit value to the hardware required SGPRn-3 and
509
/// SGPRn-4 respectively using the Finalizer ?FLAT_SCRATCH? Register.
510
///
511
/// The global segment can be accessed either using flat operations or buffer
512
/// operations. If buffer operations are used then the Global Buffer used to
513
/// access HSAIL Global/Readonly/Kernarg (which are combine) segments using a
514
/// segment address is not passed into the kernel code by CP since its base
515
/// address is always 0. Instead the Finalizer generates prolog code to
516
/// initialize 4 SGPRs with a V# that has the following properties, and then
517
/// uses that in the buffer instructions:
518
///   - base address of 0
519
///   - no swizzle
520
///   - ATC=1
521
///   - MTYPE set to support memory coherence specified in
522
///     amd_kernel_code_t.globalMemoryCoherence
523
///
524
/// When the Global Buffer is used to access the Kernarg segment, must add the
525
/// dispatch packet kernArgPtr to a kernarg segment address before using this V#.
526
/// Alternatively scalar loads can be used if the kernarg offset is uniform, as
527
/// the kernarg segment is constant for the duration of the kernel execution.
528
///
529
530
typedef struct amd_kernel_code_s {
531
  uint32_t amd_kernel_code_version_major;
532
  uint32_t amd_kernel_code_version_minor;
533
  uint16_t amd_machine_kind;
534
  uint16_t amd_machine_version_major;
535
  uint16_t amd_machine_version_minor;
536
  uint16_t amd_machine_version_stepping;
537
538
  /// Byte offset (possibly negative) from start of amd_kernel_code_t
539
  /// object to kernel's entry point instruction. The actual code for
540
  /// the kernel is required to be 256 byte aligned to match hardware
541
  /// requirements (SQ cache line is 16). The code must be position
542
  /// independent code (PIC) for AMD devices to give runtime the
543
  /// option of copying code to discrete GPU memory or APU L2
544
  /// cache. The Finalizer should endeavour to allocate all kernel
545
  /// machine code in contiguous memory pages so that a device
546
  /// pre-fetcher will tend to only pre-fetch Kernel Code objects,
547
  /// improving cache performance.
548
  int64_t kernel_code_entry_byte_offset;
549
550
  /// Range of bytes to consider prefetching expressed as an offset
551
  /// and size. The offset is from the start (possibly negative) of
552
  /// amd_kernel_code_t object. Set both to 0 if no prefetch
553
  /// information is available.
554
  int64_t kernel_code_prefetch_byte_offset;
555
  uint64_t kernel_code_prefetch_byte_size;
556
557
  /// Reserved. Must be 0.
558
  uint64_t reserved0;
559
560
  /// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and
561
  /// COMPUTE_PGM_RSRC2 registers.
562
  uint64_t compute_pgm_resource_registers;
563
564
  /// Code properties. See amd_code_property_mask_t for a full list of
565
  /// properties.
566
  uint32_t code_properties;
567
568
  /// The amount of memory required for the combined private, spill
569
  /// and arg segments for a work-item in bytes. If
570
  /// is_dynamic_callstack is 1 then additional space must be added to
571
  /// this value for the call stack.
572
  uint32_t workitem_private_segment_byte_size;
573
574
  /// The amount of group segment memory required by a work-group in
575
  /// bytes. This does not include any dynamically allocated group
576
  /// segment memory that may be added when the kernel is
577
  /// dispatched.
578
  uint32_t workgroup_group_segment_byte_size;
579
580
  /// Number of byte of GDS required by kernel dispatch. Must be 0 if
581
  /// not using GDS.
582
  uint32_t gds_segment_byte_size;
583
584
  /// The size in bytes of the kernarg segment that holds the values
585
  /// of the arguments to the kernel. This could be used by CP to
586
  /// prefetch the kernarg segment pointed to by the dispatch packet.
587
  uint64_t kernarg_segment_byte_size;
588
589
  /// Number of fbarrier's used in the kernel and all functions it
590
  /// calls. If the implementation uses group memory to allocate the
591
  /// fbarriers then that amount must already be included in the
592
  /// workgroup_group_segment_byte_size total.
593
  uint32_t workgroup_fbarrier_count;
594
595
  /// Number of scalar registers used by a wavefront. This includes
596
  /// the special SGPRs for VCC, Flat Scratch Base, Flat Scratch Size
597
  /// and XNACK (for GFX8 (VI)). It does not include the 16 SGPR added if a
598
  /// trap handler is enabled. Used to set COMPUTE_PGM_RSRC1.SGPRS.
599
  uint16_t wavefront_sgpr_count;
600
601
  /// Number of vector registers used by each work-item. Used to set
602
  /// COMPUTE_PGM_RSRC1.VGPRS.
603
  uint16_t workitem_vgpr_count;
604
605
  /// If reserved_vgpr_count is 0 then must be 0. Otherwise, this is the
606
  /// first fixed VGPR number reserved.
607
  uint16_t reserved_vgpr_first;
608
609
  /// The number of consecutive VGPRs reserved by the client. If
610
  /// is_debug_supported then this count includes VGPRs reserved
611
  /// for debugger use.
612
  uint16_t reserved_vgpr_count;
613
614
  /// If reserved_sgpr_count is 0 then must be 0. Otherwise, this is the
615
  /// first fixed SGPR number reserved.
616
  uint16_t reserved_sgpr_first;
617
618
  /// The number of consecutive SGPRs reserved by the client. If
619
  /// is_debug_supported then this count includes SGPRs reserved
620
  /// for debugger use.
621
  uint16_t reserved_sgpr_count;
622
623
  /// If is_debug_supported is 0 then must be 0. Otherwise, this is the
624
  /// fixed SGPR number used to hold the wave scratch offset for the
625
  /// entire kernel execution, or uint16_t(-1) if the register is not
626
  /// used or not known.
627
  uint16_t debug_wavefront_private_segment_offset_sgpr;
628
629
  /// If is_debug_supported is 0 then must be 0. Otherwise, this is the
630
  /// fixed SGPR number of the first of 4 SGPRs used to hold the
631
  /// scratch V# used for the entire kernel execution, or uint16_t(-1)
632
  /// if the registers are not used or not known.
633
  uint16_t debug_private_segment_buffer_sgpr;
634
635
  /// The maximum byte alignment of variables used by the kernel in
636
  /// the specified memory segment. Expressed as a power of two. Must
637
  /// be at least HSA_POWERTWO_16.
638
  uint8_t kernarg_segment_alignment;
639
  uint8_t group_segment_alignment;
640
  uint8_t private_segment_alignment;
641
642
  /// Wavefront size expressed as a power of two. Must be a power of 2
643
  /// in range 1..64 inclusive. Used to support runtime query that
644
  /// obtains wavefront size, which may be used by application to
645
  /// allocated dynamic group memory and set the dispatch work-group
646
  /// size.
647
  uint8_t wavefront_size;
648
649
  int32_t call_convention;
650
  uint8_t reserved3[12];
651
  uint64_t runtime_loader_kernel_symbol;
652
  uint64_t control_directives[16];
653
} amd_kernel_code_t;
654
655
#endif // AMDKERNELCODET_H