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