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