LLVM: lib/Target/AMDGPU/AMDKernelCodeT.h Source File (original) (raw)
1
2
3
4
5
6
7
8
9
10
11#ifndef AMDKERNELCODET_H
12#define AMDKERNELCODET_H
13
14#include
15
16
17
18
19
27
33
34
35
41
42
43#define AMD_HSA_BITS_SET(dst, mask, val) \
44 dst &= (~(1 << mask ## _SHIFT) & ~mask); \
45 dst |= (((val) << mask ## _SHIFT) & mask)
46
47
48#define AMD_HSA_BITS_GET(src, mask) \
49 ((src & mask) >> mask ## _SHIFT) \
50
51
52
59
60
61
63
64
65
66
67
68
69
70
71
74
75
76
77
78
79
80
81
82
83
84
88
92
96
100
104
108
112
116
120
124
128
132
133
134
135
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
164
165
166
167
168
169
173
174
175
176
177
178
179
180
181
185
186
190
194
198};
199
200
201
202
203
204
205
206
207
208
210
211
212
213
214
216
217
218
219
220
221
222
223
224
225
227
228
229
230
231
232
233
234
235
236
237
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
256
257
258
259
261
262
263
264
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
283
284
285
286
288
289
290
291
292
293
294
295
296
297
298
299
301
302
303
304
305
306
307
308
309
311
312
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
533
534
535
536
537
538
539
540
541
542
543
545
546
547
548
549
552
553
555
556
557
559
560
561
563
564
565
566
567
569
570
571
572
573
575
576
577
579
580
581
582
584
585
586
587
588
590
591
592
593
594
596
597
598
600
601
602
604
605
606
607
609
610
611
613
614
615
616
618
619
620
621
622
624
625
626
627
628
630
631
632
633
637
638
639
640
641
642
644
649};
650
651#endif
uint8_t hsa_powertwo8_t
Definition AMDKernelCodeT.h:20
amd_element_byte_size_t
The values used to define the number of bytes to use for the swizzle element size.
Definition AMDKernelCodeT.h:53
@ AMD_ELEMENT_8_BYTES
Definition AMDKernelCodeT.h:56
@ AMD_ELEMENT_2_BYTES
Definition AMDKernelCodeT.h:54
@ AMD_ELEMENT_16_BYTES
Definition AMDKernelCodeT.h:57
@ AMD_ELEMENT_4_BYTES
Definition AMDKernelCodeT.h:55
uint64_t hsa_ext_control_directive_present64_t
Definition AMDKernelCodeT.h:24
uint8_t hsa_ext_brig_profile8_t
Definition AMDKernelCodeT.h:22
uint8_t hsa_ext_brig_machine_model8_t
Definition AMDKernelCodeT.h:23
uint64_t amd_compute_pgm_resource_register64_t
Shader program settings for CS.
Definition AMDKernelCodeT.h:62
uint32_t amd_code_property32_t
Every amd_*_code_t has the following properties, which are composed of a number of bit fields.
Definition AMDKernelCodeT.h:72
uint32_t hsa_ext_code_kind_t
Definition AMDKernelCodeT.h:21
amd_code_version_t
Definition AMDKernelCodeT.h:37
@ AMD_CODE_VERSION_MINOR
Definition AMDKernelCodeT.h:39
@ AMD_CODE_VERSION_MAJOR
Definition AMDKernelCodeT.h:38
struct hsa_dim3_s hsa_dim3_t
amd_code_property_mask_t
Definition AMDKernelCodeT.h:73
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_SHIFT
Definition AMDKernelCodeT.h:113
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y
Definition AMDKernelCodeT.h:119
@ AMD_CODE_PROPERTY_IS_PTR64_WIDTH
Definition AMDKernelCodeT.h:171
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT
Indicate if the generated ISA is using a dynamically sized call stack.
Definition AMDKernelCodeT.h:182
@ AMD_CODE_PROPERTY_RESERVED1
Definition AMDKernelCodeT.h:131
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_WIDTH
Definition AMDKernelCodeT.h:98
@ AMD_CODE_PROPERTY_RESERVED2_WIDTH
Definition AMDKernelCodeT.h:196
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_WIDTH
Definition AMDKernelCodeT.h:162
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID
Definition AMDKernelCodeT.h:103
@ AMD_CODE_PROPERTY_RESERVED1_WIDTH
Definition AMDKernelCodeT.h:130
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE_SHIFT
The interleave (swizzle) element size in bytes required by the code for private memory.
Definition AMDKernelCodeT.h:161
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_SHIFT
Definition AMDKernelCodeT.h:117
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32
Definition AMDKernelCodeT.h:127
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_WIDTH
Definition AMDKernelCodeT.h:90
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_SHIFT
Definition AMDKernelCodeT.h:109
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH
Definition AMDKernelCodeT.h:183
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_SHIFT
Definition AMDKernelCodeT.h:101
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_WIDTH
Definition AMDKernelCodeT.h:126
@ AMD_CODE_PROPERTY_PRIVATE_ELEMENT_SIZE
Definition AMDKernelCodeT.h:163
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_SHIFT
Definition AMDKernelCodeT.h:105
@ AMD_CODE_PROPERTY_RESERVED2
Definition AMDKernelCodeT.h:197
@ AMD_CODE_PROPERTY_RESERVED2_SHIFT
Definition AMDKernelCodeT.h:195
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR
Definition AMDKernelCodeT.h:99
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR
Definition AMDKernelCodeT.h:95
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED
Definition AMDKernelCodeT.h:189
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_WIDTH
Definition AMDKernelCodeT.h:137
@ AMD_CODE_PROPERTY_IS_PTR64_SHIFT
Are global memory addresses 64 bits.
Definition AMDKernelCodeT.h:170
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE
Definition AMDKernelCodeT.h:111
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS
Definition AMDKernelCodeT.h:138
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR_SHIFT
Definition AMDKernelCodeT.h:89
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER
Definition AMDKernelCodeT.h:87
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Y_WIDTH
Definition AMDKernelCodeT.h:118
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X
Definition AMDKernelCodeT.h:115
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_SHIFT
Definition AMDKernelCodeT.h:93
@ AMD_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32_SHIFT
Definition AMDKernelCodeT.h:125
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_SHIFT
Enable the setup of the SGPR user data registers (AMD_CODE_PROPERTY_ENABLE_SGPR_*),...
Definition AMDKernelCodeT.h:85
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER_WIDTH
Definition AMDKernelCodeT.h:86
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR
Definition AMDKernelCodeT.h:91
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED
Definition AMDKernelCodeT.h:193
@ AMD_CODE_PROPERTY_ENABLE_ORDERED_APPEND_GDS_SHIFT
Control wave ID base counter for GDS ordered-append.
Definition AMDKernelCodeT.h:136
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_WIDTH
Definition AMDKernelCodeT.h:122
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z
Definition AMDKernelCodeT.h:123
@ AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR_SHIFT
Definition AMDKernelCodeT.h:97
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_WIDTH
Definition AMDKernelCodeT.h:188
@ AMD_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR_WIDTH
Definition AMDKernelCodeT.h:94
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_Z_SHIFT
Definition AMDKernelCodeT.h:121
@ AMD_CODE_PROPERTY_RESERVED1_SHIFT
Definition AMDKernelCodeT.h:129
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT_WIDTH
Definition AMDKernelCodeT.h:106
@ AMD_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT
Definition AMDKernelCodeT.h:107
@ AMD_CODE_PROPERTY_IS_PTR64
Definition AMDKernelCodeT.h:172
@ AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID_WIDTH
Definition AMDKernelCodeT.h:102
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_WIDTH
Definition AMDKernelCodeT.h:192
@ AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK
Definition AMDKernelCodeT.h:184
@ AMD_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE_WIDTH
Definition AMDKernelCodeT.h:110
@ AMD_CODE_PROPERTY_IS_XNACK_SUPPORTED_SHIFT
Definition AMDKernelCodeT.h:191
@ AMD_CODE_PROPERTY_ENABLE_SGPR_GRID_WORKGROUP_COUNT_X_WIDTH
Definition AMDKernelCodeT.h:114
@ AMD_CODE_PROPERTY_IS_DEBUG_SUPPORTED_SHIFT
Indicate if code generated has support for debugging.
Definition AMDKernelCodeT.h:187
uint32_t amd_code_version32_t
The version of the amd_*_code_t struct.
Definition AMDKernelCodeT.h:36
struct hsa_ext_control_directives_s hsa_ext_control_directives_t
The hsa_ext_control_directives_t specifies the values for the HSAIL control directives.
uint32_t hsa_ext_code_kind32_t
Definition AMDKernelCodeT.h:26
uint16_t hsa_ext_exception_kind16_t
Definition AMDKernelCodeT.h:25
AMD Kernel Code Object (amd_kernel_code_t).
Definition AMDKernelCodeT.h:526
uint16_t amd_machine_version_minor
Definition AMDKernelCodeT.h:531
uint16_t workitem_vgpr_count
Number of vector registers used by each work-item.
Definition AMDKernelCodeT.h:599
uint16_t amd_machine_kind
Definition AMDKernelCodeT.h:529
uint16_t reserved_vgpr_first
If reserved_vgpr_count is 0 then must be 0.
Definition AMDKernelCodeT.h:603
uint64_t runtime_loader_kernel_symbol
Definition AMDKernelCodeT.h:647
uint8_t group_segment_alignment
Definition AMDKernelCodeT.h:635
uint32_t code_properties
Code properties.
Definition AMDKernelCodeT.h:562
uint8_t private_segment_alignment
Definition AMDKernelCodeT.h:636
uint32_t amd_kernel_code_version_major
Definition AMDKernelCodeT.h:527
uint8_t reserved3[12]
Definition AMDKernelCodeT.h:646
uint16_t reserved_vgpr_count
The number of consecutive VGPRs reserved by the client.
Definition AMDKernelCodeT.h:608
uint8_t kernarg_segment_alignment
The maximum byte alignment of variables used by the kernel in the specified memory segment.
Definition AMDKernelCodeT.h:634
uint8_t wavefront_size
Wavefront size expressed as a power of two.
Definition AMDKernelCodeT.h:643
uint16_t reserved_sgpr_first
If reserved_sgpr_count is 0 then must be 0.
Definition AMDKernelCodeT.h:612
uint16_t amd_machine_version_major
Definition AMDKernelCodeT.h:530
uint32_t workgroup_group_segment_byte_size
The amount of group segment memory required by a work-group in bytes.
Definition AMDKernelCodeT.h:574
uint16_t wavefront_sgpr_count
Number of scalar registers used by a wavefront.
Definition AMDKernelCodeT.h:595
uint32_t gds_segment_byte_size
Number of byte of GDS required by kernel dispatch.
Definition AMDKernelCodeT.h:578
int64_t kernel_code_entry_byte_offset
Byte offset (possibly negative) from start of amd_kernel_code_t object to kernel's entry point instru...
Definition AMDKernelCodeT.h:544
int64_t kernel_code_prefetch_byte_offset
Range of bytes to consider prefetching expressed as an offset and size.
Definition AMDKernelCodeT.h:550
int32_t call_convention
Definition AMDKernelCodeT.h:645
uint16_t amd_machine_version_stepping
Definition AMDKernelCodeT.h:532
uint32_t workitem_private_segment_byte_size
The amount of memory required for the combined private, spill and arg segments for a work-item in byt...
Definition AMDKernelCodeT.h:568
uint64_t reserved0
Reserved. Must be 0.
Definition AMDKernelCodeT.h:554
uint64_t kernarg_segment_byte_size
The size in bytes of the kernarg segment that holds the values of the arguments to the kernel.
Definition AMDKernelCodeT.h:583
uint64_t kernel_code_prefetch_byte_size
Definition AMDKernelCodeT.h:551
uint64_t control_directives[16]
Definition AMDKernelCodeT.h:648
uint16_t debug_private_segment_buffer_sgpr
If is_debug_supported is 0 then must be 0.
Definition AMDKernelCodeT.h:629
uint32_t workgroup_fbarrier_count
Number of fbarrier's used in the kernel and all functions it calls.
Definition AMDKernelCodeT.h:589
uint32_t amd_kernel_code_version_minor
Definition AMDKernelCodeT.h:528
uint16_t debug_wavefront_private_segment_offset_sgpr
If is_debug_supported is 0 then must be 0.
Definition AMDKernelCodeT.h:623
uint16_t reserved_sgpr_count
The number of consecutive SGPRs reserved by the client.
Definition AMDKernelCodeT.h:617
uint64_t compute_pgm_resource_registers
Shader program settings for CS.
Definition AMDKernelCodeT.h:558
Definition AMDKernelCodeT.h:28
uint32_t x
Definition AMDKernelCodeT.h:29
uint32_t z
Definition AMDKernelCodeT.h:31
uint32_t y
Definition AMDKernelCodeT.h:30
The hsa_ext_control_directives_t specifies the values for the HSAIL control directives.
Definition AMDKernelCodeT.h:209
uint8_t reserved[75]
Reserved. Must be 0.
Definition AMDKernelCodeT.h:313
hsa_ext_exception_kind16_t enable_detect_exceptions
If enableDetectExceptions is not enabled then must be 0, otherwise must be non-0 and specifies the se...
Definition AMDKernelCodeT.h:238
hsa_ext_control_directive_present64_t enabled_control_directives
This is a bit set indicating which control directives have been specified.
Definition AMDKernelCodeT.h:215
uint32_t max_dynamic_group_size
If maxDynamicGroupSize is not enabled then must be 0, and any amount of dynamic group segment can be ...
Definition AMDKernelCodeT.h:255
uint32_t requested_workgroups_per_cu
If requestedWorkgroupsPerCu is not enabled then must be 0, and the finalizer is free to generate ISA ...
Definition AMDKernelCodeT.h:282
hsa_ext_exception_kind16_t enable_break_exceptions
If enableBreakExceptions is not enabled then must be 0, otherwise must be non-0 and specifies the set...
Definition AMDKernelCodeT.h:226
uint32_t max_flat_grid_size
If maxFlatGridSize is not enabled then must be 0, otherwise must be greater than 0.
Definition AMDKernelCodeT.h:260
hsa_dim3_t required_grid_size
If not enabled then all elements for Dim3 must be 0, otherwise every element must be greater than 0.
Definition AMDKernelCodeT.h:287
hsa_dim3_t required_workgroup_size
If requiredWorkgroupSize is not enabled then all elements for Dim3 must be 0, and the produced code c...
Definition AMDKernelCodeT.h:300
uint8_t required_dim
If requiredDim is not enabled then must be 0 and the produced kernel code can be dispatched with 1,...
Definition AMDKernelCodeT.h:310
uint32_t max_flat_workgroup_size
If maxFlatWorkgroupSize is not enabled then must be 0, otherwise must be greater than 0.
Definition AMDKernelCodeT.h:265