LLVM  13.0.0git
AMDKernelCodeT.h
Go to the documentation of this file.
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 
20 typedef uint8_t hsa_powertwo8_t;
22 typedef uint8_t hsa_ext_brig_profile8_t;
27 
28 typedef struct hsa_dim3_s {
32 } hsa_dim3_t;
33 
34 /// The version of the amd_*_code_t struct. Minor versions must be
35 /// backward compatible.
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.
58 };
59 
60 /// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and
61 /// COMPUTE_PGM_RSRC2 registers.
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)
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 
88 
92 
96 
100 
104 
108 
112 
116 
120 
124 
128 
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)
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 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.
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)+).
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.
185 
186  /// Indicate if code generated has support for debugging.
190 
194 
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.
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.
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.
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.
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.
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.
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.
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.
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.
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.
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];
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 
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.
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.
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.
559 
560  /// Code properties. See amd_code_property_mask_t for a full list of
561  /// 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.
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.
575 
576  /// Number of byte of GDS required by kernel dispatch. Must be 0 if
577  /// not using GDS.
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.
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.
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.
596 
597  /// Number of vector registers used by each work-item. Used to set
598  /// COMPUTE_PGM_RSRC1.VGPRS.
600 
601  /// If reserved_vgpr_count is 0 then must be 0. Otherwise, this is the
602  /// first fixed VGPR number reserved.
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.
609 
610  /// If reserved_sgpr_count is 0 then must be 0. Otherwise, this is the
611  /// first fixed SGPR number reserved.
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.
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.
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.
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.
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 
646  uint8_t reserved3[12];
648  uint64_t control_directives[16];
649 };
650 
651 #endif // AMDKERNELCODET_H
AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT
Definition: AMDKernelCodeT.h:105
AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR
Definition: AMDKernelCodeT.h:95
amd_compute_pgm_resource_register64_t
uint64_t amd_compute_pgm_resource_register64_t
Shader program settings for CS.
Definition: AMDKernelCodeT.h:62
amd_code_version32_t
uint32_t amd_code_version32_t
The version of the amd_*_code_t struct.
Definition: AMDKernelCodeT.h:36
AMD_ELEMENT_4_BYTES
@ AMD_ELEMENT_4_BYTES
Definition: AMDKernelCodeT.h:55
AMD_CODE_PROPERTY_RESERVED2_WIDTH
@ AMD_CODE_PROPERTY_RESERVED2_WIDTH
Definition: AMDKernelCodeT.h:196
hsa_ext_control_directives_s::required_grid_size
hsa_dim3_t required_grid_size
If not enabled then all elements for Dim3 must be 0, otherwise every element must be greater than 0.
Definition: AMDKernelCodeT.h:287
amd_kernel_code_t::reserved0
uint64_t reserved0
Reserved. Must be 0.
Definition: AMDKernelCodeT.h:554
AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK
Definition: AMDKernelCodeT.h:184
amd_kernel_code_t::reserved_sgpr_count
uint16_t reserved_sgpr_count
The number of consecutive SGPRs reserved by the client.
Definition: AMDKernelCodeT.h:617
AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT
Definition: AMDKernelCodeT.h:97
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y
Definition: AMDKernelCodeT.h:119
hsa_ext_exception_kind16_t
uint16_t hsa_ext_exception_kind16_t
Definition: AMDKernelCodeT.h:25
AMD_CODE_PROPERTY_IS_PTR64_SHIFT
@ AMD_CODE_PROPERTY_IS_PTR64_SHIFT
Are global memory addresses 64 bits.
Definition: AMDKernelCodeT.h:170
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT
Definition: AMDKernelCodeT.h:113
amd_kernel_code_t::reserved3
uint8_t reserved3[12]
Definition: AMDKernelCodeT.h:646
amd_kernel_code_t::compute_pgm_resource_registers
uint64_t compute_pgm_resource_registers
Shader program settings for CS.
Definition: AMDKernelCodeT.h:558
amd_kernel_code_t::reserved_sgpr_first
uint16_t reserved_sgpr_first
If reserved_sgpr_count is 0 then must be 0.
Definition: AMDKernelCodeT.h:612
AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT
Definition: AMDKernelCodeT.h:93
AMD_CODE_PROPERTY_RESERVED1
@ AMD_CODE_PROPERTY_RESERVED1
Definition: AMDKernelCodeT.h:131
AMD_CODE_VERSION_MINOR
@ AMD_CODE_VERSION_MINOR
Definition: AMDKernelCodeT.h:39
amd_element_byte_size_t
amd_element_byte_size_t
The values used to define the number of bytes to use for the swizzle element size.
Definition: AMDKernelCodeT.h:53
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT
Definition: AMDKernelCodeT.h:101
AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH
Definition: AMDKernelCodeT.h:106
amd_kernel_code_t::workgroup_group_segment_byte_size
uint32_t workgroup_group_segment_byte_size
The amount of group segment memory required by a work-group in bytes.
Definition: AMDKernelCodeT.h:574
AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH
Definition: AMDKernelCodeT.h:137
AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH
Definition: AMDKernelCodeT.h:188
amd_kernel_code_t::control_directives
uint64_t control_directives[16]
Definition: AMDKernelCodeT.h:648
AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT
Definition: AMDKernelCodeT.h:107
AMD_CODE_PROPERTY_RESERVED2
@ AMD_CODE_PROPERTY_RESERVED2
Definition: AMDKernelCodeT.h:197
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH
Definition: AMDKernelCodeT.h:114
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z
Definition: AMDKernelCodeT.h:123
AMD_CODE_VERSION_MAJOR
@ AMD_CODE_VERSION_MAJOR
Definition: AMDKernelCodeT.h:38
AMD_ELEMENT_8_BYTES
@ AMD_ELEMENT_8_BYTES
Definition: AMDKernelCodeT.h:56
amd_kernel_code_t::kernel_code_prefetch_byte_offset
int64_t kernel_code_prefetch_byte_offset
Range of bytes to consider prefetching expressed as an offset and size.
Definition: AMDKernelCodeT.h:550
hsa_powertwo8_t
uint8_t hsa_powertwo8_t
Definition: AMDKernelCodeT.h:20
amd_kernel_code_t::amd_kernel_code_version_major
uint32_t amd_kernel_code_version_major
Definition: AMDKernelCodeT.h:527
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER
Definition: AMDKernelCodeT.h:87
hsa_ext_control_directives_s::enabled_control_directives
hsa_ext_control_directive_present64_t enabled_control_directives
This is a bit set indicating which control directives have been specified.
Definition: AMDKernelCodeT.h:215
hsa_ext_control_directives_s::max_dynamic_group_size
uint32_t max_dynamic_group_size
If maxDynamicGroupSize is not enabled then must be 0, and any amount of dynamic group segment can be ...
Definition: AMDKernelCodeT.h:255
hsa_dim3_s::y
uint32_t y
Definition: AMDKernelCodeT.h:30
AMD_CODE_PROPERTY_RESERVED2_SHIFT
@ AMD_CODE_PROPERTY_RESERVED2_SHIFT
Definition: AMDKernelCodeT.h:195
amd_kernel_code_t::workitem_private_segment_byte_size
uint32_t workitem_private_segment_byte_size
The amount of memory required for the combined private, spill and arg segments for a work-item in byt...
Definition: AMDKernelCodeT.h:568
amd_kernel_code_t::wavefront_size
uint8_t wavefront_size
Wavefront size expressed as a power of two.
Definition: AMDKernelCodeT.h:643
AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT
Indicate if the generated ISA is using a dynamically sized call stack.
Definition: AMDKernelCodeT.h:182
amd_code_property_mask_t
amd_code_property_mask_t
Definition: AMDKernelCodeT.h:73
amd_kernel_code_t::amd_machine_version_minor
uint16_t amd_machine_version_minor
Definition: AMDKernelCodeT.h:531
AMD_CODE_PROPERTY_RESERVED1_WIDTH
@ AMD_CODE_PROPERTY_RESERVED1_WIDTH
Definition: AMDKernelCodeT.h:130
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH
Definition: AMDKernelCodeT.h:86
amd_kernel_code_t::debug_private_segment_buffer_sgpr
uint16_t debug_private_segment_buffer_sgpr
If is_debug_supported is 0 then must be 0.
Definition: AMDKernelCodeT.h:629
AMD_CODE_PROPERTY_IS_PTR64
@ AMD_CODE_PROPERTY_IS_PTR64
Definition: AMDKernelCodeT.h:172
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH
Definition: AMDKernelCodeT.h:102
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH
Definition: AMDKernelCodeT.h:122
AMD_ELEMENT_16_BYTES
@ AMD_ELEMENT_16_BYTES
Definition: AMDKernelCodeT.h:57
AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT
Definition: AMDKernelCodeT.h:125
hsa_dim3_s::x
uint32_t x
Definition: AMDKernelCodeT.h:29
hsa_ext_control_directives_s::enable_detect_exceptions
hsa_ext_exception_kind16_t enable_detect_exceptions
If enableDetectExceptions is not enabled then must be 0, otherwise must be non-0 and specifies the se...
Definition: AMDKernelCodeT.h:238
AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT
Control wave ID base counter for GDS ordered-append.
Definition: AMDKernelCodeT.h:136
hsa_ext_brig_profile8_t
uint8_t hsa_ext_brig_profile8_t
Definition: AMDKernelCodeT.h:22
amd_kernel_code_t::call_convention
int32_t call_convention
Definition: AMDKernelCodeT.h:645
AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED
Definition: AMDKernelCodeT.h:189
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X
Definition: AMDKernelCodeT.h:115
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT
Definition: AMDKernelCodeT.h:109
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE
Definition: AMDKernelCodeT.h:111
AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH
Definition: AMDKernelCodeT.h:98
amd_kernel_code_t::debug_wavefront_private_segment_offset_sgpr
uint16_t debug_wavefront_private_segment_offset_sgpr
If is_debug_supported is 0 then must be 0.
Definition: AMDKernelCodeT.h:623
amd_kernel_code_t::amd_machine_version_stepping
uint16_t amd_machine_version_stepping
Definition: AMDKernelCodeT.h:532
amd_kernel_code_t::group_segment_alignment
uint8_t group_segment_alignment
Definition: AMDKernelCodeT.h:635
amd_kernel_code_t::kernarg_segment_alignment
uint8_t kernarg_segment_alignment
The maximum byte alignment of variables used by the kernel in the specified memory segment.
Definition: AMDKernelCodeT.h:634
amd_kernel_code_t::amd_kernel_code_version_minor
uint32_t amd_kernel_code_version_minor
Definition: AMDKernelCodeT.h:528
amd_kernel_code_t::kernarg_segment_byte_size
uint64_t kernarg_segment_byte_size
The size in bytes of the kernarg segment that holds the values of the arguments to the kernel.
Definition: AMDKernelCodeT.h:583
AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH
Definition: AMDKernelCodeT.h:94
hsa_ext_control_directives_s
The hsa_ext_control_directives_t specifies the values for the HSAIL control directives.
Definition: AMDKernelCodeT.h:209
hsa_dim3_s
Definition: AMDKernelCodeT.h:28
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH
Definition: AMDKernelCodeT.h:90
hsa_ext_control_directive_present64_t
uint64_t hsa_ext_control_directive_present64_t
Definition: AMDKernelCodeT.h:24
AMD_CODE_PROPERTY_IS_PTR64_WIDTH
@ AMD_CODE_PROPERTY_IS_PTR64_WIDTH
Definition: AMDKernelCodeT.h:171
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT
Definition: AMDKernelCodeT.h:89
uint32_t
AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS
Definition: AMDKernelCodeT.h:138
amd_kernel_code_t
AMD Kernel Code Object (amd_kernel_code_t).
Definition: AMDKernelCodeT.h:526
amd_kernel_code_t::gds_segment_byte_size
uint32_t gds_segment_byte_size
Number of byte of GDS required by kernel dispatch.
Definition: AMDKernelCodeT.h:578
hsa_ext_control_directives_s::enable_break_exceptions
hsa_ext_exception_kind16_t enable_break_exceptions
If enableBreakExceptions is not enabled then must be 0, otherwise must be non-0 and specifies the set...
Definition: AMDKernelCodeT.h:226
amd_kernel_code_t::kernel_code_entry_byte_offset
int64_t kernel_code_entry_byte_offset
Byte offset (possibly negative) from start of amd_kernel_code_t object to kernel's entry point instru...
Definition: AMDKernelCodeT.h:544
hsa_ext_control_directives_s::max_flat_workgroup_size
uint32_t max_flat_workgroup_size
If maxFlatWorkgroupSize is not enabled then must be 0, otherwise must be greater than 0.
Definition: AMDKernelCodeT.h:265
amd_code_version_t
amd_code_version_t
Definition: AMDKernelCodeT.h:37
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH
Definition: AMDKernelCodeT.h:110
hsa_ext_code_kind_t
uint32_t hsa_ext_code_kind_t
Definition: AMDKernelCodeT.h:21
amd_kernel_code_t::amd_machine_kind
uint16_t amd_machine_kind
Definition: AMDKernelCodeT.h:529
hsa_ext_control_directives_s::required_dim
uint8_t required_dim
If requiredDim is not enabled then must be 0 and the produced kernel code can be dispatched with 1,...
Definition: AMDKernelCodeT.h:310
amd_code_property32_t
uint32_t amd_code_property32_t
Every amd_*_code_t has the following properties, which are composed of a number of bit fields.
Definition: AMDKernelCodeT.h:72
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH
Definition: AMDKernelCodeT.h:118
amd_kernel_code_t::amd_machine_version_major
uint16_t amd_machine_version_major
Definition: AMDKernelCodeT.h:530
amd_kernel_code_t::reserved_vgpr_count
uint16_t reserved_vgpr_count
The number of consecutive VGPRs reserved by the client.
Definition: AMDKernelCodeT.h:608
hsa_ext_control_directives_s::reserved
uint8_t reserved[75]
Reserved. Must be 0.
Definition: AMDKernelCodeT.h:313
amd_kernel_code_t::private_segment_alignment
uint8_t private_segment_alignment
Definition: AMDKernelCodeT.h:636
uint16_t
AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_WIDTH
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_WIDTH
Definition: AMDKernelCodeT.h:126
amd_kernel_code_t::code_properties
uint32_t code_properties
Code properties.
Definition: AMDKernelCodeT.h:562
hsa_ext_brig_machine_model8_t
uint8_t hsa_ext_brig_machine_model8_t
Definition: AMDKernelCodeT.h:23
AMD_ELEMENT_2_BYTES
@ AMD_ELEMENT_2_BYTES
Definition: AMDKernelCodeT.h:54
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT
Definition: AMDKernelCodeT.h:117
AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_WIDTH
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_WIDTH
Definition: AMDKernelCodeT.h:192
amd_kernel_code_t::kernel_code_prefetch_byte_size
uint64_t kernel_code_prefetch_byte_size
Definition: AMDKernelCodeT.h:551
AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32
Definition: AMDKernelCodeT.h:127
AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE
Definition: AMDKernelCodeT.h:163
AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT
The interleave (swizzle) element size in bytes required by the code for private memory.
Definition: AMDKernelCodeT.h:161
AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH
Definition: AMDKernelCodeT.h:162
amd_kernel_code_t::workitem_vgpr_count
uint16_t workitem_vgpr_count
Number of vector registers used by each work-item.
Definition: AMDKernelCodeT.h:599
AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT
Enable the setup of the SGPR user data registers (AMD_CODE_PROPERTY_ENABLE_SGPR_*),...
Definition: AMDKernelCodeT.h:85
amd_kernel_code_t::runtime_loader_kernel_symbol
uint64_t runtime_loader_kernel_symbol
Definition: AMDKernelCodeT.h:647
hsa_dim3_s::z
uint32_t z
Definition: AMDKernelCodeT.h:31
hsa_ext_control_directives_s::required_workgroup_size
hsa_dim3_t required_workgroup_size
If requiredWorkgroupSize is not enabled then all elements for Dim3 must be 0, and the produced code c...
Definition: AMDKernelCodeT.h:300
AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT
Definition: AMDKernelCodeT.h:121
hsa_dim3_t
struct hsa_dim3_s hsa_dim3_t
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID
Definition: AMDKernelCodeT.h:103
hsa_ext_control_directives_s::max_flat_grid_size
uint32_t max_flat_grid_size
If maxFlatGridSize is not enabled then must be 0, otherwise must be greater than 0.
Definition: AMDKernelCodeT.h:260
hsa_ext_control_directives_s::requested_workgroups_per_cu
uint32_t requested_workgroups_per_cu
If requestedWorkgroupsPerCu is not enabled then must be 0, and the finalizer is free to generate ISA ...
Definition: AMDKernelCodeT.h:282
amd_kernel_code_t::reserved_vgpr_first
uint16_t reserved_vgpr_first
If reserved_vgpr_count is 0 then must be 0.
Definition: AMDKernelCodeT.h:603
AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR
Definition: AMDKernelCodeT.h:99
AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT
Indicate if code generated has support for debugging.
Definition: AMDKernelCodeT.h:187
AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_SHIFT
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_SHIFT
Definition: AMDKernelCodeT.h:191
hsa_ext_code_kind32_t
uint32_t hsa_ext_code_kind32_t
Definition: AMDKernelCodeT.h:26
AMD_CODE_PROPERTY_RESERVED1_SHIFT
@ AMD_CODE_PROPERTY_RESERVED1_SHIFT
Definition: AMDKernelCodeT.h:129
AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED
Definition: AMDKernelCodeT.h:193
amd_kernel_code_t::wavefront_sgpr_count
uint16_t wavefront_sgpr_count
Number of scalar registers used by a wavefront.
Definition: AMDKernelCodeT.h:595
AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR
Definition: AMDKernelCodeT.h:91
hsa_ext_control_directives_t
struct hsa_ext_control_directives_s hsa_ext_control_directives_t
The hsa_ext_control_directives_t specifies the values for the HSAIL control directives.
AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH
Definition: AMDKernelCodeT.h:183
amd_kernel_code_t::workgroup_fbarrier_count
uint32_t workgroup_fbarrier_count
Number of fbarrier's used in the kernel and all functions it calls.
Definition: AMDKernelCodeT.h:589