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