652 lines · c
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-exception6//7//===----------------------------------------------------------------------===//8/// \file AMDKernelCodeT.h9//===----------------------------------------------------------------------===//10 11#ifndef AMDKERNELCODET_H12#define AMDKERNELCODET_H13 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 be35/// 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 = 140};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 the52/// 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 = 358};59 60/// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and61/// 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 of65/// a number of bit fields. Every bit field has a mask (AMD_CODE_PROPERTY_*),66/// bit width (AMD_CODE_PROPERTY_*_WIDTH, and bit shift amount67/// (AMD_CODE_PROPERTY_*_SHIFT) for convenient access. Unused bits must be 0.68///69/// (Note that bit fields cannot be used as their layout is70/// implementation defined in the C standard and so cannot be used to71/// 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 registers76 /// (AMD_CODE_PROPERTY_ENABLE_SGPR_*), see documentation of amd_kernel_code_t77 /// for initial register state.78 ///79 /// The total number of SGPRuser data registers requested must not80 /// exceed 16. Any requests beyond 16 will be ignored.81 ///82 /// Used to set COMPUTE_PGM_RSRC2.USER_SGPR (set to total count of83 /// 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 set134 /// COMPUTE_DISPATCH_INITIATOR.ORDERED_APPEND_ENBL. (Not sure if135 /// 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 the141 /// code for private memory. This must be 2, 4, 8 or 16. This value142 /// is provided to the finalizer when it is invoked and is recorded143 /// here. The hardware will interleave the memory requests of each144 /// lane of a wavefront by this element size to ensure each145 /// work-item gets a distinct memory location. Therefore, the146 /// finalizer ensures that all load and store operations done to147 /// private memory do not exceed this size. For example, if the148 /// element size is 4 (32-bits or dword) and a 64-bit value must be149 /// loaded, the finalizer will generate two 32-bit loads. This150 /// ensures that the interleaving will get the work-item151 /// specific dword for both halves of the 64-bit value. If it just152 /// did a 64-bit load then it would get one dword which belonged to153 /// its own work-item, but the second dword would belong to the154 /// adjacent lane work-item since the interleaving is in dwords.155 ///156 /// The value used must match the value that the runtime configures157 /// the GPU flat scratch (SH_STATIC_MEM_CONFIG.ELEMENT_SIZE). This158 /// 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 match166 /// amd_kernel_code_t.hsail_machine_model ==167 /// HSA_MACHINE_LARGE. Must also match168 /// 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 call175 /// stack. This can happen if calls are implemented using a call176 /// stack and recursion, alloca or calls to indirect functions are177 /// present. In these cases the Finalizer cannot compute the total178 /// private segment size at compile time. In this case the179 /// workitem_private_segment_byte_size only specifies the statically180 /// know private segment size, and additional space must be added181 /// 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_SHIFT198};199 200/// The hsa_ext_control_directives_t specifies the values for the HSAIL201/// control directives. These control how the finalizer generates code. This202/// struct is used both as an argument to hsaFinalizeKernel to specify values for203/// the control directives, and is used in HsaKernelCode to record the values of204/// the control directives that the finalize used when generating the code which205/// either came from the finalizer argument or explicit HSAIL control206/// directives. See the definition of the control directives in HSA Programmer's207/// Reference Manual which also defines how the values specified as finalizer208/// 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 been211 /// specified. If the value is 0 then there are no control directives specified212 /// and the rest of the fields can be ignored. The bits are accessed using the213 /// hsa_ext_control_directives_present_mask_t. Any control directive that is not214 /// 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 be218 /// non-0 and specifies the set of HSAIL exceptions that must have the BREAK219 /// policy enabled. If this set is not empty then the generated code may have220 /// lower performance than if the set is empty. If the kernel being finalized221 /// has any enablebreakexceptions control directives, then the values specified222 /// by this argument are unioned with the values in these control223 /// directives. If any of the functions the kernel calls have an224 /// enablebreakexceptions control directive, then they must be equal or a225 /// 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 be229 /// non-0 and specifies the set of HSAIL exceptions that must have the DETECT230 /// policy enabled. If this set is not empty then the generated code may have231 /// lower performance than if the set is empty. However, an implementation232 /// should endeavour to make the performance impact small. If the kernel being233 /// finalized has any enabledetectexceptions control directives, then the234 /// values specified by this argument are unioned with the values in these235 /// control directives. If any of the functions the kernel calls have an236 /// enabledetectexceptions control directive, then they must be equal or a237 /// 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 of241 /// dynamic group segment can be allocated for a dispatch, otherwise the value242 /// specifies the maximum number of bytes of dynamic group segment that can be243 /// allocated for a dispatch. If the kernel being finalized has any244 /// maxdynamicsize control directives, then the values must be the same, and245 /// must be the same as this argument if it is enabled. This value can be used246 /// by the finalizer to determine the maximum number of bytes of group memory247 /// used by each work-group by adding this value to the group memory required248 /// for all group segment variables used by the kernel and all functions it249 /// calls, and group memory used to implement other HSAIL features such as250 /// fbarriers and the detect exception operations. This can allow the finalizer251 /// to determine the expected number of work-groups that can be executed by a252 /// compute unit and allow more resources to be allocated to the work-items if253 /// it is known that fewer work-groups can be executed due to group memory254 /// limitations.255 uint32_t max_dynamic_group_size;256 257 /// If maxFlatGridSize is not enabled then must be 0, otherwise must be greater258 /// than 0. See HSA Programmer's Reference Manual description of259 /// maxflatgridsize control directive.260 uint32_t max_flat_grid_size;261 262 /// If maxFlatWorkgroupSize is not enabled then must be 0, otherwise must be263 /// greater than 0. See HSA Programmer's Reference Manual description of264 /// maxflatworkgroupsize control directive.265 uint32_t max_flat_workgroup_size;266 267 /// If requestedWorkgroupsPerCu is not enabled then must be 0, and the268 /// finalizer is free to generate ISA that may result in any number of269 /// work-groups executing on a single compute unit. Otherwise, the finalizer270 /// should attempt to generate ISA that will allow the specified number of271 /// work-groups to execute on a single compute unit. This is only a hint and272 /// can be ignored by the finalizer. If the kernel being finalized, or any of273 /// the functions it calls, has a requested control directive, then the values274 /// must be the same. This can be used to determine the number of resources275 /// 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 higher277 /// per work-item performance, as it is known there will never be more than the278 /// specified number of work-groups actually executing on the compute279 /// unit. Conversely, a high value may allocate fewer resources, resulting in280 /// lower per work-item performance, which is offset by the fact it allows more281 /// 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 every285 /// element must be greater than 0. See HSA Programmer's Reference Manual286 /// 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 be290 /// 0, and the produced code can be dispatched with any legal work-group range291 /// consistent with the dispatch dimensions. Otherwise, the code produced must292 /// always be dispatched with the specified work-group range. No element of the293 /// specified range must be 0. It must be consistent with required_dimensions294 /// and max_flat_workgroup_size. If the kernel being finalized, or any of the295 /// functions it calls, has a requiredworkgroupsize control directive, then the296 /// values must be the same. Specifying a value can allow the finalizer to297 /// optimize work-group id operations, and if the number of work-items in the298 /// work-group is less than the WAVESIZE then barrier operations can be299 /// 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 code303 /// can be dispatched with 1, 2 or 3 dimensions. If enabled then the value is304 /// 1..3 and the code produced must only be dispatched with a dimension that305 /// matches. Other values are illegal. If the kernel being finalized, or any of306 /// the functions it calls, has a requireddimsize control directive, then the307 /// values must be the same. This can be used to optimize the code generated to308 /// compute the absolute and flat work-group and work-item id, and the dim309 /// 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 Kernel317/// 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 start322/// of execution of every wavefront. This is limited by the constraints of the323/// current hardware.324///325/// The order of the SGPR registers is defined, but the Finalizer can specify326/// which ones are actually setup in the amd_kernel_code_t object using the327/// enable_sgpr_* bit fields. The register numbers used for enabled registers328/// are dense starting at SGPR0: the first enabled register is SGPR0, the next329/// enabled register is SGPR1 etc.; disabled registers do not have an SGPR330/// number.331///332/// The initial SGPRs comprise up to 16 User SRGPs that are set up by CP and333/// apply to all waves of the grid. It is possible to specify more than 16 User334/// SGPRs using the enable_sgpr_* bit fields, in which case only the first 16335/// are actually initialized. These are then immediately followed by the System336/// SGPRs that are set up by ADC/SPI and can have different values for each wave337/// 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 with343/// Scratch Wave Offset as an offset, to access the Private/Spill/Arg344/// segments using a segment address. It must be set as follows:345/// - Base address: of the scratch memory area used by the dispatch. It346/// does not include the scratch wave offset. It will be the per process347/// SH_HIDDEN_PRIVATE_BASE_VMID plus any offset from this dispatch (for348/// 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 for352/// 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, must357/// agree with amd_kernel_code_t.privateElementSize)358/// - Index_stride: SH_STATIC_MEM_CONFIG.INDEX_STRIDE (will be 64 as must359/// be number of wavefront lanes for scratch, must agree with360/// amd_kernel_code_t.wavefrontSize)361/// - Add tid enable: 1362/// - ATC: from SH_MEM_CONFIG.PRIVATE_ATC,363/// - Hash_enable: ???364/// - Heap: ???365/// - Mtype: from SH_STATIC_MEM_CONFIG.PRIVATE_MTYPE366/// - 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 packet370/// 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 for374/// 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. This378/// is directly copied from the kernargPtr in the dispatch packet. Having CP379/// 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 dispatch383/// 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_BASE390/// to base of memory for scratch for this dispatch. This is the same offset391/// used in computing the Scratch Segment Buffer base address. The value of392/// Scratch Wave Offset must be added by the kernel code and moved to393/// 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 scratch396/// memory usage. This is directly loaded from the dispatch packet Private397/// 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 in402/// flat memory instructions. Having CP load it once avoids loading it at403/// the beginning of every wavefront.404///405/// For PI:406/// This is the 64 bit base address of the scratch backing memory for407/// 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 single411/// work-item's scratch memory allocation. This is the value from the dispatch412/// 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 every417/// wavefront.418///419/// \todo [This will not be used for CI/VI since it is the same value as420/// the second SGPR of Flat Scratch Init. However, it is need for PI which421/// 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 of425/// work-groups in the X dimension for the grid being executed. Computed from426/// the fields in the HsaDispatchPacket as427/// ((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 of431/// work-groups in the Y dimension for the grid being executed. Computed from432/// the fields in the HsaDispatchPacket as433/// ((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 of439/// work-groups in the Z dimension for the grid being executed. Computed440/// from the fields in the HsaDispatchPacket as441/// ((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 dimension447/// 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 dimension451/// 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 dimension455/// of grid for wavefront. If present then Work-group Id Y will also be456/// present457///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 Offset463/// (enable_sgpr_private_segment_wave_byte_offset):464/// Number of System SGPR registers: 1. 32 bit byte offset from base of465/// dispatch scratch base. Must be used as an offset with Private/Spill/Arg466/// segment address when using Scratch Segment Buffer. It must be added to467/// 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 specify471/// which ones are actually setup in the amd_kernel_code_t object using the472/// enableVgpr* bit fields. The register numbers used for enabled registers473/// are dense starting at VGPR0: the first enabled register is VGPR0, the next474/// enabled register is VGPR1 etc.; disabled registers do not have an VGPR475/// 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-group481/// 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-group485/// 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-group489/// 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 Data494/// registers.495/// 2) Work-group Id registers X, Y, Z are set by SPI which supports any496/// combination including none.497/// 3) Scratch Wave Offset is also set by SPI which is why its value cannot498/// be added into the value Flat Scratch Offset which would avoid the499/// 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 so504/// they can be moved as a 64 bit value to the hardware required SGPRn-3 and505/// SGPRn-4 respectively using the Finalizer ?FLAT_SCRATCH? Register.506///507/// The global segment can be accessed either using flat operations or buffer508/// operations. If buffer operations are used then the Global Buffer used to509/// access HSAIL Global/Readonly/Kernarg (which are combine) segments using a510/// segment address is not passed into the kernel code by CP since its base511/// address is always 0. Instead the Finalizer generates prolog code to512/// initialize 4 SGPRs with a V# that has the following properties, and then513/// uses that in the buffer instructions:514/// - base address of 0515/// - no swizzle516/// - ATC=1517/// - MTYPE set to support memory coherence specified in518/// amd_kernel_code_t.globalMemoryCoherence519///520/// When the Global Buffer is used to access the Kernarg segment, must add the521/// 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, as523/// 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_t535 /// object to kernel's entry point instruction. The actual code for536 /// the kernel is required to be 256 byte aligned to match hardware537 /// requirements (SQ cache line is 16). The code must be position538 /// independent code (PIC) for AMD devices to give runtime the539 /// option of copying code to discrete GPU memory or APU L2540 /// cache. The Finalizer should endeavour to allocate all kernel541 /// machine code in contiguous memory pages so that a device542 /// 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 offset547 /// and size. The offset is from the start (possibly negative) of548 /// amd_kernel_code_t object. Set both to 0 if no prefetch549 /// 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 and557 /// 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 of561 /// properties.562 uint32_t code_properties;563 564 /// The amount of memory required for the combined private, spill565 /// and arg segments for a work-item in bytes. If566 /// is_dynamic_callstack is 1 then additional space must be added to567 /// 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 in571 /// bytes. This does not include any dynamically allocated group572 /// segment memory that may be added when the kernel is573 /// dispatched.574 uint32_t workgroup_group_segment_byte_size;575 576 /// Number of byte of GDS required by kernel dispatch. Must be 0 if577 /// not using GDS.578 uint32_t gds_segment_byte_size;579 580 /// The size in bytes of the kernarg segment that holds the values581 /// of the arguments to the kernel. This could be used by CP to582 /// 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 it586 /// calls. If the implementation uses group memory to allocate the587 /// fbarriers then that amount must already be included in the588 /// workgroup_group_segment_byte_size total.589 uint32_t workgroup_fbarrier_count;590 591 /// Number of scalar registers used by a wavefront. This includes592 /// the special SGPRs for VCC, Flat Scratch Base, Flat Scratch Size593 /// and XNACK (for GFX8 (VI)). It does not include the 16 SGPR added if a594 /// 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 set598 /// 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 the602 /// first fixed VGPR number reserved.603 uint16_t reserved_vgpr_first;604 605 /// The number of consecutive VGPRs reserved by the client. If606 /// is_debug_supported then this count includes VGPRs reserved607 /// for debugger use.608 uint16_t reserved_vgpr_count;609 610 /// If reserved_sgpr_count is 0 then must be 0. Otherwise, this is the611 /// first fixed SGPR number reserved.612 uint16_t reserved_sgpr_first;613 614 /// The number of consecutive SGPRs reserved by the client. If615 /// is_debug_supported then this count includes SGPRs reserved616 /// for debugger use.617 uint16_t reserved_sgpr_count;618 619 /// If is_debug_supported is 0 then must be 0. Otherwise, this is the620 /// fixed SGPR number used to hold the wave scratch offset for the621 /// entire kernel execution, or uint16_t(-1) if the register is not622 /// 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 the626 /// fixed SGPR number of the first of 4 SGPRs used to hold the627 /// 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 in632 /// the specified memory segment. Expressed as a power of two. Must633 /// 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 2639 /// in range 1..64 inclusive. Used to support runtime query that640 /// obtains wavefront size, which may be used by application to641 /// allocated dynamic group memory and set the dispatch work-group642 /// 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_H652