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