master
1//! This file is auto-generated by tools/update_cpu_features.zig.
2
3const std = @import("../std.zig");
4const CpuFeature = std.Target.Cpu.Feature;
5const CpuModel = std.Target.Cpu.Model;
6
7pub const Feature = enum {
8 @"16_bit_insts",
9 @"64_bit_literals",
10 a16,
11 add_no_carry_insts,
12 addressablelocalmemorysize163840,
13 addressablelocalmemorysize32768,
14 addressablelocalmemorysize65536,
15 agent_scope_fine_grained_remote_memory_atomics,
16 allocate1_5xvgprs,
17 aperture_regs,
18 architected_flat_scratch,
19 architected_sgprs,
20 ashr_pk_insts,
21 atomic_buffer_global_pk_add_f16_insts,
22 atomic_buffer_global_pk_add_f16_no_rtn_insts,
23 atomic_buffer_pk_add_bf16_inst,
24 atomic_csub_no_rtn_insts,
25 atomic_ds_pk_add_16_insts,
26 atomic_fadd_no_rtn_insts,
27 atomic_fadd_rtn_insts,
28 atomic_flat_pk_add_16_insts,
29 atomic_fmin_fmax_flat_f32,
30 atomic_fmin_fmax_flat_f64,
31 atomic_fmin_fmax_global_f32,
32 atomic_fmin_fmax_global_f64,
33 atomic_global_pk_add_bf16_inst,
34 auto_waitcnt_before_barrier,
35 back_off_barrier,
36 bf16_cvt_insts,
37 bf16_trans_insts,
38 bf8_cvt_scale_insts,
39 bitop3_insts,
40 block_vgpr_csr,
41 bvh_dual_bvh_8_insts,
42 ci_insts,
43 cumode,
44 cvt_fp8_vop1_bug,
45 cvt_pk_f16_f32_inst,
46 default_component_broadcast,
47 default_component_zero,
48 dl_insts,
49 dot10_insts,
50 dot11_insts,
51 dot12_insts,
52 dot13_insts,
53 dot1_insts,
54 dot2_insts,
55 dot3_insts,
56 dot4_insts,
57 dot5_insts,
58 dot6_insts,
59 dot7_insts,
60 dot8_insts,
61 dot9_insts,
62 dpp,
63 dpp8,
64 dpp_64bit,
65 dpp_src1_sgpr,
66 ds128,
67 ds_src2_insts,
68 dynamic_vgpr,
69 dynamic_vgpr_block_size_32,
70 extended_image_insts,
71 f16bf16_to_fp6bf6_cvt_scale_insts,
72 f32_to_f16bf16_cvt_sr_insts,
73 fast_denormal_f32,
74 fast_fmaf,
75 flat_address_space,
76 flat_atomic_fadd_f32_inst,
77 flat_buffer_global_fadd_f64_inst,
78 flat_for_global,
79 flat_global_insts,
80 flat_inst_offsets,
81 flat_scratch,
82 flat_scratch_insts,
83 flat_segment_offset_bug,
84 fma_mix_insts,
85 fmacf64_inst,
86 fmaf,
87 fp4_cvt_scale_insts,
88 fp64,
89 fp6bf6_cvt_scale_insts,
90 fp8_conversion_insts,
91 fp8_cvt_scale_insts,
92 fp8_insts,
93 fp8e5m3_insts,
94 full_rate_64_ops,
95 g16,
96 gcn3_encoding,
97 gds,
98 get_wave_id_inst,
99 gfx10,
100 gfx10_3_insts,
101 gfx10_a_encoding,
102 gfx10_b_encoding,
103 gfx10_insts,
104 gfx11,
105 gfx11_insts,
106 gfx12,
107 gfx1250_insts,
108 gfx12_insts,
109 gfx7_gfx8_gfx9_insts,
110 gfx8_insts,
111 gfx9,
112 gfx90a_insts,
113 gfx940_insts,
114 gfx950_insts,
115 gfx9_insts,
116 gws,
117 half_rate_64_ops,
118 ieee_minimum_maximum_insts,
119 image_gather4_d16_bug,
120 image_insts,
121 image_store_d16_bug,
122 inst_fwd_prefetch_bug,
123 int_clamp_insts,
124 inv_2pi_inline_imm,
125 kernarg_preload,
126 lds_barrier_arrive_atomic,
127 lds_branch_vmem_war_hazard,
128 lds_misaligned_bug,
129 ldsbankcount16,
130 ldsbankcount32,
131 load_store_opt,
132 lshl_add_u64_inst,
133 mad_intra_fwd_bug,
134 mad_mac_f32_insts,
135 mad_mix_insts,
136 mai_insts,
137 max_hard_clause_length_32,
138 max_hard_clause_length_63,
139 max_private_element_size_16,
140 max_private_element_size_4,
141 max_private_element_size_8,
142 memory_atomic_fadd_f32_denormal_support,
143 mfma_inline_literal_bug,
144 mimg_r128,
145 minimum3_maximum3_f16,
146 minimum3_maximum3_f32,
147 minimum3_maximum3_pkf16,
148 movrel,
149 msaa_load_dst_sel_bug,
150 negative_scratch_offset_bug,
151 negative_unaligned_scratch_offset_bug,
152 no_data_dep_hazard,
153 no_sdst_cmpx,
154 nsa_clause_bug,
155 nsa_encoding,
156 nsa_to_vmem_bug,
157 offset_3f_bug,
158 packed_fp32_ops,
159 packed_tid,
160 partial_nsa_encoding,
161 permlane16_swap,
162 permlane32_swap,
163 pk_fmac_f16_inst,
164 point_sample_accel,
165 precise_memory,
166 priv_enabled_trap2_nop_bug,
167 prng_inst,
168 promote_alloca,
169 prt_strict_null,
170 pseudo_scalar_trans,
171 r128_a16,
172 real_true16,
173 relaxed_buffer_oob_mode,
174 required_export_priority,
175 requires_cov6,
176 restricted_soffset,
177 s_memrealtime,
178 s_memtime_inst,
179 safe_smem_prefetch,
180 salu_float,
181 scalar_atomics,
182 scalar_dwordx3_loads,
183 scalar_flat_scratch_insts,
184 scalar_stores,
185 sdwa,
186 sdwa_mav,
187 sdwa_omod,
188 sdwa_out_mods_vopc,
189 sdwa_scalar,
190 sdwa_sdst,
191 sea_islands,
192 setprio_inc_wg_inst,
193 sgpr_init_bug,
194 shader_cycles_hi_lo_registers,
195 shader_cycles_register,
196 si_scheduler,
197 smem_to_vector_write_hazard,
198 southern_islands,
199 sramecc,
200 sramecc_support,
201 tgsplit,
202 transpose_load_f4f6_insts,
203 trap_handler,
204 trig_reduced_range,
205 true16,
206 unaligned_access_mode,
207 unaligned_buffer_access,
208 unaligned_ds_access,
209 unaligned_scratch_access,
210 unpacked_d16_vmem,
211 unsafe_ds_offset_folding,
212 user_sgpr_init16_bug,
213 valu_trans_use_hazard,
214 vcmpx_exec_war_hazard,
215 vcmpx_permlane_hazard,
216 vgpr_index_mode,
217 vmem_to_lds_load_insts,
218 vmem_to_scalar_write_hazard,
219 vmem_write_vgpr_in_order,
220 volcanic_islands,
221 vop3_literal,
222 vop3p,
223 vopd,
224 vscnt,
225 wait_xcnt,
226 wavefrontsize16,
227 wavefrontsize32,
228 wavefrontsize64,
229 xf32_insts,
230 xnack,
231 xnack_support,
232};
233
234pub const featureSet = CpuFeature.FeatureSetFns(Feature).featureSet;
235pub const featureSetHas = CpuFeature.FeatureSetFns(Feature).featureSetHas;
236pub const featureSetHasAny = CpuFeature.FeatureSetFns(Feature).featureSetHasAny;
237pub const featureSetHasAll = CpuFeature.FeatureSetFns(Feature).featureSetHasAll;
238
239pub const all_features = blk: {
240 @setEvalBranchQuota(2000);
241 const len = @typeInfo(Feature).@"enum".fields.len;
242 std.debug.assert(len <= CpuFeature.Set.needed_bit_count);
243 var result: [len]CpuFeature = undefined;
244 result[@intFromEnum(Feature.@"16_bit_insts")] = .{
245 .llvm_name = "16-bit-insts",
246 .description = "Has i16/f16 instructions",
247 .dependencies = featureSet(&[_]Feature{}),
248 };
249 result[@intFromEnum(Feature.@"64_bit_literals")] = .{
250 .llvm_name = "64-bit-literals",
251 .description = "Can use 64-bit literals with single DWORD instructions",
252 .dependencies = featureSet(&[_]Feature{}),
253 };
254 result[@intFromEnum(Feature.a16)] = .{
255 .llvm_name = "a16",
256 .description = "Support A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands",
257 .dependencies = featureSet(&[_]Feature{}),
258 };
259 result[@intFromEnum(Feature.add_no_carry_insts)] = .{
260 .llvm_name = "add-no-carry-insts",
261 .description = "Have VALU add/sub instructions without carry out",
262 .dependencies = featureSet(&[_]Feature{}),
263 };
264 result[@intFromEnum(Feature.addressablelocalmemorysize163840)] = .{
265 .llvm_name = "addressablelocalmemorysize163840",
266 .description = "The size of local memory in bytes",
267 .dependencies = featureSet(&[_]Feature{}),
268 };
269 result[@intFromEnum(Feature.addressablelocalmemorysize32768)] = .{
270 .llvm_name = "addressablelocalmemorysize32768",
271 .description = "The size of local memory in bytes",
272 .dependencies = featureSet(&[_]Feature{}),
273 };
274 result[@intFromEnum(Feature.addressablelocalmemorysize65536)] = .{
275 .llvm_name = "addressablelocalmemorysize65536",
276 .description = "The size of local memory in bytes",
277 .dependencies = featureSet(&[_]Feature{}),
278 };
279 result[@intFromEnum(Feature.agent_scope_fine_grained_remote_memory_atomics)] = .{
280 .llvm_name = "agent-scope-fine-grained-remote-memory-atomics",
281 .description = "Agent (device) scoped atomic operations, excluding those directly supported by PCIe (i.e. integer atomic add, exchange, and compare-and-swap), are functional for allocations in host or peer device memory.",
282 .dependencies = featureSet(&[_]Feature{}),
283 };
284 result[@intFromEnum(Feature.allocate1_5xvgprs)] = .{
285 .llvm_name = "allocate1_5xvgprs",
286 .description = "Has 50% more physical VGPRs and 50% larger allocation granule",
287 .dependencies = featureSet(&[_]Feature{}),
288 };
289 result[@intFromEnum(Feature.aperture_regs)] = .{
290 .llvm_name = "aperture-regs",
291 .description = "Has Memory Aperture Base and Size Registers",
292 .dependencies = featureSet(&[_]Feature{}),
293 };
294 result[@intFromEnum(Feature.architected_flat_scratch)] = .{
295 .llvm_name = "architected-flat-scratch",
296 .description = "Flat Scratch register is a readonly SPI initialized architected register",
297 .dependencies = featureSet(&[_]Feature{}),
298 };
299 result[@intFromEnum(Feature.architected_sgprs)] = .{
300 .llvm_name = "architected-sgprs",
301 .description = "Enable the architected SGPRs",
302 .dependencies = featureSet(&[_]Feature{}),
303 };
304 result[@intFromEnum(Feature.ashr_pk_insts)] = .{
305 .llvm_name = "ashr-pk-insts",
306 .description = "Has Arithmetic Shift Pack instructions",
307 .dependencies = featureSet(&[_]Feature{}),
308 };
309 result[@intFromEnum(Feature.atomic_buffer_global_pk_add_f16_insts)] = .{
310 .llvm_name = "atomic-buffer-global-pk-add-f16-insts",
311 .description = "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that can return original value",
312 .dependencies = featureSet(&[_]Feature{
313 .flat_global_insts,
314 }),
315 };
316 result[@intFromEnum(Feature.atomic_buffer_global_pk_add_f16_no_rtn_insts)] = .{
317 .llvm_name = "atomic-buffer-global-pk-add-f16-no-rtn-insts",
318 .description = "Has buffer_atomic_pk_add_f16 and global_atomic_pk_add_f16 instructions that don't return original value",
319 .dependencies = featureSet(&[_]Feature{
320 .flat_global_insts,
321 }),
322 };
323 result[@intFromEnum(Feature.atomic_buffer_pk_add_bf16_inst)] = .{
324 .llvm_name = "atomic-buffer-pk-add-bf16-inst",
325 .description = "Has buffer_atomic_pk_add_bf16 instruction",
326 .dependencies = featureSet(&[_]Feature{}),
327 };
328 result[@intFromEnum(Feature.atomic_csub_no_rtn_insts)] = .{
329 .llvm_name = "atomic-csub-no-rtn-insts",
330 .description = "Has buffer_atomic_csub and global_atomic_csub instructions that don't return original value",
331 .dependencies = featureSet(&[_]Feature{}),
332 };
333 result[@intFromEnum(Feature.atomic_ds_pk_add_16_insts)] = .{
334 .llvm_name = "atomic-ds-pk-add-16-insts",
335 .description = "Has ds_pk_add_bf16, ds_pk_add_f16, ds_pk_add_rtn_bf16, ds_pk_add_rtn_f16 instructions",
336 .dependencies = featureSet(&[_]Feature{}),
337 };
338 result[@intFromEnum(Feature.atomic_fadd_no_rtn_insts)] = .{
339 .llvm_name = "atomic-fadd-no-rtn-insts",
340 .description = "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that don't return original value",
341 .dependencies = featureSet(&[_]Feature{
342 .flat_global_insts,
343 }),
344 };
345 result[@intFromEnum(Feature.atomic_fadd_rtn_insts)] = .{
346 .llvm_name = "atomic-fadd-rtn-insts",
347 .description = "Has buffer_atomic_add_f32 and global_atomic_add_f32 instructions that return original value",
348 .dependencies = featureSet(&[_]Feature{
349 .flat_global_insts,
350 }),
351 };
352 result[@intFromEnum(Feature.atomic_flat_pk_add_16_insts)] = .{
353 .llvm_name = "atomic-flat-pk-add-16-insts",
354 .description = "Has flat_atomic_pk_add_f16 and flat_atomic_pk_add_bf16 instructions",
355 .dependencies = featureSet(&[_]Feature{}),
356 };
357 result[@intFromEnum(Feature.atomic_fmin_fmax_flat_f32)] = .{
358 .llvm_name = "atomic-fmin-fmax-flat-f32",
359 .description = "Has flat memory instructions for atomicrmw fmin/fmax for float",
360 .dependencies = featureSet(&[_]Feature{}),
361 };
362 result[@intFromEnum(Feature.atomic_fmin_fmax_flat_f64)] = .{
363 .llvm_name = "atomic-fmin-fmax-flat-f64",
364 .description = "Has flat memory instructions for atomicrmw fmin/fmax for double",
365 .dependencies = featureSet(&[_]Feature{}),
366 };
367 result[@intFromEnum(Feature.atomic_fmin_fmax_global_f32)] = .{
368 .llvm_name = "atomic-fmin-fmax-global-f32",
369 .description = "Has global/buffer instructions for atomicrmw fmin/fmax for float",
370 .dependencies = featureSet(&[_]Feature{}),
371 };
372 result[@intFromEnum(Feature.atomic_fmin_fmax_global_f64)] = .{
373 .llvm_name = "atomic-fmin-fmax-global-f64",
374 .description = "Has global/buffer instructions for atomicrmw fmin/fmax for float",
375 .dependencies = featureSet(&[_]Feature{}),
376 };
377 result[@intFromEnum(Feature.atomic_global_pk_add_bf16_inst)] = .{
378 .llvm_name = "atomic-global-pk-add-bf16-inst",
379 .description = "Has global_atomic_pk_add_bf16 instruction",
380 .dependencies = featureSet(&[_]Feature{
381 .flat_global_insts,
382 }),
383 };
384 result[@intFromEnum(Feature.auto_waitcnt_before_barrier)] = .{
385 .llvm_name = "auto-waitcnt-before-barrier",
386 .description = "Hardware automatically inserts waitcnt before barrier",
387 .dependencies = featureSet(&[_]Feature{}),
388 };
389 result[@intFromEnum(Feature.back_off_barrier)] = .{
390 .llvm_name = "back-off-barrier",
391 .description = "Hardware supports backing off s_barrier if an exception occurs",
392 .dependencies = featureSet(&[_]Feature{}),
393 };
394 result[@intFromEnum(Feature.bf16_cvt_insts)] = .{
395 .llvm_name = "bf16-cvt-insts",
396 .description = "Has bf16 conversion instructions",
397 .dependencies = featureSet(&[_]Feature{}),
398 };
399 result[@intFromEnum(Feature.bf16_trans_insts)] = .{
400 .llvm_name = "bf16-trans-insts",
401 .description = "Has bf16 transcendental instructions",
402 .dependencies = featureSet(&[_]Feature{}),
403 };
404 result[@intFromEnum(Feature.bf8_cvt_scale_insts)] = .{
405 .llvm_name = "bf8-cvt-scale-insts",
406 .description = "Has bf8 conversion scale instructions",
407 .dependencies = featureSet(&[_]Feature{}),
408 };
409 result[@intFromEnum(Feature.bitop3_insts)] = .{
410 .llvm_name = "bitop3-insts",
411 .description = "Has v_bitop3_b32/v_bitop3_b16 instructions",
412 .dependencies = featureSet(&[_]Feature{}),
413 };
414 result[@intFromEnum(Feature.block_vgpr_csr)] = .{
415 .llvm_name = "block-vgpr-csr",
416 .description = "Use block load/store for VGPR callee saved registers",
417 .dependencies = featureSet(&[_]Feature{}),
418 };
419 result[@intFromEnum(Feature.bvh_dual_bvh_8_insts)] = .{
420 .llvm_name = "bvh-dual-bvh-8-insts",
421 .description = "Has image_bvh_dual_intersect_ray and image_bvh8_intersect_ray instructions",
422 .dependencies = featureSet(&[_]Feature{}),
423 };
424 result[@intFromEnum(Feature.ci_insts)] = .{
425 .llvm_name = "ci-insts",
426 .description = "Additional instructions for CI+",
427 .dependencies = featureSet(&[_]Feature{}),
428 };
429 result[@intFromEnum(Feature.cumode)] = .{
430 .llvm_name = "cumode",
431 .description = "Enable CU wavefront execution mode",
432 .dependencies = featureSet(&[_]Feature{}),
433 };
434 result[@intFromEnum(Feature.cvt_fp8_vop1_bug)] = .{
435 .llvm_name = "cvt-fp8-vop1-bug",
436 .description = "FP8/BF8 VOP1 form of conversion to F32 is unreliable",
437 .dependencies = featureSet(&[_]Feature{
438 .fp8_conversion_insts,
439 }),
440 };
441 result[@intFromEnum(Feature.cvt_pk_f16_f32_inst)] = .{
442 .llvm_name = "cvt-pk-f16-f32-inst",
443 .description = "Has cvt_pk_f16_f32 instruction",
444 .dependencies = featureSet(&[_]Feature{}),
445 };
446 result[@intFromEnum(Feature.default_component_broadcast)] = .{
447 .llvm_name = "default-component-broadcast",
448 .description = "BUFFER/IMAGE store instructions set unspecified components to x component (GFX12)",
449 .dependencies = featureSet(&[_]Feature{}),
450 };
451 result[@intFromEnum(Feature.default_component_zero)] = .{
452 .llvm_name = "default-component-zero",
453 .description = "BUFFER/IMAGE store instructions set unspecified components to zero (before GFX12)",
454 .dependencies = featureSet(&[_]Feature{}),
455 };
456 result[@intFromEnum(Feature.dl_insts)] = .{
457 .llvm_name = "dl-insts",
458 .description = "Has v_fmac_f32 and v_xnor_b32 instructions",
459 .dependencies = featureSet(&[_]Feature{}),
460 };
461 result[@intFromEnum(Feature.dot10_insts)] = .{
462 .llvm_name = "dot10-insts",
463 .description = "Has v_dot2_f32_f16 instruction",
464 .dependencies = featureSet(&[_]Feature{}),
465 };
466 result[@intFromEnum(Feature.dot11_insts)] = .{
467 .llvm_name = "dot11-insts",
468 .description = "Has v_dot4_f32_fp8_fp8, v_dot4_f32_fp8_bf8, v_dot4_f32_bf8_fp8, v_dot4_f32_bf8_bf8 instructions",
469 .dependencies = featureSet(&[_]Feature{}),
470 };
471 result[@intFromEnum(Feature.dot12_insts)] = .{
472 .llvm_name = "dot12-insts",
473 .description = "Has v_dot2_f32_bf16 instructions",
474 .dependencies = featureSet(&[_]Feature{}),
475 };
476 result[@intFromEnum(Feature.dot13_insts)] = .{
477 .llvm_name = "dot13-insts",
478 .description = "Has v_dot2c_f32_bf16 instructions",
479 .dependencies = featureSet(&[_]Feature{}),
480 };
481 result[@intFromEnum(Feature.dot1_insts)] = .{
482 .llvm_name = "dot1-insts",
483 .description = "Has v_dot4_i32_i8 and v_dot8_i32_i4 instructions",
484 .dependencies = featureSet(&[_]Feature{}),
485 };
486 result[@intFromEnum(Feature.dot2_insts)] = .{
487 .llvm_name = "dot2-insts",
488 .description = "Has v_dot2_i32_i16, v_dot2_u32_u16 instructions",
489 .dependencies = featureSet(&[_]Feature{}),
490 };
491 result[@intFromEnum(Feature.dot3_insts)] = .{
492 .llvm_name = "dot3-insts",
493 .description = "Has v_dot8c_i32_i4 instruction",
494 .dependencies = featureSet(&[_]Feature{}),
495 };
496 result[@intFromEnum(Feature.dot4_insts)] = .{
497 .llvm_name = "dot4-insts",
498 .description = "Has v_dot2c_i32_i16 instruction",
499 .dependencies = featureSet(&[_]Feature{}),
500 };
501 result[@intFromEnum(Feature.dot5_insts)] = .{
502 .llvm_name = "dot5-insts",
503 .description = "Has v_dot2c_f32_f16 instruction",
504 .dependencies = featureSet(&[_]Feature{}),
505 };
506 result[@intFromEnum(Feature.dot6_insts)] = .{
507 .llvm_name = "dot6-insts",
508 .description = "Has v_dot4c_i32_i8 instruction",
509 .dependencies = featureSet(&[_]Feature{}),
510 };
511 result[@intFromEnum(Feature.dot7_insts)] = .{
512 .llvm_name = "dot7-insts",
513 .description = "Has v_dot4_u32_u8, v_dot8_u32_u4 instructions",
514 .dependencies = featureSet(&[_]Feature{}),
515 };
516 result[@intFromEnum(Feature.dot8_insts)] = .{
517 .llvm_name = "dot8-insts",
518 .description = "Has v_dot4_i32_iu8, v_dot8_i32_iu4 instructions",
519 .dependencies = featureSet(&[_]Feature{}),
520 };
521 result[@intFromEnum(Feature.dot9_insts)] = .{
522 .llvm_name = "dot9-insts",
523 .description = "Has v_dot2_f16_f16, v_dot2_bf16_bf16 instructions",
524 .dependencies = featureSet(&[_]Feature{}),
525 };
526 result[@intFromEnum(Feature.dpp)] = .{
527 .llvm_name = "dpp",
528 .description = "Support DPP (Data Parallel Primitives) extension",
529 .dependencies = featureSet(&[_]Feature{}),
530 };
531 result[@intFromEnum(Feature.dpp8)] = .{
532 .llvm_name = "dpp8",
533 .description = "Support DPP8 (Data Parallel Primitives) extension",
534 .dependencies = featureSet(&[_]Feature{}),
535 };
536 result[@intFromEnum(Feature.dpp_64bit)] = .{
537 .llvm_name = "dpp-64bit",
538 .description = "Support DPP (Data Parallel Primitives) extension in DP ALU",
539 .dependencies = featureSet(&[_]Feature{}),
540 };
541 result[@intFromEnum(Feature.dpp_src1_sgpr)] = .{
542 .llvm_name = "dpp-src1-sgpr",
543 .description = "Support SGPR for Src1 of DPP instructions",
544 .dependencies = featureSet(&[_]Feature{}),
545 };
546 result[@intFromEnum(Feature.ds128)] = .{
547 .llvm_name = "enable-ds128",
548 .description = "Use ds_{read|write}_b128",
549 .dependencies = featureSet(&[_]Feature{}),
550 };
551 result[@intFromEnum(Feature.ds_src2_insts)] = .{
552 .llvm_name = "ds-src2-insts",
553 .description = "Has ds_*_src2 instructions",
554 .dependencies = featureSet(&[_]Feature{}),
555 };
556 result[@intFromEnum(Feature.dynamic_vgpr)] = .{
557 .llvm_name = "dynamic-vgpr",
558 .description = "Enable dynamic VGPR mode",
559 .dependencies = featureSet(&[_]Feature{}),
560 };
561 result[@intFromEnum(Feature.dynamic_vgpr_block_size_32)] = .{
562 .llvm_name = "dynamic-vgpr-block-size-32",
563 .description = "Use a block size of 32 for dynamic VGPR allocation (default is 16)",
564 .dependencies = featureSet(&[_]Feature{}),
565 };
566 result[@intFromEnum(Feature.extended_image_insts)] = .{
567 .llvm_name = "extended-image-insts",
568 .description = "Support mips != 0, lod != 0, gather4, and get_lod",
569 .dependencies = featureSet(&[_]Feature{}),
570 };
571 result[@intFromEnum(Feature.f16bf16_to_fp6bf6_cvt_scale_insts)] = .{
572 .llvm_name = "f16bf16-to-fp6bf6-cvt-scale-insts",
573 .description = "Has f16bf16 to fp6bf6 conversion scale instructions",
574 .dependencies = featureSet(&[_]Feature{}),
575 };
576 result[@intFromEnum(Feature.f32_to_f16bf16_cvt_sr_insts)] = .{
577 .llvm_name = "f32-to-f16bf16-cvt-sr-insts",
578 .description = "Has f32 to f16bf16 conversion scale instructions",
579 .dependencies = featureSet(&[_]Feature{}),
580 };
581 result[@intFromEnum(Feature.fast_denormal_f32)] = .{
582 .llvm_name = "fast-denormal-f32",
583 .description = "Enabling denormals does not cause f32 instructions to run at f64 rates",
584 .dependencies = featureSet(&[_]Feature{}),
585 };
586 result[@intFromEnum(Feature.fast_fmaf)] = .{
587 .llvm_name = "fast-fmaf",
588 .description = "Assuming f32 fma is at least as fast as mul + add",
589 .dependencies = featureSet(&[_]Feature{}),
590 };
591 result[@intFromEnum(Feature.flat_address_space)] = .{
592 .llvm_name = "flat-address-space",
593 .description = "Support flat address space",
594 .dependencies = featureSet(&[_]Feature{}),
595 };
596 result[@intFromEnum(Feature.flat_atomic_fadd_f32_inst)] = .{
597 .llvm_name = "flat-atomic-fadd-f32-inst",
598 .description = "Has flat_atomic_add_f32 instruction",
599 .dependencies = featureSet(&[_]Feature{}),
600 };
601 result[@intFromEnum(Feature.flat_buffer_global_fadd_f64_inst)] = .{
602 .llvm_name = "flat-buffer-global-fadd-f64-inst",
603 .description = "Has flat, buffer, and global instructions for f64 atomic fadd",
604 .dependencies = featureSet(&[_]Feature{}),
605 };
606 result[@intFromEnum(Feature.flat_for_global)] = .{
607 .llvm_name = "flat-for-global",
608 .description = "Force to generate flat instruction for global",
609 .dependencies = featureSet(&[_]Feature{}),
610 };
611 result[@intFromEnum(Feature.flat_global_insts)] = .{
612 .llvm_name = "flat-global-insts",
613 .description = "Have global_* flat memory instructions",
614 .dependencies = featureSet(&[_]Feature{}),
615 };
616 result[@intFromEnum(Feature.flat_inst_offsets)] = .{
617 .llvm_name = "flat-inst-offsets",
618 .description = "Flat instructions have immediate offset addressing mode",
619 .dependencies = featureSet(&[_]Feature{}),
620 };
621 result[@intFromEnum(Feature.flat_scratch)] = .{
622 .llvm_name = "enable-flat-scratch",
623 .description = "Use scratch_* flat memory instructions to access scratch",
624 .dependencies = featureSet(&[_]Feature{}),
625 };
626 result[@intFromEnum(Feature.flat_scratch_insts)] = .{
627 .llvm_name = "flat-scratch-insts",
628 .description = "Have scratch_* flat memory instructions",
629 .dependencies = featureSet(&[_]Feature{}),
630 };
631 result[@intFromEnum(Feature.flat_segment_offset_bug)] = .{
632 .llvm_name = "flat-segment-offset-bug",
633 .description = "GFX10 bug where inst_offset is ignored when flat instructions access global memory",
634 .dependencies = featureSet(&[_]Feature{}),
635 };
636 result[@intFromEnum(Feature.fma_mix_insts)] = .{
637 .llvm_name = "fma-mix-insts",
638 .description = "Has v_fma_mix_f32, v_fma_mixlo_f16, v_fma_mixhi_f16 instructions",
639 .dependencies = featureSet(&[_]Feature{}),
640 };
641 result[@intFromEnum(Feature.fmacf64_inst)] = .{
642 .llvm_name = "fmacf64-inst",
643 .description = "Has v_fmac_f64 instruction",
644 .dependencies = featureSet(&[_]Feature{}),
645 };
646 result[@intFromEnum(Feature.fmaf)] = .{
647 .llvm_name = "fmaf",
648 .description = "Enable single precision FMA (not as fast as mul+add, but fused)",
649 .dependencies = featureSet(&[_]Feature{}),
650 };
651 result[@intFromEnum(Feature.fp4_cvt_scale_insts)] = .{
652 .llvm_name = "fp4-cvt-scale-insts",
653 .description = "Has fp4 conversion scale instructions",
654 .dependencies = featureSet(&[_]Feature{}),
655 };
656 result[@intFromEnum(Feature.fp64)] = .{
657 .llvm_name = "fp64",
658 .description = "Enable double precision operations",
659 .dependencies = featureSet(&[_]Feature{}),
660 };
661 result[@intFromEnum(Feature.fp6bf6_cvt_scale_insts)] = .{
662 .llvm_name = "fp6bf6-cvt-scale-insts",
663 .description = "Has fp6 and bf6 conversion scale instructions",
664 .dependencies = featureSet(&[_]Feature{}),
665 };
666 result[@intFromEnum(Feature.fp8_conversion_insts)] = .{
667 .llvm_name = "fp8-conversion-insts",
668 .description = "Has fp8 and bf8 conversion instructions",
669 .dependencies = featureSet(&[_]Feature{}),
670 };
671 result[@intFromEnum(Feature.fp8_cvt_scale_insts)] = .{
672 .llvm_name = "fp8-cvt-scale-insts",
673 .description = "Has fp8 conversion scale instructions",
674 .dependencies = featureSet(&[_]Feature{}),
675 };
676 result[@intFromEnum(Feature.fp8_insts)] = .{
677 .llvm_name = "fp8-insts",
678 .description = "Has fp8 and bf8 instructions",
679 .dependencies = featureSet(&[_]Feature{}),
680 };
681 result[@intFromEnum(Feature.fp8e5m3_insts)] = .{
682 .llvm_name = "fp8e5m3-insts",
683 .description = "Has fp8 e5m3 format support",
684 .dependencies = featureSet(&[_]Feature{}),
685 };
686 result[@intFromEnum(Feature.full_rate_64_ops)] = .{
687 .llvm_name = "full-rate-64-ops",
688 .description = "Most fp64 instructions are full rate",
689 .dependencies = featureSet(&[_]Feature{}),
690 };
691 result[@intFromEnum(Feature.g16)] = .{
692 .llvm_name = "g16",
693 .description = "Support G16 for 16-bit gradient image operands",
694 .dependencies = featureSet(&[_]Feature{}),
695 };
696 result[@intFromEnum(Feature.gcn3_encoding)] = .{
697 .llvm_name = "gcn3-encoding",
698 .description = "Encoding format for VI",
699 .dependencies = featureSet(&[_]Feature{}),
700 };
701 result[@intFromEnum(Feature.gds)] = .{
702 .llvm_name = "gds",
703 .description = "Has Global Data Share",
704 .dependencies = featureSet(&[_]Feature{}),
705 };
706 result[@intFromEnum(Feature.get_wave_id_inst)] = .{
707 .llvm_name = "get-wave-id-inst",
708 .description = "Has s_get_waveid_in_workgroup instruction",
709 .dependencies = featureSet(&[_]Feature{}),
710 };
711 result[@intFromEnum(Feature.gfx10)] = .{
712 .llvm_name = "gfx10",
713 .description = "GFX10 GPU generation",
714 .dependencies = featureSet(&[_]Feature{
715 .@"16_bit_insts",
716 .a16,
717 .add_no_carry_insts,
718 .addressablelocalmemorysize65536,
719 .aperture_regs,
720 .atomic_fmin_fmax_flat_f32,
721 .atomic_fmin_fmax_flat_f64,
722 .atomic_fmin_fmax_global_f32,
723 .atomic_fmin_fmax_global_f64,
724 .ci_insts,
725 .default_component_zero,
726 .dpp,
727 .dpp8,
728 .extended_image_insts,
729 .fast_denormal_f32,
730 .fast_fmaf,
731 .flat_address_space,
732 .flat_global_insts,
733 .flat_inst_offsets,
734 .flat_scratch_insts,
735 .fma_mix_insts,
736 .fp64,
737 .g16,
738 .gds,
739 .gfx10_insts,
740 .gfx8_insts,
741 .gfx9_insts,
742 .gws,
743 .image_insts,
744 .int_clamp_insts,
745 .inv_2pi_inline_imm,
746 .max_hard_clause_length_63,
747 .mimg_r128,
748 .movrel,
749 .no_data_dep_hazard,
750 .no_sdst_cmpx,
751 .pk_fmac_f16_inst,
752 .s_memrealtime,
753 .s_memtime_inst,
754 .sdwa,
755 .sdwa_omod,
756 .sdwa_scalar,
757 .sdwa_sdst,
758 .unaligned_buffer_access,
759 .unaligned_ds_access,
760 .unaligned_scratch_access,
761 .vmem_to_lds_load_insts,
762 .vmem_write_vgpr_in_order,
763 .vop3_literal,
764 .vop3p,
765 .vscnt,
766 }),
767 };
768 result[@intFromEnum(Feature.gfx10_3_insts)] = .{
769 .llvm_name = "gfx10-3-insts",
770 .description = "Additional instructions for GFX10.3",
771 .dependencies = featureSet(&[_]Feature{}),
772 };
773 result[@intFromEnum(Feature.gfx10_a_encoding)] = .{
774 .llvm_name = "gfx10_a-encoding",
775 .description = "Has BVH ray tracing instructions",
776 .dependencies = featureSet(&[_]Feature{}),
777 };
778 result[@intFromEnum(Feature.gfx10_b_encoding)] = .{
779 .llvm_name = "gfx10_b-encoding",
780 .description = "Encoding format GFX10_B",
781 .dependencies = featureSet(&[_]Feature{}),
782 };
783 result[@intFromEnum(Feature.gfx10_insts)] = .{
784 .llvm_name = "gfx10-insts",
785 .description = "Additional instructions for GFX10+",
786 .dependencies = featureSet(&[_]Feature{}),
787 };
788 result[@intFromEnum(Feature.gfx11)] = .{
789 .llvm_name = "gfx11",
790 .description = "GFX11 GPU generation",
791 .dependencies = featureSet(&[_]Feature{
792 .@"16_bit_insts",
793 .a16,
794 .add_no_carry_insts,
795 .addressablelocalmemorysize65536,
796 .aperture_regs,
797 .atomic_fmin_fmax_flat_f32,
798 .atomic_fmin_fmax_global_f32,
799 .ci_insts,
800 .default_component_zero,
801 .dpp,
802 .dpp8,
803 .extended_image_insts,
804 .fast_denormal_f32,
805 .fast_fmaf,
806 .flat_address_space,
807 .flat_global_insts,
808 .flat_inst_offsets,
809 .flat_scratch_insts,
810 .fma_mix_insts,
811 .fp64,
812 .g16,
813 .gds,
814 .gfx10_3_insts,
815 .gfx10_a_encoding,
816 .gfx10_b_encoding,
817 .gfx10_insts,
818 .gfx11_insts,
819 .gfx8_insts,
820 .gfx9_insts,
821 .gws,
822 .int_clamp_insts,
823 .inv_2pi_inline_imm,
824 .max_hard_clause_length_32,
825 .mimg_r128,
826 .movrel,
827 .no_data_dep_hazard,
828 .no_sdst_cmpx,
829 .pk_fmac_f16_inst,
830 .true16,
831 .unaligned_buffer_access,
832 .unaligned_ds_access,
833 .unaligned_scratch_access,
834 .vmem_write_vgpr_in_order,
835 .vop3_literal,
836 .vop3p,
837 .vopd,
838 .vscnt,
839 }),
840 };
841 result[@intFromEnum(Feature.gfx11_insts)] = .{
842 .llvm_name = "gfx11-insts",
843 .description = "Additional instructions for GFX11+",
844 .dependencies = featureSet(&[_]Feature{}),
845 };
846 result[@intFromEnum(Feature.gfx12)] = .{
847 .llvm_name = "gfx12",
848 .description = "GFX12 GPU generation",
849 .dependencies = featureSet(&[_]Feature{
850 .@"16_bit_insts",
851 .a16,
852 .add_no_carry_insts,
853 .addressablelocalmemorysize65536,
854 .agent_scope_fine_grained_remote_memory_atomics,
855 .aperture_regs,
856 .atomic_fmin_fmax_flat_f32,
857 .atomic_fmin_fmax_global_f32,
858 .ci_insts,
859 .default_component_broadcast,
860 .dpp,
861 .dpp8,
862 .fast_denormal_f32,
863 .fast_fmaf,
864 .flat_address_space,
865 .flat_global_insts,
866 .flat_inst_offsets,
867 .flat_scratch_insts,
868 .fma_mix_insts,
869 .fp64,
870 .g16,
871 .gfx10_3_insts,
872 .gfx10_a_encoding,
873 .gfx10_b_encoding,
874 .gfx10_insts,
875 .gfx11_insts,
876 .gfx12_insts,
877 .gfx8_insts,
878 .gfx9_insts,
879 .ieee_minimum_maximum_insts,
880 .int_clamp_insts,
881 .inv_2pi_inline_imm,
882 .max_hard_clause_length_32,
883 .mimg_r128,
884 .minimum3_maximum3_f16,
885 .minimum3_maximum3_f32,
886 .movrel,
887 .no_data_dep_hazard,
888 .no_sdst_cmpx,
889 .pk_fmac_f16_inst,
890 .true16,
891 .unaligned_buffer_access,
892 .unaligned_ds_access,
893 .unaligned_scratch_access,
894 .vop3_literal,
895 .vop3p,
896 .vopd,
897 .vscnt,
898 }),
899 };
900 result[@intFromEnum(Feature.gfx1250_insts)] = .{
901 .llvm_name = "gfx1250-insts",
902 .description = "Additional instructions for GFX1250+",
903 .dependencies = featureSet(&[_]Feature{}),
904 };
905 result[@intFromEnum(Feature.gfx12_insts)] = .{
906 .llvm_name = "gfx12-insts",
907 .description = "Additional instructions for GFX12+",
908 .dependencies = featureSet(&[_]Feature{}),
909 };
910 result[@intFromEnum(Feature.gfx7_gfx8_gfx9_insts)] = .{
911 .llvm_name = "gfx7-gfx8-gfx9-insts",
912 .description = "Instructions shared in GFX7, GFX8, GFX9",
913 .dependencies = featureSet(&[_]Feature{}),
914 };
915 result[@intFromEnum(Feature.gfx8_insts)] = .{
916 .llvm_name = "gfx8-insts",
917 .description = "Additional instructions for GFX8+",
918 .dependencies = featureSet(&[_]Feature{}),
919 };
920 result[@intFromEnum(Feature.gfx9)] = .{
921 .llvm_name = "gfx9",
922 .description = "GFX9 GPU generation",
923 .dependencies = featureSet(&[_]Feature{
924 .@"16_bit_insts",
925 .a16,
926 .add_no_carry_insts,
927 .aperture_regs,
928 .ci_insts,
929 .default_component_zero,
930 .dpp,
931 .fast_denormal_f32,
932 .fast_fmaf,
933 .flat_address_space,
934 .flat_global_insts,
935 .flat_inst_offsets,
936 .flat_scratch_insts,
937 .fp64,
938 .gcn3_encoding,
939 .gfx7_gfx8_gfx9_insts,
940 .gfx8_insts,
941 .gfx9_insts,
942 .gws,
943 .int_clamp_insts,
944 .inv_2pi_inline_imm,
945 .negative_scratch_offset_bug,
946 .r128_a16,
947 .s_memrealtime,
948 .s_memtime_inst,
949 .scalar_atomics,
950 .scalar_flat_scratch_insts,
951 .scalar_stores,
952 .sdwa,
953 .sdwa_omod,
954 .sdwa_scalar,
955 .sdwa_sdst,
956 .unaligned_buffer_access,
957 .unaligned_ds_access,
958 .unaligned_scratch_access,
959 .vgpr_index_mode,
960 .vmem_to_lds_load_insts,
961 .vmem_write_vgpr_in_order,
962 .vop3p,
963 .wavefrontsize64,
964 .xnack_support,
965 }),
966 };
967 result[@intFromEnum(Feature.gfx90a_insts)] = .{
968 .llvm_name = "gfx90a-insts",
969 .description = "Additional instructions for GFX90A+",
970 .dependencies = featureSet(&[_]Feature{}),
971 };
972 result[@intFromEnum(Feature.gfx940_insts)] = .{
973 .llvm_name = "gfx940-insts",
974 .description = "Additional instructions for GFX940+",
975 .dependencies = featureSet(&[_]Feature{}),
976 };
977 result[@intFromEnum(Feature.gfx950_insts)] = .{
978 .llvm_name = "gfx950-insts",
979 .description = "Additional instructions for GFX950+",
980 .dependencies = featureSet(&[_]Feature{
981 .ashr_pk_insts,
982 .bf8_cvt_scale_insts,
983 .cvt_pk_f16_f32_inst,
984 .f16bf16_to_fp6bf6_cvt_scale_insts,
985 .f32_to_f16bf16_cvt_sr_insts,
986 .fp4_cvt_scale_insts,
987 .fp6bf6_cvt_scale_insts,
988 .fp8_cvt_scale_insts,
989 .minimum3_maximum3_f32,
990 .minimum3_maximum3_pkf16,
991 .permlane16_swap,
992 .permlane32_swap,
993 }),
994 };
995 result[@intFromEnum(Feature.gfx9_insts)] = .{
996 .llvm_name = "gfx9-insts",
997 .description = "Additional instructions for GFX9+",
998 .dependencies = featureSet(&[_]Feature{}),
999 };
1000 result[@intFromEnum(Feature.gws)] = .{
1001 .llvm_name = "gws",
1002 .description = "Has Global Wave Sync",
1003 .dependencies = featureSet(&[_]Feature{}),
1004 };
1005 result[@intFromEnum(Feature.half_rate_64_ops)] = .{
1006 .llvm_name = "half-rate-64-ops",
1007 .description = "Most fp64 instructions are half rate instead of quarter",
1008 .dependencies = featureSet(&[_]Feature{}),
1009 };
1010 result[@intFromEnum(Feature.ieee_minimum_maximum_insts)] = .{
1011 .llvm_name = "ieee-minimum-maximum-insts",
1012 .description = "Has v_minimum/maximum_f16/f32/f64, v_minimummaximum/maximumminimum_f16/f32 and v_pk_minimum/maximum_f16 instructions",
1013 .dependencies = featureSet(&[_]Feature{}),
1014 };
1015 result[@intFromEnum(Feature.image_gather4_d16_bug)] = .{
1016 .llvm_name = "image-gather4-d16-bug",
1017 .description = "Image Gather4 D16 hardware bug",
1018 .dependencies = featureSet(&[_]Feature{}),
1019 };
1020 result[@intFromEnum(Feature.image_insts)] = .{
1021 .llvm_name = "image-insts",
1022 .description = "Support image instructions",
1023 .dependencies = featureSet(&[_]Feature{}),
1024 };
1025 result[@intFromEnum(Feature.image_store_d16_bug)] = .{
1026 .llvm_name = "image-store-d16-bug",
1027 .description = "Image Store D16 hardware bug",
1028 .dependencies = featureSet(&[_]Feature{}),
1029 };
1030 result[@intFromEnum(Feature.inst_fwd_prefetch_bug)] = .{
1031 .llvm_name = "inst-fwd-prefetch-bug",
1032 .description = "S_INST_PREFETCH instruction causes shader to hang",
1033 .dependencies = featureSet(&[_]Feature{}),
1034 };
1035 result[@intFromEnum(Feature.int_clamp_insts)] = .{
1036 .llvm_name = "int-clamp-insts",
1037 .description = "Support clamp for integer destination",
1038 .dependencies = featureSet(&[_]Feature{}),
1039 };
1040 result[@intFromEnum(Feature.inv_2pi_inline_imm)] = .{
1041 .llvm_name = "inv-2pi-inline-imm",
1042 .description = "Has 1 / (2 * pi) as inline immediate",
1043 .dependencies = featureSet(&[_]Feature{}),
1044 };
1045 result[@intFromEnum(Feature.kernarg_preload)] = .{
1046 .llvm_name = "kernarg-preload",
1047 .description = "Hardware supports preloading of kernel arguments in user SGPRs.",
1048 .dependencies = featureSet(&[_]Feature{}),
1049 };
1050 result[@intFromEnum(Feature.lds_barrier_arrive_atomic)] = .{
1051 .llvm_name = "lds-barrier-arrive-atomic",
1052 .description = "Has LDS barrier-arrive atomic instructions",
1053 .dependencies = featureSet(&[_]Feature{}),
1054 };
1055 result[@intFromEnum(Feature.lds_branch_vmem_war_hazard)] = .{
1056 .llvm_name = "lds-branch-vmem-war-hazard",
1057 .description = "Switching between LDS and VMEM-tex not waiting VM_VSRC=0",
1058 .dependencies = featureSet(&[_]Feature{}),
1059 };
1060 result[@intFromEnum(Feature.lds_misaligned_bug)] = .{
1061 .llvm_name = "lds-misaligned-bug",
1062 .description = "Some GFX10 bug with multi-dword LDS and flat access that is not naturally aligned in WGP mode",
1063 .dependencies = featureSet(&[_]Feature{}),
1064 };
1065 result[@intFromEnum(Feature.ldsbankcount16)] = .{
1066 .llvm_name = "ldsbankcount16",
1067 .description = "The number of LDS banks per compute unit.",
1068 .dependencies = featureSet(&[_]Feature{}),
1069 };
1070 result[@intFromEnum(Feature.ldsbankcount32)] = .{
1071 .llvm_name = "ldsbankcount32",
1072 .description = "The number of LDS banks per compute unit.",
1073 .dependencies = featureSet(&[_]Feature{}),
1074 };
1075 result[@intFromEnum(Feature.load_store_opt)] = .{
1076 .llvm_name = "load-store-opt",
1077 .description = "Enable SI load/store optimizer pass",
1078 .dependencies = featureSet(&[_]Feature{}),
1079 };
1080 result[@intFromEnum(Feature.lshl_add_u64_inst)] = .{
1081 .llvm_name = "lshl-add-u64-inst",
1082 .description = "Has v_lshl_add_u64 instruction",
1083 .dependencies = featureSet(&[_]Feature{}),
1084 };
1085 result[@intFromEnum(Feature.mad_intra_fwd_bug)] = .{
1086 .llvm_name = "mad-intra-fwd-bug",
1087 .description = "MAD_U64/I64 intra instruction forwarding bug",
1088 .dependencies = featureSet(&[_]Feature{}),
1089 };
1090 result[@intFromEnum(Feature.mad_mac_f32_insts)] = .{
1091 .llvm_name = "mad-mac-f32-insts",
1092 .description = "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions",
1093 .dependencies = featureSet(&[_]Feature{}),
1094 };
1095 result[@intFromEnum(Feature.mad_mix_insts)] = .{
1096 .llvm_name = "mad-mix-insts",
1097 .description = "Has v_mad_mix_f32, v_mad_mixlo_f16, v_mad_mixhi_f16 instructions",
1098 .dependencies = featureSet(&[_]Feature{}),
1099 };
1100 result[@intFromEnum(Feature.mai_insts)] = .{
1101 .llvm_name = "mai-insts",
1102 .description = "Has mAI instructions",
1103 .dependencies = featureSet(&[_]Feature{}),
1104 };
1105 result[@intFromEnum(Feature.max_hard_clause_length_32)] = .{
1106 .llvm_name = "max-hard-clause-length-32",
1107 .description = "Maximum number of instructions in an explicit S_CLAUSE is 32",
1108 .dependencies = featureSet(&[_]Feature{}),
1109 };
1110 result[@intFromEnum(Feature.max_hard_clause_length_63)] = .{
1111 .llvm_name = "max-hard-clause-length-63",
1112 .description = "Maximum number of instructions in an explicit S_CLAUSE is 63",
1113 .dependencies = featureSet(&[_]Feature{}),
1114 };
1115 result[@intFromEnum(Feature.max_private_element_size_16)] = .{
1116 .llvm_name = "max-private-element-size-16",
1117 .description = "Maximum private access size may be 16",
1118 .dependencies = featureSet(&[_]Feature{}),
1119 };
1120 result[@intFromEnum(Feature.max_private_element_size_4)] = .{
1121 .llvm_name = "max-private-element-size-4",
1122 .description = "Maximum private access size may be 4",
1123 .dependencies = featureSet(&[_]Feature{}),
1124 };
1125 result[@intFromEnum(Feature.max_private_element_size_8)] = .{
1126 .llvm_name = "max-private-element-size-8",
1127 .description = "Maximum private access size may be 8",
1128 .dependencies = featureSet(&[_]Feature{}),
1129 };
1130 result[@intFromEnum(Feature.memory_atomic_fadd_f32_denormal_support)] = .{
1131 .llvm_name = "memory-atomic-fadd-f32-denormal-support",
1132 .description = "global/flat/buffer atomic fadd for float supports denormal handling",
1133 .dependencies = featureSet(&[_]Feature{}),
1134 };
1135 result[@intFromEnum(Feature.mfma_inline_literal_bug)] = .{
1136 .llvm_name = "mfma-inline-literal-bug",
1137 .description = "MFMA cannot use inline literal as SrcC",
1138 .dependencies = featureSet(&[_]Feature{}),
1139 };
1140 result[@intFromEnum(Feature.mimg_r128)] = .{
1141 .llvm_name = "mimg-r128",
1142 .description = "Support 128-bit texture resources",
1143 .dependencies = featureSet(&[_]Feature{}),
1144 };
1145 result[@intFromEnum(Feature.minimum3_maximum3_f16)] = .{
1146 .llvm_name = "minimum3-maximum3-f16",
1147 .description = "Has v_minimum3_f16 and v_maximum3_f16 instructions",
1148 .dependencies = featureSet(&[_]Feature{}),
1149 };
1150 result[@intFromEnum(Feature.minimum3_maximum3_f32)] = .{
1151 .llvm_name = "minimum3-maximum3-f32",
1152 .description = "Has v_minimum3_f32 and v_maximum3_f32 instructions",
1153 .dependencies = featureSet(&[_]Feature{}),
1154 };
1155 result[@intFromEnum(Feature.minimum3_maximum3_pkf16)] = .{
1156 .llvm_name = "minimum3-maximum3-pkf16",
1157 .description = "Has v_pk_minimum3_f16 and v_pk_maximum3_f16 instructions",
1158 .dependencies = featureSet(&[_]Feature{}),
1159 };
1160 result[@intFromEnum(Feature.movrel)] = .{
1161 .llvm_name = "movrel",
1162 .description = "Has v_movrel*_b32 instructions",
1163 .dependencies = featureSet(&[_]Feature{}),
1164 };
1165 result[@intFromEnum(Feature.msaa_load_dst_sel_bug)] = .{
1166 .llvm_name = "msaa-load-dst-sel-bug",
1167 .description = "MSAA loads not honoring dst_sel bug",
1168 .dependencies = featureSet(&[_]Feature{}),
1169 };
1170 result[@intFromEnum(Feature.negative_scratch_offset_bug)] = .{
1171 .llvm_name = "negative-scratch-offset-bug",
1172 .description = "Negative immediate offsets in scratch instructions with an SGPR offset page fault on GFX9",
1173 .dependencies = featureSet(&[_]Feature{}),
1174 };
1175 result[@intFromEnum(Feature.negative_unaligned_scratch_offset_bug)] = .{
1176 .llvm_name = "negative-unaligned-scratch-offset-bug",
1177 .description = "Scratch instructions with a VGPR offset and a negative immediate offset that is not a multiple of 4 read wrong memory on GFX10",
1178 .dependencies = featureSet(&[_]Feature{}),
1179 };
1180 result[@intFromEnum(Feature.no_data_dep_hazard)] = .{
1181 .llvm_name = "no-data-dep-hazard",
1182 .description = "Does not need SW waitstates",
1183 .dependencies = featureSet(&[_]Feature{}),
1184 };
1185 result[@intFromEnum(Feature.no_sdst_cmpx)] = .{
1186 .llvm_name = "no-sdst-cmpx",
1187 .description = "V_CMPX does not write VCC/SGPR in addition to EXEC",
1188 .dependencies = featureSet(&[_]Feature{}),
1189 };
1190 result[@intFromEnum(Feature.nsa_clause_bug)] = .{
1191 .llvm_name = "nsa-clause-bug",
1192 .description = "MIMG-NSA in a hard clause has unpredictable results on GFX10.1",
1193 .dependencies = featureSet(&[_]Feature{}),
1194 };
1195 result[@intFromEnum(Feature.nsa_encoding)] = .{
1196 .llvm_name = "nsa-encoding",
1197 .description = "Support NSA encoding for image instructions",
1198 .dependencies = featureSet(&[_]Feature{}),
1199 };
1200 result[@intFromEnum(Feature.nsa_to_vmem_bug)] = .{
1201 .llvm_name = "nsa-to-vmem-bug",
1202 .description = "MIMG-NSA followed by VMEM fail if EXEC_LO or EXEC_HI equals zero",
1203 .dependencies = featureSet(&[_]Feature{}),
1204 };
1205 result[@intFromEnum(Feature.offset_3f_bug)] = .{
1206 .llvm_name = "offset-3f-bug",
1207 .description = "Branch offset of 3f hardware bug",
1208 .dependencies = featureSet(&[_]Feature{}),
1209 };
1210 result[@intFromEnum(Feature.packed_fp32_ops)] = .{
1211 .llvm_name = "packed-fp32-ops",
1212 .description = "Support packed fp32 instructions",
1213 .dependencies = featureSet(&[_]Feature{}),
1214 };
1215 result[@intFromEnum(Feature.packed_tid)] = .{
1216 .llvm_name = "packed-tid",
1217 .description = "Workitem IDs are packed into v0 at kernel launch",
1218 .dependencies = featureSet(&[_]Feature{}),
1219 };
1220 result[@intFromEnum(Feature.partial_nsa_encoding)] = .{
1221 .llvm_name = "partial-nsa-encoding",
1222 .description = "Support partial NSA encoding for image instructions",
1223 .dependencies = featureSet(&[_]Feature{}),
1224 };
1225 result[@intFromEnum(Feature.permlane16_swap)] = .{
1226 .llvm_name = "permlane16-swap",
1227 .description = "Has v_permlane16_swap_b32 instructions",
1228 .dependencies = featureSet(&[_]Feature{}),
1229 };
1230 result[@intFromEnum(Feature.permlane32_swap)] = .{
1231 .llvm_name = "permlane32-swap",
1232 .description = "Has v_permlane32_swap_b32 instructions",
1233 .dependencies = featureSet(&[_]Feature{}),
1234 };
1235 result[@intFromEnum(Feature.pk_fmac_f16_inst)] = .{
1236 .llvm_name = "pk-fmac-f16-inst",
1237 .description = "Has v_pk_fmac_f16 instruction",
1238 .dependencies = featureSet(&[_]Feature{}),
1239 };
1240 result[@intFromEnum(Feature.point_sample_accel)] = .{
1241 .llvm_name = "point-sample-accel",
1242 .description = "Has point sample acceleration feature",
1243 .dependencies = featureSet(&[_]Feature{}),
1244 };
1245 result[@intFromEnum(Feature.precise_memory)] = .{
1246 .llvm_name = "precise-memory",
1247 .description = "Enable precise memory mode",
1248 .dependencies = featureSet(&[_]Feature{}),
1249 };
1250 result[@intFromEnum(Feature.priv_enabled_trap2_nop_bug)] = .{
1251 .llvm_name = "priv-enabled-trap2-nop-bug",
1252 .description = "Hardware that runs with PRIV=1 interpreting 's_trap 2' as a nop bug",
1253 .dependencies = featureSet(&[_]Feature{}),
1254 };
1255 result[@intFromEnum(Feature.prng_inst)] = .{
1256 .llvm_name = "prng-inst",
1257 .description = "Has v_prng_b32 instruction",
1258 .dependencies = featureSet(&[_]Feature{}),
1259 };
1260 result[@intFromEnum(Feature.promote_alloca)] = .{
1261 .llvm_name = "promote-alloca",
1262 .description = "Enable promote alloca pass",
1263 .dependencies = featureSet(&[_]Feature{}),
1264 };
1265 result[@intFromEnum(Feature.prt_strict_null)] = .{
1266 .llvm_name = "enable-prt-strict-null",
1267 .description = "Enable zeroing of result registers for sparse texture fetches",
1268 .dependencies = featureSet(&[_]Feature{}),
1269 };
1270 result[@intFromEnum(Feature.pseudo_scalar_trans)] = .{
1271 .llvm_name = "pseudo-scalar-trans",
1272 .description = "Has Pseudo Scalar Transcendental instructions",
1273 .dependencies = featureSet(&[_]Feature{}),
1274 };
1275 result[@intFromEnum(Feature.r128_a16)] = .{
1276 .llvm_name = "r128-a16",
1277 .description = "Support gfx9-style A16 for 16-bit coordinates/gradients/lod/clamp/mip image operands, where a16 is aliased with r128",
1278 .dependencies = featureSet(&[_]Feature{}),
1279 };
1280 result[@intFromEnum(Feature.real_true16)] = .{
1281 .llvm_name = "real-true16",
1282 .description = "Use true 16-bit registers",
1283 .dependencies = featureSet(&[_]Feature{}),
1284 };
1285 result[@intFromEnum(Feature.relaxed_buffer_oob_mode)] = .{
1286 .llvm_name = "relaxed-buffer-oob-mode",
1287 .description = "Disable strict out-of-bounds buffer guarantees. An OOB access may potentially cause an adjacent access to be treated as if it were also OOB",
1288 .dependencies = featureSet(&[_]Feature{}),
1289 };
1290 result[@intFromEnum(Feature.required_export_priority)] = .{
1291 .llvm_name = "required-export-priority",
1292 .description = "Export priority must be explicitly manipulated on GFX11.5",
1293 .dependencies = featureSet(&[_]Feature{}),
1294 };
1295 result[@intFromEnum(Feature.requires_cov6)] = .{
1296 .llvm_name = "requires-cov6",
1297 .description = "Target Requires Code Object V6",
1298 .dependencies = featureSet(&[_]Feature{}),
1299 };
1300 result[@intFromEnum(Feature.restricted_soffset)] = .{
1301 .llvm_name = "restricted-soffset",
1302 .description = "Has restricted SOffset (immediate not supported).",
1303 .dependencies = featureSet(&[_]Feature{}),
1304 };
1305 result[@intFromEnum(Feature.s_memrealtime)] = .{
1306 .llvm_name = "s-memrealtime",
1307 .description = "Has s_memrealtime instruction",
1308 .dependencies = featureSet(&[_]Feature{}),
1309 };
1310 result[@intFromEnum(Feature.s_memtime_inst)] = .{
1311 .llvm_name = "s-memtime-inst",
1312 .description = "Has s_memtime instruction",
1313 .dependencies = featureSet(&[_]Feature{}),
1314 };
1315 result[@intFromEnum(Feature.safe_smem_prefetch)] = .{
1316 .llvm_name = "safe-smem-prefetch",
1317 .description = "SMEM prefetches do not fail on illegal address",
1318 .dependencies = featureSet(&[_]Feature{}),
1319 };
1320 result[@intFromEnum(Feature.salu_float)] = .{
1321 .llvm_name = "salu-float",
1322 .description = "Has SALU floating point instructions",
1323 .dependencies = featureSet(&[_]Feature{}),
1324 };
1325 result[@intFromEnum(Feature.scalar_atomics)] = .{
1326 .llvm_name = "scalar-atomics",
1327 .description = "Has atomic scalar memory instructions",
1328 .dependencies = featureSet(&[_]Feature{}),
1329 };
1330 result[@intFromEnum(Feature.scalar_dwordx3_loads)] = .{
1331 .llvm_name = "scalar-dwordx3-loads",
1332 .description = "Has 96-bit scalar load instructions",
1333 .dependencies = featureSet(&[_]Feature{}),
1334 };
1335 result[@intFromEnum(Feature.scalar_flat_scratch_insts)] = .{
1336 .llvm_name = "scalar-flat-scratch-insts",
1337 .description = "Have s_scratch_* flat memory instructions",
1338 .dependencies = featureSet(&[_]Feature{}),
1339 };
1340 result[@intFromEnum(Feature.scalar_stores)] = .{
1341 .llvm_name = "scalar-stores",
1342 .description = "Has store scalar memory instructions",
1343 .dependencies = featureSet(&[_]Feature{}),
1344 };
1345 result[@intFromEnum(Feature.sdwa)] = .{
1346 .llvm_name = "sdwa",
1347 .description = "Support SDWA (Sub-DWORD Addressing) extension",
1348 .dependencies = featureSet(&[_]Feature{}),
1349 };
1350 result[@intFromEnum(Feature.sdwa_mav)] = .{
1351 .llvm_name = "sdwa-mav",
1352 .description = "Support v_mac_f32/f16 with SDWA (Sub-DWORD Addressing) extension",
1353 .dependencies = featureSet(&[_]Feature{}),
1354 };
1355 result[@intFromEnum(Feature.sdwa_omod)] = .{
1356 .llvm_name = "sdwa-omod",
1357 .description = "Support OMod with SDWA (Sub-DWORD Addressing) extension",
1358 .dependencies = featureSet(&[_]Feature{}),
1359 };
1360 result[@intFromEnum(Feature.sdwa_out_mods_vopc)] = .{
1361 .llvm_name = "sdwa-out-mods-vopc",
1362 .description = "Support clamp for VOPC with SDWA (Sub-DWORD Addressing) extension",
1363 .dependencies = featureSet(&[_]Feature{}),
1364 };
1365 result[@intFromEnum(Feature.sdwa_scalar)] = .{
1366 .llvm_name = "sdwa-scalar",
1367 .description = "Support scalar register with SDWA (Sub-DWORD Addressing) extension",
1368 .dependencies = featureSet(&[_]Feature{}),
1369 };
1370 result[@intFromEnum(Feature.sdwa_sdst)] = .{
1371 .llvm_name = "sdwa-sdst",
1372 .description = "Support scalar dst for VOPC with SDWA (Sub-DWORD Addressing) extension",
1373 .dependencies = featureSet(&[_]Feature{}),
1374 };
1375 result[@intFromEnum(Feature.sea_islands)] = .{
1376 .llvm_name = "sea-islands",
1377 .description = "SEA_ISLANDS GPU generation",
1378 .dependencies = featureSet(&[_]Feature{
1379 .addressablelocalmemorysize65536,
1380 .atomic_fmin_fmax_flat_f32,
1381 .atomic_fmin_fmax_flat_f64,
1382 .atomic_fmin_fmax_global_f32,
1383 .atomic_fmin_fmax_global_f64,
1384 .ci_insts,
1385 .default_component_zero,
1386 .ds_src2_insts,
1387 .extended_image_insts,
1388 .flat_address_space,
1389 .fp64,
1390 .gds,
1391 .gfx7_gfx8_gfx9_insts,
1392 .gws,
1393 .image_insts,
1394 .mad_mac_f32_insts,
1395 .mimg_r128,
1396 .movrel,
1397 .s_memtime_inst,
1398 .trig_reduced_range,
1399 .unaligned_buffer_access,
1400 .vmem_write_vgpr_in_order,
1401 .wavefrontsize64,
1402 }),
1403 };
1404 result[@intFromEnum(Feature.setprio_inc_wg_inst)] = .{
1405 .llvm_name = "setprio-inc-wg-inst",
1406 .description = "Has s_setprio_inc_wg instruction.",
1407 .dependencies = featureSet(&[_]Feature{}),
1408 };
1409 result[@intFromEnum(Feature.sgpr_init_bug)] = .{
1410 .llvm_name = "sgpr-init-bug",
1411 .description = "VI SGPR initialization bug requiring a fixed SGPR allocation size",
1412 .dependencies = featureSet(&[_]Feature{}),
1413 };
1414 result[@intFromEnum(Feature.shader_cycles_hi_lo_registers)] = .{
1415 .llvm_name = "shader-cycles-hi-lo-registers",
1416 .description = "Has SHADER_CYCLES_HI/LO hardware registers",
1417 .dependencies = featureSet(&[_]Feature{}),
1418 };
1419 result[@intFromEnum(Feature.shader_cycles_register)] = .{
1420 .llvm_name = "shader-cycles-register",
1421 .description = "Has SHADER_CYCLES hardware register",
1422 .dependencies = featureSet(&[_]Feature{}),
1423 };
1424 result[@intFromEnum(Feature.si_scheduler)] = .{
1425 .llvm_name = "si-scheduler",
1426 .description = "Enable SI Machine Scheduler",
1427 .dependencies = featureSet(&[_]Feature{}),
1428 };
1429 result[@intFromEnum(Feature.smem_to_vector_write_hazard)] = .{
1430 .llvm_name = "smem-to-vector-write-hazard",
1431 .description = "s_load_dword followed by v_cmp page faults",
1432 .dependencies = featureSet(&[_]Feature{}),
1433 };
1434 result[@intFromEnum(Feature.southern_islands)] = .{
1435 .llvm_name = "southern-islands",
1436 .description = "SOUTHERN_ISLANDS GPU generation",
1437 .dependencies = featureSet(&[_]Feature{
1438 .addressablelocalmemorysize32768,
1439 .atomic_fmin_fmax_global_f32,
1440 .atomic_fmin_fmax_global_f64,
1441 .default_component_zero,
1442 .ds_src2_insts,
1443 .extended_image_insts,
1444 .fp64,
1445 .gds,
1446 .gws,
1447 .image_insts,
1448 .ldsbankcount32,
1449 .mad_mac_f32_insts,
1450 .mimg_r128,
1451 .movrel,
1452 .s_memtime_inst,
1453 .trig_reduced_range,
1454 .vmem_write_vgpr_in_order,
1455 .wavefrontsize64,
1456 }),
1457 };
1458 result[@intFromEnum(Feature.sramecc)] = .{
1459 .llvm_name = "sramecc",
1460 .description = "Enable SRAMECC",
1461 .dependencies = featureSet(&[_]Feature{}),
1462 };
1463 result[@intFromEnum(Feature.sramecc_support)] = .{
1464 .llvm_name = "sramecc-support",
1465 .description = "Hardware supports SRAMECC",
1466 .dependencies = featureSet(&[_]Feature{}),
1467 };
1468 result[@intFromEnum(Feature.tgsplit)] = .{
1469 .llvm_name = "tgsplit",
1470 .description = "Enable threadgroup split execution",
1471 .dependencies = featureSet(&[_]Feature{}),
1472 };
1473 result[@intFromEnum(Feature.transpose_load_f4f6_insts)] = .{
1474 .llvm_name = "transpose-load-f4f6-insts",
1475 .description = "Has ds_load_tr4/tr6 and global_load_tr4/tr6 instructions",
1476 .dependencies = featureSet(&[_]Feature{}),
1477 };
1478 result[@intFromEnum(Feature.trap_handler)] = .{
1479 .llvm_name = "trap-handler",
1480 .description = "Trap handler support",
1481 .dependencies = featureSet(&[_]Feature{}),
1482 };
1483 result[@intFromEnum(Feature.trig_reduced_range)] = .{
1484 .llvm_name = "trig-reduced-range",
1485 .description = "Requires use of fract on arguments to trig instructions",
1486 .dependencies = featureSet(&[_]Feature{}),
1487 };
1488 result[@intFromEnum(Feature.true16)] = .{
1489 .llvm_name = "true16",
1490 .description = "True 16-bit operand instructions",
1491 .dependencies = featureSet(&[_]Feature{}),
1492 };
1493 result[@intFromEnum(Feature.unaligned_access_mode)] = .{
1494 .llvm_name = "unaligned-access-mode",
1495 .description = "Enable unaligned global, local and region loads and stores if the hardware supports it",
1496 .dependencies = featureSet(&[_]Feature{}),
1497 };
1498 result[@intFromEnum(Feature.unaligned_buffer_access)] = .{
1499 .llvm_name = "unaligned-buffer-access",
1500 .description = "Hardware supports unaligned global loads and stores",
1501 .dependencies = featureSet(&[_]Feature{}),
1502 };
1503 result[@intFromEnum(Feature.unaligned_ds_access)] = .{
1504 .llvm_name = "unaligned-ds-access",
1505 .description = "Hardware supports unaligned local and region loads and stores",
1506 .dependencies = featureSet(&[_]Feature{}),
1507 };
1508 result[@intFromEnum(Feature.unaligned_scratch_access)] = .{
1509 .llvm_name = "unaligned-scratch-access",
1510 .description = "Support unaligned scratch loads and stores",
1511 .dependencies = featureSet(&[_]Feature{}),
1512 };
1513 result[@intFromEnum(Feature.unpacked_d16_vmem)] = .{
1514 .llvm_name = "unpacked-d16-vmem",
1515 .description = "Has unpacked d16 vmem instructions",
1516 .dependencies = featureSet(&[_]Feature{}),
1517 };
1518 result[@intFromEnum(Feature.unsafe_ds_offset_folding)] = .{
1519 .llvm_name = "unsafe-ds-offset-folding",
1520 .description = "Force using DS instruction immediate offsets on SI",
1521 .dependencies = featureSet(&[_]Feature{}),
1522 };
1523 result[@intFromEnum(Feature.user_sgpr_init16_bug)] = .{
1524 .llvm_name = "user-sgpr-init16-bug",
1525 .description = "Bug requiring at least 16 user+system SGPRs to be enabled",
1526 .dependencies = featureSet(&[_]Feature{}),
1527 };
1528 result[@intFromEnum(Feature.valu_trans_use_hazard)] = .{
1529 .llvm_name = "valu-trans-use-hazard",
1530 .description = "Hazard when TRANS instructions are closely followed by a use of the result",
1531 .dependencies = featureSet(&[_]Feature{}),
1532 };
1533 result[@intFromEnum(Feature.vcmpx_exec_war_hazard)] = .{
1534 .llvm_name = "vcmpx-exec-war-hazard",
1535 .description = "V_CMPX WAR hazard on EXEC (V_CMPX issue ONLY)",
1536 .dependencies = featureSet(&[_]Feature{}),
1537 };
1538 result[@intFromEnum(Feature.vcmpx_permlane_hazard)] = .{
1539 .llvm_name = "vcmpx-permlane-hazard",
1540 .description = "TODO: describe me",
1541 .dependencies = featureSet(&[_]Feature{}),
1542 };
1543 result[@intFromEnum(Feature.vgpr_index_mode)] = .{
1544 .llvm_name = "vgpr-index-mode",
1545 .description = "Has VGPR mode register indexing",
1546 .dependencies = featureSet(&[_]Feature{}),
1547 };
1548 result[@intFromEnum(Feature.vmem_to_lds_load_insts)] = .{
1549 .llvm_name = "vmem-to-lds-load-insts",
1550 .description = "The platform has memory to lds instructions (global_load w/lds bit set, buffer_load w/lds bit set or global_load_lds. This does not include scratch_load_lds.",
1551 .dependencies = featureSet(&[_]Feature{}),
1552 };
1553 result[@intFromEnum(Feature.vmem_to_scalar_write_hazard)] = .{
1554 .llvm_name = "vmem-to-scalar-write-hazard",
1555 .description = "VMEM instruction followed by scalar writing to EXEC mask, M0 or SGPR leads to incorrect execution.",
1556 .dependencies = featureSet(&[_]Feature{}),
1557 };
1558 result[@intFromEnum(Feature.vmem_write_vgpr_in_order)] = .{
1559 .llvm_name = "vmem-write-vgpr-in-order",
1560 .description = "VMEM instructions of the same type write VGPR results in order",
1561 .dependencies = featureSet(&[_]Feature{}),
1562 };
1563 result[@intFromEnum(Feature.volcanic_islands)] = .{
1564 .llvm_name = "volcanic-islands",
1565 .description = "VOLCANIC_ISLANDS GPU generation",
1566 .dependencies = featureSet(&[_]Feature{
1567 .@"16_bit_insts",
1568 .addressablelocalmemorysize65536,
1569 .ci_insts,
1570 .default_component_zero,
1571 .dpp,
1572 .ds_src2_insts,
1573 .extended_image_insts,
1574 .fast_denormal_f32,
1575 .flat_address_space,
1576 .fp64,
1577 .gcn3_encoding,
1578 .gds,
1579 .gfx7_gfx8_gfx9_insts,
1580 .gfx8_insts,
1581 .gws,
1582 .image_insts,
1583 .int_clamp_insts,
1584 .inv_2pi_inline_imm,
1585 .mad_mac_f32_insts,
1586 .mimg_r128,
1587 .movrel,
1588 .s_memrealtime,
1589 .s_memtime_inst,
1590 .scalar_stores,
1591 .sdwa,
1592 .sdwa_mav,
1593 .sdwa_out_mods_vopc,
1594 .trig_reduced_range,
1595 .unaligned_buffer_access,
1596 .vgpr_index_mode,
1597 .vmem_write_vgpr_in_order,
1598 .wavefrontsize64,
1599 }),
1600 };
1601 result[@intFromEnum(Feature.vop3_literal)] = .{
1602 .llvm_name = "vop3-literal",
1603 .description = "Can use one literal in VOP3",
1604 .dependencies = featureSet(&[_]Feature{}),
1605 };
1606 result[@intFromEnum(Feature.vop3p)] = .{
1607 .llvm_name = "vop3p",
1608 .description = "Has VOP3P packed instructions",
1609 .dependencies = featureSet(&[_]Feature{}),
1610 };
1611 result[@intFromEnum(Feature.vopd)] = .{
1612 .llvm_name = "vopd",
1613 .description = "Has VOPD dual issue wave32 instructions",
1614 .dependencies = featureSet(&[_]Feature{}),
1615 };
1616 result[@intFromEnum(Feature.vscnt)] = .{
1617 .llvm_name = "vscnt",
1618 .description = "Has separate store vscnt counter",
1619 .dependencies = featureSet(&[_]Feature{}),
1620 };
1621 result[@intFromEnum(Feature.wait_xcnt)] = .{
1622 .llvm_name = "wait-xcnt",
1623 .description = "Has s_wait_xcnt instruction",
1624 .dependencies = featureSet(&[_]Feature{}),
1625 };
1626 result[@intFromEnum(Feature.wavefrontsize16)] = .{
1627 .llvm_name = "wavefrontsize16",
1628 .description = "The number of threads per wavefront",
1629 .dependencies = featureSet(&[_]Feature{}),
1630 };
1631 result[@intFromEnum(Feature.wavefrontsize32)] = .{
1632 .llvm_name = "wavefrontsize32",
1633 .description = "The number of threads per wavefront",
1634 .dependencies = featureSet(&[_]Feature{}),
1635 };
1636 result[@intFromEnum(Feature.wavefrontsize64)] = .{
1637 .llvm_name = "wavefrontsize64",
1638 .description = "The number of threads per wavefront",
1639 .dependencies = featureSet(&[_]Feature{}),
1640 };
1641 result[@intFromEnum(Feature.xf32_insts)] = .{
1642 .llvm_name = "xf32-insts",
1643 .description = "Has instructions that support xf32 format, such as v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32",
1644 .dependencies = featureSet(&[_]Feature{}),
1645 };
1646 result[@intFromEnum(Feature.xnack)] = .{
1647 .llvm_name = "xnack",
1648 .description = "Enable XNACK support",
1649 .dependencies = featureSet(&[_]Feature{}),
1650 };
1651 result[@intFromEnum(Feature.xnack_support)] = .{
1652 .llvm_name = "xnack-support",
1653 .description = "Hardware supports XNACK",
1654 .dependencies = featureSet(&[_]Feature{}),
1655 };
1656 const ti = @typeInfo(Feature);
1657 for (&result, 0..) |*elem, i| {
1658 elem.index = i;
1659 elem.name = ti.@"enum".fields[i].name;
1660 }
1661 break :blk result;
1662};
1663
1664pub const cpu = struct {
1665 pub const bonaire: CpuModel = .{
1666 .name = "bonaire",
1667 .llvm_name = "bonaire",
1668 .features = featureSet(&[_]Feature{
1669 .ldsbankcount32,
1670 .sea_islands,
1671 }),
1672 };
1673 pub const carrizo: CpuModel = .{
1674 .name = "carrizo",
1675 .llvm_name = "carrizo",
1676 .features = featureSet(&[_]Feature{
1677 .fast_fmaf,
1678 .half_rate_64_ops,
1679 .ldsbankcount32,
1680 .unpacked_d16_vmem,
1681 .volcanic_islands,
1682 .xnack_support,
1683 }),
1684 };
1685 pub const fiji: CpuModel = .{
1686 .name = "fiji",
1687 .llvm_name = "fiji",
1688 .features = featureSet(&[_]Feature{
1689 .ldsbankcount32,
1690 .unpacked_d16_vmem,
1691 .volcanic_islands,
1692 }),
1693 };
1694 pub const generic: CpuModel = .{
1695 .name = "generic",
1696 .llvm_name = "generic",
1697 .features = featureSet(&[_]Feature{}),
1698 };
1699 pub const generic_hsa: CpuModel = .{
1700 .name = "generic_hsa",
1701 .llvm_name = "generic-hsa",
1702 .features = featureSet(&[_]Feature{
1703 .flat_address_space,
1704 }),
1705 };
1706 pub const gfx1010: CpuModel = .{
1707 .name = "gfx1010",
1708 .llvm_name = "gfx1010",
1709 .features = featureSet(&[_]Feature{
1710 .back_off_barrier,
1711 .dl_insts,
1712 .ds_src2_insts,
1713 .flat_segment_offset_bug,
1714 .get_wave_id_inst,
1715 .gfx10,
1716 .inst_fwd_prefetch_bug,
1717 .lds_branch_vmem_war_hazard,
1718 .lds_misaligned_bug,
1719 .ldsbankcount32,
1720 .mad_mac_f32_insts,
1721 .negative_unaligned_scratch_offset_bug,
1722 .nsa_clause_bug,
1723 .nsa_encoding,
1724 .nsa_to_vmem_bug,
1725 .offset_3f_bug,
1726 .scalar_atomics,
1727 .scalar_flat_scratch_insts,
1728 .scalar_stores,
1729 .smem_to_vector_write_hazard,
1730 .vcmpx_exec_war_hazard,
1731 .vcmpx_permlane_hazard,
1732 .vmem_to_scalar_write_hazard,
1733 .xnack_support,
1734 }),
1735 };
1736 pub const gfx1011: CpuModel = .{
1737 .name = "gfx1011",
1738 .llvm_name = "gfx1011",
1739 .features = featureSet(&[_]Feature{
1740 .back_off_barrier,
1741 .dl_insts,
1742 .dot10_insts,
1743 .dot1_insts,
1744 .dot2_insts,
1745 .dot5_insts,
1746 .dot6_insts,
1747 .dot7_insts,
1748 .ds_src2_insts,
1749 .flat_segment_offset_bug,
1750 .get_wave_id_inst,
1751 .gfx10,
1752 .inst_fwd_prefetch_bug,
1753 .lds_branch_vmem_war_hazard,
1754 .lds_misaligned_bug,
1755 .ldsbankcount32,
1756 .mad_mac_f32_insts,
1757 .negative_unaligned_scratch_offset_bug,
1758 .nsa_clause_bug,
1759 .nsa_encoding,
1760 .nsa_to_vmem_bug,
1761 .offset_3f_bug,
1762 .scalar_atomics,
1763 .scalar_flat_scratch_insts,
1764 .scalar_stores,
1765 .smem_to_vector_write_hazard,
1766 .vcmpx_exec_war_hazard,
1767 .vcmpx_permlane_hazard,
1768 .vmem_to_scalar_write_hazard,
1769 .xnack_support,
1770 }),
1771 };
1772 pub const gfx1012: CpuModel = .{
1773 .name = "gfx1012",
1774 .llvm_name = "gfx1012",
1775 .features = featureSet(&[_]Feature{
1776 .back_off_barrier,
1777 .dl_insts,
1778 .dot10_insts,
1779 .dot1_insts,
1780 .dot2_insts,
1781 .dot5_insts,
1782 .dot6_insts,
1783 .dot7_insts,
1784 .ds_src2_insts,
1785 .flat_segment_offset_bug,
1786 .get_wave_id_inst,
1787 .gfx10,
1788 .inst_fwd_prefetch_bug,
1789 .lds_branch_vmem_war_hazard,
1790 .lds_misaligned_bug,
1791 .ldsbankcount32,
1792 .mad_mac_f32_insts,
1793 .negative_unaligned_scratch_offset_bug,
1794 .nsa_clause_bug,
1795 .nsa_encoding,
1796 .nsa_to_vmem_bug,
1797 .offset_3f_bug,
1798 .scalar_atomics,
1799 .scalar_flat_scratch_insts,
1800 .scalar_stores,
1801 .smem_to_vector_write_hazard,
1802 .vcmpx_exec_war_hazard,
1803 .vcmpx_permlane_hazard,
1804 .vmem_to_scalar_write_hazard,
1805 .xnack_support,
1806 }),
1807 };
1808 pub const gfx1013: CpuModel = .{
1809 .name = "gfx1013",
1810 .llvm_name = "gfx1013",
1811 .features = featureSet(&[_]Feature{
1812 .back_off_barrier,
1813 .dl_insts,
1814 .ds_src2_insts,
1815 .flat_segment_offset_bug,
1816 .get_wave_id_inst,
1817 .gfx10,
1818 .gfx10_a_encoding,
1819 .inst_fwd_prefetch_bug,
1820 .lds_branch_vmem_war_hazard,
1821 .lds_misaligned_bug,
1822 .ldsbankcount32,
1823 .mad_mac_f32_insts,
1824 .negative_unaligned_scratch_offset_bug,
1825 .nsa_clause_bug,
1826 .nsa_encoding,
1827 .nsa_to_vmem_bug,
1828 .offset_3f_bug,
1829 .scalar_atomics,
1830 .scalar_flat_scratch_insts,
1831 .scalar_stores,
1832 .smem_to_vector_write_hazard,
1833 .vcmpx_exec_war_hazard,
1834 .vcmpx_permlane_hazard,
1835 .vmem_to_scalar_write_hazard,
1836 .xnack_support,
1837 }),
1838 };
1839 pub const gfx1030: CpuModel = .{
1840 .name = "gfx1030",
1841 .llvm_name = "gfx1030",
1842 .features = featureSet(&[_]Feature{
1843 .back_off_barrier,
1844 .dl_insts,
1845 .dot10_insts,
1846 .dot1_insts,
1847 .dot2_insts,
1848 .dot5_insts,
1849 .dot6_insts,
1850 .dot7_insts,
1851 .gfx10,
1852 .gfx10_3_insts,
1853 .gfx10_a_encoding,
1854 .gfx10_b_encoding,
1855 .ldsbankcount32,
1856 .nsa_encoding,
1857 .shader_cycles_register,
1858 }),
1859 };
1860 pub const gfx1031: CpuModel = .{
1861 .name = "gfx1031",
1862 .llvm_name = "gfx1031",
1863 .features = featureSet(&[_]Feature{
1864 .back_off_barrier,
1865 .dl_insts,
1866 .dot10_insts,
1867 .dot1_insts,
1868 .dot2_insts,
1869 .dot5_insts,
1870 .dot6_insts,
1871 .dot7_insts,
1872 .gfx10,
1873 .gfx10_3_insts,
1874 .gfx10_a_encoding,
1875 .gfx10_b_encoding,
1876 .ldsbankcount32,
1877 .nsa_encoding,
1878 .shader_cycles_register,
1879 }),
1880 };
1881 pub const gfx1032: CpuModel = .{
1882 .name = "gfx1032",
1883 .llvm_name = "gfx1032",
1884 .features = featureSet(&[_]Feature{
1885 .back_off_barrier,
1886 .dl_insts,
1887 .dot10_insts,
1888 .dot1_insts,
1889 .dot2_insts,
1890 .dot5_insts,
1891 .dot6_insts,
1892 .dot7_insts,
1893 .gfx10,
1894 .gfx10_3_insts,
1895 .gfx10_a_encoding,
1896 .gfx10_b_encoding,
1897 .ldsbankcount32,
1898 .nsa_encoding,
1899 .shader_cycles_register,
1900 }),
1901 };
1902 pub const gfx1033: CpuModel = .{
1903 .name = "gfx1033",
1904 .llvm_name = "gfx1033",
1905 .features = featureSet(&[_]Feature{
1906 .back_off_barrier,
1907 .dl_insts,
1908 .dot10_insts,
1909 .dot1_insts,
1910 .dot2_insts,
1911 .dot5_insts,
1912 .dot6_insts,
1913 .dot7_insts,
1914 .gfx10,
1915 .gfx10_3_insts,
1916 .gfx10_a_encoding,
1917 .gfx10_b_encoding,
1918 .ldsbankcount32,
1919 .nsa_encoding,
1920 .shader_cycles_register,
1921 }),
1922 };
1923 pub const gfx1034: CpuModel = .{
1924 .name = "gfx1034",
1925 .llvm_name = "gfx1034",
1926 .features = featureSet(&[_]Feature{
1927 .back_off_barrier,
1928 .dl_insts,
1929 .dot10_insts,
1930 .dot1_insts,
1931 .dot2_insts,
1932 .dot5_insts,
1933 .dot6_insts,
1934 .dot7_insts,
1935 .gfx10,
1936 .gfx10_3_insts,
1937 .gfx10_a_encoding,
1938 .gfx10_b_encoding,
1939 .ldsbankcount32,
1940 .nsa_encoding,
1941 .shader_cycles_register,
1942 }),
1943 };
1944 pub const gfx1035: CpuModel = .{
1945 .name = "gfx1035",
1946 .llvm_name = "gfx1035",
1947 .features = featureSet(&[_]Feature{
1948 .back_off_barrier,
1949 .dl_insts,
1950 .dot10_insts,
1951 .dot1_insts,
1952 .dot2_insts,
1953 .dot5_insts,
1954 .dot6_insts,
1955 .dot7_insts,
1956 .gfx10,
1957 .gfx10_3_insts,
1958 .gfx10_a_encoding,
1959 .gfx10_b_encoding,
1960 .ldsbankcount32,
1961 .nsa_encoding,
1962 .shader_cycles_register,
1963 }),
1964 };
1965 pub const gfx1036: CpuModel = .{
1966 .name = "gfx1036",
1967 .llvm_name = "gfx1036",
1968 .features = featureSet(&[_]Feature{
1969 .back_off_barrier,
1970 .dl_insts,
1971 .dot10_insts,
1972 .dot1_insts,
1973 .dot2_insts,
1974 .dot5_insts,
1975 .dot6_insts,
1976 .dot7_insts,
1977 .gfx10,
1978 .gfx10_3_insts,
1979 .gfx10_a_encoding,
1980 .gfx10_b_encoding,
1981 .ldsbankcount32,
1982 .nsa_encoding,
1983 .shader_cycles_register,
1984 }),
1985 };
1986 pub const gfx10_1_generic: CpuModel = .{
1987 .name = "gfx10_1_generic",
1988 .llvm_name = "gfx10-1-generic",
1989 .features = featureSet(&[_]Feature{
1990 .back_off_barrier,
1991 .dl_insts,
1992 .ds_src2_insts,
1993 .flat_segment_offset_bug,
1994 .get_wave_id_inst,
1995 .gfx10,
1996 .inst_fwd_prefetch_bug,
1997 .lds_branch_vmem_war_hazard,
1998 .lds_misaligned_bug,
1999 .ldsbankcount32,
2000 .mad_mac_f32_insts,
2001 .negative_unaligned_scratch_offset_bug,
2002 .nsa_clause_bug,
2003 .nsa_encoding,
2004 .nsa_to_vmem_bug,
2005 .offset_3f_bug,
2006 .requires_cov6,
2007 .scalar_atomics,
2008 .scalar_flat_scratch_insts,
2009 .scalar_stores,
2010 .smem_to_vector_write_hazard,
2011 .vcmpx_exec_war_hazard,
2012 .vcmpx_permlane_hazard,
2013 .vmem_to_scalar_write_hazard,
2014 .xnack_support,
2015 }),
2016 };
2017 pub const gfx10_3_generic: CpuModel = .{
2018 .name = "gfx10_3_generic",
2019 .llvm_name = "gfx10-3-generic",
2020 .features = featureSet(&[_]Feature{
2021 .back_off_barrier,
2022 .dl_insts,
2023 .dot10_insts,
2024 .dot1_insts,
2025 .dot2_insts,
2026 .dot5_insts,
2027 .dot6_insts,
2028 .dot7_insts,
2029 .gfx10,
2030 .gfx10_3_insts,
2031 .gfx10_a_encoding,
2032 .gfx10_b_encoding,
2033 .ldsbankcount32,
2034 .nsa_encoding,
2035 .requires_cov6,
2036 .shader_cycles_register,
2037 }),
2038 };
2039 pub const gfx1100: CpuModel = .{
2040 .name = "gfx1100",
2041 .llvm_name = "gfx1100",
2042 .features = featureSet(&[_]Feature{
2043 .allocate1_5xvgprs,
2044 .architected_flat_scratch,
2045 .atomic_fadd_no_rtn_insts,
2046 .atomic_fadd_rtn_insts,
2047 .dl_insts,
2048 .dot10_insts,
2049 .dot12_insts,
2050 .dot5_insts,
2051 .dot7_insts,
2052 .dot8_insts,
2053 .dot9_insts,
2054 .flat_atomic_fadd_f32_inst,
2055 .gfx11,
2056 .image_insts,
2057 .ldsbankcount32,
2058 .mad_intra_fwd_bug,
2059 .memory_atomic_fadd_f32_denormal_support,
2060 .msaa_load_dst_sel_bug,
2061 .nsa_encoding,
2062 .packed_tid,
2063 .partial_nsa_encoding,
2064 .priv_enabled_trap2_nop_bug,
2065 .real_true16,
2066 .shader_cycles_register,
2067 .user_sgpr_init16_bug,
2068 .valu_trans_use_hazard,
2069 .vcmpx_permlane_hazard,
2070 }),
2071 };
2072 pub const gfx1101: CpuModel = .{
2073 .name = "gfx1101",
2074 .llvm_name = "gfx1101",
2075 .features = featureSet(&[_]Feature{
2076 .allocate1_5xvgprs,
2077 .architected_flat_scratch,
2078 .atomic_fadd_no_rtn_insts,
2079 .atomic_fadd_rtn_insts,
2080 .dl_insts,
2081 .dot10_insts,
2082 .dot12_insts,
2083 .dot5_insts,
2084 .dot7_insts,
2085 .dot8_insts,
2086 .dot9_insts,
2087 .flat_atomic_fadd_f32_inst,
2088 .gfx11,
2089 .image_insts,
2090 .ldsbankcount32,
2091 .mad_intra_fwd_bug,
2092 .memory_atomic_fadd_f32_denormal_support,
2093 .msaa_load_dst_sel_bug,
2094 .nsa_encoding,
2095 .packed_tid,
2096 .partial_nsa_encoding,
2097 .priv_enabled_trap2_nop_bug,
2098 .real_true16,
2099 .shader_cycles_register,
2100 .valu_trans_use_hazard,
2101 .vcmpx_permlane_hazard,
2102 }),
2103 };
2104 pub const gfx1102: CpuModel = .{
2105 .name = "gfx1102",
2106 .llvm_name = "gfx1102",
2107 .features = featureSet(&[_]Feature{
2108 .architected_flat_scratch,
2109 .atomic_fadd_no_rtn_insts,
2110 .atomic_fadd_rtn_insts,
2111 .dl_insts,
2112 .dot10_insts,
2113 .dot12_insts,
2114 .dot5_insts,
2115 .dot7_insts,
2116 .dot8_insts,
2117 .dot9_insts,
2118 .flat_atomic_fadd_f32_inst,
2119 .gfx11,
2120 .image_insts,
2121 .ldsbankcount32,
2122 .mad_intra_fwd_bug,
2123 .memory_atomic_fadd_f32_denormal_support,
2124 .msaa_load_dst_sel_bug,
2125 .nsa_encoding,
2126 .packed_tid,
2127 .partial_nsa_encoding,
2128 .priv_enabled_trap2_nop_bug,
2129 .real_true16,
2130 .shader_cycles_register,
2131 .user_sgpr_init16_bug,
2132 .valu_trans_use_hazard,
2133 .vcmpx_permlane_hazard,
2134 }),
2135 };
2136 pub const gfx1103: CpuModel = .{
2137 .name = "gfx1103",
2138 .llvm_name = "gfx1103",
2139 .features = featureSet(&[_]Feature{
2140 .architected_flat_scratch,
2141 .atomic_fadd_no_rtn_insts,
2142 .atomic_fadd_rtn_insts,
2143 .dl_insts,
2144 .dot10_insts,
2145 .dot12_insts,
2146 .dot5_insts,
2147 .dot7_insts,
2148 .dot8_insts,
2149 .dot9_insts,
2150 .flat_atomic_fadd_f32_inst,
2151 .gfx11,
2152 .image_insts,
2153 .ldsbankcount32,
2154 .mad_intra_fwd_bug,
2155 .memory_atomic_fadd_f32_denormal_support,
2156 .msaa_load_dst_sel_bug,
2157 .nsa_encoding,
2158 .packed_tid,
2159 .partial_nsa_encoding,
2160 .priv_enabled_trap2_nop_bug,
2161 .real_true16,
2162 .shader_cycles_register,
2163 .valu_trans_use_hazard,
2164 .vcmpx_permlane_hazard,
2165 }),
2166 };
2167 pub const gfx1150: CpuModel = .{
2168 .name = "gfx1150",
2169 .llvm_name = "gfx1150",
2170 .features = featureSet(&[_]Feature{
2171 .architected_flat_scratch,
2172 .atomic_fadd_no_rtn_insts,
2173 .atomic_fadd_rtn_insts,
2174 .dl_insts,
2175 .dot10_insts,
2176 .dot12_insts,
2177 .dot5_insts,
2178 .dot7_insts,
2179 .dot8_insts,
2180 .dot9_insts,
2181 .dpp_src1_sgpr,
2182 .flat_atomic_fadd_f32_inst,
2183 .gfx11,
2184 .image_insts,
2185 .ldsbankcount32,
2186 .memory_atomic_fadd_f32_denormal_support,
2187 .nsa_encoding,
2188 .packed_tid,
2189 .partial_nsa_encoding,
2190 .point_sample_accel,
2191 .required_export_priority,
2192 .salu_float,
2193 .shader_cycles_register,
2194 .vcmpx_permlane_hazard,
2195 }),
2196 };
2197 pub const gfx1151: CpuModel = .{
2198 .name = "gfx1151",
2199 .llvm_name = "gfx1151",
2200 .features = featureSet(&[_]Feature{
2201 .allocate1_5xvgprs,
2202 .architected_flat_scratch,
2203 .atomic_fadd_no_rtn_insts,
2204 .atomic_fadd_rtn_insts,
2205 .dl_insts,
2206 .dot10_insts,
2207 .dot12_insts,
2208 .dot5_insts,
2209 .dot7_insts,
2210 .dot8_insts,
2211 .dot9_insts,
2212 .dpp_src1_sgpr,
2213 .flat_atomic_fadd_f32_inst,
2214 .gfx11,
2215 .image_insts,
2216 .ldsbankcount32,
2217 .memory_atomic_fadd_f32_denormal_support,
2218 .nsa_encoding,
2219 .packed_tid,
2220 .partial_nsa_encoding,
2221 .point_sample_accel,
2222 .required_export_priority,
2223 .salu_float,
2224 .shader_cycles_register,
2225 .vcmpx_permlane_hazard,
2226 }),
2227 };
2228 pub const gfx1152: CpuModel = .{
2229 .name = "gfx1152",
2230 .llvm_name = "gfx1152",
2231 .features = featureSet(&[_]Feature{
2232 .architected_flat_scratch,
2233 .atomic_fadd_no_rtn_insts,
2234 .atomic_fadd_rtn_insts,
2235 .dl_insts,
2236 .dot10_insts,
2237 .dot12_insts,
2238 .dot5_insts,
2239 .dot7_insts,
2240 .dot8_insts,
2241 .dot9_insts,
2242 .dpp_src1_sgpr,
2243 .flat_atomic_fadd_f32_inst,
2244 .gfx11,
2245 .image_insts,
2246 .ldsbankcount32,
2247 .memory_atomic_fadd_f32_denormal_support,
2248 .nsa_encoding,
2249 .packed_tid,
2250 .partial_nsa_encoding,
2251 .point_sample_accel,
2252 .required_export_priority,
2253 .salu_float,
2254 .shader_cycles_register,
2255 .vcmpx_permlane_hazard,
2256 }),
2257 };
2258 pub const gfx1153: CpuModel = .{
2259 .name = "gfx1153",
2260 .llvm_name = "gfx1153",
2261 .features = featureSet(&[_]Feature{
2262 .architected_flat_scratch,
2263 .atomic_fadd_no_rtn_insts,
2264 .atomic_fadd_rtn_insts,
2265 .dl_insts,
2266 .dot10_insts,
2267 .dot12_insts,
2268 .dot5_insts,
2269 .dot7_insts,
2270 .dot8_insts,
2271 .dot9_insts,
2272 .dpp_src1_sgpr,
2273 .flat_atomic_fadd_f32_inst,
2274 .gfx11,
2275 .image_insts,
2276 .ldsbankcount32,
2277 .memory_atomic_fadd_f32_denormal_support,
2278 .nsa_encoding,
2279 .packed_tid,
2280 .partial_nsa_encoding,
2281 .required_export_priority,
2282 .salu_float,
2283 .shader_cycles_register,
2284 .vcmpx_permlane_hazard,
2285 }),
2286 };
2287 pub const gfx11_generic: CpuModel = .{
2288 .name = "gfx11_generic",
2289 .llvm_name = "gfx11-generic",
2290 .features = featureSet(&[_]Feature{
2291 .architected_flat_scratch,
2292 .atomic_fadd_no_rtn_insts,
2293 .atomic_fadd_rtn_insts,
2294 .dl_insts,
2295 .dot10_insts,
2296 .dot12_insts,
2297 .dot5_insts,
2298 .dot7_insts,
2299 .dot8_insts,
2300 .dot9_insts,
2301 .flat_atomic_fadd_f32_inst,
2302 .gfx11,
2303 .image_insts,
2304 .ldsbankcount32,
2305 .mad_intra_fwd_bug,
2306 .memory_atomic_fadd_f32_denormal_support,
2307 .msaa_load_dst_sel_bug,
2308 .nsa_encoding,
2309 .packed_tid,
2310 .partial_nsa_encoding,
2311 .priv_enabled_trap2_nop_bug,
2312 .required_export_priority,
2313 .requires_cov6,
2314 .shader_cycles_register,
2315 .user_sgpr_init16_bug,
2316 .valu_trans_use_hazard,
2317 .vcmpx_permlane_hazard,
2318 }),
2319 };
2320 pub const gfx1200: CpuModel = .{
2321 .name = "gfx1200",
2322 .llvm_name = "gfx1200",
2323 .features = featureSet(&[_]Feature{
2324 .allocate1_5xvgprs,
2325 .architected_flat_scratch,
2326 .architected_sgprs,
2327 .atomic_buffer_global_pk_add_f16_insts,
2328 .atomic_buffer_pk_add_bf16_inst,
2329 .atomic_ds_pk_add_16_insts,
2330 .atomic_fadd_no_rtn_insts,
2331 .atomic_fadd_rtn_insts,
2332 .atomic_flat_pk_add_16_insts,
2333 .atomic_global_pk_add_bf16_inst,
2334 .bvh_dual_bvh_8_insts,
2335 .dl_insts,
2336 .dot10_insts,
2337 .dot11_insts,
2338 .dot12_insts,
2339 .dot7_insts,
2340 .dot8_insts,
2341 .dot9_insts,
2342 .dpp_src1_sgpr,
2343 .extended_image_insts,
2344 .flat_atomic_fadd_f32_inst,
2345 .fp8_conversion_insts,
2346 .gfx12,
2347 .image_insts,
2348 .ldsbankcount32,
2349 .memory_atomic_fadd_f32_denormal_support,
2350 .nsa_encoding,
2351 .packed_tid,
2352 .partial_nsa_encoding,
2353 .pseudo_scalar_trans,
2354 .restricted_soffset,
2355 .salu_float,
2356 .scalar_dwordx3_loads,
2357 .shader_cycles_hi_lo_registers,
2358 .vcmpx_permlane_hazard,
2359 }),
2360 };
2361 pub const gfx1201: CpuModel = .{
2362 .name = "gfx1201",
2363 .llvm_name = "gfx1201",
2364 .features = featureSet(&[_]Feature{
2365 .allocate1_5xvgprs,
2366 .architected_flat_scratch,
2367 .architected_sgprs,
2368 .atomic_buffer_global_pk_add_f16_insts,
2369 .atomic_buffer_pk_add_bf16_inst,
2370 .atomic_ds_pk_add_16_insts,
2371 .atomic_fadd_no_rtn_insts,
2372 .atomic_fadd_rtn_insts,
2373 .atomic_flat_pk_add_16_insts,
2374 .atomic_global_pk_add_bf16_inst,
2375 .bvh_dual_bvh_8_insts,
2376 .dl_insts,
2377 .dot10_insts,
2378 .dot11_insts,
2379 .dot12_insts,
2380 .dot7_insts,
2381 .dot8_insts,
2382 .dot9_insts,
2383 .dpp_src1_sgpr,
2384 .extended_image_insts,
2385 .flat_atomic_fadd_f32_inst,
2386 .fp8_conversion_insts,
2387 .gfx12,
2388 .image_insts,
2389 .ldsbankcount32,
2390 .memory_atomic_fadd_f32_denormal_support,
2391 .nsa_encoding,
2392 .packed_tid,
2393 .partial_nsa_encoding,
2394 .pseudo_scalar_trans,
2395 .restricted_soffset,
2396 .salu_float,
2397 .scalar_dwordx3_loads,
2398 .shader_cycles_hi_lo_registers,
2399 .vcmpx_permlane_hazard,
2400 }),
2401 };
2402 pub const gfx1250: CpuModel = .{
2403 .name = "gfx1250",
2404 .llvm_name = "gfx1250",
2405 .features = featureSet(&[_]Feature{
2406 .@"64_bit_literals",
2407 .architected_flat_scratch,
2408 .architected_sgprs,
2409 .ashr_pk_insts,
2410 .atomic_buffer_global_pk_add_f16_insts,
2411 .atomic_buffer_pk_add_bf16_inst,
2412 .atomic_ds_pk_add_16_insts,
2413 .atomic_fadd_no_rtn_insts,
2414 .atomic_fadd_rtn_insts,
2415 .atomic_flat_pk_add_16_insts,
2416 .atomic_fmin_fmax_flat_f64,
2417 .atomic_fmin_fmax_global_f64,
2418 .atomic_global_pk_add_bf16_inst,
2419 .bf16_cvt_insts,
2420 .bf16_trans_insts,
2421 .bitop3_insts,
2422 .cumode,
2423 .cvt_pk_f16_f32_inst,
2424 .dl_insts,
2425 .dot7_insts,
2426 .dot8_insts,
2427 .dpp_src1_sgpr,
2428 .flat_atomic_fadd_f32_inst,
2429 .flat_buffer_global_fadd_f64_inst,
2430 .fmacf64_inst,
2431 .fp8_conversion_insts,
2432 .fp8e5m3_insts,
2433 .gfx12,
2434 .gfx1250_insts,
2435 .kernarg_preload,
2436 .lds_barrier_arrive_atomic,
2437 .ldsbankcount32,
2438 .lshl_add_u64_inst,
2439 .max_hard_clause_length_63,
2440 .memory_atomic_fadd_f32_denormal_support,
2441 .minimum3_maximum3_pkf16,
2442 .packed_fp32_ops,
2443 .packed_tid,
2444 .permlane16_swap,
2445 .prng_inst,
2446 .pseudo_scalar_trans,
2447 .restricted_soffset,
2448 .salu_float,
2449 .scalar_dwordx3_loads,
2450 .setprio_inc_wg_inst,
2451 .shader_cycles_hi_lo_registers,
2452 .sramecc_support,
2453 .transpose_load_f4f6_insts,
2454 .vcmpx_permlane_hazard,
2455 .wait_xcnt,
2456 .wavefrontsize32,
2457 }),
2458 };
2459 pub const gfx12_generic: CpuModel = .{
2460 .name = "gfx12_generic",
2461 .llvm_name = "gfx12-generic",
2462 .features = featureSet(&[_]Feature{
2463 .allocate1_5xvgprs,
2464 .architected_flat_scratch,
2465 .architected_sgprs,
2466 .atomic_buffer_global_pk_add_f16_insts,
2467 .atomic_buffer_pk_add_bf16_inst,
2468 .atomic_ds_pk_add_16_insts,
2469 .atomic_fadd_no_rtn_insts,
2470 .atomic_fadd_rtn_insts,
2471 .atomic_flat_pk_add_16_insts,
2472 .atomic_global_pk_add_bf16_inst,
2473 .bvh_dual_bvh_8_insts,
2474 .dl_insts,
2475 .dot10_insts,
2476 .dot11_insts,
2477 .dot12_insts,
2478 .dot7_insts,
2479 .dot8_insts,
2480 .dot9_insts,
2481 .dpp_src1_sgpr,
2482 .extended_image_insts,
2483 .flat_atomic_fadd_f32_inst,
2484 .fp8_conversion_insts,
2485 .gfx12,
2486 .image_insts,
2487 .ldsbankcount32,
2488 .memory_atomic_fadd_f32_denormal_support,
2489 .nsa_encoding,
2490 .packed_tid,
2491 .partial_nsa_encoding,
2492 .pseudo_scalar_trans,
2493 .requires_cov6,
2494 .restricted_soffset,
2495 .salu_float,
2496 .scalar_dwordx3_loads,
2497 .shader_cycles_hi_lo_registers,
2498 .vcmpx_permlane_hazard,
2499 }),
2500 };
2501 pub const gfx600: CpuModel = .{
2502 .name = "gfx600",
2503 .llvm_name = "gfx600",
2504 .features = featureSet(&[_]Feature{
2505 .fast_fmaf,
2506 .half_rate_64_ops,
2507 .southern_islands,
2508 }),
2509 };
2510 pub const gfx601: CpuModel = .{
2511 .name = "gfx601",
2512 .llvm_name = "gfx601",
2513 .features = featureSet(&[_]Feature{
2514 .southern_islands,
2515 }),
2516 };
2517 pub const gfx602: CpuModel = .{
2518 .name = "gfx602",
2519 .llvm_name = "gfx602",
2520 .features = featureSet(&[_]Feature{
2521 .southern_islands,
2522 }),
2523 };
2524 pub const gfx700: CpuModel = .{
2525 .name = "gfx700",
2526 .llvm_name = "gfx700",
2527 .features = featureSet(&[_]Feature{
2528 .ldsbankcount32,
2529 .sea_islands,
2530 }),
2531 };
2532 pub const gfx701: CpuModel = .{
2533 .name = "gfx701",
2534 .llvm_name = "gfx701",
2535 .features = featureSet(&[_]Feature{
2536 .fast_fmaf,
2537 .half_rate_64_ops,
2538 .ldsbankcount32,
2539 .sea_islands,
2540 }),
2541 };
2542 pub const gfx702: CpuModel = .{
2543 .name = "gfx702",
2544 .llvm_name = "gfx702",
2545 .features = featureSet(&[_]Feature{
2546 .fast_fmaf,
2547 .ldsbankcount16,
2548 .sea_islands,
2549 }),
2550 };
2551 pub const gfx703: CpuModel = .{
2552 .name = "gfx703",
2553 .llvm_name = "gfx703",
2554 .features = featureSet(&[_]Feature{
2555 .ldsbankcount16,
2556 .sea_islands,
2557 }),
2558 };
2559 pub const gfx704: CpuModel = .{
2560 .name = "gfx704",
2561 .llvm_name = "gfx704",
2562 .features = featureSet(&[_]Feature{
2563 .ldsbankcount32,
2564 .sea_islands,
2565 }),
2566 };
2567 pub const gfx705: CpuModel = .{
2568 .name = "gfx705",
2569 .llvm_name = "gfx705",
2570 .features = featureSet(&[_]Feature{
2571 .ldsbankcount16,
2572 .sea_islands,
2573 }),
2574 };
2575 pub const gfx801: CpuModel = .{
2576 .name = "gfx801",
2577 .llvm_name = "gfx801",
2578 .features = featureSet(&[_]Feature{
2579 .fast_fmaf,
2580 .half_rate_64_ops,
2581 .ldsbankcount32,
2582 .unpacked_d16_vmem,
2583 .volcanic_islands,
2584 .xnack_support,
2585 }),
2586 };
2587 pub const gfx802: CpuModel = .{
2588 .name = "gfx802",
2589 .llvm_name = "gfx802",
2590 .features = featureSet(&[_]Feature{
2591 .ldsbankcount32,
2592 .sgpr_init_bug,
2593 .unpacked_d16_vmem,
2594 .volcanic_islands,
2595 }),
2596 };
2597 pub const gfx803: CpuModel = .{
2598 .name = "gfx803",
2599 .llvm_name = "gfx803",
2600 .features = featureSet(&[_]Feature{
2601 .ldsbankcount32,
2602 .unpacked_d16_vmem,
2603 .volcanic_islands,
2604 }),
2605 };
2606 pub const gfx805: CpuModel = .{
2607 .name = "gfx805",
2608 .llvm_name = "gfx805",
2609 .features = featureSet(&[_]Feature{
2610 .ldsbankcount32,
2611 .sgpr_init_bug,
2612 .unpacked_d16_vmem,
2613 .volcanic_islands,
2614 }),
2615 };
2616 pub const gfx810: CpuModel = .{
2617 .name = "gfx810",
2618 .llvm_name = "gfx810",
2619 .features = featureSet(&[_]Feature{
2620 .image_gather4_d16_bug,
2621 .image_store_d16_bug,
2622 .ldsbankcount16,
2623 .volcanic_islands,
2624 .xnack_support,
2625 }),
2626 };
2627 pub const gfx900: CpuModel = .{
2628 .name = "gfx900",
2629 .llvm_name = "gfx900",
2630 .features = featureSet(&[_]Feature{
2631 .addressablelocalmemorysize65536,
2632 .ds_src2_insts,
2633 .extended_image_insts,
2634 .gds,
2635 .gfx9,
2636 .image_gather4_d16_bug,
2637 .image_insts,
2638 .ldsbankcount32,
2639 .mad_mac_f32_insts,
2640 .mad_mix_insts,
2641 }),
2642 };
2643 pub const gfx902: CpuModel = .{
2644 .name = "gfx902",
2645 .llvm_name = "gfx902",
2646 .features = featureSet(&[_]Feature{
2647 .addressablelocalmemorysize65536,
2648 .ds_src2_insts,
2649 .extended_image_insts,
2650 .gds,
2651 .gfx9,
2652 .image_gather4_d16_bug,
2653 .image_insts,
2654 .ldsbankcount32,
2655 .mad_mac_f32_insts,
2656 .mad_mix_insts,
2657 }),
2658 };
2659 pub const gfx904: CpuModel = .{
2660 .name = "gfx904",
2661 .llvm_name = "gfx904",
2662 .features = featureSet(&[_]Feature{
2663 .addressablelocalmemorysize65536,
2664 .ds_src2_insts,
2665 .extended_image_insts,
2666 .fma_mix_insts,
2667 .gds,
2668 .gfx9,
2669 .image_gather4_d16_bug,
2670 .image_insts,
2671 .ldsbankcount32,
2672 .mad_mac_f32_insts,
2673 }),
2674 };
2675 pub const gfx906: CpuModel = .{
2676 .name = "gfx906",
2677 .llvm_name = "gfx906",
2678 .features = featureSet(&[_]Feature{
2679 .addressablelocalmemorysize65536,
2680 .dl_insts,
2681 .dot10_insts,
2682 .dot1_insts,
2683 .dot2_insts,
2684 .dot7_insts,
2685 .ds_src2_insts,
2686 .extended_image_insts,
2687 .fma_mix_insts,
2688 .gds,
2689 .gfx9,
2690 .half_rate_64_ops,
2691 .image_gather4_d16_bug,
2692 .image_insts,
2693 .ldsbankcount32,
2694 .mad_mac_f32_insts,
2695 .sramecc_support,
2696 }),
2697 };
2698 pub const gfx908: CpuModel = .{
2699 .name = "gfx908",
2700 .llvm_name = "gfx908",
2701 .features = featureSet(&[_]Feature{
2702 .addressablelocalmemorysize65536,
2703 .atomic_buffer_global_pk_add_f16_no_rtn_insts,
2704 .atomic_fadd_no_rtn_insts,
2705 .dl_insts,
2706 .dot10_insts,
2707 .dot1_insts,
2708 .dot2_insts,
2709 .dot3_insts,
2710 .dot4_insts,
2711 .dot5_insts,
2712 .dot6_insts,
2713 .dot7_insts,
2714 .ds_src2_insts,
2715 .extended_image_insts,
2716 .fma_mix_insts,
2717 .gds,
2718 .gfx9,
2719 .half_rate_64_ops,
2720 .image_gather4_d16_bug,
2721 .image_insts,
2722 .ldsbankcount32,
2723 .mad_mac_f32_insts,
2724 .mai_insts,
2725 .mfma_inline_literal_bug,
2726 .pk_fmac_f16_inst,
2727 .sramecc_support,
2728 }),
2729 };
2730 pub const gfx909: CpuModel = .{
2731 .name = "gfx909",
2732 .llvm_name = "gfx909",
2733 .features = featureSet(&[_]Feature{
2734 .addressablelocalmemorysize65536,
2735 .ds_src2_insts,
2736 .extended_image_insts,
2737 .gds,
2738 .gfx9,
2739 .image_gather4_d16_bug,
2740 .image_insts,
2741 .ldsbankcount32,
2742 .mad_mac_f32_insts,
2743 .mad_mix_insts,
2744 }),
2745 };
2746 pub const gfx90a: CpuModel = .{
2747 .name = "gfx90a",
2748 .llvm_name = "gfx90a",
2749 .features = featureSet(&[_]Feature{
2750 .addressablelocalmemorysize65536,
2751 .atomic_buffer_global_pk_add_f16_insts,
2752 .atomic_fadd_no_rtn_insts,
2753 .atomic_fadd_rtn_insts,
2754 .atomic_fmin_fmax_flat_f64,
2755 .atomic_fmin_fmax_global_f64,
2756 .back_off_barrier,
2757 .dl_insts,
2758 .dot10_insts,
2759 .dot1_insts,
2760 .dot2_insts,
2761 .dot3_insts,
2762 .dot4_insts,
2763 .dot5_insts,
2764 .dot6_insts,
2765 .dot7_insts,
2766 .dpp_64bit,
2767 .flat_buffer_global_fadd_f64_inst,
2768 .fma_mix_insts,
2769 .fmacf64_inst,
2770 .full_rate_64_ops,
2771 .gfx9,
2772 .gfx90a_insts,
2773 .image_insts,
2774 .kernarg_preload,
2775 .ldsbankcount32,
2776 .mad_mac_f32_insts,
2777 .mai_insts,
2778 .packed_fp32_ops,
2779 .packed_tid,
2780 .pk_fmac_f16_inst,
2781 .sramecc_support,
2782 }),
2783 };
2784 pub const gfx90c: CpuModel = .{
2785 .name = "gfx90c",
2786 .llvm_name = "gfx90c",
2787 .features = featureSet(&[_]Feature{
2788 .addressablelocalmemorysize65536,
2789 .ds_src2_insts,
2790 .extended_image_insts,
2791 .gds,
2792 .gfx9,
2793 .image_gather4_d16_bug,
2794 .image_insts,
2795 .ldsbankcount32,
2796 .mad_mac_f32_insts,
2797 .mad_mix_insts,
2798 }),
2799 };
2800 pub const gfx942: CpuModel = .{
2801 .name = "gfx942",
2802 .llvm_name = "gfx942",
2803 .features = featureSet(&[_]Feature{
2804 .addressablelocalmemorysize65536,
2805 .agent_scope_fine_grained_remote_memory_atomics,
2806 .architected_flat_scratch,
2807 .atomic_buffer_global_pk_add_f16_insts,
2808 .atomic_ds_pk_add_16_insts,
2809 .atomic_fadd_no_rtn_insts,
2810 .atomic_fadd_rtn_insts,
2811 .atomic_flat_pk_add_16_insts,
2812 .atomic_fmin_fmax_flat_f64,
2813 .atomic_fmin_fmax_global_f64,
2814 .atomic_global_pk_add_bf16_inst,
2815 .back_off_barrier,
2816 .cvt_fp8_vop1_bug,
2817 .dl_insts,
2818 .dot10_insts,
2819 .dot1_insts,
2820 .dot2_insts,
2821 .dot3_insts,
2822 .dot4_insts,
2823 .dot5_insts,
2824 .dot6_insts,
2825 .dot7_insts,
2826 .dpp_64bit,
2827 .flat_atomic_fadd_f32_inst,
2828 .flat_buffer_global_fadd_f64_inst,
2829 .fma_mix_insts,
2830 .fmacf64_inst,
2831 .fp8_insts,
2832 .full_rate_64_ops,
2833 .gfx9,
2834 .gfx90a_insts,
2835 .gfx940_insts,
2836 .kernarg_preload,
2837 .ldsbankcount32,
2838 .lshl_add_u64_inst,
2839 .mai_insts,
2840 .memory_atomic_fadd_f32_denormal_support,
2841 .packed_fp32_ops,
2842 .packed_tid,
2843 .pk_fmac_f16_inst,
2844 .sramecc_support,
2845 .xf32_insts,
2846 }),
2847 };
2848 pub const gfx950: CpuModel = .{
2849 .name = "gfx950",
2850 .llvm_name = "gfx950",
2851 .features = featureSet(&[_]Feature{
2852 .addressablelocalmemorysize163840,
2853 .agent_scope_fine_grained_remote_memory_atomics,
2854 .architected_flat_scratch,
2855 .atomic_buffer_global_pk_add_f16_insts,
2856 .atomic_buffer_pk_add_bf16_inst,
2857 .atomic_ds_pk_add_16_insts,
2858 .atomic_fadd_no_rtn_insts,
2859 .atomic_fadd_rtn_insts,
2860 .atomic_flat_pk_add_16_insts,
2861 .atomic_fmin_fmax_flat_f64,
2862 .atomic_fmin_fmax_global_f64,
2863 .atomic_global_pk_add_bf16_inst,
2864 .back_off_barrier,
2865 .bf16_cvt_insts,
2866 .bitop3_insts,
2867 .dl_insts,
2868 .dot10_insts,
2869 .dot12_insts,
2870 .dot13_insts,
2871 .dot1_insts,
2872 .dot2_insts,
2873 .dot3_insts,
2874 .dot4_insts,
2875 .dot5_insts,
2876 .dot6_insts,
2877 .dot7_insts,
2878 .dpp_64bit,
2879 .flat_atomic_fadd_f32_inst,
2880 .flat_buffer_global_fadd_f64_inst,
2881 .fma_mix_insts,
2882 .fmacf64_inst,
2883 .fp8_conversion_insts,
2884 .fp8_insts,
2885 .full_rate_64_ops,
2886 .gfx9,
2887 .gfx90a_insts,
2888 .gfx940_insts,
2889 .gfx950_insts,
2890 .kernarg_preload,
2891 .ldsbankcount32,
2892 .lshl_add_u64_inst,
2893 .mai_insts,
2894 .memory_atomic_fadd_f32_denormal_support,
2895 .packed_fp32_ops,
2896 .packed_tid,
2897 .pk_fmac_f16_inst,
2898 .prng_inst,
2899 .sramecc_support,
2900 }),
2901 };
2902 pub const gfx9_4_generic: CpuModel = .{
2903 .name = "gfx9_4_generic",
2904 .llvm_name = "gfx9-4-generic",
2905 .features = featureSet(&[_]Feature{
2906 .addressablelocalmemorysize65536,
2907 .agent_scope_fine_grained_remote_memory_atomics,
2908 .architected_flat_scratch,
2909 .atomic_buffer_global_pk_add_f16_insts,
2910 .atomic_ds_pk_add_16_insts,
2911 .atomic_fadd_no_rtn_insts,
2912 .atomic_fadd_rtn_insts,
2913 .atomic_flat_pk_add_16_insts,
2914 .atomic_fmin_fmax_flat_f64,
2915 .atomic_fmin_fmax_global_f64,
2916 .atomic_global_pk_add_bf16_inst,
2917 .back_off_barrier,
2918 .dl_insts,
2919 .dot10_insts,
2920 .dot1_insts,
2921 .dot2_insts,
2922 .dot3_insts,
2923 .dot4_insts,
2924 .dot5_insts,
2925 .dot6_insts,
2926 .dot7_insts,
2927 .dpp_64bit,
2928 .flat_atomic_fadd_f32_inst,
2929 .flat_buffer_global_fadd_f64_inst,
2930 .fma_mix_insts,
2931 .fmacf64_inst,
2932 .full_rate_64_ops,
2933 .gfx9,
2934 .gfx90a_insts,
2935 .gfx940_insts,
2936 .kernarg_preload,
2937 .ldsbankcount32,
2938 .lshl_add_u64_inst,
2939 .mai_insts,
2940 .memory_atomic_fadd_f32_denormal_support,
2941 .packed_fp32_ops,
2942 .packed_tid,
2943 .pk_fmac_f16_inst,
2944 .requires_cov6,
2945 .sramecc_support,
2946 }),
2947 };
2948 pub const gfx9_generic: CpuModel = .{
2949 .name = "gfx9_generic",
2950 .llvm_name = "gfx9-generic",
2951 .features = featureSet(&[_]Feature{
2952 .addressablelocalmemorysize65536,
2953 .ds_src2_insts,
2954 .extended_image_insts,
2955 .gds,
2956 .gfx9,
2957 .image_gather4_d16_bug,
2958 .image_insts,
2959 .ldsbankcount32,
2960 .mad_mac_f32_insts,
2961 .requires_cov6,
2962 }),
2963 };
2964 pub const hainan: CpuModel = .{
2965 .name = "hainan",
2966 .llvm_name = "hainan",
2967 .features = featureSet(&[_]Feature{
2968 .southern_islands,
2969 }),
2970 };
2971 pub const hawaii: CpuModel = .{
2972 .name = "hawaii",
2973 .llvm_name = "hawaii",
2974 .features = featureSet(&[_]Feature{
2975 .fast_fmaf,
2976 .half_rate_64_ops,
2977 .ldsbankcount32,
2978 .sea_islands,
2979 }),
2980 };
2981 pub const iceland: CpuModel = .{
2982 .name = "iceland",
2983 .llvm_name = "iceland",
2984 .features = featureSet(&[_]Feature{
2985 .ldsbankcount32,
2986 .sgpr_init_bug,
2987 .unpacked_d16_vmem,
2988 .volcanic_islands,
2989 }),
2990 };
2991 pub const kabini: CpuModel = .{
2992 .name = "kabini",
2993 .llvm_name = "kabini",
2994 .features = featureSet(&[_]Feature{
2995 .ldsbankcount16,
2996 .sea_islands,
2997 }),
2998 };
2999 pub const kaveri: CpuModel = .{
3000 .name = "kaveri",
3001 .llvm_name = "kaveri",
3002 .features = featureSet(&[_]Feature{
3003 .ldsbankcount32,
3004 .sea_islands,
3005 }),
3006 };
3007 pub const mullins: CpuModel = .{
3008 .name = "mullins",
3009 .llvm_name = "mullins",
3010 .features = featureSet(&[_]Feature{
3011 .ldsbankcount16,
3012 .sea_islands,
3013 }),
3014 };
3015 pub const oland: CpuModel = .{
3016 .name = "oland",
3017 .llvm_name = "oland",
3018 .features = featureSet(&[_]Feature{
3019 .southern_islands,
3020 }),
3021 };
3022 pub const pitcairn: CpuModel = .{
3023 .name = "pitcairn",
3024 .llvm_name = "pitcairn",
3025 .features = featureSet(&[_]Feature{
3026 .southern_islands,
3027 }),
3028 };
3029 pub const polaris10: CpuModel = .{
3030 .name = "polaris10",
3031 .llvm_name = "polaris10",
3032 .features = featureSet(&[_]Feature{
3033 .ldsbankcount32,
3034 .unpacked_d16_vmem,
3035 .volcanic_islands,
3036 }),
3037 };
3038 pub const polaris11: CpuModel = .{
3039 .name = "polaris11",
3040 .llvm_name = "polaris11",
3041 .features = featureSet(&[_]Feature{
3042 .ldsbankcount32,
3043 .unpacked_d16_vmem,
3044 .volcanic_islands,
3045 }),
3046 };
3047 pub const stoney: CpuModel = .{
3048 .name = "stoney",
3049 .llvm_name = "stoney",
3050 .features = featureSet(&[_]Feature{
3051 .image_gather4_d16_bug,
3052 .image_store_d16_bug,
3053 .ldsbankcount16,
3054 .volcanic_islands,
3055 .xnack_support,
3056 }),
3057 };
3058 pub const tahiti: CpuModel = .{
3059 .name = "tahiti",
3060 .llvm_name = "tahiti",
3061 .features = featureSet(&[_]Feature{
3062 .fast_fmaf,
3063 .half_rate_64_ops,
3064 .southern_islands,
3065 }),
3066 };
3067 pub const tonga: CpuModel = .{
3068 .name = "tonga",
3069 .llvm_name = "tonga",
3070 .features = featureSet(&[_]Feature{
3071 .ldsbankcount32,
3072 .sgpr_init_bug,
3073 .unpacked_d16_vmem,
3074 .volcanic_islands,
3075 }),
3076 };
3077 pub const tongapro: CpuModel = .{
3078 .name = "tongapro",
3079 .llvm_name = "tongapro",
3080 .features = featureSet(&[_]Feature{
3081 .ldsbankcount32,
3082 .sgpr_init_bug,
3083 .unpacked_d16_vmem,
3084 .volcanic_islands,
3085 }),
3086 };
3087 pub const verde: CpuModel = .{
3088 .name = "verde",
3089 .llvm_name = "verde",
3090 .features = featureSet(&[_]Feature{
3091 .southern_islands,
3092 }),
3093 };
3094};