//===-- AMDGPUKernelCodeT.h - Print AMDGPU assembly code ---------*- C++ -*-===// // // The LLVM Compiler Infrastructure // // This file is distributed under the University of Illinois Open Source // License. See LICENSE.TXT for details. // //===----------------------------------------------------------------------===// /// \file AMDKernelCodeT.h //===----------------------------------------------------------------------===// #ifndef AMDKERNELCODET_H #define AMDKERNELCODET_H #include <cstddef> #include <cstdint> //---------------------------------------------------------------------------// // AMD Kernel Code, and its dependencies // //---------------------------------------------------------------------------// typedef uint8_t hsa_powertwo8_t; typedef uint32_t hsa_ext_code_kind_t; typedef uint8_t hsa_ext_brig_profile8_t; typedef uint8_t hsa_ext_brig_machine_model8_t; typedef uint64_t hsa_ext_control_directive_present64_t; typedef uint16_t hsa_ext_exception_kind16_t; typedef uint32_t hsa_ext_code_kind32_t; typedef struct hsa_dim3_s { uint32_t x; uint32_t y; uint32_t z; } hsa_dim3_t; /// The version of the amd_*_code_t struct. Minor versions must be /// backward compatible. typedef uint32_t amd_code_version32_t; enum amd_code_version_t { AMD_CODE_VERSION_MAJOR = 0, AMD_CODE_VERSION_MINOR = 1 }; /// The values used to define the number of bytes to use for the /// swizzle element size. enum amd_element_byte_size_t { AMD_ELEMENT_2_BYTES = 0, AMD_ELEMENT_4_BYTES = 1, AMD_ELEMENT_8_BYTES = 2, AMD_ELEMENT_16_BYTES = 3 }; /// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and /// COMPUTE_PGM_RSRC2 registers. typedef uint64_t amd_compute_pgm_resource_register64_t; /// Every amd_*_code_t has the following properties, which are composed of /// a number of bit fields. Every bit field has a mask (AMD_CODE_PROPERTY_*), /// bit width (AMD_CODE_PROPERTY_*_WIDTH, and bit shift amount /// (AMD_CODE_PROPERTY_*_SHIFT) for convenient access. Unused bits must be 0. /// /// (Note that bit fields cannot be used as their layout is /// implementation defined in the C standard and so cannot be used to /// specify an ABI) typedef uint32_t amd_code_property32_t; enum amd_code_property_mask_t { /// Enable the setup of the SGPR user data registers /// (AMD_CODE_PROPERTY_ENABLE_SGPR_*), see documentation of amd_kernel_code_t /// for initial register state. /// /// The total number of SGPRuser data registers requested must not /// exceed 16. Any requests beyond 16 will be ignored. /// /// Used to set COMPUTE_PGM_RSRC2.USER_SGPR (set to total count of /// SGPR user data registers enabled up to 16). AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT = 0, AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT = 1, AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT = 2, AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT = 3, AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT = 4, AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT = 5, AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT = 6, AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT = 7, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT = 8, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH = 1, 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, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT = 9, AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH = 1, 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, /// Control wave ID base counter for GDS ordered-append. Used to set /// COMPUTE_DISPATCH_INITIATOR.ORDERED_APPEND_ENBL. (Not sure if /// ORDERED_APPEND_MODE also needs to be settable) AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT = 10, AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH = 1, 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, /// The interleave (swizzle) element size in bytes required by the /// code for private memory. This must be 2, 4, 8 or 16. This value /// is provided to the finalizer when it is invoked and is recorded /// here. The hardware will interleave the memory requests of each /// lane of a wavefront by this element size to ensure each /// work-item gets a distinct memory memory location. Therefore, the /// finalizer ensures that all load and store operations done to /// private memory do not exceed this size. For example, if the /// element size is 4 (32-bits or dword) and a 64-bit value must be /// loaded, the finalizer will generate two 32-bit loads. This /// ensures that the interleaving will get the the work-item /// specific dword for both halves of the 64-bit value. If it just /// did a 64-bit load then it would get one dword which belonged to /// its own work-item, but the second dword would belong to the /// adjacent lane work-item since the interleaving is in dwords. /// /// The value used must match the value that the runtime configures /// the GPU flat scratch (SH_STATIC_MEM_CONFIG.ELEMENT_SIZE). This /// is generally DWORD. /// /// Use values from the amd_element_byte_size_t enum. AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT = 11, AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH = 2, AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE = ((1 << AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH) - 1) << AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT, /// Are global memory addresses 64 bits. Must match /// amd_kernel_code_t.hsail_machine_model == /// HSA_MACHINE_LARGE. Must also match /// SH_MEM_CONFIG.PTR32 (GFX6 (SI)/GFX7 (CI)), /// SH_MEM_CONFIG.ADDRESS_MODE (GFX8 (VI)+). AMD_CODE_PROPERTY_IS_PTR64_SHIFT = 13, AMD_CODE_PROPERTY_IS_PTR64_WIDTH = 1, AMD_CODE_PROPERTY_IS_PTR64 = ((1 << AMD_CODE_PROPERTY_IS_PTR64_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_PTR64_SHIFT, /// Indicate if the generated ISA is using a dynamically sized call /// stack. This can happen if calls are implemented using a call /// stack and recursion, alloca or calls to indirect functions are /// present. In these cases the Finalizer cannot compute the total /// private segment size at compile time. In this case the /// workitem_private_segment_byte_size only specifies the statically /// know private segment size, and additional space must be added /// for the call stack. AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT = 14, AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH = 1, AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK = ((1 << AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT, /// Indicate if code generated has support for debugging. AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT = 15, AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH = 1, AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED = ((1 << AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH) - 1) << AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT }; /// @brief The hsa_ext_control_directives_t specifies the values for the HSAIL /// control directives. These control how the finalizer generates code. This /// struct is used both as an argument to hsaFinalizeKernel to specify values for /// the control directives, and is used in HsaKernelCode to record the values of /// the control directives that the finalize used when generating the code which /// either came from the finalizer argument or explicit HSAIL control /// directives. See the definition of the control directives in HSA Programmer's /// Reference Manual which also defines how the values specified as finalizer /// arguments have to agree with the control directives in the HSAIL code. typedef struct hsa_ext_control_directives_s { /// This is a bit set indicating which control directives have been /// specified. If the value is 0 then there are no control directives specified /// and the rest of the fields can be ignored. The bits are accessed using the /// hsa_ext_control_directives_present_mask_t. Any control directive that is not /// enabled in this bit set must have the value of all 0s. hsa_ext_control_directive_present64_t enabled_control_directives; /// If enableBreakExceptions is not enabled then must be 0, otherwise must be /// non-0 and specifies the set of HSAIL exceptions that must have the BREAK /// policy enabled. If this set is not empty then the generated code may have /// lower performance than if the set is empty. If the kernel being finalized /// has any enablebreakexceptions control directives, then the values specified /// by this argument are unioned with the values in these control /// directives. If any of the functions the kernel calls have an /// enablebreakexceptions control directive, then they must be equal or a /// subset of, this union. hsa_ext_exception_kind16_t enable_break_exceptions; /// If enableDetectExceptions is not enabled then must be 0, otherwise must be /// non-0 and specifies the set of HSAIL exceptions that must have the DETECT /// policy enabled. If this set is not empty then the generated code may have /// lower performance than if the set is empty. However, an implementation /// should endeavour to make the performance impact small. If the kernel being /// finalized has any enabledetectexceptions control directives, then the /// values specified by this argument are unioned with the values in these /// control directives. If any of the functions the kernel calls have an /// enabledetectexceptions control directive, then they must be equal or a /// subset of, this union. hsa_ext_exception_kind16_t enable_detect_exceptions; /// If maxDynamicGroupSize is not enabled then must be 0, and any amount of /// dynamic group segment can be allocated for a dispatch, otherwise the value /// specifies the maximum number of bytes of dynamic group segment that can be /// allocated for a dispatch. If the kernel being finalized has any /// maxdynamicsize control directives, then the values must be the same, and /// must be the same as this argument if it is enabled. This value can be used /// by the finalizer to determine the maximum number of bytes of group memory /// used by each work-group by adding this value to the group memory required /// for all group segment variables used by the kernel and all functions it /// calls, and group memory used to implement other HSAIL features such as /// fbarriers and the detect exception operations. This can allow the finalizer /// to determine the expected number of work-groups that can be executed by a /// compute unit and allow more resources to be allocated to the work-items if /// it is known that fewer work-groups can be executed due to group memory /// limitations. uint32_t max_dynamic_group_size; /// If maxFlatGridSize is not enabled then must be 0, otherwise must be greater /// than 0. See HSA Programmer's Reference Manual description of /// maxflatgridsize control directive. uint32_t max_flat_grid_size; /// If maxFlatWorkgroupSize is not enabled then must be 0, otherwise must be /// greater than 0. See HSA Programmer's Reference Manual description of /// maxflatworkgroupsize control directive. uint32_t max_flat_workgroup_size; /// If requestedWorkgroupsPerCu is not enabled then must be 0, and the /// finalizer is free to generate ISA that may result in any number of /// work-groups executing on a single compute unit. Otherwise, the finalizer /// should attempt to generate ISA that will allow the specified number of /// work-groups to execute on a single compute unit. This is only a hint and /// can be ignored by the finalizer. If the kernel being finalized, or any of /// the functions it calls, has a requested control directive, then the values /// must be the same. This can be used to determine the number of resources /// that should be allocated to a single work-group and work-item. For example, /// a low value may allow more resources to be allocated, resulting in higher /// per work-item performance, as it is known there will never be more than the /// specified number of work-groups actually executing on the compute /// unit. Conversely, a high value may allocate fewer resources, resulting in /// lower per work-item performance, which is offset by the fact it allows more /// work-groups to actually execute on the compute unit. uint32_t requested_workgroups_per_cu; /// If not enabled then all elements for Dim3 must be 0, otherwise every /// element must be greater than 0. See HSA Programmer's Reference Manual /// description of requiredgridsize control directive. hsa_dim3_t required_grid_size; /// If requiredWorkgroupSize is not enabled then all elements for Dim3 must be /// 0, and the produced code can be dispatched with any legal work-group range /// consistent with the dispatch dimensions. Otherwise, the code produced must /// always be dispatched with the specified work-group range. No element of the /// specified range must be 0. It must be consistent with required_dimensions /// and max_flat_workgroup_size. If the kernel being finalized, or any of the /// functions it calls, has a requiredworkgroupsize control directive, then the /// values must be the same. Specifying a value can allow the finalizer to /// optimize work-group id operations, and if the number of work-items in the /// work-group is less than the WAVESIZE then barrier operations can be /// optimized to just a memory fence. hsa_dim3_t required_workgroup_size; /// If requiredDim is not enabled then must be 0 and the produced kernel code /// can be dispatched with 1, 2 or 3 dimensions. If enabled then the value is /// 1..3 and the code produced must only be dispatched with a dimension that /// matches. Other values are illegal. If the kernel being finalized, or any of /// the functions it calls, has a requireddimsize control directive, then the /// values must be the same. This can be used to optimize the code generated to /// compute the absolute and flat work-group and work-item id, and the dim /// HSAIL operations. uint8_t required_dim; /// Reserved. Must be 0. uint8_t reserved[75]; } hsa_ext_control_directives_t; /// AMD Kernel Code Object (amd_kernel_code_t). GPU CP uses the AMD Kernel /// Code Object to set up the hardware to execute the kernel dispatch. /// /// Initial Kernel Register State. /// /// Initial kernel register state will be set up by CP/SPI prior to the start /// of execution of every wavefront. This is limited by the constraints of the /// current hardware. /// /// The order of the SGPR registers is defined, but the Finalizer can specify /// which ones are actually setup in the amd_kernel_code_t object using the /// enable_sgpr_* bit fields. The register numbers used for enabled registers /// are dense starting at SGPR0: the first enabled register is SGPR0, the next /// enabled register is SGPR1 etc.; disabled registers do not have an SGPR /// number. /// /// The initial SGPRs comprise up to 16 User SRGPs that are set up by CP and /// apply to all waves of the grid. It is possible to specify more than 16 User /// SGPRs using the enable_sgpr_* bit fields, in which case only the first 16 /// are actually initialized. These are then immediately followed by the System /// SGPRs that are set up by ADC/SPI and can have different values for each wave /// of the grid dispatch. /// /// SGPR register initial state is defined as follows: /// /// Private Segment Buffer (enable_sgpr_private_segment_buffer): /// Number of User SGPR registers: 4. V# that can be used, together with /// Scratch Wave Offset as an offset, to access the Private/Spill/Arg /// segments using a segment address. It must be set as follows: /// - Base address: of the scratch memory area used by the dispatch. It /// does not include the scratch wave offset. It will be the per process /// SH_HIDDEN_PRIVATE_BASE_VMID plus any offset from this dispatch (for /// example there may be a per pipe offset, or per AQL Queue offset). /// - Stride + data_format: Element Size * Index Stride (???) /// - Cache swizzle: ??? /// - Swizzle enable: SH_STATIC_MEM_CONFIG.SWIZZLE_ENABLE (must be 1 for /// scratch) /// - Num records: Flat Scratch Work Item Size / Element Size (???) /// - Dst_sel_*: ??? /// - Num_format: ??? /// - Element_size: SH_STATIC_MEM_CONFIG.ELEMENT_SIZE (will be DWORD, must /// agree with amd_kernel_code_t.privateElementSize) /// - Index_stride: SH_STATIC_MEM_CONFIG.INDEX_STRIDE (will be 64 as must /// be number of wavefront lanes for scratch, must agree with /// amd_kernel_code_t.wavefrontSize) /// - Add tid enable: 1 /// - ATC: from SH_MEM_CONFIG.PRIVATE_ATC, /// - Hash_enable: ??? /// - Heap: ??? /// - Mtype: from SH_STATIC_MEM_CONFIG.PRIVATE_MTYPE /// - Type: 0 (a buffer) (???) /// /// Dispatch Ptr (enable_sgpr_dispatch_ptr): /// Number of User SGPR registers: 2. 64 bit address of AQL dispatch packet /// for kernel actually executing. /// /// Queue Ptr (enable_sgpr_queue_ptr): /// Number of User SGPR registers: 2. 64 bit address of AmdQueue object for /// AQL queue on which the dispatch packet was queued. /// /// Kernarg Segment Ptr (enable_sgpr_kernarg_segment_ptr): /// Number of User SGPR registers: 2. 64 bit address of Kernarg segment. This /// is directly copied from the kernargPtr in the dispatch packet. Having CP /// load it once avoids loading it at the beginning of every wavefront. /// /// Dispatch Id (enable_sgpr_dispatch_id): /// Number of User SGPR registers: 2. 64 bit Dispatch ID of the dispatch /// packet being executed. /// /// Flat Scratch Init (enable_sgpr_flat_scratch_init): /// Number of User SGPR registers: 2. This is 2 SGPRs. /// /// For CI/VI: /// The first SGPR is a 32 bit byte offset from SH_MEM_HIDDEN_PRIVATE_BASE /// to base of memory for scratch for this dispatch. This is the same offset /// used in computing the Scratch Segment Buffer base address. The value of /// Scratch Wave Offset must be added by the kernel code and moved to /// SGPRn-4 for use as the FLAT SCRATCH BASE in flat memory instructions. /// /// The second SGPR is 32 bit byte size of a single work-item’s scratch /// memory usage. This is directly loaded from the dispatch packet Private /// Segment Byte Size and rounded up to a multiple of DWORD. /// /// \todo [Does CP need to round this to >4 byte alignment?] /// /// The kernel code must move to SGPRn-3 for use as the FLAT SCRATCH SIZE in /// flat memory instructions. Having CP load it once avoids loading it at /// the beginning of every wavefront. /// /// For PI: /// This is the 64 bit base address of the scratch backing memory for /// allocated by CP for this dispatch. /// /// Private Segment Size (enable_sgpr_private_segment_size): /// Number of User SGPR registers: 1. The 32 bit byte size of a single /// work-item’s scratch memory allocation. This is the value from the dispatch /// packet. Private Segment Byte Size rounded up by CP to a multiple of DWORD. /// /// \todo [Does CP need to round this to >4 byte alignment?] /// /// Having CP load it once avoids loading it at the beginning of every /// wavefront. /// /// \todo [This will not be used for CI/VI since it is the same value as /// the second SGPR of Flat Scratch Init. However, it is need for PI which /// changes meaning of Flat Scratchg Init..] /// /// Grid Work-Group Count X (enable_sgpr_grid_workgroup_count_x): /// Number of User SGPR registers: 1. 32 bit count of the number of /// work-groups in the X dimension for the grid being executed. Computed from /// the fields in the HsaDispatchPacket as /// ((gridSize.x+workgroupSize.x-1)/workgroupSize.x). /// /// Grid Work-Group Count Y (enable_sgpr_grid_workgroup_count_y): /// Number of User SGPR registers: 1. 32 bit count of the number of /// work-groups in the Y dimension for the grid being executed. Computed from /// the fields in the HsaDispatchPacket as /// ((gridSize.y+workgroupSize.y-1)/workgroupSize.y). /// /// Only initialized if <16 previous SGPRs initialized. /// /// Grid Work-Group Count Z (enable_sgpr_grid_workgroup_count_z): /// Number of User SGPR registers: 1. 32 bit count of the number of /// work-groups in the Z dimension for the grid being executed. Computed /// from the fields in the HsaDispatchPacket as /// ((gridSize.z+workgroupSize.z-1)/workgroupSize.z). /// /// Only initialized if <16 previous SGPRs initialized. /// /// Work-Group Id X (enable_sgpr_workgroup_id_x): /// Number of System SGPR registers: 1. 32 bit work group id in X dimension /// of grid for wavefront. Always present. /// /// Work-Group Id Y (enable_sgpr_workgroup_id_y): /// Number of System SGPR registers: 1. 32 bit work group id in Y dimension /// of grid for wavefront. /// /// Work-Group Id Z (enable_sgpr_workgroup_id_z): /// Number of System SGPR registers: 1. 32 bit work group id in Z dimension /// of grid for wavefront. If present then Work-group Id Y will also be /// present /// /// Work-Group Info (enable_sgpr_workgroup_info): /// Number of System SGPR registers: 1. {first_wave, 14’b0000, /// ordered_append_term[10:0], threadgroup_size_in_waves[5:0]} /// /// Private Segment Wave Byte Offset /// (enable_sgpr_private_segment_wave_byte_offset): /// Number of System SGPR registers: 1. 32 bit byte offset from base of /// dispatch scratch base. Must be used as an offset with Private/Spill/Arg /// segment address when using Scratch Segment Buffer. It must be added to /// Flat Scratch Offset if setting up FLAT SCRATCH for flat addressing. /// /// /// The order of the VGPR registers is defined, but the Finalizer can specify /// which ones are actually setup in the amd_kernel_code_t object using the /// enableVgpr* bit fields. The register numbers used for enabled registers /// are dense starting at VGPR0: the first enabled register is VGPR0, the next /// enabled register is VGPR1 etc.; disabled registers do not have an VGPR /// number. /// /// VGPR register initial state is defined as follows: /// /// Work-Item Id X (always initialized): /// Number of registers: 1. 32 bit work item id in X dimension of work-group /// for wavefront lane. /// /// Work-Item Id X (enable_vgpr_workitem_id > 0): /// Number of registers: 1. 32 bit work item id in Y dimension of work-group /// for wavefront lane. /// /// Work-Item Id X (enable_vgpr_workitem_id > 0): /// Number of registers: 1. 32 bit work item id in Z dimension of work-group /// for wavefront lane. /// /// /// The setting of registers is being done by existing GPU hardware as follows: /// 1) SGPRs before the Work-Group Ids are set by CP using the 16 User Data /// registers. /// 2) Work-group Id registers X, Y, Z are set by SPI which supports any /// combination including none. /// 3) Scratch Wave Offset is also set by SPI which is why its value cannot /// be added into the value Flat Scratch Offset which would avoid the /// Finalizer generated prolog having to do the add. /// 4) The VGPRs are set by SPI which only supports specifying either (X), /// (X, Y) or (X, Y, Z). /// /// Flat Scratch Dispatch Offset and Flat Scratch Size are adjacent SGRRs so /// they can be moved as a 64 bit value to the hardware required SGPRn-3 and /// SGPRn-4 respectively using the Finalizer ?FLAT_SCRATCH? Register. /// /// The global segment can be accessed either using flat operations or buffer /// operations. If buffer operations are used then the Global Buffer used to /// access HSAIL Global/Readonly/Kernarg (which are combine) segments using a /// segment address is not passed into the kernel code by CP since its base /// address is always 0. Instead the Finalizer generates prolog code to /// initialize 4 SGPRs with a V# that has the following properties, and then /// uses that in the buffer instructions: /// - base address of 0 /// - no swizzle /// - ATC=1 /// - MTYPE set to support memory coherence specified in /// amd_kernel_code_t.globalMemoryCoherence /// /// When the Global Buffer is used to access the Kernarg segment, must add the /// dispatch packet kernArgPtr to a kernarg segment address before using this V#. /// Alternatively scalar loads can be used if the kernarg offset is uniform, as /// the kernarg segment is constant for the duration of the kernel execution. /// typedef struct amd_kernel_code_s { /// The AMD major version of the Code Object. Must be the value /// AMD_CODE_VERSION_MAJOR. amd_code_version32_t amd_code_version_major; /// The AMD minor version of the Code Object. Minor versions must be /// backward compatible. Must be the value /// AMD_CODE_VERSION_MINOR. amd_code_version32_t amd_code_version_minor; /// The byte size of this struct. Must be set to /// sizeof(amd_kernel_code_t). Used for backward /// compatibility. uint32_t struct_byte_size; /// The target chip instruction set for which code has been /// generated. Values are from the E_SC_INSTRUCTION_SET enumeration /// in sc/Interface/SCCommon.h. uint32_t target_chip; /// Byte offset (possibly negative) from start of amd_kernel_code_t /// object to kernel's entry point instruction. The actual code for /// the kernel is required to be 256 byte aligned to match hardware /// requirements (SQ cache line is 16). The code must be position /// independent code (PIC) for AMD devices to give runtime the /// option of copying code to discrete GPU memory or APU L2 /// cache. The Finalizer should endeavour to allocate all kernel /// machine code in contiguous memory pages so that a device /// pre-fetcher will tend to only pre-fetch Kernel Code objects, /// improving cache performance. int64_t kernel_code_entry_byte_offset; /// Range of bytes to consider prefetching expressed as an offset /// and size. The offset is from the start (possibly negative) of /// amd_kernel_code_t object. Set both to 0 if no prefetch /// information is available. /// /// \todo ttye 11/15/2013 Is the prefetch definition we want? Did /// not make the size a uint64_t as prefetching more than 4GiB seems /// excessive. int64_t kernel_code_prefetch_byte_offset; uint64_t kernel_code_prefetch_byte_size; /// Number of bytes of scratch backing memory required for full /// occupancy of target chip. This takes into account the number of /// bytes of scratch per work-item, the wavefront size, the maximum /// number of wavefronts per CU, and the number of CUs. This is an /// upper limit on scratch. If the grid being dispatched is small it /// may only need less than this. If the kernel uses no scratch, or /// the Finalizer has not computed this value, it must be 0. uint64_t max_scratch_backing_memory_byte_size; /// Shader program settings for CS. Contains COMPUTE_PGM_RSRC1 and /// COMPUTE_PGM_RSRC2 registers. amd_compute_pgm_resource_register64_t compute_pgm_resource_registers; /// Code properties. See amd_code_property_mask_t for a full list of /// properties. amd_code_property32_t code_properties; /// The amount of memory required for the combined private, spill /// and arg segments for a work-item in bytes. If /// is_dynamic_callstack is 1 then additional space must be added to /// this value for the call stack. uint32_t workitem_private_segment_byte_size; /// The amount of group segment memory required by a work-group in /// bytes. This does not include any dynamically allocated group /// segment memory that may be added when the kernel is /// dispatched. uint32_t workgroup_group_segment_byte_size; /// Number of byte of GDS required by kernel dispatch. Must be 0 if /// not using GDS. uint32_t gds_segment_byte_size; /// The size in bytes of the kernarg segment that holds the values /// of the arguments to the kernel. This could be used by CP to /// prefetch the kernarg segment pointed to by the dispatch packet. uint64_t kernarg_segment_byte_size; /// Number of fbarrier's used in the kernel and all functions it /// calls. If the implementation uses group memory to allocate the /// fbarriers then that amount must already be included in the /// workgroup_group_segment_byte_size total. uint32_t workgroup_fbarrier_count; /// Number of scalar registers used by a wavefront. This includes /// the special SGPRs for VCC, Flat Scratch Base, Flat Scratch Size /// and XNACK (for GFX8 (VI)). It does not include the 16 SGPR added if a /// trap handler is enabled. Used to set COMPUTE_PGM_RSRC1.SGPRS. uint16_t wavefront_sgpr_count; /// Number of vector registers used by each work-item. Used to set /// COMPUTE_PGM_RSRC1.VGPRS. uint16_t workitem_vgpr_count; /// If reserved_vgpr_count is 0 then must be 0. Otherwise, this is the /// first fixed VGPR number reserved. uint16_t reserved_vgpr_first; /// The number of consecutive VGPRs reserved by the client. If /// is_debug_supported then this count includes VGPRs reserved /// for debugger use. uint16_t reserved_vgpr_count; /// If reserved_sgpr_count is 0 then must be 0. Otherwise, this is the /// first fixed SGPR number reserved. uint16_t reserved_sgpr_first; /// The number of consecutive SGPRs reserved by the client. If /// is_debug_supported then this count includes SGPRs reserved /// for debugger use. uint16_t reserved_sgpr_count; /// If is_debug_supported is 0 then must be 0. Otherwise, this is the /// fixed SGPR number used to hold the wave scratch offset for the /// entire kernel execution, or uint16_t(-1) if the register is not /// used or not known. uint16_t debug_wavefront_private_segment_offset_sgpr; /// If is_debug_supported is 0 then must be 0. Otherwise, this is the /// fixed SGPR number of the first of 4 SGPRs used to hold the /// scratch V# used for the entire kernel execution, or uint16_t(-1) /// if the registers are not used or not known. uint16_t debug_private_segment_buffer_sgpr; /// The maximum byte alignment of variables used by the kernel in /// the specified memory segment. Expressed as a power of two. Must /// be at least HSA_POWERTWO_16. hsa_powertwo8_t kernarg_segment_alignment; hsa_powertwo8_t group_segment_alignment; hsa_powertwo8_t private_segment_alignment; uint8_t reserved3; /// Type of code object. hsa_ext_code_kind32_t code_type; /// Reserved for code properties if any are defined in the future. /// There are currently no code properties so this field must be 0. uint32_t reserved4; /// Wavefront size expressed as a power of two. Must be a power of 2 /// in range 1..64 inclusive. Used to support runtime query that /// obtains wavefront size, which may be used by application to /// allocated dynamic group memory and set the dispatch work-group /// size. hsa_powertwo8_t wavefront_size; /// The optimization level specified when the kernel was /// finalized. uint8_t optimization_level; /// The HSAIL profile defines which features are used. This /// information is from the HSAIL version directive. If this /// amd_kernel_code_t is not generated from an HSAIL compilation /// unit then must be 0. hsa_ext_brig_profile8_t hsail_profile; /// The HSAIL machine model gives the address sizes used by the /// code. This information is from the HSAIL version directive. If /// not generated from an HSAIL compilation unit then must still /// indicate for what machine mode the code is generated. hsa_ext_brig_machine_model8_t hsail_machine_model; /// The HSAIL major version. This information is from the HSAIL /// version directive. If this amd_kernel_code_t is not /// generated from an HSAIL compilation unit then must be 0. uint32_t hsail_version_major; /// The HSAIL minor version. This information is from the HSAIL /// version directive. If this amd_kernel_code_t is not /// generated from an HSAIL compilation unit then must be 0. uint32_t hsail_version_minor; /// Reserved for HSAIL target options if any are defined in the /// future. There are currently no target options so this field /// must be 0. uint16_t reserved5; /// Reserved. Must be 0. uint16_t reserved6; /// The values should be the actually values used by the finalizer /// in generating the code. This may be the union of values /// specified as finalizer arguments and explicit HSAIL control /// directives. If the finalizer chooses to ignore a control /// directive, and not generate constrained code, then the control /// directive should not be marked as enabled even though it was /// present in the HSAIL or finalizer argument. The values are /// intended to reflect the constraints that the code actually /// requires to correctly execute, not the values that were /// actually specified at finalize time. hsa_ext_control_directives_t control_directive; /// The code can immediately follow the amd_kernel_code_t, or can /// come after subsequent amd_kernel_code_t structs when there are /// multiple kernels in the compilation unit. } amd_kernel_code_t; #endif // AMDKERNELCODET_H