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