ptx_parser/parser/instruction/
mod.rs

1// Auto-generated module declarations
2// DO NOT EDIT MANUALLY
3#![allow(unused)]
4
5use crate::parser::{PtxParser, PtxParseError, PtxTokenStream, Span};
6use crate::r#type::instruction::Inst;
7use crate::lexer::PtxToken;
8
9pub mod abs;
10pub mod activemask;
11pub mod add_cc;
12pub mod add;
13pub mod addc;
14pub mod alloca;
15pub mod and;
16pub mod applypriority;
17pub mod atom;
18pub mod bar;
19pub mod bar_warp_sync;
20pub mod barrier_cluster;
21pub mod bfe;
22pub mod bfi;
23pub mod bfind;
24pub mod bmsk;
25pub mod bra;
26pub mod brev;
27pub mod brkpt;
28pub mod brx_idx;
29pub mod call;
30pub mod clusterlaunchcontrol_query_cancel;
31pub mod clusterlaunchcontrol_try_cancel;
32pub mod clz;
33pub mod cnot;
34pub mod copysign;
35pub mod cos;
36pub mod cp_async_bulk_commit_group;
37pub mod cp_async_bulk_prefetch_tensor;
38pub mod cp_async_bulk_prefetch;
39pub mod cp_async_bulk_tensor;
40pub mod cp_async_bulk;
41pub mod cp_async_bulk_wait_group;
42pub mod cp_async_commit_group;
43pub mod cp_async_mbarrier_arrive;
44pub mod cp_async;
45pub mod cp_async_wait_group;
46pub mod cp_reduce_async_bulk_tensor;
47pub mod cp_reduce_async_bulk;
48pub mod createpolicy;
49pub mod cvt_pack;
50pub mod cvt;
51pub mod cvta;
52pub mod discard;
53pub mod div;
54pub mod dp2a;
55pub mod dp4a;
56pub mod elect_sync;
57pub mod ex2;
58pub mod exit;
59pub mod fma;
60pub mod fns;
61pub mod getctarank;
62pub mod griddepcontrol;
63pub mod isspacep;
64pub mod istypep;
65pub mod ld_global_nc;
66pub mod ld;
67pub mod ldmatrix;
68pub mod ldu;
69pub mod lg2;
70pub mod lop3;
71pub mod mad_cc;
72pub mod mad;
73pub mod mad24;
74pub mod madc;
75pub mod mapa;
76pub mod match_sync;
77pub mod max;
78pub mod mbarrier_arrive;
79pub mod mbarrier_arrive_drop;
80pub mod mbarrier_complete_tx;
81pub mod mbarrier_expect_tx;
82pub mod mbarrier_init;
83pub mod mbarrier_inval;
84pub mod mbarrier_pending_count;
85pub mod mbarrier_test_wait;
86pub mod membar;
87pub mod min;
88pub mod mma_sp;
89pub mod mma;
90pub mod mov;
91pub mod movmatrix;
92pub mod mul;
93pub mod mul24;
94pub mod multimem_ld_reduce;
95pub mod nanosleep;
96pub mod neg;
97pub mod not;
98pub mod or;
99pub mod pmevent;
100pub mod popc;
101pub mod prefetch;
102pub mod prmt;
103pub mod rcp_approx_ftz_f64;
104pub mod rcp;
105pub mod red_async;
106pub mod red;
107pub mod redux_sync;
108pub mod rem;
109pub mod ret;
110pub mod rsqrt_approx_ftz_f64;
111pub mod rsqrt;
112pub mod sad;
113pub mod selp;
114pub mod set;
115pub mod setmaxnreg;
116pub mod setp;
117pub mod shf;
118pub mod shfl_sync;
119pub mod shfl;
120pub mod shl;
121pub mod shr;
122pub mod sin;
123pub mod slct;
124pub mod sqrt;
125pub mod st_async;
126pub mod st_bulk;
127pub mod st;
128pub mod stackrestore;
129pub mod stacksave;
130pub mod stmatrix;
131pub mod sub_cc;
132pub mod sub;
133pub mod subc;
134pub mod suld;
135pub mod suq;
136pub mod sured;
137pub mod sust;
138pub mod szext;
139pub mod tanh;
140pub mod tcgen05_alloc;
141pub mod tcgen05_commit;
142pub mod tcgen05_cp;
143pub mod tcgen05_fence;
144pub mod tcgen05_ld;
145pub mod tcgen05_mma_sp;
146pub mod tcgen05_mma;
147pub mod tcgen05_mma_ws_sp;
148pub mod tcgen05_mma_ws;
149pub mod tcgen05_shift;
150pub mod tcgen05_st;
151pub mod tcgen05_wait;
152pub mod tensormap_cp_fenceproxy;
153pub mod tensormap_replace;
154pub mod testp;
155pub mod tex;
156pub mod tld4;
157pub mod trap;
158pub mod txq;
159pub mod vmad;
160pub mod vop;
161pub mod vop2;
162pub mod vop4;
163pub mod vote_sync;
164pub mod vote;
165pub mod vset;
166pub mod vset2;
167pub mod vset4;
168pub mod vsh;
169pub mod wgmma_commit_group;
170pub mod wgmma_fence;
171pub mod wgmma_mma_async_sp;
172pub mod wgmma_mma_async;
173pub mod wgmma_wait_group;
174pub mod wmma_load;
175pub mod wmma_mma;
176pub mod wmma_store;
177pub mod xor;
178
179/// Parse instruction without label or predicate
180pub(crate) fn parse_instruction_inner(stream: &mut PtxTokenStream) -> Result<Inst, PtxParseError> {
181    let start_pos = stream.position();
182    
183    // Peek at the opcode to determine which parser to try
184    let opcode = if let Ok((PtxToken::Identifier(name), _)) = stream.peek() {
185        name.as_str()
186    } else {
187        let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(0..0);
188        return Err(crate::parser::unexpected_value(span, &["instruction opcode"], "not an identifier"));
189    };
190    
191    // Dispatch based on opcode
192    match opcode {
193        "abs" => {
194            stream.set_position(start_pos);
195            match <crate::r#type::instruction::abs::section_0::AbsType as PtxParser>::parse(stream) {
196                Ok(inst) => return Ok(Inst::AbsType(inst)),
197                Err(_) => {}
198            }
199            stream.set_position(start_pos);
200            match <crate::r#type::instruction::abs::section_0::AbsFtzF32 as PtxParser>::parse(stream) {
201                Ok(inst) => return Ok(Inst::AbsFtzF32(inst)),
202                Err(_) => {}
203            }
204            stream.set_position(start_pos);
205            match <crate::r#type::instruction::abs::section_0::AbsF64 as PtxParser>::parse(stream) {
206                Ok(inst) => return Ok(Inst::AbsF64(inst)),
207                Err(_) => {}
208            }
209            stream.set_position(start_pos);
210            match <crate::r#type::instruction::abs::section_0::AbsFtzF16 as PtxParser>::parse(stream) {
211                Ok(inst) => return Ok(Inst::AbsFtzF16(inst)),
212                Err(_) => {}
213            }
214            stream.set_position(start_pos);
215            match <crate::r#type::instruction::abs::section_0::AbsFtzF16x2 as PtxParser>::parse(stream) {
216                Ok(inst) => return Ok(Inst::AbsFtzF16x2(inst)),
217                Err(_) => {}
218            }
219            stream.set_position(start_pos);
220            match <crate::r#type::instruction::abs::section_0::AbsBf16 as PtxParser>::parse(stream) {
221                Ok(inst) => return Ok(Inst::AbsBf16(inst)),
222                Err(_) => {}
223            }
224            stream.set_position(start_pos);
225            match <crate::r#type::instruction::abs::section_0::AbsBf16x2 as PtxParser>::parse(stream) {
226                Ok(inst) => return Ok(Inst::AbsBf16x2(inst)),
227                Err(_) => {}
228            }
229        }
230        "activemask" => {
231            stream.set_position(start_pos);
232            match <crate::r#type::instruction::activemask::section_0::ActivemaskB32 as PtxParser>::parse(stream) {
233                Ok(inst) => return Ok(Inst::ActivemaskB32(inst)),
234                Err(_) => {}
235            }
236        }
237        "add" => {
238            stream.set_position(start_pos);
239            match <crate::r#type::instruction::add_cc::section_0::AddCcType as PtxParser>::parse(stream) {
240                Ok(inst) => return Ok(Inst::AddCcType(inst)),
241                Err(_) => {}
242            }
243            stream.set_position(start_pos);
244            match <crate::r#type::instruction::add::section_0::AddType as PtxParser>::parse(stream) {
245                Ok(inst) => return Ok(Inst::AddType(inst)),
246                Err(_) => {}
247            }
248            stream.set_position(start_pos);
249            match <crate::r#type::instruction::add::section_0::AddSatS32 as PtxParser>::parse(stream) {
250                Ok(inst) => return Ok(Inst::AddSatS32(inst)),
251                Err(_) => {}
252            }
253            stream.set_position(start_pos);
254            match <crate::r#type::instruction::add::section_1::AddRndFtzSatF32 as PtxParser>::parse(stream) {
255                Ok(inst) => return Ok(Inst::AddRndFtzSatF32(inst)),
256                Err(_) => {}
257            }
258            stream.set_position(start_pos);
259            match <crate::r#type::instruction::add::section_1::AddRndFtzF32x2 as PtxParser>::parse(stream) {
260                Ok(inst) => return Ok(Inst::AddRndFtzF32x2(inst)),
261                Err(_) => {}
262            }
263            stream.set_position(start_pos);
264            match <crate::r#type::instruction::add::section_1::AddRndF64 as PtxParser>::parse(stream) {
265                Ok(inst) => return Ok(Inst::AddRndF64(inst)),
266                Err(_) => {}
267            }
268            stream.set_position(start_pos);
269            match <crate::r#type::instruction::add::section_2::AddRndFtzSatF16 as PtxParser>::parse(stream) {
270                Ok(inst) => return Ok(Inst::AddRndFtzSatF16(inst)),
271                Err(_) => {}
272            }
273            stream.set_position(start_pos);
274            match <crate::r#type::instruction::add::section_2::AddRndFtzSatF16x2 as PtxParser>::parse(stream) {
275                Ok(inst) => return Ok(Inst::AddRndFtzSatF16x2(inst)),
276                Err(_) => {}
277            }
278            stream.set_position(start_pos);
279            match <crate::r#type::instruction::add::section_2::AddRndBf16 as PtxParser>::parse(stream) {
280                Ok(inst) => return Ok(Inst::AddRndBf16(inst)),
281                Err(_) => {}
282            }
283            stream.set_position(start_pos);
284            match <crate::r#type::instruction::add::section_2::AddRndBf16x2 as PtxParser>::parse(stream) {
285                Ok(inst) => return Ok(Inst::AddRndBf16x2(inst)),
286                Err(_) => {}
287            }
288            stream.set_position(start_pos);
289            match <crate::r#type::instruction::add::section_3::AddRndSatF32Atype as PtxParser>::parse(stream) {
290                Ok(inst) => return Ok(Inst::AddRndSatF32Atype(inst)),
291                Err(_) => {}
292            }
293        }
294        "addc" => {
295            stream.set_position(start_pos);
296            match <crate::r#type::instruction::addc::section_0::AddcCcType as PtxParser>::parse(stream) {
297                Ok(inst) => return Ok(Inst::AddcCcType(inst)),
298                Err(_) => {}
299            }
300        }
301        "alloca" => {
302            stream.set_position(start_pos);
303            match <crate::r#type::instruction::alloca::section_0::AllocaType as PtxParser>::parse(stream) {
304                Ok(inst) => return Ok(Inst::AllocaType(inst)),
305                Err(_) => {}
306            }
307        }
308        "and" => {
309            stream.set_position(start_pos);
310            match <crate::r#type::instruction::and::section_0::AndType as PtxParser>::parse(stream) {
311                Ok(inst) => return Ok(Inst::AndType(inst)),
312                Err(_) => {}
313            }
314        }
315        "applypriority" => {
316            stream.set_position(start_pos);
317            match <crate::r#type::instruction::applypriority::section_0::ApplypriorityGlobalLevelEvictionPriority as PtxParser>::parse(stream) {
318                Ok(inst) => return Ok(Inst::ApplypriorityGlobalLevelEvictionPriority(inst)),
319                Err(_) => {}
320            }
321        }
322        "atom" => {
323            stream.set_position(start_pos);
324            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceOpLevelCacheHintType as PtxParser>::parse(stream) {
325                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceOpLevelCacheHintType(inst)),
326                Err(_) => {}
327            }
328            stream.set_position(start_pos);
329            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceOpType as PtxParser>::parse(stream) {
330                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceOpType(inst)),
331                Err(_) => {}
332            }
333            stream.set_position(start_pos);
334            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceCasB16 as PtxParser>::parse(stream) {
335                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceCasB16(inst)),
336                Err(_) => {}
337            }
338            stream.set_position(start_pos);
339            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceCasB128 as PtxParser>::parse(stream) {
340                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceCasB128(inst)),
341                Err(_) => {}
342            }
343            stream.set_position(start_pos);
344            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceExchLevelCacheHintB128 as PtxParser>::parse(stream) {
345                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceExchLevelCacheHintB128(inst)),
346                Err(_) => {}
347            }
348            stream.set_position(start_pos);
349            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceAddNoftzLevelCacheHintF16 as PtxParser>::parse(stream) {
350                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceAddNoftzLevelCacheHintF16(inst)),
351                Err(_) => {}
352            }
353            stream.set_position(start_pos);
354            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceAddNoftzLevelCacheHintF16x2 as PtxParser>::parse(stream) {
355                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceAddNoftzLevelCacheHintF16x2(inst)),
356                Err(_) => {}
357            }
358            stream.set_position(start_pos);
359            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16 as PtxParser>::parse(stream) {
360                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16(inst)),
361                Err(_) => {}
362            }
363            stream.set_position(start_pos);
364            match <crate::r#type::instruction::atom::section_0::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16x2 as PtxParser>::parse(stream) {
365                Ok(inst) => return Ok(Inst::AtomSemScopeSpaceAddNoftzLevelCacheHintBf16x2(inst)),
366                Err(_) => {}
367            }
368            stream.set_position(start_pos);
369            match <crate::r#type::instruction::atom::section_1::AtomSemScopeGlobalAddLevelCacheHintVec32BitF32 as PtxParser>::parse(stream) {
370                Ok(inst) => return Ok(Inst::AtomSemScopeGlobalAddLevelCacheHintVec32BitF32(inst)),
371                Err(_) => {}
372            }
373            stream.set_position(start_pos);
374            match <crate::r#type::instruction::atom::section_1::AtomSemScopeGlobalOpNoftzLevelCacheHintVec16BitHalfWordType as PtxParser>::parse(stream) {
375                Ok(inst) => return Ok(Inst::AtomSemScopeGlobalOpNoftzLevelCacheHintVec16BitHalfWordType(inst)),
376                Err(_) => {}
377            }
378            stream.set_position(start_pos);
379            match <crate::r#type::instruction::atom::section_1::AtomSemScopeGlobalOpNoftzLevelCacheHintVec32BitPackedType as PtxParser>::parse(stream) {
380                Ok(inst) => return Ok(Inst::AtomSemScopeGlobalOpNoftzLevelCacheHintVec32BitPackedType(inst)),
381                Err(_) => {}
382            }
383        }
384        "bar" => {
385            stream.set_position(start_pos);
386            match <crate::r#type::instruction::bar::section_0::BarrierCtaSyncAligned as PtxParser>::parse(stream) {
387                Ok(inst) => return Ok(Inst::BarrierCtaSyncAligned(inst)),
388                Err(_) => {}
389            }
390            stream.set_position(start_pos);
391            match <crate::r#type::instruction::bar::section_0::BarrierCtaArriveAligned as PtxParser>::parse(stream) {
392                Ok(inst) => return Ok(Inst::BarrierCtaArriveAligned(inst)),
393                Err(_) => {}
394            }
395            stream.set_position(start_pos);
396            match <crate::r#type::instruction::bar::section_0::BarrierCtaRedPopcAlignedU32 as PtxParser>::parse(stream) {
397                Ok(inst) => return Ok(Inst::BarrierCtaRedPopcAlignedU32(inst)),
398                Err(_) => {}
399            }
400            stream.set_position(start_pos);
401            match <crate::r#type::instruction::bar::section_0::BarrierCtaRedOpAlignedPred as PtxParser>::parse(stream) {
402                Ok(inst) => return Ok(Inst::BarrierCtaRedOpAlignedPred(inst)),
403                Err(_) => {}
404            }
405            stream.set_position(start_pos);
406            match <crate::r#type::instruction::bar::section_0::BarCtaSync as PtxParser>::parse(stream) {
407                Ok(inst) => return Ok(Inst::BarCtaSync(inst)),
408                Err(_) => {}
409            }
410            stream.set_position(start_pos);
411            match <crate::r#type::instruction::bar::section_0::BarCtaArrive as PtxParser>::parse(stream) {
412                Ok(inst) => return Ok(Inst::BarCtaArrive(inst)),
413                Err(_) => {}
414            }
415            stream.set_position(start_pos);
416            match <crate::r#type::instruction::bar::section_0::BarCtaRedPopcU32 as PtxParser>::parse(stream) {
417                Ok(inst) => return Ok(Inst::BarCtaRedPopcU32(inst)),
418                Err(_) => {}
419            }
420            stream.set_position(start_pos);
421            match <crate::r#type::instruction::bar::section_0::BarCtaRedOpPred as PtxParser>::parse(stream) {
422                Ok(inst) => return Ok(Inst::BarCtaRedOpPred(inst)),
423                Err(_) => {}
424            }
425            stream.set_position(start_pos);
426            match <crate::r#type::instruction::bar_warp_sync::section_0::BarWarpSync as PtxParser>::parse(stream) {
427                Ok(inst) => return Ok(Inst::BarWarpSync(inst)),
428                Err(_) => {}
429            }
430        }
431        "barrier" => {
432            stream.set_position(start_pos);
433            match <crate::r#type::instruction::bar::section_0::BarrierCtaSyncAligned as PtxParser>::parse(stream) {
434                Ok(inst) => return Ok(Inst::BarrierCtaSyncAligned(inst)),
435                Err(_) => {}
436            }
437            stream.set_position(start_pos);
438            match <crate::r#type::instruction::bar::section_0::BarrierCtaArriveAligned as PtxParser>::parse(stream) {
439                Ok(inst) => return Ok(Inst::BarrierCtaArriveAligned(inst)),
440                Err(_) => {}
441            }
442            stream.set_position(start_pos);
443            match <crate::r#type::instruction::bar::section_0::BarrierCtaRedPopcAlignedU32 as PtxParser>::parse(stream) {
444                Ok(inst) => return Ok(Inst::BarrierCtaRedPopcAlignedU32(inst)),
445                Err(_) => {}
446            }
447            stream.set_position(start_pos);
448            match <crate::r#type::instruction::bar::section_0::BarrierCtaRedOpAlignedPred as PtxParser>::parse(stream) {
449                Ok(inst) => return Ok(Inst::BarrierCtaRedOpAlignedPred(inst)),
450                Err(_) => {}
451            }
452            stream.set_position(start_pos);
453            match <crate::r#type::instruction::bar::section_0::BarCtaSync as PtxParser>::parse(stream) {
454                Ok(inst) => return Ok(Inst::BarCtaSync(inst)),
455                Err(_) => {}
456            }
457            stream.set_position(start_pos);
458            match <crate::r#type::instruction::bar::section_0::BarCtaArrive as PtxParser>::parse(stream) {
459                Ok(inst) => return Ok(Inst::BarCtaArrive(inst)),
460                Err(_) => {}
461            }
462            stream.set_position(start_pos);
463            match <crate::r#type::instruction::bar::section_0::BarCtaRedPopcU32 as PtxParser>::parse(stream) {
464                Ok(inst) => return Ok(Inst::BarCtaRedPopcU32(inst)),
465                Err(_) => {}
466            }
467            stream.set_position(start_pos);
468            match <crate::r#type::instruction::bar::section_0::BarCtaRedOpPred as PtxParser>::parse(stream) {
469                Ok(inst) => return Ok(Inst::BarCtaRedOpPred(inst)),
470                Err(_) => {}
471            }
472            stream.set_position(start_pos);
473            match <crate::r#type::instruction::barrier_cluster::section_0::BarrierClusterArriveSemAligned as PtxParser>::parse(stream) {
474                Ok(inst) => return Ok(Inst::BarrierClusterArriveSemAligned(inst)),
475                Err(_) => {}
476            }
477            stream.set_position(start_pos);
478            match <crate::r#type::instruction::barrier_cluster::section_0::BarrierClusterWaitAcquireAligned as PtxParser>::parse(stream) {
479                Ok(inst) => return Ok(Inst::BarrierClusterWaitAcquireAligned(inst)),
480                Err(_) => {}
481            }
482        }
483        "bfe" => {
484            stream.set_position(start_pos);
485            match <crate::r#type::instruction::bfe::section_0::BfeType as PtxParser>::parse(stream) {
486                Ok(inst) => return Ok(Inst::BfeType(inst)),
487                Err(_) => {}
488            }
489        }
490        "bfi" => {
491            stream.set_position(start_pos);
492            match <crate::r#type::instruction::bfi::section_0::BfiType as PtxParser>::parse(stream) {
493                Ok(inst) => return Ok(Inst::BfiType(inst)),
494                Err(_) => {}
495            }
496        }
497        "bfind" => {
498            stream.set_position(start_pos);
499            match <crate::r#type::instruction::bfind::section_0::BfindType as PtxParser>::parse(stream) {
500                Ok(inst) => return Ok(Inst::BfindType(inst)),
501                Err(_) => {}
502            }
503            stream.set_position(start_pos);
504            match <crate::r#type::instruction::bfind::section_0::BfindShiftamtType as PtxParser>::parse(stream) {
505                Ok(inst) => return Ok(Inst::BfindShiftamtType(inst)),
506                Err(_) => {}
507            }
508        }
509        "bmsk" => {
510            stream.set_position(start_pos);
511            match <crate::r#type::instruction::bmsk::section_0::BmskModeB32 as PtxParser>::parse(stream) {
512                Ok(inst) => return Ok(Inst::BmskModeB32(inst)),
513                Err(_) => {}
514            }
515        }
516        "bra" => {
517            stream.set_position(start_pos);
518            match <crate::r#type::instruction::bra::section_0::BraUni as PtxParser>::parse(stream) {
519                Ok(inst) => return Ok(Inst::BraUni(inst)),
520                Err(_) => {}
521            }
522            stream.set_position(start_pos);
523            match <crate::r#type::instruction::bra::section_0::BraUni1 as PtxParser>::parse(stream) {
524                Ok(inst) => return Ok(Inst::BraUni1(inst)),
525                Err(_) => {}
526            }
527        }
528        "brev" => {
529            stream.set_position(start_pos);
530            match <crate::r#type::instruction::brev::section_0::BrevType as PtxParser>::parse(stream) {
531                Ok(inst) => return Ok(Inst::BrevType(inst)),
532                Err(_) => {}
533            }
534        }
535        "brkpt" => {
536            stream.set_position(start_pos);
537            match <crate::r#type::instruction::brkpt::section_0::Brkpt as PtxParser>::parse(stream) {
538                Ok(inst) => return Ok(Inst::Brkpt(inst)),
539                Err(_) => {}
540            }
541        }
542        "brx" => {
543            stream.set_position(start_pos);
544            match <crate::r#type::instruction::brx_idx::section_0::BrxIdxUni as PtxParser>::parse(stream) {
545                Ok(inst) => return Ok(Inst::BrxIdxUni(inst)),
546                Err(_) => {}
547            }
548            stream.set_position(start_pos);
549            match <crate::r#type::instruction::brx_idx::section_0::BrxIdxUni1 as PtxParser>::parse(stream) {
550                Ok(inst) => return Ok(Inst::BrxIdxUni1(inst)),
551                Err(_) => {}
552            }
553        }
554        "call" => {
555            stream.set_position(start_pos);
556            match <crate::r#type::instruction::call::section_0::CallUni as PtxParser>::parse(stream) {
557                Ok(inst) => return Ok(Inst::CallUni(inst)),
558                Err(_) => {}
559            }
560            stream.set_position(start_pos);
561            match <crate::r#type::instruction::call::section_0::CallUni1 as PtxParser>::parse(stream) {
562                Ok(inst) => return Ok(Inst::CallUni1(inst)),
563                Err(_) => {}
564            }
565            stream.set_position(start_pos);
566            match <crate::r#type::instruction::call::section_0::CallUni2 as PtxParser>::parse(stream) {
567                Ok(inst) => return Ok(Inst::CallUni2(inst)),
568                Err(_) => {}
569            }
570            stream.set_position(start_pos);
571            match <crate::r#type::instruction::call::section_0::CallUni3 as PtxParser>::parse(stream) {
572                Ok(inst) => return Ok(Inst::CallUni3(inst)),
573                Err(_) => {}
574            }
575            stream.set_position(start_pos);
576            match <crate::r#type::instruction::call::section_0::CallUni4 as PtxParser>::parse(stream) {
577                Ok(inst) => return Ok(Inst::CallUni4(inst)),
578                Err(_) => {}
579            }
580            stream.set_position(start_pos);
581            match <crate::r#type::instruction::call::section_0::CallUni5 as PtxParser>::parse(stream) {
582                Ok(inst) => return Ok(Inst::CallUni5(inst)),
583                Err(_) => {}
584            }
585            stream.set_position(start_pos);
586            match <crate::r#type::instruction::call::section_0::CallUni6 as PtxParser>::parse(stream) {
587                Ok(inst) => return Ok(Inst::CallUni6(inst)),
588                Err(_) => {}
589            }
590            stream.set_position(start_pos);
591            match <crate::r#type::instruction::call::section_0::CallUni7 as PtxParser>::parse(stream) {
592                Ok(inst) => return Ok(Inst::CallUni7(inst)),
593                Err(_) => {}
594            }
595            stream.set_position(start_pos);
596            match <crate::r#type::instruction::call::section_0::CallUni8 as PtxParser>::parse(stream) {
597                Ok(inst) => return Ok(Inst::CallUni8(inst)),
598                Err(_) => {}
599            }
600        }
601        "clusterlaunchcontrol" => {
602            stream.set_position(start_pos);
603            match <crate::r#type::instruction::clusterlaunchcontrol_query_cancel::section_0::ClusterlaunchcontrolQueryCancelIsCanceledPredB128 as PtxParser>::parse(stream) {
604                Ok(inst) => return Ok(Inst::ClusterlaunchcontrolQueryCancelIsCanceledPredB128(inst)),
605                Err(_) => {}
606            }
607            stream.set_position(start_pos);
608            match <crate::r#type::instruction::clusterlaunchcontrol_query_cancel::section_0::ClusterlaunchcontrolQueryCancelGetFirstCtaidV4B32B128 as PtxParser>::parse(stream) {
609                Ok(inst) => return Ok(Inst::ClusterlaunchcontrolQueryCancelGetFirstCtaidV4B32B128(inst)),
610                Err(_) => {}
611            }
612            stream.set_position(start_pos);
613            match <crate::r#type::instruction::clusterlaunchcontrol_query_cancel::section_0::ClusterlaunchcontrolQueryCancelGetFirstCtaidDimensionB32B128 as PtxParser>::parse(stream) {
614                Ok(inst) => return Ok(Inst::ClusterlaunchcontrolQueryCancelGetFirstCtaidDimensionB32B128(inst)),
615                Err(_) => {}
616            }
617            stream.set_position(start_pos);
618            match <crate::r#type::instruction::clusterlaunchcontrol_try_cancel::section_0::ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128 as PtxParser>::parse(stream) {
619                Ok(inst) => return Ok(Inst::ClusterlaunchcontrolTryCancelAsyncSpaceCompletionMechanismMulticastClusterAllB128(inst)),
620                Err(_) => {}
621            }
622        }
623        "clz" => {
624            stream.set_position(start_pos);
625            match <crate::r#type::instruction::clz::section_0::ClzType as PtxParser>::parse(stream) {
626                Ok(inst) => return Ok(Inst::ClzType(inst)),
627                Err(_) => {}
628            }
629        }
630        "cnot" => {
631            stream.set_position(start_pos);
632            match <crate::r#type::instruction::cnot::section_0::CnotType as PtxParser>::parse(stream) {
633                Ok(inst) => return Ok(Inst::CnotType(inst)),
634                Err(_) => {}
635            }
636        }
637        "copysign" => {
638            stream.set_position(start_pos);
639            match <crate::r#type::instruction::copysign::section_0::CopysignType as PtxParser>::parse(stream) {
640                Ok(inst) => return Ok(Inst::CopysignType(inst)),
641                Err(_) => {}
642            }
643        }
644        "cos" => {
645            stream.set_position(start_pos);
646            match <crate::r#type::instruction::cos::section_0::CosApproxFtzF32 as PtxParser>::parse(stream) {
647                Ok(inst) => return Ok(Inst::CosApproxFtzF32(inst)),
648                Err(_) => {}
649            }
650        }
651        "cp" => {
652            stream.set_position(start_pos);
653            match <crate::r#type::instruction::cp_async_bulk_commit_group::section_0::CpAsyncBulkCommitGroup as PtxParser>::parse(stream) {
654                Ok(inst) => return Ok(Inst::CpAsyncBulkCommitGroup(inst)),
655                Err(_) => {}
656            }
657            stream.set_position(start_pos);
658            match <crate::r#type::instruction::cp_async_bulk_prefetch_tensor::section_0::CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint as PtxParser>::parse(stream) {
659                Ok(inst) => return Ok(Inst::CpAsyncBulkPrefetchTensorDimL2SrcLoadModeLevelCacheHint(inst)),
660                Err(_) => {}
661            }
662            stream.set_position(start_pos);
663            match <crate::r#type::instruction::cp_async_bulk_prefetch::section_0::CpAsyncBulkPrefetchL2SrcLevelCacheHint as PtxParser>::parse(stream) {
664                Ok(inst) => return Ok(Inst::CpAsyncBulkPrefetchL2SrcLevelCacheHint(inst)),
665                Err(_) => {}
666            }
667            stream.set_position(start_pos);
668            match <crate::r#type::instruction::cp_async_bulk_tensor::section_0::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint as PtxParser>::parse(stream) {
669                Ok(inst) => return Ok(Inst::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismCtaGroupLevelCacheHint(inst)),
670                Err(_) => {}
671            }
672            stream.set_position(start_pos);
673            match <crate::r#type::instruction::cp_async_bulk_tensor::section_1::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint as PtxParser>::parse(stream) {
674                Ok(inst) => return Ok(Inst::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismMulticastCtaGroupLevelCacheHint(inst)),
675                Err(_) => {}
676            }
677            stream.set_position(start_pos);
678            match <crate::r#type::instruction::cp_async_bulk_tensor::section_2::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint as PtxParser>::parse(stream) {
679                Ok(inst) => return Ok(Inst::CpAsyncBulkTensorDimDstSrcLoadModeCompletionMechanismLevelCacheHint(inst)),
680                Err(_) => {}
681            }
682            stream.set_position(start_pos);
683            match <crate::r#type::instruction::cp_async_bulk::section_0::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint as PtxParser>::parse(stream) {
684                Ok(inst) => return Ok(Inst::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHint(inst)),
685                Err(_) => {}
686            }
687            stream.set_position(start_pos);
688            match <crate::r#type::instruction::cp_async_bulk::section_1::CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint as PtxParser>::parse(stream) {
689                Ok(inst) => return Ok(Inst::CpAsyncBulkDstSrcCompletionMechanismMulticastLevelCacheHint(inst)),
690                Err(_) => {}
691            }
692            stream.set_position(start_pos);
693            match <crate::r#type::instruction::cp_async_bulk::section_2::CpAsyncBulkDstSrcCompletionMechanism as PtxParser>::parse(stream) {
694                Ok(inst) => return Ok(Inst::CpAsyncBulkDstSrcCompletionMechanism(inst)),
695                Err(_) => {}
696            }
697            stream.set_position(start_pos);
698            match <crate::r#type::instruction::cp_async_bulk::section_3::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask as PtxParser>::parse(stream) {
699                Ok(inst) => return Ok(Inst::CpAsyncBulkDstSrcCompletionMechanismLevelCacheHintCpMask(inst)),
700                Err(_) => {}
701            }
702            stream.set_position(start_pos);
703            match <crate::r#type::instruction::cp_async_bulk_wait_group::section_0::CpAsyncBulkWaitGroupRead as PtxParser>::parse(stream) {
704                Ok(inst) => return Ok(Inst::CpAsyncBulkWaitGroupRead(inst)),
705                Err(_) => {}
706            }
707            stream.set_position(start_pos);
708            match <crate::r#type::instruction::cp_async_commit_group::section_0::CpAsyncCommitGroup as PtxParser>::parse(stream) {
709                Ok(inst) => return Ok(Inst::CpAsyncCommitGroup(inst)),
710                Err(_) => {}
711            }
712            stream.set_position(start_pos);
713            match <crate::r#type::instruction::cp_async_mbarrier_arrive::section_0::CpAsyncMbarrierArriveNoincStateB64 as PtxParser>::parse(stream) {
714                Ok(inst) => return Ok(Inst::CpAsyncMbarrierArriveNoincStateB64(inst)),
715                Err(_) => {}
716            }
717            stream.set_position(start_pos);
718            match <crate::r#type::instruction::cp_async::section_0::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize as PtxParser>::parse(stream) {
719                Ok(inst) => return Ok(Inst::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize(inst)),
720                Err(_) => {}
721            }
722            stream.set_position(start_pos);
723            match <crate::r#type::instruction::cp_async::section_0::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize as PtxParser>::parse(stream) {
724                Ok(inst) => return Ok(Inst::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize(inst)),
725                Err(_) => {}
726            }
727            stream.set_position(start_pos);
728            match <crate::r#type::instruction::cp_async::section_0::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1 as PtxParser>::parse(stream) {
729                Ok(inst) => return Ok(Inst::CpAsyncCaStateGlobalLevelCacheHintLevelPrefetchSize1(inst)),
730                Err(_) => {}
731            }
732            stream.set_position(start_pos);
733            match <crate::r#type::instruction::cp_async::section_0::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1 as PtxParser>::parse(stream) {
734                Ok(inst) => return Ok(Inst::CpAsyncCgStateGlobalLevelCacheHintLevelPrefetchSize1(inst)),
735                Err(_) => {}
736            }
737            stream.set_position(start_pos);
738            match <crate::r#type::instruction::cp_async_wait_group::section_0::CpAsyncWaitGroup as PtxParser>::parse(stream) {
739                Ok(inst) => return Ok(Inst::CpAsyncWaitGroup(inst)),
740                Err(_) => {}
741            }
742            stream.set_position(start_pos);
743            match <crate::r#type::instruction::cp_async_wait_group::section_0::CpAsyncWaitAll as PtxParser>::parse(stream) {
744                Ok(inst) => return Ok(Inst::CpAsyncWaitAll(inst)),
745                Err(_) => {}
746            }
747            stream.set_position(start_pos);
748            match <crate::r#type::instruction::cp_reduce_async_bulk_tensor::section_0::CpReduceAsyncBulkTensorDimDstSrcRedopLoadModeCompletionMechanismLevelCacheHint as PtxParser>::parse(stream) {
749                Ok(inst) => return Ok(Inst::CpReduceAsyncBulkTensorDimDstSrcRedopLoadModeCompletionMechanismLevelCacheHint(inst)),
750                Err(_) => {}
751            }
752            stream.set_position(start_pos);
753            match <crate::r#type::instruction::cp_reduce_async_bulk::section_0::CpReduceAsyncBulkDstSrcCompletionMechanismRedopType as PtxParser>::parse(stream) {
754                Ok(inst) => return Ok(Inst::CpReduceAsyncBulkDstSrcCompletionMechanismRedopType(inst)),
755                Err(_) => {}
756            }
757            stream.set_position(start_pos);
758            match <crate::r#type::instruction::cp_reduce_async_bulk::section_1::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType as PtxParser>::parse(stream) {
759                Ok(inst) => return Ok(Inst::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintRedopType(inst)),
760                Err(_) => {}
761            }
762            stream.set_position(start_pos);
763            match <crate::r#type::instruction::cp_reduce_async_bulk::section_2::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType as PtxParser>::parse(stream) {
764                Ok(inst) => return Ok(Inst::CpReduceAsyncBulkDstSrcCompletionMechanismLevelCacheHintAddNoftzType(inst)),
765                Err(_) => {}
766            }
767        }
768        "createpolicy" => {
769            stream.set_position(start_pos);
770            match <crate::r#type::instruction::createpolicy::section_0::CreatepolicyRangeGlobalLevelPrimaryPriorityLevelSecondaryPriorityB64 as PtxParser>::parse(stream) {
771                Ok(inst) => return Ok(Inst::CreatepolicyRangeGlobalLevelPrimaryPriorityLevelSecondaryPriorityB64(inst)),
772                Err(_) => {}
773            }
774            stream.set_position(start_pos);
775            match <crate::r#type::instruction::createpolicy::section_0::CreatepolicyFractionalLevelPrimaryPriorityLevelSecondaryPriorityB64 as PtxParser>::parse(stream) {
776                Ok(inst) => return Ok(Inst::CreatepolicyFractionalLevelPrimaryPriorityLevelSecondaryPriorityB64(inst)),
777                Err(_) => {}
778            }
779            stream.set_position(start_pos);
780            match <crate::r#type::instruction::createpolicy::section_0::CreatepolicyCvtL2B64 as PtxParser>::parse(stream) {
781                Ok(inst) => return Ok(Inst::CreatepolicyCvtL2B64(inst)),
782                Err(_) => {}
783            }
784        }
785        "cvt" => {
786            stream.set_position(start_pos);
787            match <crate::r#type::instruction::cvt_pack::section_0::CvtPackSatConverttypeAbtype as PtxParser>::parse(stream) {
788                Ok(inst) => return Ok(Inst::CvtPackSatConverttypeAbtype(inst)),
789                Err(_) => {}
790            }
791            stream.set_position(start_pos);
792            match <crate::r#type::instruction::cvt_pack::section_1::CvtPackSatConverttypeAbtypeCtype as PtxParser>::parse(stream) {
793                Ok(inst) => return Ok(Inst::CvtPackSatConverttypeAbtypeCtype(inst)),
794                Err(_) => {}
795            }
796            stream.set_position(start_pos);
797            match <crate::r#type::instruction::cvt::section_0::CvtIrndFtzSatDtypeAtype as PtxParser>::parse(stream) {
798                Ok(inst) => return Ok(Inst::CvtIrndFtzSatDtypeAtype(inst)),
799                Err(_) => {}
800            }
801            stream.set_position(start_pos);
802            match <crate::r#type::instruction::cvt::section_0::CvtFrndFtzSatDtypeAtype as PtxParser>::parse(stream) {
803                Ok(inst) => return Ok(Inst::CvtFrndFtzSatDtypeAtype(inst)),
804                Err(_) => {}
805            }
806            stream.set_position(start_pos);
807            match <crate::r#type::instruction::cvt::section_0::CvtFrnd2ReluSatfiniteF16F32 as PtxParser>::parse(stream) {
808                Ok(inst) => return Ok(Inst::CvtFrnd2ReluSatfiniteF16F32(inst)),
809                Err(_) => {}
810            }
811            stream.set_position(start_pos);
812            match <crate::r#type::instruction::cvt::section_0::CvtFrnd2ReluSatfiniteF16x2F32 as PtxParser>::parse(stream) {
813                Ok(inst) => return Ok(Inst::CvtFrnd2ReluSatfiniteF16x2F32(inst)),
814                Err(_) => {}
815            }
816            stream.set_position(start_pos);
817            match <crate::r#type::instruction::cvt::section_0::CvtRsReluSatfiniteF16x2F32 as PtxParser>::parse(stream) {
818                Ok(inst) => return Ok(Inst::CvtRsReluSatfiniteF16x2F32(inst)),
819                Err(_) => {}
820            }
821            stream.set_position(start_pos);
822            match <crate::r#type::instruction::cvt::section_0::CvtFrnd2ReluSatfiniteBf16F32 as PtxParser>::parse(stream) {
823                Ok(inst) => return Ok(Inst::CvtFrnd2ReluSatfiniteBf16F32(inst)),
824                Err(_) => {}
825            }
826            stream.set_position(start_pos);
827            match <crate::r#type::instruction::cvt::section_0::CvtFrnd2ReluSatfiniteBf16x2F32 as PtxParser>::parse(stream) {
828                Ok(inst) => return Ok(Inst::CvtFrnd2ReluSatfiniteBf16x2F32(inst)),
829                Err(_) => {}
830            }
831            stream.set_position(start_pos);
832            match <crate::r#type::instruction::cvt::section_0::CvtRsReluSatfiniteBf16x2F32 as PtxParser>::parse(stream) {
833                Ok(inst) => return Ok(Inst::CvtRsReluSatfiniteBf16x2F32(inst)),
834                Err(_) => {}
835            }
836            stream.set_position(start_pos);
837            match <crate::r#type::instruction::cvt::section_0::CvtRnaSatfiniteTf32F32 as PtxParser>::parse(stream) {
838                Ok(inst) => return Ok(Inst::CvtRnaSatfiniteTf32F32(inst)),
839                Err(_) => {}
840            }
841            stream.set_position(start_pos);
842            match <crate::r#type::instruction::cvt::section_0::CvtFrnd2SatfiniteReluTf32F32 as PtxParser>::parse(stream) {
843                Ok(inst) => return Ok(Inst::CvtFrnd2SatfiniteReluTf32F32(inst)),
844                Err(_) => {}
845            }
846            stream.set_position(start_pos);
847            match <crate::r#type::instruction::cvt::section_0::CvtRnSatfiniteReluF8x2typeF32 as PtxParser>::parse(stream) {
848                Ok(inst) => return Ok(Inst::CvtRnSatfiniteReluF8x2typeF32(inst)),
849                Err(_) => {}
850            }
851            stream.set_position(start_pos);
852            match <crate::r#type::instruction::cvt::section_0::CvtRnSatfiniteReluF8x2typeF16x2 as PtxParser>::parse(stream) {
853                Ok(inst) => return Ok(Inst::CvtRnSatfiniteReluF8x2typeF16x2(inst)),
854                Err(_) => {}
855            }
856            stream.set_position(start_pos);
857            match <crate::r#type::instruction::cvt::section_0::CvtRnReluF16x2F8x2type as PtxParser>::parse(stream) {
858                Ok(inst) => return Ok(Inst::CvtRnReluF16x2F8x2type(inst)),
859                Err(_) => {}
860            }
861            stream.set_position(start_pos);
862            match <crate::r#type::instruction::cvt::section_0::CvtRsReluSatfiniteF8x4typeF32 as PtxParser>::parse(stream) {
863                Ok(inst) => return Ok(Inst::CvtRsReluSatfiniteF8x4typeF32(inst)),
864                Err(_) => {}
865            }
866            stream.set_position(start_pos);
867            match <crate::r#type::instruction::cvt::section_0::CvtRnSatfiniteReluF4x2typeF32 as PtxParser>::parse(stream) {
868                Ok(inst) => return Ok(Inst::CvtRnSatfiniteReluF4x2typeF32(inst)),
869                Err(_) => {}
870            }
871            stream.set_position(start_pos);
872            match <crate::r#type::instruction::cvt::section_0::CvtRnReluF16x2F4x2type as PtxParser>::parse(stream) {
873                Ok(inst) => return Ok(Inst::CvtRnReluF16x2F4x2type(inst)),
874                Err(_) => {}
875            }
876            stream.set_position(start_pos);
877            match <crate::r#type::instruction::cvt::section_0::CvtRsReluSatfiniteF4x4typeF32 as PtxParser>::parse(stream) {
878                Ok(inst) => return Ok(Inst::CvtRsReluSatfiniteF4x4typeF32(inst)),
879                Err(_) => {}
880            }
881            stream.set_position(start_pos);
882            match <crate::r#type::instruction::cvt::section_0::CvtRnSatfiniteReluF6x2typeF32 as PtxParser>::parse(stream) {
883                Ok(inst) => return Ok(Inst::CvtRnSatfiniteReluF6x2typeF32(inst)),
884                Err(_) => {}
885            }
886            stream.set_position(start_pos);
887            match <crate::r#type::instruction::cvt::section_0::CvtRnReluF16x2F6x2type as PtxParser>::parse(stream) {
888                Ok(inst) => return Ok(Inst::CvtRnReluF16x2F6x2type(inst)),
889                Err(_) => {}
890            }
891            stream.set_position(start_pos);
892            match <crate::r#type::instruction::cvt::section_0::CvtRsReluSatfiniteF6x4typeF32 as PtxParser>::parse(stream) {
893                Ok(inst) => return Ok(Inst::CvtRsReluSatfiniteF6x4typeF32(inst)),
894                Err(_) => {}
895            }
896            stream.set_position(start_pos);
897            match <crate::r#type::instruction::cvt::section_0::CvtFrnd3SatfiniteUe8m0x2F32 as PtxParser>::parse(stream) {
898                Ok(inst) => return Ok(Inst::CvtFrnd3SatfiniteUe8m0x2F32(inst)),
899                Err(_) => {}
900            }
901            stream.set_position(start_pos);
902            match <crate::r#type::instruction::cvt::section_0::CvtFrnd3SatfiniteUe8m0x2Bf16x2 as PtxParser>::parse(stream) {
903                Ok(inst) => return Ok(Inst::CvtFrnd3SatfiniteUe8m0x2Bf16x2(inst)),
904                Err(_) => {}
905            }
906            stream.set_position(start_pos);
907            match <crate::r#type::instruction::cvt::section_0::CvtRnBf16x2Ue8m0x2 as PtxParser>::parse(stream) {
908                Ok(inst) => return Ok(Inst::CvtRnBf16x2Ue8m0x2(inst)),
909                Err(_) => {}
910            }
911        }
912        "cvta" => {
913            stream.set_position(start_pos);
914            match <crate::r#type::instruction::cvta::section_0::CvtaSpaceSize as PtxParser>::parse(stream) {
915                Ok(inst) => return Ok(Inst::CvtaSpaceSize(inst)),
916                Err(_) => {}
917            }
918            stream.set_position(start_pos);
919            match <crate::r#type::instruction::cvta::section_0::CvtaToSpaceSize as PtxParser>::parse(stream) {
920                Ok(inst) => return Ok(Inst::CvtaToSpaceSize(inst)),
921                Err(_) => {}
922            }
923        }
924        "discard" => {
925            stream.set_position(start_pos);
926            match <crate::r#type::instruction::discard::section_0::DiscardGlobalLevel as PtxParser>::parse(stream) {
927                Ok(inst) => return Ok(Inst::DiscardGlobalLevel(inst)),
928                Err(_) => {}
929            }
930        }
931        "div" => {
932            stream.set_position(start_pos);
933            match <crate::r#type::instruction::div::section_0::DivType as PtxParser>::parse(stream) {
934                Ok(inst) => return Ok(Inst::DivType(inst)),
935                Err(_) => {}
936            }
937            stream.set_position(start_pos);
938            match <crate::r#type::instruction::div::section_0::DivApproxFtzF32 as PtxParser>::parse(stream) {
939                Ok(inst) => return Ok(Inst::DivApproxFtzF32(inst)),
940                Err(_) => {}
941            }
942            stream.set_position(start_pos);
943            match <crate::r#type::instruction::div::section_0::DivFullFtzF32 as PtxParser>::parse(stream) {
944                Ok(inst) => return Ok(Inst::DivFullFtzF32(inst)),
945                Err(_) => {}
946            }
947            stream.set_position(start_pos);
948            match <crate::r#type::instruction::div::section_0::DivRndFtzF32 as PtxParser>::parse(stream) {
949                Ok(inst) => return Ok(Inst::DivRndFtzF32(inst)),
950                Err(_) => {}
951            }
952            stream.set_position(start_pos);
953            match <crate::r#type::instruction::div::section_0::DivRndF64 as PtxParser>::parse(stream) {
954                Ok(inst) => return Ok(Inst::DivRndF64(inst)),
955                Err(_) => {}
956            }
957        }
958        "dp2a" => {
959            stream.set_position(start_pos);
960            match <crate::r#type::instruction::dp2a::section_0::Dp2aModeAtypeBtype as PtxParser>::parse(stream) {
961                Ok(inst) => return Ok(Inst::Dp2aModeAtypeBtype(inst)),
962                Err(_) => {}
963            }
964        }
965        "dp4a" => {
966            stream.set_position(start_pos);
967            match <crate::r#type::instruction::dp4a::section_0::Dp4aAtypeBtype as PtxParser>::parse(stream) {
968                Ok(inst) => return Ok(Inst::Dp4aAtypeBtype(inst)),
969                Err(_) => {}
970            }
971        }
972        "elect" => {
973            stream.set_position(start_pos);
974            match <crate::r#type::instruction::elect_sync::section_0::ElectSync as PtxParser>::parse(stream) {
975                Ok(inst) => return Ok(Inst::ElectSync(inst)),
976                Err(_) => {}
977            }
978        }
979        "ex2" => {
980            stream.set_position(start_pos);
981            match <crate::r#type::instruction::ex2::section_0::Ex2ApproxFtzF32 as PtxParser>::parse(stream) {
982                Ok(inst) => return Ok(Inst::Ex2ApproxFtzF32(inst)),
983                Err(_) => {}
984            }
985            stream.set_position(start_pos);
986            match <crate::r#type::instruction::ex2::section_0::Ex2ApproxAtype as PtxParser>::parse(stream) {
987                Ok(inst) => return Ok(Inst::Ex2ApproxAtype(inst)),
988                Err(_) => {}
989            }
990            stream.set_position(start_pos);
991            match <crate::r#type::instruction::ex2::section_0::Ex2ApproxFtzBtype as PtxParser>::parse(stream) {
992                Ok(inst) => return Ok(Inst::Ex2ApproxFtzBtype(inst)),
993                Err(_) => {}
994            }
995        }
996        "exit" => {
997            stream.set_position(start_pos);
998            match <crate::r#type::instruction::exit::section_0::Exit as PtxParser>::parse(stream) {
999                Ok(inst) => return Ok(Inst::Exit(inst)),
1000                Err(_) => {}
1001            }
1002        }
1003        "fence" => {
1004            stream.set_position(start_pos);
1005            match <crate::r#type::instruction::membar::section_0::FenceSemScope as PtxParser>::parse(stream) {
1006                Ok(inst) => return Ok(Inst::FenceSemScope(inst)),
1007                Err(_) => {}
1008            }
1009            stream.set_position(start_pos);
1010            match <crate::r#type::instruction::membar::section_0::FenceAcquireSyncRestrictSharedClusterCluster as PtxParser>::parse(stream) {
1011                Ok(inst) => return Ok(Inst::FenceAcquireSyncRestrictSharedClusterCluster(inst)),
1012                Err(_) => {}
1013            }
1014            stream.set_position(start_pos);
1015            match <crate::r#type::instruction::membar::section_0::FenceReleaseSyncRestrictSharedCtaCluster as PtxParser>::parse(stream) {
1016                Ok(inst) => return Ok(Inst::FenceReleaseSyncRestrictSharedCtaCluster(inst)),
1017                Err(_) => {}
1018            }
1019            stream.set_position(start_pos);
1020            match <crate::r#type::instruction::membar::section_0::FenceOpRestrictReleaseCluster as PtxParser>::parse(stream) {
1021                Ok(inst) => return Ok(Inst::FenceOpRestrictReleaseCluster(inst)),
1022                Err(_) => {}
1023            }
1024            stream.set_position(start_pos);
1025            match <crate::r#type::instruction::membar::section_0::FenceProxyProxykind as PtxParser>::parse(stream) {
1026                Ok(inst) => return Ok(Inst::FenceProxyProxykind(inst)),
1027                Err(_) => {}
1028            }
1029            stream.set_position(start_pos);
1030            match <crate::r#type::instruction::membar::section_0::FenceProxyToProxykindFromProxykindReleaseScope as PtxParser>::parse(stream) {
1031                Ok(inst) => return Ok(Inst::FenceProxyToProxykindFromProxykindReleaseScope(inst)),
1032                Err(_) => {}
1033            }
1034            stream.set_position(start_pos);
1035            match <crate::r#type::instruction::membar::section_0::FenceProxyToProxykindFromProxykindAcquireScope as PtxParser>::parse(stream) {
1036                Ok(inst) => return Ok(Inst::FenceProxyToProxykindFromProxykindAcquireScope(inst)),
1037                Err(_) => {}
1038            }
1039            stream.set_position(start_pos);
1040            match <crate::r#type::instruction::membar::section_0::FenceProxyAsyncGenericAcquireSyncRestrictSharedClusterCluster as PtxParser>::parse(stream) {
1041                Ok(inst) => return Ok(Inst::FenceProxyAsyncGenericAcquireSyncRestrictSharedClusterCluster(inst)),
1042                Err(_) => {}
1043            }
1044            stream.set_position(start_pos);
1045            match <crate::r#type::instruction::membar::section_0::FenceProxyAsyncGenericReleaseSyncRestrictSharedCtaCluster as PtxParser>::parse(stream) {
1046                Ok(inst) => return Ok(Inst::FenceProxyAsyncGenericReleaseSyncRestrictSharedCtaCluster(inst)),
1047                Err(_) => {}
1048            }
1049            stream.set_position(start_pos);
1050            match <crate::r#type::instruction::membar::section_0::MembarLevel as PtxParser>::parse(stream) {
1051                Ok(inst) => return Ok(Inst::MembarLevel(inst)),
1052                Err(_) => {}
1053            }
1054            stream.set_position(start_pos);
1055            match <crate::r#type::instruction::membar::section_0::MembarProxyProxykind as PtxParser>::parse(stream) {
1056                Ok(inst) => return Ok(Inst::MembarProxyProxykind(inst)),
1057                Err(_) => {}
1058            }
1059        }
1060        "fma" => {
1061            stream.set_position(start_pos);
1062            match <crate::r#type::instruction::fma::section_0::FmaRndFtzSatF32 as PtxParser>::parse(stream) {
1063                Ok(inst) => return Ok(Inst::FmaRndFtzSatF32(inst)),
1064                Err(_) => {}
1065            }
1066            stream.set_position(start_pos);
1067            match <crate::r#type::instruction::fma::section_0::FmaRndFtzF32x2 as PtxParser>::parse(stream) {
1068                Ok(inst) => return Ok(Inst::FmaRndFtzF32x2(inst)),
1069                Err(_) => {}
1070            }
1071            stream.set_position(start_pos);
1072            match <crate::r#type::instruction::fma::section_0::FmaRndF64 as PtxParser>::parse(stream) {
1073                Ok(inst) => return Ok(Inst::FmaRndF64(inst)),
1074                Err(_) => {}
1075            }
1076            stream.set_position(start_pos);
1077            match <crate::r#type::instruction::fma::section_1::FmaRndFtzSatF16 as PtxParser>::parse(stream) {
1078                Ok(inst) => return Ok(Inst::FmaRndFtzSatF16(inst)),
1079                Err(_) => {}
1080            }
1081            stream.set_position(start_pos);
1082            match <crate::r#type::instruction::fma::section_1::FmaRndFtzSatF16x2 as PtxParser>::parse(stream) {
1083                Ok(inst) => return Ok(Inst::FmaRndFtzSatF16x2(inst)),
1084                Err(_) => {}
1085            }
1086            stream.set_position(start_pos);
1087            match <crate::r#type::instruction::fma::section_1::FmaRndFtzReluF16 as PtxParser>::parse(stream) {
1088                Ok(inst) => return Ok(Inst::FmaRndFtzReluF16(inst)),
1089                Err(_) => {}
1090            }
1091            stream.set_position(start_pos);
1092            match <crate::r#type::instruction::fma::section_1::FmaRndFtzReluF16x2 as PtxParser>::parse(stream) {
1093                Ok(inst) => return Ok(Inst::FmaRndFtzReluF16x2(inst)),
1094                Err(_) => {}
1095            }
1096            stream.set_position(start_pos);
1097            match <crate::r#type::instruction::fma::section_1::FmaRndReluBf16 as PtxParser>::parse(stream) {
1098                Ok(inst) => return Ok(Inst::FmaRndReluBf16(inst)),
1099                Err(_) => {}
1100            }
1101            stream.set_position(start_pos);
1102            match <crate::r#type::instruction::fma::section_1::FmaRndReluBf16x2 as PtxParser>::parse(stream) {
1103                Ok(inst) => return Ok(Inst::FmaRndReluBf16x2(inst)),
1104                Err(_) => {}
1105            }
1106            stream.set_position(start_pos);
1107            match <crate::r#type::instruction::fma::section_1::FmaRndOobReluType as PtxParser>::parse(stream) {
1108                Ok(inst) => return Ok(Inst::FmaRndOobReluType(inst)),
1109                Err(_) => {}
1110            }
1111            stream.set_position(start_pos);
1112            match <crate::r#type::instruction::fma::section_2::FmaRndSatF32Abtype as PtxParser>::parse(stream) {
1113                Ok(inst) => return Ok(Inst::FmaRndSatF32Abtype(inst)),
1114                Err(_) => {}
1115            }
1116        }
1117        "fns" => {
1118            stream.set_position(start_pos);
1119            match <crate::r#type::instruction::fns::section_0::FnsB32 as PtxParser>::parse(stream) {
1120                Ok(inst) => return Ok(Inst::FnsB32(inst)),
1121                Err(_) => {}
1122            }
1123        }
1124        "getctarank" => {
1125            stream.set_position(start_pos);
1126            match <crate::r#type::instruction::getctarank::section_0::GetctarankSpaceType as PtxParser>::parse(stream) {
1127                Ok(inst) => return Ok(Inst::GetctarankSpaceType(inst)),
1128                Err(_) => {}
1129            }
1130            stream.set_position(start_pos);
1131            match <crate::r#type::instruction::getctarank::section_0::GetctarankSharedClusterType as PtxParser>::parse(stream) {
1132                Ok(inst) => return Ok(Inst::GetctarankSharedClusterType(inst)),
1133                Err(_) => {}
1134            }
1135            stream.set_position(start_pos);
1136            match <crate::r#type::instruction::getctarank::section_0::GetctarankType as PtxParser>::parse(stream) {
1137                Ok(inst) => return Ok(Inst::GetctarankType(inst)),
1138                Err(_) => {}
1139            }
1140        }
1141        "griddepcontrol" => {
1142            stream.set_position(start_pos);
1143            match <crate::r#type::instruction::griddepcontrol::section_0::GriddepcontrolAction as PtxParser>::parse(stream) {
1144                Ok(inst) => return Ok(Inst::GriddepcontrolAction(inst)),
1145                Err(_) => {}
1146            }
1147        }
1148        "isspacep" => {
1149            stream.set_position(start_pos);
1150            match <crate::r#type::instruction::isspacep::section_0::IsspacepSpace as PtxParser>::parse(stream) {
1151                Ok(inst) => return Ok(Inst::IsspacepSpace(inst)),
1152                Err(_) => {}
1153            }
1154        }
1155        "istypep" => {
1156            stream.set_position(start_pos);
1157            match <crate::r#type::instruction::istypep::section_0::IstypepType as PtxParser>::parse(stream) {
1158                Ok(inst) => return Ok(Inst::IstypepType(inst)),
1159                Err(_) => {}
1160            }
1161        }
1162        "ld" => {
1163            stream.set_position(start_pos);
1164            match <crate::r#type::instruction::ld_global_nc::section_0::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType as PtxParser>::parse(stream) {
1165                Ok(inst) => return Ok(Inst::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeType(inst)),
1166                Err(_) => {}
1167            }
1168            stream.set_position(start_pos);
1169            match <crate::r#type::instruction::ld_global_nc::section_0::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1170                Ok(inst) => return Ok(Inst::LdGlobalCopNcLevelCacheHintLevelPrefetchSizeVecType(inst)),
1171                Err(_) => {}
1172            }
1173            stream.set_position(start_pos);
1174            match <crate::r#type::instruction::ld_global_nc::section_0::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType as PtxParser>::parse(stream) {
1175                Ok(inst) => return Ok(Inst::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeType(inst)),
1176                Err(_) => {}
1177            }
1178            stream.set_position(start_pos);
1179            match <crate::r#type::instruction::ld_global_nc::section_0::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1180                Ok(inst) => return Ok(Inst::LdGlobalNcLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(inst)),
1181                Err(_) => {}
1182            }
1183            stream.set_position(start_pos);
1184            match <crate::r#type::instruction::ld::section_0::LdWeakSsCopLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1185                Ok(inst) => return Ok(Inst::LdWeakSsCopLevelCacheHintLevelPrefetchSizeVecType(inst)),
1186                Err(_) => {}
1187            }
1188            stream.set_position(start_pos);
1189            match <crate::r#type::instruction::ld::section_0::LdWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1190                Ok(inst) => return Ok(Inst::LdWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(inst)),
1191                Err(_) => {}
1192            }
1193            stream.set_position(start_pos);
1194            match <crate::r#type::instruction::ld::section_0::LdVolatileSsLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1195                Ok(inst) => return Ok(Inst::LdVolatileSsLevelPrefetchSizeVecType(inst)),
1196                Err(_) => {}
1197            }
1198            stream.set_position(start_pos);
1199            match <crate::r#type::instruction::ld::section_0::LdRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1200                Ok(inst) => return Ok(Inst::LdRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(inst)),
1201                Err(_) => {}
1202            }
1203            stream.set_position(start_pos);
1204            match <crate::r#type::instruction::ld::section_0::LdAcquireScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType as PtxParser>::parse(stream) {
1205                Ok(inst) => return Ok(Inst::LdAcquireScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintLevelPrefetchSizeVecType(inst)),
1206                Err(_) => {}
1207            }
1208            stream.set_position(start_pos);
1209            match <crate::r#type::instruction::ld::section_0::LdMmioRelaxedSysGlobalType as PtxParser>::parse(stream) {
1210                Ok(inst) => return Ok(Inst::LdMmioRelaxedSysGlobalType(inst)),
1211                Err(_) => {}
1212            }
1213        }
1214        "ldmatrix" => {
1215            stream.set_position(start_pos);
1216            match <crate::r#type::instruction::ldmatrix::section_0::LdmatrixSyncAlignedShapeNumTransSsType as PtxParser>::parse(stream) {
1217                Ok(inst) => return Ok(Inst::LdmatrixSyncAlignedShapeNumTransSsType(inst)),
1218                Err(_) => {}
1219            }
1220            stream.set_position(start_pos);
1221            match <crate::r#type::instruction::ldmatrix::section_0::LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt as PtxParser>::parse(stream) {
1222                Ok(inst) => return Ok(Inst::LdmatrixSyncAlignedM8n16NumSsDstFmtSrcFmt(inst)),
1223                Err(_) => {}
1224            }
1225            stream.set_position(start_pos);
1226            match <crate::r#type::instruction::ldmatrix::section_0::LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt as PtxParser>::parse(stream) {
1227                Ok(inst) => return Ok(Inst::LdmatrixSyncAlignedM16n16NumTransSsDstFmtSrcFmt(inst)),
1228                Err(_) => {}
1229            }
1230        }
1231        "ldu" => {
1232            stream.set_position(start_pos);
1233            match <crate::r#type::instruction::ldu::section_0::LduSsType as PtxParser>::parse(stream) {
1234                Ok(inst) => return Ok(Inst::LduSsType(inst)),
1235                Err(_) => {}
1236            }
1237            stream.set_position(start_pos);
1238            match <crate::r#type::instruction::ldu::section_0::LduSsVecType as PtxParser>::parse(stream) {
1239                Ok(inst) => return Ok(Inst::LduSsVecType(inst)),
1240                Err(_) => {}
1241            }
1242        }
1243        "lg2" => {
1244            stream.set_position(start_pos);
1245            match <crate::r#type::instruction::lg2::section_0::Lg2ApproxFtzF32 as PtxParser>::parse(stream) {
1246                Ok(inst) => return Ok(Inst::Lg2ApproxFtzF32(inst)),
1247                Err(_) => {}
1248            }
1249        }
1250        "lop3" => {
1251            stream.set_position(start_pos);
1252            match <crate::r#type::instruction::lop3::section_0::Lop3B32 as PtxParser>::parse(stream) {
1253                Ok(inst) => return Ok(Inst::Lop3B32(inst)),
1254                Err(_) => {}
1255            }
1256            stream.set_position(start_pos);
1257            match <crate::r#type::instruction::lop3::section_0::Lop3BoolopB32 as PtxParser>::parse(stream) {
1258                Ok(inst) => return Ok(Inst::Lop3BoolopB32(inst)),
1259                Err(_) => {}
1260            }
1261        }
1262        "mad" => {
1263            stream.set_position(start_pos);
1264            match <crate::r#type::instruction::mad_cc::section_0::MadHiloCcType as PtxParser>::parse(stream) {
1265                Ok(inst) => return Ok(Inst::MadHiloCcType(inst)),
1266                Err(_) => {}
1267            }
1268            stream.set_position(start_pos);
1269            match <crate::r#type::instruction::mad::section_0::MadModeType as PtxParser>::parse(stream) {
1270                Ok(inst) => return Ok(Inst::MadModeType(inst)),
1271                Err(_) => {}
1272            }
1273            stream.set_position(start_pos);
1274            match <crate::r#type::instruction::mad::section_0::MadHiSatS32 as PtxParser>::parse(stream) {
1275                Ok(inst) => return Ok(Inst::MadHiSatS32(inst)),
1276                Err(_) => {}
1277            }
1278            stream.set_position(start_pos);
1279            match <crate::r#type::instruction::mad::section_0::MadFtzSatF32 as PtxParser>::parse(stream) {
1280                Ok(inst) => return Ok(Inst::MadFtzSatF32(inst)),
1281                Err(_) => {}
1282            }
1283            stream.set_position(start_pos);
1284            match <crate::r#type::instruction::mad::section_0::MadRndFtzSatF32 as PtxParser>::parse(stream) {
1285                Ok(inst) => return Ok(Inst::MadRndFtzSatF32(inst)),
1286                Err(_) => {}
1287            }
1288            stream.set_position(start_pos);
1289            match <crate::r#type::instruction::mad::section_0::MadRndF64 as PtxParser>::parse(stream) {
1290                Ok(inst) => return Ok(Inst::MadRndF64(inst)),
1291                Err(_) => {}
1292            }
1293        }
1294        "mad24" => {
1295            stream.set_position(start_pos);
1296            match <crate::r#type::instruction::mad24::section_0::Mad24ModeType as PtxParser>::parse(stream) {
1297                Ok(inst) => return Ok(Inst::Mad24ModeType(inst)),
1298                Err(_) => {}
1299            }
1300            stream.set_position(start_pos);
1301            match <crate::r#type::instruction::mad24::section_0::Mad24HiSatS32 as PtxParser>::parse(stream) {
1302                Ok(inst) => return Ok(Inst::Mad24HiSatS32(inst)),
1303                Err(_) => {}
1304            }
1305        }
1306        "madc" => {
1307            stream.set_position(start_pos);
1308            match <crate::r#type::instruction::madc::section_0::MadcHiloCcType as PtxParser>::parse(stream) {
1309                Ok(inst) => return Ok(Inst::MadcHiloCcType(inst)),
1310                Err(_) => {}
1311            }
1312        }
1313        "mapa" => {
1314            stream.set_position(start_pos);
1315            match <crate::r#type::instruction::mapa::section_0::MapaSpaceType as PtxParser>::parse(stream) {
1316                Ok(inst) => return Ok(Inst::MapaSpaceType(inst)),
1317                Err(_) => {}
1318            }
1319        }
1320        "match" => {
1321            stream.set_position(start_pos);
1322            match <crate::r#type::instruction::match_sync::section_0::MatchAnySyncType as PtxParser>::parse(stream) {
1323                Ok(inst) => return Ok(Inst::MatchAnySyncType(inst)),
1324                Err(_) => {}
1325            }
1326            stream.set_position(start_pos);
1327            match <crate::r#type::instruction::match_sync::section_0::MatchAllSyncType as PtxParser>::parse(stream) {
1328                Ok(inst) => return Ok(Inst::MatchAllSyncType(inst)),
1329                Err(_) => {}
1330            }
1331        }
1332        "max" => {
1333            stream.set_position(start_pos);
1334            match <crate::r#type::instruction::max::section_0::MaxAtype as PtxParser>::parse(stream) {
1335                Ok(inst) => return Ok(Inst::MaxAtype(inst)),
1336                Err(_) => {}
1337            }
1338            stream.set_position(start_pos);
1339            match <crate::r#type::instruction::max::section_0::MaxReluBtype as PtxParser>::parse(stream) {
1340                Ok(inst) => return Ok(Inst::MaxReluBtype(inst)),
1341                Err(_) => {}
1342            }
1343            stream.set_position(start_pos);
1344            match <crate::r#type::instruction::max::section_0::MaxFtzNanXorsignAbsF32 as PtxParser>::parse(stream) {
1345                Ok(inst) => return Ok(Inst::MaxFtzNanXorsignAbsF32(inst)),
1346                Err(_) => {}
1347            }
1348            stream.set_position(start_pos);
1349            match <crate::r#type::instruction::max::section_0::MaxFtzNanAbsF32 as PtxParser>::parse(stream) {
1350                Ok(inst) => return Ok(Inst::MaxFtzNanAbsF32(inst)),
1351                Err(_) => {}
1352            }
1353            stream.set_position(start_pos);
1354            match <crate::r#type::instruction::max::section_0::MaxF64 as PtxParser>::parse(stream) {
1355                Ok(inst) => return Ok(Inst::MaxF64(inst)),
1356                Err(_) => {}
1357            }
1358            stream.set_position(start_pos);
1359            match <crate::r#type::instruction::max::section_0::MaxFtzNanXorsignAbsF16 as PtxParser>::parse(stream) {
1360                Ok(inst) => return Ok(Inst::MaxFtzNanXorsignAbsF16(inst)),
1361                Err(_) => {}
1362            }
1363            stream.set_position(start_pos);
1364            match <crate::r#type::instruction::max::section_0::MaxFtzNanXorsignAbsF16x2 as PtxParser>::parse(stream) {
1365                Ok(inst) => return Ok(Inst::MaxFtzNanXorsignAbsF16x2(inst)),
1366                Err(_) => {}
1367            }
1368            stream.set_position(start_pos);
1369            match <crate::r#type::instruction::max::section_0::MaxNanXorsignAbsBf16 as PtxParser>::parse(stream) {
1370                Ok(inst) => return Ok(Inst::MaxNanXorsignAbsBf16(inst)),
1371                Err(_) => {}
1372            }
1373            stream.set_position(start_pos);
1374            match <crate::r#type::instruction::max::section_0::MaxNanXorsignAbsBf16x2 as PtxParser>::parse(stream) {
1375                Ok(inst) => return Ok(Inst::MaxNanXorsignAbsBf16x2(inst)),
1376                Err(_) => {}
1377            }
1378        }
1379        "mbarrier" => {
1380            stream.set_position(start_pos);
1381            match <crate::r#type::instruction::mbarrier_arrive::section_0::MbarrierArriveSemScopeStateB64 as PtxParser>::parse(stream) {
1382                Ok(inst) => return Ok(Inst::MbarrierArriveSemScopeStateB64(inst)),
1383                Err(_) => {}
1384            }
1385            stream.set_position(start_pos);
1386            match <crate::r#type::instruction::mbarrier_arrive::section_0::MbarrierArriveSemScopeSharedClusterB64 as PtxParser>::parse(stream) {
1387                Ok(inst) => return Ok(Inst::MbarrierArriveSemScopeSharedClusterB64(inst)),
1388                Err(_) => {}
1389            }
1390            stream.set_position(start_pos);
1391            match <crate::r#type::instruction::mbarrier_arrive::section_0::MbarrierArriveExpectTxSemScopeStateB64 as PtxParser>::parse(stream) {
1392                Ok(inst) => return Ok(Inst::MbarrierArriveExpectTxSemScopeStateB64(inst)),
1393                Err(_) => {}
1394            }
1395            stream.set_position(start_pos);
1396            match <crate::r#type::instruction::mbarrier_arrive::section_0::MbarrierArriveExpectTxSemScopeSharedClusterB64 as PtxParser>::parse(stream) {
1397                Ok(inst) => return Ok(Inst::MbarrierArriveExpectTxSemScopeSharedClusterB64(inst)),
1398                Err(_) => {}
1399            }
1400            stream.set_position(start_pos);
1401            match <crate::r#type::instruction::mbarrier_arrive::section_0::MbarrierArriveNocompleteReleaseCtaStateB64 as PtxParser>::parse(stream) {
1402                Ok(inst) => return Ok(Inst::MbarrierArriveNocompleteReleaseCtaStateB64(inst)),
1403                Err(_) => {}
1404            }
1405            stream.set_position(start_pos);
1406            match <crate::r#type::instruction::mbarrier_arrive_drop::section_0::MbarrierArriveDropSemScopeStateB64 as PtxParser>::parse(stream) {
1407                Ok(inst) => return Ok(Inst::MbarrierArriveDropSemScopeStateB64(inst)),
1408                Err(_) => {}
1409            }
1410            stream.set_position(start_pos);
1411            match <crate::r#type::instruction::mbarrier_arrive_drop::section_0::MbarrierArriveDropSemScopeSharedClusterB64 as PtxParser>::parse(stream) {
1412                Ok(inst) => return Ok(Inst::MbarrierArriveDropSemScopeSharedClusterB64(inst)),
1413                Err(_) => {}
1414            }
1415            stream.set_position(start_pos);
1416            match <crate::r#type::instruction::mbarrier_arrive_drop::section_0::MbarrierArriveDropExpectTxStateSemScopeB64 as PtxParser>::parse(stream) {
1417                Ok(inst) => return Ok(Inst::MbarrierArriveDropExpectTxStateSemScopeB64(inst)),
1418                Err(_) => {}
1419            }
1420            stream.set_position(start_pos);
1421            match <crate::r#type::instruction::mbarrier_arrive_drop::section_0::MbarrierArriveDropExpectTxSharedClusterSemScopeB64 as PtxParser>::parse(stream) {
1422                Ok(inst) => return Ok(Inst::MbarrierArriveDropExpectTxSharedClusterSemScopeB64(inst)),
1423                Err(_) => {}
1424            }
1425            stream.set_position(start_pos);
1426            match <crate::r#type::instruction::mbarrier_arrive_drop::section_0::MbarrierArriveDropNocompleteReleaseCtaStateB64 as PtxParser>::parse(stream) {
1427                Ok(inst) => return Ok(Inst::MbarrierArriveDropNocompleteReleaseCtaStateB64(inst)),
1428                Err(_) => {}
1429            }
1430            stream.set_position(start_pos);
1431            match <crate::r#type::instruction::mbarrier_complete_tx::section_0::MbarrierCompleteTxSemScopeSpaceB64 as PtxParser>::parse(stream) {
1432                Ok(inst) => return Ok(Inst::MbarrierCompleteTxSemScopeSpaceB64(inst)),
1433                Err(_) => {}
1434            }
1435            stream.set_position(start_pos);
1436            match <crate::r#type::instruction::mbarrier_expect_tx::section_0::MbarrierExpectTxSemScopeSpaceB64 as PtxParser>::parse(stream) {
1437                Ok(inst) => return Ok(Inst::MbarrierExpectTxSemScopeSpaceB64(inst)),
1438                Err(_) => {}
1439            }
1440            stream.set_position(start_pos);
1441            match <crate::r#type::instruction::mbarrier_init::section_0::MbarrierInitStateB64 as PtxParser>::parse(stream) {
1442                Ok(inst) => return Ok(Inst::MbarrierInitStateB64(inst)),
1443                Err(_) => {}
1444            }
1445            stream.set_position(start_pos);
1446            match <crate::r#type::instruction::mbarrier_inval::section_0::MbarrierInvalStateB64 as PtxParser>::parse(stream) {
1447                Ok(inst) => return Ok(Inst::MbarrierInvalStateB64(inst)),
1448                Err(_) => {}
1449            }
1450            stream.set_position(start_pos);
1451            match <crate::r#type::instruction::mbarrier_pending_count::section_0::MbarrierPendingCountB64 as PtxParser>::parse(stream) {
1452                Ok(inst) => return Ok(Inst::MbarrierPendingCountB64(inst)),
1453                Err(_) => {}
1454            }
1455            stream.set_position(start_pos);
1456            match <crate::r#type::instruction::mbarrier_test_wait::section_0::MbarrierTestWaitSemScopeStateB64 as PtxParser>::parse(stream) {
1457                Ok(inst) => return Ok(Inst::MbarrierTestWaitSemScopeStateB64(inst)),
1458                Err(_) => {}
1459            }
1460            stream.set_position(start_pos);
1461            match <crate::r#type::instruction::mbarrier_test_wait::section_0::MbarrierTestWaitParitySemScopeStateB64 as PtxParser>::parse(stream) {
1462                Ok(inst) => return Ok(Inst::MbarrierTestWaitParitySemScopeStateB64(inst)),
1463                Err(_) => {}
1464            }
1465            stream.set_position(start_pos);
1466            match <crate::r#type::instruction::mbarrier_test_wait::section_0::MbarrierTryWaitSemScopeStateB64 as PtxParser>::parse(stream) {
1467                Ok(inst) => return Ok(Inst::MbarrierTryWaitSemScopeStateB64(inst)),
1468                Err(_) => {}
1469            }
1470            stream.set_position(start_pos);
1471            match <crate::r#type::instruction::mbarrier_test_wait::section_0::MbarrierTryWaitParitySemScopeStateB64 as PtxParser>::parse(stream) {
1472                Ok(inst) => return Ok(Inst::MbarrierTryWaitParitySemScopeStateB64(inst)),
1473                Err(_) => {}
1474            }
1475        }
1476        "membar" => {
1477            stream.set_position(start_pos);
1478            match <crate::r#type::instruction::membar::section_0::FenceSemScope as PtxParser>::parse(stream) {
1479                Ok(inst) => return Ok(Inst::FenceSemScope(inst)),
1480                Err(_) => {}
1481            }
1482            stream.set_position(start_pos);
1483            match <crate::r#type::instruction::membar::section_0::FenceAcquireSyncRestrictSharedClusterCluster as PtxParser>::parse(stream) {
1484                Ok(inst) => return Ok(Inst::FenceAcquireSyncRestrictSharedClusterCluster(inst)),
1485                Err(_) => {}
1486            }
1487            stream.set_position(start_pos);
1488            match <crate::r#type::instruction::membar::section_0::FenceReleaseSyncRestrictSharedCtaCluster as PtxParser>::parse(stream) {
1489                Ok(inst) => return Ok(Inst::FenceReleaseSyncRestrictSharedCtaCluster(inst)),
1490                Err(_) => {}
1491            }
1492            stream.set_position(start_pos);
1493            match <crate::r#type::instruction::membar::section_0::FenceOpRestrictReleaseCluster as PtxParser>::parse(stream) {
1494                Ok(inst) => return Ok(Inst::FenceOpRestrictReleaseCluster(inst)),
1495                Err(_) => {}
1496            }
1497            stream.set_position(start_pos);
1498            match <crate::r#type::instruction::membar::section_0::FenceProxyProxykind as PtxParser>::parse(stream) {
1499                Ok(inst) => return Ok(Inst::FenceProxyProxykind(inst)),
1500                Err(_) => {}
1501            }
1502            stream.set_position(start_pos);
1503            match <crate::r#type::instruction::membar::section_0::FenceProxyToProxykindFromProxykindReleaseScope as PtxParser>::parse(stream) {
1504                Ok(inst) => return Ok(Inst::FenceProxyToProxykindFromProxykindReleaseScope(inst)),
1505                Err(_) => {}
1506            }
1507            stream.set_position(start_pos);
1508            match <crate::r#type::instruction::membar::section_0::FenceProxyToProxykindFromProxykindAcquireScope as PtxParser>::parse(stream) {
1509                Ok(inst) => return Ok(Inst::FenceProxyToProxykindFromProxykindAcquireScope(inst)),
1510                Err(_) => {}
1511            }
1512            stream.set_position(start_pos);
1513            match <crate::r#type::instruction::membar::section_0::FenceProxyAsyncGenericAcquireSyncRestrictSharedClusterCluster as PtxParser>::parse(stream) {
1514                Ok(inst) => return Ok(Inst::FenceProxyAsyncGenericAcquireSyncRestrictSharedClusterCluster(inst)),
1515                Err(_) => {}
1516            }
1517            stream.set_position(start_pos);
1518            match <crate::r#type::instruction::membar::section_0::FenceProxyAsyncGenericReleaseSyncRestrictSharedCtaCluster as PtxParser>::parse(stream) {
1519                Ok(inst) => return Ok(Inst::FenceProxyAsyncGenericReleaseSyncRestrictSharedCtaCluster(inst)),
1520                Err(_) => {}
1521            }
1522            stream.set_position(start_pos);
1523            match <crate::r#type::instruction::membar::section_0::MembarLevel as PtxParser>::parse(stream) {
1524                Ok(inst) => return Ok(Inst::MembarLevel(inst)),
1525                Err(_) => {}
1526            }
1527            stream.set_position(start_pos);
1528            match <crate::r#type::instruction::membar::section_0::MembarProxyProxykind as PtxParser>::parse(stream) {
1529                Ok(inst) => return Ok(Inst::MembarProxyProxykind(inst)),
1530                Err(_) => {}
1531            }
1532        }
1533        "min" => {
1534            stream.set_position(start_pos);
1535            match <crate::r#type::instruction::min::section_0::MinAtype as PtxParser>::parse(stream) {
1536                Ok(inst) => return Ok(Inst::MinAtype(inst)),
1537                Err(_) => {}
1538            }
1539            stream.set_position(start_pos);
1540            match <crate::r#type::instruction::min::section_0::MinReluBtype as PtxParser>::parse(stream) {
1541                Ok(inst) => return Ok(Inst::MinReluBtype(inst)),
1542                Err(_) => {}
1543            }
1544            stream.set_position(start_pos);
1545            match <crate::r#type::instruction::min::section_0::MinFtzNanXorsignAbsF32 as PtxParser>::parse(stream) {
1546                Ok(inst) => return Ok(Inst::MinFtzNanXorsignAbsF32(inst)),
1547                Err(_) => {}
1548            }
1549            stream.set_position(start_pos);
1550            match <crate::r#type::instruction::min::section_0::MinFtzNanAbsF32 as PtxParser>::parse(stream) {
1551                Ok(inst) => return Ok(Inst::MinFtzNanAbsF32(inst)),
1552                Err(_) => {}
1553            }
1554            stream.set_position(start_pos);
1555            match <crate::r#type::instruction::min::section_0::MinF64 as PtxParser>::parse(stream) {
1556                Ok(inst) => return Ok(Inst::MinF64(inst)),
1557                Err(_) => {}
1558            }
1559            stream.set_position(start_pos);
1560            match <crate::r#type::instruction::min::section_0::MinFtzNanXorsignAbsF16 as PtxParser>::parse(stream) {
1561                Ok(inst) => return Ok(Inst::MinFtzNanXorsignAbsF16(inst)),
1562                Err(_) => {}
1563            }
1564            stream.set_position(start_pos);
1565            match <crate::r#type::instruction::min::section_0::MinFtzNanXorsignAbsF16x2 as PtxParser>::parse(stream) {
1566                Ok(inst) => return Ok(Inst::MinFtzNanXorsignAbsF16x2(inst)),
1567                Err(_) => {}
1568            }
1569            stream.set_position(start_pos);
1570            match <crate::r#type::instruction::min::section_0::MinNanXorsignAbsBf16 as PtxParser>::parse(stream) {
1571                Ok(inst) => return Ok(Inst::MinNanXorsignAbsBf16(inst)),
1572                Err(_) => {}
1573            }
1574            stream.set_position(start_pos);
1575            match <crate::r#type::instruction::min::section_0::MinNanXorsignAbsBf16x2 as PtxParser>::parse(stream) {
1576                Ok(inst) => return Ok(Inst::MinNanXorsignAbsBf16x2(inst)),
1577                Err(_) => {}
1578            }
1579        }
1580        "mma" => {
1581            stream.set_position(start_pos);
1582            match <crate::r#type::instruction::mma_sp::section_0::MmaSpvariantSyncAlignedM16n8k16RowColDtypeF16F16Ctype as PtxParser>::parse(stream) {
1583                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k16RowColDtypeF16F16Ctype(inst)),
1584                Err(_) => {}
1585            }
1586            stream.set_position(start_pos);
1587            match <crate::r#type::instruction::mma_sp::section_0::MmaSpvariantSyncAlignedM16n8k32RowColDtypeF16F16Ctype as PtxParser>::parse(stream) {
1588                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k32RowColDtypeF16F16Ctype(inst)),
1589                Err(_) => {}
1590            }
1591            stream.set_position(start_pos);
1592            match <crate::r#type::instruction::mma_sp::section_1::MmaSpvariantSyncAlignedM16n8k16RowColF32Bf16Bf16F32 as PtxParser>::parse(stream) {
1593                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k16RowColF32Bf16Bf16F32(inst)),
1594                Err(_) => {}
1595            }
1596            stream.set_position(start_pos);
1597            match <crate::r#type::instruction::mma_sp::section_1::MmaSpvariantSyncAlignedM16n8k32RowColF32Bf16Bf16F32 as PtxParser>::parse(stream) {
1598                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k32RowColF32Bf16Bf16F32(inst)),
1599                Err(_) => {}
1600            }
1601            stream.set_position(start_pos);
1602            match <crate::r#type::instruction::mma_sp::section_1::MmaSpvariantSyncAlignedM16n8k8RowColF32Tf32Tf32F32 as PtxParser>::parse(stream) {
1603                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k8RowColF32Tf32Tf32F32(inst)),
1604                Err(_) => {}
1605            }
1606            stream.set_position(start_pos);
1607            match <crate::r#type::instruction::mma_sp::section_1::MmaSpvariantSyncAlignedM16n8k16RowColF32Tf32Tf32F32 as PtxParser>::parse(stream) {
1608                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k16RowColF32Tf32Tf32F32(inst)),
1609                Err(_) => {}
1610            }
1611            stream.set_position(start_pos);
1612            match <crate::r#type::instruction::mma_sp::section_1::MmaSpvariantSyncAlignedM16n8k64RowColF32F8typeF8typeF32 as PtxParser>::parse(stream) {
1613                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k64RowColF32F8typeF8typeF32(inst)),
1614                Err(_) => {}
1615            }
1616            stream.set_position(start_pos);
1617            match <crate::r#type::instruction::mma_sp::section_1::MmaSpOrderedMetadataSyncAlignedM16n8k64RowColKindDtypeF8f6f4typeF8f6f4typeCtype as PtxParser>::parse(stream) {
1618                Ok(inst) => return Ok(Inst::MmaSpOrderedMetadataSyncAlignedM16n8k64RowColKindDtypeF8f6f4typeF8f6f4typeCtype(inst)),
1619                Err(_) => {}
1620            }
1621            stream.set_position(start_pos);
1622            match <crate::r#type::instruction::mma_sp::section_2::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype as PtxParser>::parse(stream) {
1623                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype(inst)),
1624                Err(_) => {}
1625            }
1626            stream.set_position(start_pos);
1627            match <crate::r#type::instruction::mma_sp::section_3::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 as PtxParser>::parse(stream) {
1628                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k128RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1(inst)),
1629                Err(_) => {}
1630            }
1631            stream.set_position(start_pos);
1632            match <crate::r#type::instruction::mma_sp::section_4::MmaSpvariantSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype as PtxParser>::parse(stream) {
1633                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype(inst)),
1634                Err(_) => {}
1635            }
1636            stream.set_position(start_pos);
1637            match <crate::r#type::instruction::mma_sp::section_5::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 as PtxParser>::parse(stream) {
1638                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32(inst)),
1639                Err(_) => {}
1640            }
1641            stream.set_position(start_pos);
1642            match <crate::r#type::instruction::mma_sp::section_6::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 as PtxParser>::parse(stream) {
1643                Ok(inst) => return Ok(Inst::MmaSpvariantSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321(inst)),
1644                Err(_) => {}
1645            }
1646            stream.set_position(start_pos);
1647            match <crate::r#type::instruction::mma::section_0::MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype as PtxParser>::parse(stream) {
1648                Ok(inst) => return Ok(Inst::MmaSyncAlignedM8n8k4AlayoutBlayoutDtypeF16F16Ctype(inst)),
1649                Err(_) => {}
1650            }
1651            stream.set_position(start_pos);
1652            match <crate::r#type::instruction::mma::section_0::MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype as PtxParser>::parse(stream) {
1653                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k8RowColDtypeF16F16Ctype(inst)),
1654                Err(_) => {}
1655            }
1656            stream.set_position(start_pos);
1657            match <crate::r#type::instruction::mma::section_0::MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype as PtxParser>::parse(stream) {
1658                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k16RowColDtypeF16F16Ctype(inst)),
1659                Err(_) => {}
1660            }
1661            stream.set_position(start_pos);
1662            match <crate::r#type::instruction::mma::section_1::MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32 as PtxParser>::parse(stream) {
1663                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k4RowColF32Tf32Tf32F32(inst)),
1664                Err(_) => {}
1665            }
1666            stream.set_position(start_pos);
1667            match <crate::r#type::instruction::mma::section_1::MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32 as PtxParser>::parse(stream) {
1668                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k8RowColF32AtypeBtypeF32(inst)),
1669                Err(_) => {}
1670            }
1671            stream.set_position(start_pos);
1672            match <crate::r#type::instruction::mma::section_1::MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32 as PtxParser>::parse(stream) {
1673                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k16RowColF32Bf16Bf16F32(inst)),
1674                Err(_) => {}
1675            }
1676            stream.set_position(start_pos);
1677            match <crate::r#type::instruction::mma::section_1::MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype as PtxParser>::parse(stream) {
1678                Ok(inst) => return Ok(Inst::MmaSyncAlignedShapeRowColDtypeF8typeF8typeCtype(inst)),
1679                Err(_) => {}
1680            }
1681            stream.set_position(start_pos);
1682            match <crate::r#type::instruction::mma::section_1::MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype as PtxParser>::parse(stream) {
1683                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k32RowColKindDtypeF8f6f4typeF8f6f4typeCtype(inst)),
1684                Err(_) => {}
1685            }
1686            stream.set_position(start_pos);
1687            match <crate::r#type::instruction::mma::section_2::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype as PtxParser>::parse(stream) {
1688                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype(inst)),
1689                Err(_) => {}
1690            }
1691            stream.set_position(start_pos);
1692            match <crate::r#type::instruction::mma::section_3::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1 as PtxParser>::parse(stream) {
1693                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k64RowColKindBlockScaleScaleVecSizeF32E2m1E2m1F32Stype1(inst)),
1694                Err(_) => {}
1695            }
1696            stream.set_position(start_pos);
1697            match <crate::r#type::instruction::mma::section_4::MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype as PtxParser>::parse(stream) {
1698                Ok(inst) => return Ok(Inst::MmaSyncAlignedM16n8k32RowColKindBlockScaleScaleVecSizeF32F8f6f4typeF8f6f4typeF32Stype(inst)),
1699                Err(_) => {}
1700            }
1701            stream.set_position(start_pos);
1702            match <crate::r#type::instruction::mma::section_5::MmaSyncAlignedShapeRowColF64F64F64F64 as PtxParser>::parse(stream) {
1703                Ok(inst) => return Ok(Inst::MmaSyncAlignedShapeRowColF64F64F64F64(inst)),
1704                Err(_) => {}
1705            }
1706            stream.set_position(start_pos);
1707            match <crate::r#type::instruction::mma::section_6::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32 as PtxParser>::parse(stream) {
1708                Ok(inst) => return Ok(Inst::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS32(inst)),
1709                Err(_) => {}
1710            }
1711            stream.set_position(start_pos);
1712            match <crate::r#type::instruction::mma::section_7::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321 as PtxParser>::parse(stream) {
1713                Ok(inst) => return Ok(Inst::MmaSyncAlignedShapeRowColSatfiniteS32AtypeBtypeS321(inst)),
1714                Err(_) => {}
1715            }
1716            stream.set_position(start_pos);
1717            match <crate::r#type::instruction::mma::section_8::MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc as PtxParser>::parse(stream) {
1718                Ok(inst) => return Ok(Inst::MmaSyncAlignedShapeRowColS32B1B1S32BitopPopc(inst)),
1719                Err(_) => {}
1720            }
1721        }
1722        "mov" => {
1723            stream.set_position(start_pos);
1724            match <crate::r#type::instruction::mov::section_0::MovType as PtxParser>::parse(stream) {
1725                Ok(inst) => return Ok(Inst::MovType(inst)),
1726                Err(_) => {}
1727            }
1728            stream.set_position(start_pos);
1729            match <crate::r#type::instruction::mov::section_0::MovU32 as PtxParser>::parse(stream) {
1730                Ok(inst) => return Ok(Inst::MovU32(inst)),
1731                Err(_) => {}
1732            }
1733            stream.set_position(start_pos);
1734            match <crate::r#type::instruction::mov::section_0::MovU64 as PtxParser>::parse(stream) {
1735                Ok(inst) => return Ok(Inst::MovU64(inst)),
1736                Err(_) => {}
1737            }
1738            stream.set_position(start_pos);
1739            match <crate::r#type::instruction::mov::section_0::MovU321 as PtxParser>::parse(stream) {
1740                Ok(inst) => return Ok(Inst::MovU321(inst)),
1741                Err(_) => {}
1742            }
1743            stream.set_position(start_pos);
1744            match <crate::r#type::instruction::mov::section_0::MovU641 as PtxParser>::parse(stream) {
1745                Ok(inst) => return Ok(Inst::MovU641(inst)),
1746                Err(_) => {}
1747            }
1748            stream.set_position(start_pos);
1749            match <crate::r#type::instruction::mov::section_1::MovType1 as PtxParser>::parse(stream) {
1750                Ok(inst) => return Ok(Inst::MovType1(inst)),
1751                Err(_) => {}
1752            }
1753        }
1754        "movmatrix" => {
1755            stream.set_position(start_pos);
1756            match <crate::r#type::instruction::movmatrix::section_0::MovmatrixSyncAlignedShapeTransType as PtxParser>::parse(stream) {
1757                Ok(inst) => return Ok(Inst::MovmatrixSyncAlignedShapeTransType(inst)),
1758                Err(_) => {}
1759            }
1760        }
1761        "mul" => {
1762            stream.set_position(start_pos);
1763            match <crate::r#type::instruction::mul::section_0::MulModeType as PtxParser>::parse(stream) {
1764                Ok(inst) => return Ok(Inst::MulModeType(inst)),
1765                Err(_) => {}
1766            }
1767            stream.set_position(start_pos);
1768            match <crate::r#type::instruction::mul::section_1::MulRndFtzSatF32 as PtxParser>::parse(stream) {
1769                Ok(inst) => return Ok(Inst::MulRndFtzSatF32(inst)),
1770                Err(_) => {}
1771            }
1772            stream.set_position(start_pos);
1773            match <crate::r#type::instruction::mul::section_1::MulRndFtzF32x2 as PtxParser>::parse(stream) {
1774                Ok(inst) => return Ok(Inst::MulRndFtzF32x2(inst)),
1775                Err(_) => {}
1776            }
1777            stream.set_position(start_pos);
1778            match <crate::r#type::instruction::mul::section_1::MulRndF64 as PtxParser>::parse(stream) {
1779                Ok(inst) => return Ok(Inst::MulRndF64(inst)),
1780                Err(_) => {}
1781            }
1782            stream.set_position(start_pos);
1783            match <crate::r#type::instruction::mul::section_2::MulRndFtzSatF16 as PtxParser>::parse(stream) {
1784                Ok(inst) => return Ok(Inst::MulRndFtzSatF16(inst)),
1785                Err(_) => {}
1786            }
1787            stream.set_position(start_pos);
1788            match <crate::r#type::instruction::mul::section_2::MulRndFtzSatF16x2 as PtxParser>::parse(stream) {
1789                Ok(inst) => return Ok(Inst::MulRndFtzSatF16x2(inst)),
1790                Err(_) => {}
1791            }
1792            stream.set_position(start_pos);
1793            match <crate::r#type::instruction::mul::section_2::MulRndBf16 as PtxParser>::parse(stream) {
1794                Ok(inst) => return Ok(Inst::MulRndBf16(inst)),
1795                Err(_) => {}
1796            }
1797            stream.set_position(start_pos);
1798            match <crate::r#type::instruction::mul::section_2::MulRndBf16x2 as PtxParser>::parse(stream) {
1799                Ok(inst) => return Ok(Inst::MulRndBf16x2(inst)),
1800                Err(_) => {}
1801            }
1802        }
1803        "mul24" => {
1804            stream.set_position(start_pos);
1805            match <crate::r#type::instruction::mul24::section_0::Mul24ModeType as PtxParser>::parse(stream) {
1806                Ok(inst) => return Ok(Inst::Mul24ModeType(inst)),
1807                Err(_) => {}
1808            }
1809        }
1810        "multimem" => {
1811            stream.set_position(start_pos);
1812            match <crate::r#type::instruction::multimem_ld_reduce::section_0::MultimemLdReduceLdsemScopeSsOpType as PtxParser>::parse(stream) {
1813                Ok(inst) => return Ok(Inst::MultimemLdReduceLdsemScopeSsOpType(inst)),
1814                Err(_) => {}
1815            }
1816            stream.set_position(start_pos);
1817            match <crate::r#type::instruction::multimem_ld_reduce::section_0::MultimemLdReduceWeakSsOpType as PtxParser>::parse(stream) {
1818                Ok(inst) => return Ok(Inst::MultimemLdReduceWeakSsOpType(inst)),
1819                Err(_) => {}
1820            }
1821            stream.set_position(start_pos);
1822            match <crate::r#type::instruction::multimem_ld_reduce::section_0::MultimemStStsemScopeSsType as PtxParser>::parse(stream) {
1823                Ok(inst) => return Ok(Inst::MultimemStStsemScopeSsType(inst)),
1824                Err(_) => {}
1825            }
1826            stream.set_position(start_pos);
1827            match <crate::r#type::instruction::multimem_ld_reduce::section_0::MultimemStWeakSsType as PtxParser>::parse(stream) {
1828                Ok(inst) => return Ok(Inst::MultimemStWeakSsType(inst)),
1829                Err(_) => {}
1830            }
1831            stream.set_position(start_pos);
1832            match <crate::r#type::instruction::multimem_ld_reduce::section_0::MultimemRedRedsemScopeSsOpType as PtxParser>::parse(stream) {
1833                Ok(inst) => return Ok(Inst::MultimemRedRedsemScopeSsOpType(inst)),
1834                Err(_) => {}
1835            }
1836            stream.set_position(start_pos);
1837            match <crate::r#type::instruction::multimem_ld_reduce::section_1::MultimemLdReduceLdsemScopeSsOpAccPrecVecType as PtxParser>::parse(stream) {
1838                Ok(inst) => return Ok(Inst::MultimemLdReduceLdsemScopeSsOpAccPrecVecType(inst)),
1839                Err(_) => {}
1840            }
1841            stream.set_position(start_pos);
1842            match <crate::r#type::instruction::multimem_ld_reduce::section_1::MultimemLdReduceWeakSsOpAccPrecVecType as PtxParser>::parse(stream) {
1843                Ok(inst) => return Ok(Inst::MultimemLdReduceWeakSsOpAccPrecVecType(inst)),
1844                Err(_) => {}
1845            }
1846            stream.set_position(start_pos);
1847            match <crate::r#type::instruction::multimem_ld_reduce::section_1::MultimemStStsemScopeSsVecType as PtxParser>::parse(stream) {
1848                Ok(inst) => return Ok(Inst::MultimemStStsemScopeSsVecType(inst)),
1849                Err(_) => {}
1850            }
1851            stream.set_position(start_pos);
1852            match <crate::r#type::instruction::multimem_ld_reduce::section_1::MultimemStWeakSsVecType as PtxParser>::parse(stream) {
1853                Ok(inst) => return Ok(Inst::MultimemStWeakSsVecType(inst)),
1854                Err(_) => {}
1855            }
1856            stream.set_position(start_pos);
1857            match <crate::r#type::instruction::multimem_ld_reduce::section_1::MultimemRedRedsemScopeSsRedopVecRedtype as PtxParser>::parse(stream) {
1858                Ok(inst) => return Ok(Inst::MultimemRedRedsemScopeSsRedopVecRedtype(inst)),
1859                Err(_) => {}
1860            }
1861        }
1862        "nanosleep" => {
1863            stream.set_position(start_pos);
1864            match <crate::r#type::instruction::nanosleep::section_0::NanosleepU32 as PtxParser>::parse(stream) {
1865                Ok(inst) => return Ok(Inst::NanosleepU32(inst)),
1866                Err(_) => {}
1867            }
1868        }
1869        "neg" => {
1870            stream.set_position(start_pos);
1871            match <crate::r#type::instruction::neg::section_0::NegType as PtxParser>::parse(stream) {
1872                Ok(inst) => return Ok(Inst::NegType(inst)),
1873                Err(_) => {}
1874            }
1875            stream.set_position(start_pos);
1876            match <crate::r#type::instruction::neg::section_0::NegFtzF32 as PtxParser>::parse(stream) {
1877                Ok(inst) => return Ok(Inst::NegFtzF32(inst)),
1878                Err(_) => {}
1879            }
1880            stream.set_position(start_pos);
1881            match <crate::r#type::instruction::neg::section_0::NegF64 as PtxParser>::parse(stream) {
1882                Ok(inst) => return Ok(Inst::NegF64(inst)),
1883                Err(_) => {}
1884            }
1885            stream.set_position(start_pos);
1886            match <crate::r#type::instruction::neg::section_0::NegFtzF16 as PtxParser>::parse(stream) {
1887                Ok(inst) => return Ok(Inst::NegFtzF16(inst)),
1888                Err(_) => {}
1889            }
1890            stream.set_position(start_pos);
1891            match <crate::r#type::instruction::neg::section_0::NegFtzF16x2 as PtxParser>::parse(stream) {
1892                Ok(inst) => return Ok(Inst::NegFtzF16x2(inst)),
1893                Err(_) => {}
1894            }
1895            stream.set_position(start_pos);
1896            match <crate::r#type::instruction::neg::section_0::NegBf16 as PtxParser>::parse(stream) {
1897                Ok(inst) => return Ok(Inst::NegBf16(inst)),
1898                Err(_) => {}
1899            }
1900            stream.set_position(start_pos);
1901            match <crate::r#type::instruction::neg::section_0::NegBf16x2 as PtxParser>::parse(stream) {
1902                Ok(inst) => return Ok(Inst::NegBf16x2(inst)),
1903                Err(_) => {}
1904            }
1905        }
1906        "not" => {
1907            stream.set_position(start_pos);
1908            match <crate::r#type::instruction::not::section_0::NotType as PtxParser>::parse(stream) {
1909                Ok(inst) => return Ok(Inst::NotType(inst)),
1910                Err(_) => {}
1911            }
1912        }
1913        "or" => {
1914            stream.set_position(start_pos);
1915            match <crate::r#type::instruction::or::section_0::OrType as PtxParser>::parse(stream) {
1916                Ok(inst) => return Ok(Inst::OrType(inst)),
1917                Err(_) => {}
1918            }
1919        }
1920        "pmevent" => {
1921            stream.set_position(start_pos);
1922            match <crate::r#type::instruction::pmevent::section_0::Pmevent as PtxParser>::parse(stream) {
1923                Ok(inst) => return Ok(Inst::Pmevent(inst)),
1924                Err(_) => {}
1925            }
1926            stream.set_position(start_pos);
1927            match <crate::r#type::instruction::pmevent::section_0::PmeventMask as PtxParser>::parse(stream) {
1928                Ok(inst) => return Ok(Inst::PmeventMask(inst)),
1929                Err(_) => {}
1930            }
1931        }
1932        "popc" => {
1933            stream.set_position(start_pos);
1934            match <crate::r#type::instruction::popc::section_0::PopcType as PtxParser>::parse(stream) {
1935                Ok(inst) => return Ok(Inst::PopcType(inst)),
1936                Err(_) => {}
1937            }
1938        }
1939        "prefetch" => {
1940            stream.set_position(start_pos);
1941            match <crate::r#type::instruction::prefetch::section_0::PrefetchSpaceLevel as PtxParser>::parse(stream) {
1942                Ok(inst) => return Ok(Inst::PrefetchSpaceLevel(inst)),
1943                Err(_) => {}
1944            }
1945            stream.set_position(start_pos);
1946            match <crate::r#type::instruction::prefetch::section_0::PrefetchGlobalLevelEvictionPriority as PtxParser>::parse(stream) {
1947                Ok(inst) => return Ok(Inst::PrefetchGlobalLevelEvictionPriority(inst)),
1948                Err(_) => {}
1949            }
1950            stream.set_position(start_pos);
1951            match <crate::r#type::instruction::prefetch::section_0::PrefetchuL1 as PtxParser>::parse(stream) {
1952                Ok(inst) => return Ok(Inst::PrefetchuL1(inst)),
1953                Err(_) => {}
1954            }
1955            stream.set_position(start_pos);
1956            match <crate::r#type::instruction::prefetch::section_0::PrefetchTensormapSpaceTensormap as PtxParser>::parse(stream) {
1957                Ok(inst) => return Ok(Inst::PrefetchTensormapSpaceTensormap(inst)),
1958                Err(_) => {}
1959            }
1960        }
1961        "prefetchu" => {
1962            stream.set_position(start_pos);
1963            match <crate::r#type::instruction::prefetch::section_0::PrefetchSpaceLevel as PtxParser>::parse(stream) {
1964                Ok(inst) => return Ok(Inst::PrefetchSpaceLevel(inst)),
1965                Err(_) => {}
1966            }
1967            stream.set_position(start_pos);
1968            match <crate::r#type::instruction::prefetch::section_0::PrefetchGlobalLevelEvictionPriority as PtxParser>::parse(stream) {
1969                Ok(inst) => return Ok(Inst::PrefetchGlobalLevelEvictionPriority(inst)),
1970                Err(_) => {}
1971            }
1972            stream.set_position(start_pos);
1973            match <crate::r#type::instruction::prefetch::section_0::PrefetchuL1 as PtxParser>::parse(stream) {
1974                Ok(inst) => return Ok(Inst::PrefetchuL1(inst)),
1975                Err(_) => {}
1976            }
1977            stream.set_position(start_pos);
1978            match <crate::r#type::instruction::prefetch::section_0::PrefetchTensormapSpaceTensormap as PtxParser>::parse(stream) {
1979                Ok(inst) => return Ok(Inst::PrefetchTensormapSpaceTensormap(inst)),
1980                Err(_) => {}
1981            }
1982        }
1983        "prmt" => {
1984            stream.set_position(start_pos);
1985            match <crate::r#type::instruction::prmt::section_0::PrmtB32Mode as PtxParser>::parse(stream) {
1986                Ok(inst) => return Ok(Inst::PrmtB32Mode(inst)),
1987                Err(_) => {}
1988            }
1989        }
1990        "rcp" => {
1991            stream.set_position(start_pos);
1992            match <crate::r#type::instruction::rcp_approx_ftz_f64::section_0::RcpApproxFtzF64 as PtxParser>::parse(stream) {
1993                Ok(inst) => return Ok(Inst::RcpApproxFtzF64(inst)),
1994                Err(_) => {}
1995            }
1996            stream.set_position(start_pos);
1997            match <crate::r#type::instruction::rcp::section_0::RcpApproxFtzF32 as PtxParser>::parse(stream) {
1998                Ok(inst) => return Ok(Inst::RcpApproxFtzF32(inst)),
1999                Err(_) => {}
2000            }
2001            stream.set_position(start_pos);
2002            match <crate::r#type::instruction::rcp::section_0::RcpRndFtzF32 as PtxParser>::parse(stream) {
2003                Ok(inst) => return Ok(Inst::RcpRndFtzF32(inst)),
2004                Err(_) => {}
2005            }
2006            stream.set_position(start_pos);
2007            match <crate::r#type::instruction::rcp::section_0::RcpRndF64 as PtxParser>::parse(stream) {
2008                Ok(inst) => return Ok(Inst::RcpRndF64(inst)),
2009                Err(_) => {}
2010            }
2011        }
2012        "red" => {
2013            stream.set_position(start_pos);
2014            match <crate::r#type::instruction::red_async::section_0::RedAsyncSemScopeSsCompletionMechanismOpType as PtxParser>::parse(stream) {
2015                Ok(inst) => return Ok(Inst::RedAsyncSemScopeSsCompletionMechanismOpType(inst)),
2016                Err(_) => {}
2017            }
2018            stream.set_position(start_pos);
2019            match <crate::r#type::instruction::red_async::section_1::RedAsyncSemScopeSsCompletionMechanismOpType1 as PtxParser>::parse(stream) {
2020                Ok(inst) => return Ok(Inst::RedAsyncSemScopeSsCompletionMechanismOpType1(inst)),
2021                Err(_) => {}
2022            }
2023            stream.set_position(start_pos);
2024            match <crate::r#type::instruction::red_async::section_2::RedAsyncSemScopeSsCompletionMechanismOpType2 as PtxParser>::parse(stream) {
2025                Ok(inst) => return Ok(Inst::RedAsyncSemScopeSsCompletionMechanismOpType2(inst)),
2026                Err(_) => {}
2027            }
2028            stream.set_position(start_pos);
2029            match <crate::r#type::instruction::red_async::section_3::RedAsyncSemScopeSsCompletionMechanismAddType as PtxParser>::parse(stream) {
2030                Ok(inst) => return Ok(Inst::RedAsyncSemScopeSsCompletionMechanismAddType(inst)),
2031                Err(_) => {}
2032            }
2033            stream.set_position(start_pos);
2034            match <crate::r#type::instruction::red_async::section_4::RedAsyncMmioSemScopeSsAddType as PtxParser>::parse(stream) {
2035                Ok(inst) => return Ok(Inst::RedAsyncMmioSemScopeSsAddType(inst)),
2036                Err(_) => {}
2037            }
2038            stream.set_position(start_pos);
2039            match <crate::r#type::instruction::red::section_0::RedOpSpaceSemScopeLevelCacheHintType as PtxParser>::parse(stream) {
2040                Ok(inst) => return Ok(Inst::RedOpSpaceSemScopeLevelCacheHintType(inst)),
2041                Err(_) => {}
2042            }
2043            stream.set_position(start_pos);
2044            match <crate::r#type::instruction::red::section_0::RedAddSpaceSemScopeNoftzLevelCacheHintF16 as PtxParser>::parse(stream) {
2045                Ok(inst) => return Ok(Inst::RedAddSpaceSemScopeNoftzLevelCacheHintF16(inst)),
2046                Err(_) => {}
2047            }
2048            stream.set_position(start_pos);
2049            match <crate::r#type::instruction::red::section_0::RedAddSpaceSemScopeNoftzLevelCacheHintF16x2 as PtxParser>::parse(stream) {
2050                Ok(inst) => return Ok(Inst::RedAddSpaceSemScopeNoftzLevelCacheHintF16x2(inst)),
2051                Err(_) => {}
2052            }
2053            stream.set_position(start_pos);
2054            match <crate::r#type::instruction::red::section_0::RedAddSpaceSemScopeNoftzLevelCacheHintBf16 as PtxParser>::parse(stream) {
2055                Ok(inst) => return Ok(Inst::RedAddSpaceSemScopeNoftzLevelCacheHintBf16(inst)),
2056                Err(_) => {}
2057            }
2058            stream.set_position(start_pos);
2059            match <crate::r#type::instruction::red::section_0::RedAddSpaceSemScopeNoftzLevelCacheHintBf16x2 as PtxParser>::parse(stream) {
2060                Ok(inst) => return Ok(Inst::RedAddSpaceSemScopeNoftzLevelCacheHintBf16x2(inst)),
2061                Err(_) => {}
2062            }
2063            stream.set_position(start_pos);
2064            match <crate::r#type::instruction::red::section_1::RedAddSpaceSemScopeLevelCacheHintVec32BitF32 as PtxParser>::parse(stream) {
2065                Ok(inst) => return Ok(Inst::RedAddSpaceSemScopeLevelCacheHintVec32BitF32(inst)),
2066                Err(_) => {}
2067            }
2068            stream.set_position(start_pos);
2069            match <crate::r#type::instruction::red::section_1::RedOpSpaceSemScopeNoftzLevelCacheHintVec16BitHalfWordType as PtxParser>::parse(stream) {
2070                Ok(inst) => return Ok(Inst::RedOpSpaceSemScopeNoftzLevelCacheHintVec16BitHalfWordType(inst)),
2071                Err(_) => {}
2072            }
2073            stream.set_position(start_pos);
2074            match <crate::r#type::instruction::red::section_1::RedOpSpaceSemScopeNoftzLevelCacheHintVec32BitPackedType as PtxParser>::parse(stream) {
2075                Ok(inst) => return Ok(Inst::RedOpSpaceSemScopeNoftzLevelCacheHintVec32BitPackedType(inst)),
2076                Err(_) => {}
2077            }
2078        }
2079        "redux" => {
2080            stream.set_position(start_pos);
2081            match <crate::r#type::instruction::redux_sync::section_0::ReduxSyncOpType as PtxParser>::parse(stream) {
2082                Ok(inst) => return Ok(Inst::ReduxSyncOpType(inst)),
2083                Err(_) => {}
2084            }
2085            stream.set_position(start_pos);
2086            match <crate::r#type::instruction::redux_sync::section_1::ReduxSyncOpB32 as PtxParser>::parse(stream) {
2087                Ok(inst) => return Ok(Inst::ReduxSyncOpB32(inst)),
2088                Err(_) => {}
2089            }
2090            stream.set_position(start_pos);
2091            match <crate::r#type::instruction::redux_sync::section_2::ReduxSyncOpAbsNanF32 as PtxParser>::parse(stream) {
2092                Ok(inst) => return Ok(Inst::ReduxSyncOpAbsNanF32(inst)),
2093                Err(_) => {}
2094            }
2095        }
2096        "rem" => {
2097            stream.set_position(start_pos);
2098            match <crate::r#type::instruction::rem::section_0::RemType as PtxParser>::parse(stream) {
2099                Ok(inst) => return Ok(Inst::RemType(inst)),
2100                Err(_) => {}
2101            }
2102        }
2103        "ret" => {
2104            stream.set_position(start_pos);
2105            match <crate::r#type::instruction::ret::section_0::RetUni as PtxParser>::parse(stream) {
2106                Ok(inst) => return Ok(Inst::RetUni(inst)),
2107                Err(_) => {}
2108            }
2109        }
2110        "rsqrt" => {
2111            stream.set_position(start_pos);
2112            match <crate::r#type::instruction::rsqrt_approx_ftz_f64::section_0::RsqrtApproxFtzF64 as PtxParser>::parse(stream) {
2113                Ok(inst) => return Ok(Inst::RsqrtApproxFtzF64(inst)),
2114                Err(_) => {}
2115            }
2116            stream.set_position(start_pos);
2117            match <crate::r#type::instruction::rsqrt::section_0::RsqrtApproxFtzF32 as PtxParser>::parse(stream) {
2118                Ok(inst) => return Ok(Inst::RsqrtApproxFtzF32(inst)),
2119                Err(_) => {}
2120            }
2121            stream.set_position(start_pos);
2122            match <crate::r#type::instruction::rsqrt::section_0::RsqrtApproxF64 as PtxParser>::parse(stream) {
2123                Ok(inst) => return Ok(Inst::RsqrtApproxF64(inst)),
2124                Err(_) => {}
2125            }
2126        }
2127        "sad" => {
2128            stream.set_position(start_pos);
2129            match <crate::r#type::instruction::sad::section_0::SadType as PtxParser>::parse(stream) {
2130                Ok(inst) => return Ok(Inst::SadType(inst)),
2131                Err(_) => {}
2132            }
2133        }
2134        "selp" => {
2135            stream.set_position(start_pos);
2136            match <crate::r#type::instruction::selp::section_0::SelpType as PtxParser>::parse(stream) {
2137                Ok(inst) => return Ok(Inst::SelpType(inst)),
2138                Err(_) => {}
2139            }
2140        }
2141        "set" => {
2142            stream.set_position(start_pos);
2143            match <crate::r#type::instruction::set::section_0::SetCmpopFtzDtypeStype as PtxParser>::parse(stream) {
2144                Ok(inst) => return Ok(Inst::SetCmpopFtzDtypeStype(inst)),
2145                Err(_) => {}
2146            }
2147            stream.set_position(start_pos);
2148            match <crate::r#type::instruction::set::section_0::SetCmpopBoolopFtzDtypeStype as PtxParser>::parse(stream) {
2149                Ok(inst) => return Ok(Inst::SetCmpopBoolopFtzDtypeStype(inst)),
2150                Err(_) => {}
2151            }
2152            stream.set_position(start_pos);
2153            match <crate::r#type::instruction::set::section_1::SetCmpopFtzF16Stype as PtxParser>::parse(stream) {
2154                Ok(inst) => return Ok(Inst::SetCmpopFtzF16Stype(inst)),
2155                Err(_) => {}
2156            }
2157            stream.set_position(start_pos);
2158            match <crate::r#type::instruction::set::section_1::SetCmpopBoolopFtzF16Stype as PtxParser>::parse(stream) {
2159                Ok(inst) => return Ok(Inst::SetCmpopBoolopFtzF16Stype(inst)),
2160                Err(_) => {}
2161            }
2162            stream.set_position(start_pos);
2163            match <crate::r#type::instruction::set::section_1::SetCmpopBf16Stype as PtxParser>::parse(stream) {
2164                Ok(inst) => return Ok(Inst::SetCmpopBf16Stype(inst)),
2165                Err(_) => {}
2166            }
2167            stream.set_position(start_pos);
2168            match <crate::r#type::instruction::set::section_1::SetCmpopBoolopBf16Stype as PtxParser>::parse(stream) {
2169                Ok(inst) => return Ok(Inst::SetCmpopBoolopBf16Stype(inst)),
2170                Err(_) => {}
2171            }
2172            stream.set_position(start_pos);
2173            match <crate::r#type::instruction::set::section_1::SetCmpopFtzDtypeF16 as PtxParser>::parse(stream) {
2174                Ok(inst) => return Ok(Inst::SetCmpopFtzDtypeF16(inst)),
2175                Err(_) => {}
2176            }
2177            stream.set_position(start_pos);
2178            match <crate::r#type::instruction::set::section_1::SetCmpopBoolopFtzDtypeF16 as PtxParser>::parse(stream) {
2179                Ok(inst) => return Ok(Inst::SetCmpopBoolopFtzDtypeF16(inst)),
2180                Err(_) => {}
2181            }
2182            stream.set_position(start_pos);
2183            match <crate::r#type::instruction::set::section_2::SetCmpopDtypeBf16 as PtxParser>::parse(stream) {
2184                Ok(inst) => return Ok(Inst::SetCmpopDtypeBf16(inst)),
2185                Err(_) => {}
2186            }
2187            stream.set_position(start_pos);
2188            match <crate::r#type::instruction::set::section_2::SetCmpopBoolopDtypeBf16 as PtxParser>::parse(stream) {
2189                Ok(inst) => return Ok(Inst::SetCmpopBoolopDtypeBf16(inst)),
2190                Err(_) => {}
2191            }
2192            stream.set_position(start_pos);
2193            match <crate::r#type::instruction::set::section_3::SetCmpopFtzDtypeF16x2 as PtxParser>::parse(stream) {
2194                Ok(inst) => return Ok(Inst::SetCmpopFtzDtypeF16x2(inst)),
2195                Err(_) => {}
2196            }
2197            stream.set_position(start_pos);
2198            match <crate::r#type::instruction::set::section_3::SetCmpopBoolopFtzDtypeF16x2 as PtxParser>::parse(stream) {
2199                Ok(inst) => return Ok(Inst::SetCmpopBoolopFtzDtypeF16x2(inst)),
2200                Err(_) => {}
2201            }
2202            stream.set_position(start_pos);
2203            match <crate::r#type::instruction::set::section_4::SetCmpopDtypeBf16x2 as PtxParser>::parse(stream) {
2204                Ok(inst) => return Ok(Inst::SetCmpopDtypeBf16x2(inst)),
2205                Err(_) => {}
2206            }
2207            stream.set_position(start_pos);
2208            match <crate::r#type::instruction::set::section_4::SetCmpopBoolopDtypeBf16x2 as PtxParser>::parse(stream) {
2209                Ok(inst) => return Ok(Inst::SetCmpopBoolopDtypeBf16x2(inst)),
2210                Err(_) => {}
2211            }
2212        }
2213        "setmaxnreg" => {
2214            stream.set_position(start_pos);
2215            match <crate::r#type::instruction::setmaxnreg::section_0::SetmaxnregActionSyncAlignedU32 as PtxParser>::parse(stream) {
2216                Ok(inst) => return Ok(Inst::SetmaxnregActionSyncAlignedU32(inst)),
2217                Err(_) => {}
2218            }
2219        }
2220        "setp" => {
2221            stream.set_position(start_pos);
2222            match <crate::r#type::instruction::setp::section_0::SetpCmpopFtzType as PtxParser>::parse(stream) {
2223                Ok(inst) => return Ok(Inst::SetpCmpopFtzType(inst)),
2224                Err(_) => {}
2225            }
2226            stream.set_position(start_pos);
2227            match <crate::r#type::instruction::setp::section_0::SetpCmpopBoolopFtzType as PtxParser>::parse(stream) {
2228                Ok(inst) => return Ok(Inst::SetpCmpopBoolopFtzType(inst)),
2229                Err(_) => {}
2230            }
2231            stream.set_position(start_pos);
2232            match <crate::r#type::instruction::setp::section_1::SetpCmpopFtzF16 as PtxParser>::parse(stream) {
2233                Ok(inst) => return Ok(Inst::SetpCmpopFtzF16(inst)),
2234                Err(_) => {}
2235            }
2236            stream.set_position(start_pos);
2237            match <crate::r#type::instruction::setp::section_1::SetpCmpopBoolopFtzF16 as PtxParser>::parse(stream) {
2238                Ok(inst) => return Ok(Inst::SetpCmpopBoolopFtzF16(inst)),
2239                Err(_) => {}
2240            }
2241            stream.set_position(start_pos);
2242            match <crate::r#type::instruction::setp::section_1::SetpCmpopFtzF16x2 as PtxParser>::parse(stream) {
2243                Ok(inst) => return Ok(Inst::SetpCmpopFtzF16x2(inst)),
2244                Err(_) => {}
2245            }
2246            stream.set_position(start_pos);
2247            match <crate::r#type::instruction::setp::section_1::SetpCmpopBoolopFtzF16x2 as PtxParser>::parse(stream) {
2248                Ok(inst) => return Ok(Inst::SetpCmpopBoolopFtzF16x2(inst)),
2249                Err(_) => {}
2250            }
2251            stream.set_position(start_pos);
2252            match <crate::r#type::instruction::setp::section_1::SetpCmpopBf16 as PtxParser>::parse(stream) {
2253                Ok(inst) => return Ok(Inst::SetpCmpopBf16(inst)),
2254                Err(_) => {}
2255            }
2256            stream.set_position(start_pos);
2257            match <crate::r#type::instruction::setp::section_1::SetpCmpopBoolopBf16 as PtxParser>::parse(stream) {
2258                Ok(inst) => return Ok(Inst::SetpCmpopBoolopBf16(inst)),
2259                Err(_) => {}
2260            }
2261            stream.set_position(start_pos);
2262            match <crate::r#type::instruction::setp::section_1::SetpCmpopBf16x2 as PtxParser>::parse(stream) {
2263                Ok(inst) => return Ok(Inst::SetpCmpopBf16x2(inst)),
2264                Err(_) => {}
2265            }
2266            stream.set_position(start_pos);
2267            match <crate::r#type::instruction::setp::section_1::SetpCmpopBoolopBf16x2 as PtxParser>::parse(stream) {
2268                Ok(inst) => return Ok(Inst::SetpCmpopBoolopBf16x2(inst)),
2269                Err(_) => {}
2270            }
2271        }
2272        "shf" => {
2273            stream.set_position(start_pos);
2274            match <crate::r#type::instruction::shf::section_0::ShfLModeB32 as PtxParser>::parse(stream) {
2275                Ok(inst) => return Ok(Inst::ShfLModeB32(inst)),
2276                Err(_) => {}
2277            }
2278            stream.set_position(start_pos);
2279            match <crate::r#type::instruction::shf::section_0::ShfRModeB32 as PtxParser>::parse(stream) {
2280                Ok(inst) => return Ok(Inst::ShfRModeB32(inst)),
2281                Err(_) => {}
2282            }
2283        }
2284        "shfl" => {
2285            stream.set_position(start_pos);
2286            match <crate::r#type::instruction::shfl_sync::section_0::ShflSyncModeB32 as PtxParser>::parse(stream) {
2287                Ok(inst) => return Ok(Inst::ShflSyncModeB32(inst)),
2288                Err(_) => {}
2289            }
2290            stream.set_position(start_pos);
2291            match <crate::r#type::instruction::shfl::section_0::ShflModeB32 as PtxParser>::parse(stream) {
2292                Ok(inst) => return Ok(Inst::ShflModeB32(inst)),
2293                Err(_) => {}
2294            }
2295        }
2296        "shl" => {
2297            stream.set_position(start_pos);
2298            match <crate::r#type::instruction::shl::section_0::ShlType as PtxParser>::parse(stream) {
2299                Ok(inst) => return Ok(Inst::ShlType(inst)),
2300                Err(_) => {}
2301            }
2302        }
2303        "shr" => {
2304            stream.set_position(start_pos);
2305            match <crate::r#type::instruction::shr::section_0::ShrType as PtxParser>::parse(stream) {
2306                Ok(inst) => return Ok(Inst::ShrType(inst)),
2307                Err(_) => {}
2308            }
2309        }
2310        "sin" => {
2311            stream.set_position(start_pos);
2312            match <crate::r#type::instruction::sin::section_0::SinApproxFtzF32 as PtxParser>::parse(stream) {
2313                Ok(inst) => return Ok(Inst::SinApproxFtzF32(inst)),
2314                Err(_) => {}
2315            }
2316        }
2317        "slct" => {
2318            stream.set_position(start_pos);
2319            match <crate::r#type::instruction::slct::section_0::SlctDtypeS32 as PtxParser>::parse(stream) {
2320                Ok(inst) => return Ok(Inst::SlctDtypeS32(inst)),
2321                Err(_) => {}
2322            }
2323            stream.set_position(start_pos);
2324            match <crate::r#type::instruction::slct::section_0::SlctFtzDtypeF32 as PtxParser>::parse(stream) {
2325                Ok(inst) => return Ok(Inst::SlctFtzDtypeF32(inst)),
2326                Err(_) => {}
2327            }
2328        }
2329        "sqrt" => {
2330            stream.set_position(start_pos);
2331            match <crate::r#type::instruction::sqrt::section_0::SqrtApproxFtzF32 as PtxParser>::parse(stream) {
2332                Ok(inst) => return Ok(Inst::SqrtApproxFtzF32(inst)),
2333                Err(_) => {}
2334            }
2335            stream.set_position(start_pos);
2336            match <crate::r#type::instruction::sqrt::section_0::SqrtRndFtzF32 as PtxParser>::parse(stream) {
2337                Ok(inst) => return Ok(Inst::SqrtRndFtzF32(inst)),
2338                Err(_) => {}
2339            }
2340            stream.set_position(start_pos);
2341            match <crate::r#type::instruction::sqrt::section_0::SqrtRndF64 as PtxParser>::parse(stream) {
2342                Ok(inst) => return Ok(Inst::SqrtRndF64(inst)),
2343                Err(_) => {}
2344            }
2345        }
2346        "st" => {
2347            stream.set_position(start_pos);
2348            match <crate::r#type::instruction::st_async::section_0::StAsyncSemScopeSsCompletionMechanismVecType as PtxParser>::parse(stream) {
2349                Ok(inst) => return Ok(Inst::StAsyncSemScopeSsCompletionMechanismVecType(inst)),
2350                Err(_) => {}
2351            }
2352            stream.set_position(start_pos);
2353            match <crate::r#type::instruction::st_async::section_1::StAsyncMmioSemScopeSsType as PtxParser>::parse(stream) {
2354                Ok(inst) => return Ok(Inst::StAsyncMmioSemScopeSsType(inst)),
2355                Err(_) => {}
2356            }
2357            stream.set_position(start_pos);
2358            match <crate::r#type::instruction::st_bulk::section_0::StBulkWeakSharedCta as PtxParser>::parse(stream) {
2359                Ok(inst) => return Ok(Inst::StBulkWeakSharedCta(inst)),
2360                Err(_) => {}
2361            }
2362            stream.set_position(start_pos);
2363            match <crate::r#type::instruction::st::section_0::StWeakSsCopLevelCacheHintVecType as PtxParser>::parse(stream) {
2364                Ok(inst) => return Ok(Inst::StWeakSsCopLevelCacheHintVecType(inst)),
2365                Err(_) => {}
2366            }
2367            stream.set_position(start_pos);
2368            match <crate::r#type::instruction::st::section_0::StWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType as PtxParser>::parse(stream) {
2369                Ok(inst) => return Ok(Inst::StWeakSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(inst)),
2370                Err(_) => {}
2371            }
2372            stream.set_position(start_pos);
2373            match <crate::r#type::instruction::st::section_0::StVolatileSsVecType as PtxParser>::parse(stream) {
2374                Ok(inst) => return Ok(Inst::StVolatileSsVecType(inst)),
2375                Err(_) => {}
2376            }
2377            stream.set_position(start_pos);
2378            match <crate::r#type::instruction::st::section_0::StRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType as PtxParser>::parse(stream) {
2379                Ok(inst) => return Ok(Inst::StRelaxedScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(inst)),
2380                Err(_) => {}
2381            }
2382            stream.set_position(start_pos);
2383            match <crate::r#type::instruction::st::section_0::StReleaseScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType as PtxParser>::parse(stream) {
2384                Ok(inst) => return Ok(Inst::StReleaseScopeSsLevel1EvictionPriorityLevel2EvictionPriorityLevelCacheHintVecType(inst)),
2385                Err(_) => {}
2386            }
2387            stream.set_position(start_pos);
2388            match <crate::r#type::instruction::st::section_0::StMmioRelaxedSysGlobalType as PtxParser>::parse(stream) {
2389                Ok(inst) => return Ok(Inst::StMmioRelaxedSysGlobalType(inst)),
2390                Err(_) => {}
2391            }
2392        }
2393        "stackrestore" => {
2394            stream.set_position(start_pos);
2395            match <crate::r#type::instruction::stackrestore::section_0::StackrestoreType as PtxParser>::parse(stream) {
2396                Ok(inst) => return Ok(Inst::StackrestoreType(inst)),
2397                Err(_) => {}
2398            }
2399        }
2400        "stacksave" => {
2401            stream.set_position(start_pos);
2402            match <crate::r#type::instruction::stacksave::section_0::StacksaveType as PtxParser>::parse(stream) {
2403                Ok(inst) => return Ok(Inst::StacksaveType(inst)),
2404                Err(_) => {}
2405            }
2406        }
2407        "stmatrix" => {
2408            stream.set_position(start_pos);
2409            match <crate::r#type::instruction::stmatrix::section_0::StmatrixSyncAlignedShapeNumTransSsType as PtxParser>::parse(stream) {
2410                Ok(inst) => return Ok(Inst::StmatrixSyncAlignedShapeNumTransSsType(inst)),
2411                Err(_) => {}
2412            }
2413        }
2414        "sub" => {
2415            stream.set_position(start_pos);
2416            match <crate::r#type::instruction::sub_cc::section_0::SubCcType as PtxParser>::parse(stream) {
2417                Ok(inst) => return Ok(Inst::SubCcType(inst)),
2418                Err(_) => {}
2419            }
2420            stream.set_position(start_pos);
2421            match <crate::r#type::instruction::sub::section_0::SubType as PtxParser>::parse(stream) {
2422                Ok(inst) => return Ok(Inst::SubType(inst)),
2423                Err(_) => {}
2424            }
2425            stream.set_position(start_pos);
2426            match <crate::r#type::instruction::sub::section_0::SubSatS32 as PtxParser>::parse(stream) {
2427                Ok(inst) => return Ok(Inst::SubSatS32(inst)),
2428                Err(_) => {}
2429            }
2430            stream.set_position(start_pos);
2431            match <crate::r#type::instruction::sub::section_1::SubRndFtzSatF32 as PtxParser>::parse(stream) {
2432                Ok(inst) => return Ok(Inst::SubRndFtzSatF32(inst)),
2433                Err(_) => {}
2434            }
2435            stream.set_position(start_pos);
2436            match <crate::r#type::instruction::sub::section_1::SubRndFtzF32x2 as PtxParser>::parse(stream) {
2437                Ok(inst) => return Ok(Inst::SubRndFtzF32x2(inst)),
2438                Err(_) => {}
2439            }
2440            stream.set_position(start_pos);
2441            match <crate::r#type::instruction::sub::section_1::SubRndF64 as PtxParser>::parse(stream) {
2442                Ok(inst) => return Ok(Inst::SubRndF64(inst)),
2443                Err(_) => {}
2444            }
2445            stream.set_position(start_pos);
2446            match <crate::r#type::instruction::sub::section_2::SubRndFtzSatF16 as PtxParser>::parse(stream) {
2447                Ok(inst) => return Ok(Inst::SubRndFtzSatF16(inst)),
2448                Err(_) => {}
2449            }
2450            stream.set_position(start_pos);
2451            match <crate::r#type::instruction::sub::section_2::SubRndFtzSatF16x2 as PtxParser>::parse(stream) {
2452                Ok(inst) => return Ok(Inst::SubRndFtzSatF16x2(inst)),
2453                Err(_) => {}
2454            }
2455            stream.set_position(start_pos);
2456            match <crate::r#type::instruction::sub::section_2::SubRndBf16 as PtxParser>::parse(stream) {
2457                Ok(inst) => return Ok(Inst::SubRndBf16(inst)),
2458                Err(_) => {}
2459            }
2460            stream.set_position(start_pos);
2461            match <crate::r#type::instruction::sub::section_2::SubRndBf16x2 as PtxParser>::parse(stream) {
2462                Ok(inst) => return Ok(Inst::SubRndBf16x2(inst)),
2463                Err(_) => {}
2464            }
2465            stream.set_position(start_pos);
2466            match <crate::r#type::instruction::sub::section_3::SubRndSatF32Atype as PtxParser>::parse(stream) {
2467                Ok(inst) => return Ok(Inst::SubRndSatF32Atype(inst)),
2468                Err(_) => {}
2469            }
2470        }
2471        "subc" => {
2472            stream.set_position(start_pos);
2473            match <crate::r#type::instruction::subc::section_0::SubcCcType as PtxParser>::parse(stream) {
2474                Ok(inst) => return Ok(Inst::SubcCcType(inst)),
2475                Err(_) => {}
2476            }
2477        }
2478        "suld" => {
2479            stream.set_position(start_pos);
2480            match <crate::r#type::instruction::suld::section_0::SuldBGeomCopVecDtypeMode as PtxParser>::parse(stream) {
2481                Ok(inst) => return Ok(Inst::SuldBGeomCopVecDtypeMode(inst)),
2482                Err(_) => {}
2483            }
2484        }
2485        "suq" => {
2486            stream.set_position(start_pos);
2487            match <crate::r#type::instruction::suq::section_0::SuqQueryB32 as PtxParser>::parse(stream) {
2488                Ok(inst) => return Ok(Inst::SuqQueryB32(inst)),
2489                Err(_) => {}
2490            }
2491        }
2492        "sured" => {
2493            stream.set_position(start_pos);
2494            match <crate::r#type::instruction::sured::section_0::SuredBOpGeomCtypeMode as PtxParser>::parse(stream) {
2495                Ok(inst) => return Ok(Inst::SuredBOpGeomCtypeMode(inst)),
2496                Err(_) => {}
2497            }
2498            stream.set_position(start_pos);
2499            match <crate::r#type::instruction::sured::section_1::SuredPOpGeomCtypeMode as PtxParser>::parse(stream) {
2500                Ok(inst) => return Ok(Inst::SuredPOpGeomCtypeMode(inst)),
2501                Err(_) => {}
2502            }
2503        }
2504        "sust" => {
2505            stream.set_position(start_pos);
2506            match <crate::r#type::instruction::sust::section_0::SustBDimCopVecCtypeMode as PtxParser>::parse(stream) {
2507                Ok(inst) => return Ok(Inst::SustBDimCopVecCtypeMode(inst)),
2508                Err(_) => {}
2509            }
2510            stream.set_position(start_pos);
2511            match <crate::r#type::instruction::sust::section_0::SustPDimVecB32Mode as PtxParser>::parse(stream) {
2512                Ok(inst) => return Ok(Inst::SustPDimVecB32Mode(inst)),
2513                Err(_) => {}
2514            }
2515            stream.set_position(start_pos);
2516            match <crate::r#type::instruction::sust::section_0::SustBAdimCopVecCtypeMode as PtxParser>::parse(stream) {
2517                Ok(inst) => return Ok(Inst::SustBAdimCopVecCtypeMode(inst)),
2518                Err(_) => {}
2519            }
2520        }
2521        "szext" => {
2522            stream.set_position(start_pos);
2523            match <crate::r#type::instruction::szext::section_0::SzextModeType as PtxParser>::parse(stream) {
2524                Ok(inst) => return Ok(Inst::SzextModeType(inst)),
2525                Err(_) => {}
2526            }
2527        }
2528        "tanh" => {
2529            stream.set_position(start_pos);
2530            match <crate::r#type::instruction::tanh::section_0::TanhApproxType as PtxParser>::parse(stream) {
2531                Ok(inst) => return Ok(Inst::TanhApproxType(inst)),
2532                Err(_) => {}
2533            }
2534        }
2535        "tcgen05" => {
2536            stream.set_position(start_pos);
2537            match <crate::r#type::instruction::tcgen05_alloc::section_0::Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32 as PtxParser>::parse(stream) {
2538                Ok(inst) => return Ok(Inst::Tcgen05AllocCtaGroupSyncAlignedSharedCtaB32(inst)),
2539                Err(_) => {}
2540            }
2541            stream.set_position(start_pos);
2542            match <crate::r#type::instruction::tcgen05_alloc::section_0::Tcgen05DeallocCtaGroupSyncAlignedB32 as PtxParser>::parse(stream) {
2543                Ok(inst) => return Ok(Inst::Tcgen05DeallocCtaGroupSyncAlignedB32(inst)),
2544                Err(_) => {}
2545            }
2546            stream.set_position(start_pos);
2547            match <crate::r#type::instruction::tcgen05_alloc::section_0::Tcgen05RelinquishAllocPermitCtaGroupSyncAligned as PtxParser>::parse(stream) {
2548                Ok(inst) => return Ok(Inst::Tcgen05RelinquishAllocPermitCtaGroupSyncAligned(inst)),
2549                Err(_) => {}
2550            }
2551            stream.set_position(start_pos);
2552            match <crate::r#type::instruction::tcgen05_commit::section_0::Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64 as PtxParser>::parse(stream) {
2553                Ok(inst) => return Ok(Inst::Tcgen05CommitCtaGroupCompletionMechanismSharedClusterMulticastB64(inst)),
2554                Err(_) => {}
2555            }
2556            stream.set_position(start_pos);
2557            match <crate::r#type::instruction::tcgen05_cp::section_0::Tcgen05CpCtaGroupShapeMulticastDstSrcFmt as PtxParser>::parse(stream) {
2558                Ok(inst) => return Ok(Inst::Tcgen05CpCtaGroupShapeMulticastDstSrcFmt(inst)),
2559                Err(_) => {}
2560            }
2561            stream.set_position(start_pos);
2562            match <crate::r#type::instruction::tcgen05_fence::section_0::Tcgen05FenceBeforeThreadSync as PtxParser>::parse(stream) {
2563                Ok(inst) => return Ok(Inst::Tcgen05FenceBeforeThreadSync(inst)),
2564                Err(_) => {}
2565            }
2566            stream.set_position(start_pos);
2567            match <crate::r#type::instruction::tcgen05_fence::section_0::Tcgen05FenceAfterThreadSync as PtxParser>::parse(stream) {
2568                Ok(inst) => return Ok(Inst::Tcgen05FenceAfterThreadSync(inst)),
2569                Err(_) => {}
2570            }
2571            stream.set_position(start_pos);
2572            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdSyncAlignedShape1NumPackB32 as PtxParser>::parse(stream) {
2573                Ok(inst) => return Ok(Inst::Tcgen05LdSyncAlignedShape1NumPackB32(inst)),
2574                Err(_) => {}
2575            }
2576            stream.set_position(start_pos);
2577            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdSyncAlignedShape2NumPackB32 as PtxParser>::parse(stream) {
2578                Ok(inst) => return Ok(Inst::Tcgen05LdSyncAlignedShape2NumPackB32(inst)),
2579                Err(_) => {}
2580            }
2581            stream.set_position(start_pos);
2582            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32 as PtxParser>::parse(stream) {
2583                Ok(inst) => return Ok(Inst::Tcgen05LdRedSyncAlignedShape3NumRedopAbsNanF32(inst)),
2584                Err(_) => {}
2585            }
2586            stream.set_position(start_pos);
2587            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32 as PtxParser>::parse(stream) {
2588                Ok(inst) => return Ok(Inst::Tcgen05LdRedSyncAlignedShape4NumRedopAbsNanF32(inst)),
2589                Err(_) => {}
2590            }
2591            stream.set_position(start_pos);
2592            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdRedSyncAlignedShape3NumRedopType as PtxParser>::parse(stream) {
2593                Ok(inst) => return Ok(Inst::Tcgen05LdRedSyncAlignedShape3NumRedopType(inst)),
2594                Err(_) => {}
2595            }
2596            stream.set_position(start_pos);
2597            match <crate::r#type::instruction::tcgen05_ld::section_0::Tcgen05LdRedSyncAlignedShape4NumRedopType as PtxParser>::parse(stream) {
2598                Ok(inst) => return Ok(Inst::Tcgen05LdRedSyncAlignedShape4NumRedopType(inst)),
2599                Err(_) => {}
2600            }
2601            stream.set_position(start_pos);
2602            match <crate::r#type::instruction::tcgen05_mma_sp::section_0::Tcgen05MmaSpCtaGroupKind as PtxParser>::parse(stream) {
2603                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKind(inst)),
2604                Err(_) => {}
2605            }
2606            stream.set_position(start_pos);
2607            match <crate::r#type::instruction::tcgen05_mma_sp::section_0::Tcgen05MmaSpCtaGroupKind1 as PtxParser>::parse(stream) {
2608                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKind1(inst)),
2609                Err(_) => {}
2610            }
2611            stream.set_position(start_pos);
2612            match <crate::r#type::instruction::tcgen05_mma_sp::section_1::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize as PtxParser>::parse(stream) {
2613                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize(inst)),
2614                Err(_) => {}
2615            }
2616            stream.set_position(start_pos);
2617            match <crate::r#type::instruction::tcgen05_mma_sp::section_1::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize1 as PtxParser>::parse(stream) {
2618                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsize1(inst)),
2619                Err(_) => {}
2620            }
2621            stream.set_position(start_pos);
2622            match <crate::r#type::instruction::tcgen05_mma_sp::section_2::Tcgen05MmaSpCtaGroupKindCollectorUsage as PtxParser>::parse(stream) {
2623                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindCollectorUsage(inst)),
2624                Err(_) => {}
2625            }
2626            stream.set_position(start_pos);
2627            match <crate::r#type::instruction::tcgen05_mma_sp::section_2::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage as PtxParser>::parse(stream) {
2628                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage(inst)),
2629                Err(_) => {}
2630            }
2631            stream.set_position(start_pos);
2632            match <crate::r#type::instruction::tcgen05_mma_sp::section_2::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage1 as PtxParser>::parse(stream) {
2633                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindAshiftCollectorUsage1(inst)),
2634                Err(_) => {}
2635            }
2636            stream.set_position(start_pos);
2637            match <crate::r#type::instruction::tcgen05_mma_sp::section_3::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage as PtxParser>::parse(stream) {
2638                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage(inst)),
2639                Err(_) => {}
2640            }
2641            stream.set_position(start_pos);
2642            match <crate::r#type::instruction::tcgen05_mma_sp::section_3::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 as PtxParser>::parse(stream) {
2643                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1(inst)),
2644                Err(_) => {}
2645            }
2646            stream.set_position(start_pos);
2647            match <crate::r#type::instruction::tcgen05_mma_sp::section_4::Tcgen05MmaSpCtaGroupKindI8 as PtxParser>::parse(stream) {
2648                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindI8(inst)),
2649                Err(_) => {}
2650            }
2651            stream.set_position(start_pos);
2652            match <crate::r#type::instruction::tcgen05_mma_sp::section_4::Tcgen05MmaSpCtaGroupKindI81 as PtxParser>::parse(stream) {
2653                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindI81(inst)),
2654                Err(_) => {}
2655            }
2656            stream.set_position(start_pos);
2657            match <crate::r#type::instruction::tcgen05_mma_sp::section_5::Tcgen05MmaSpCtaGroupKindI8CollectorUsage as PtxParser>::parse(stream) {
2658                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindI8CollectorUsage(inst)),
2659                Err(_) => {}
2660            }
2661            stream.set_position(start_pos);
2662            match <crate::r#type::instruction::tcgen05_mma_sp::section_5::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage as PtxParser>::parse(stream) {
2663                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage(inst)),
2664                Err(_) => {}
2665            }
2666            stream.set_position(start_pos);
2667            match <crate::r#type::instruction::tcgen05_mma_sp::section_5::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage1 as PtxParser>::parse(stream) {
2668                Ok(inst) => return Ok(Inst::Tcgen05MmaSpCtaGroupKindI8AshiftCollectorUsage1(inst)),
2669                Err(_) => {}
2670            }
2671            stream.set_position(start_pos);
2672            match <crate::r#type::instruction::tcgen05_mma::section_0::Tcgen05MmaCtaGroupKind as PtxParser>::parse(stream) {
2673                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKind(inst)),
2674                Err(_) => {}
2675            }
2676            stream.set_position(start_pos);
2677            match <crate::r#type::instruction::tcgen05_mma::section_0::Tcgen05MmaCtaGroupKind1 as PtxParser>::parse(stream) {
2678                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKind1(inst)),
2679                Err(_) => {}
2680            }
2681            stream.set_position(start_pos);
2682            match <crate::r#type::instruction::tcgen05_mma::section_1::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize as PtxParser>::parse(stream) {
2683                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize(inst)),
2684                Err(_) => {}
2685            }
2686            stream.set_position(start_pos);
2687            match <crate::r#type::instruction::tcgen05_mma::section_1::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1 as PtxParser>::parse(stream) {
2688                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsize1(inst)),
2689                Err(_) => {}
2690            }
2691            stream.set_position(start_pos);
2692            match <crate::r#type::instruction::tcgen05_mma::section_2::Tcgen05MmaCtaGroupKindCollectorUsage as PtxParser>::parse(stream) {
2693                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindCollectorUsage(inst)),
2694                Err(_) => {}
2695            }
2696            stream.set_position(start_pos);
2697            match <crate::r#type::instruction::tcgen05_mma::section_2::Tcgen05MmaCtaGroupKindAshiftCollectorUsage as PtxParser>::parse(stream) {
2698                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindAshiftCollectorUsage(inst)),
2699                Err(_) => {}
2700            }
2701            stream.set_position(start_pos);
2702            match <crate::r#type::instruction::tcgen05_mma::section_2::Tcgen05MmaCtaGroupKindAshiftCollectorUsage1 as PtxParser>::parse(stream) {
2703                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindAshiftCollectorUsage1(inst)),
2704                Err(_) => {}
2705            }
2706            stream.set_position(start_pos);
2707            match <crate::r#type::instruction::tcgen05_mma::section_3::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage as PtxParser>::parse(stream) {
2708                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage(inst)),
2709                Err(_) => {}
2710            }
2711            stream.set_position(start_pos);
2712            match <crate::r#type::instruction::tcgen05_mma::section_3::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1 as PtxParser>::parse(stream) {
2713                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindBlockScaleScaleVectorsizeCollectorUsage1(inst)),
2714                Err(_) => {}
2715            }
2716            stream.set_position(start_pos);
2717            match <crate::r#type::instruction::tcgen05_mma::section_4::Tcgen05MmaCtaGroupKindI8 as PtxParser>::parse(stream) {
2718                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindI8(inst)),
2719                Err(_) => {}
2720            }
2721            stream.set_position(start_pos);
2722            match <crate::r#type::instruction::tcgen05_mma::section_4::Tcgen05MmaCtaGroupKindI81 as PtxParser>::parse(stream) {
2723                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindI81(inst)),
2724                Err(_) => {}
2725            }
2726            stream.set_position(start_pos);
2727            match <crate::r#type::instruction::tcgen05_mma::section_5::Tcgen05MmaCtaGroupKindI8CollectorUsage as PtxParser>::parse(stream) {
2728                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindI8CollectorUsage(inst)),
2729                Err(_) => {}
2730            }
2731            stream.set_position(start_pos);
2732            match <crate::r#type::instruction::tcgen05_mma::section_5::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage as PtxParser>::parse(stream) {
2733                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage(inst)),
2734                Err(_) => {}
2735            }
2736            stream.set_position(start_pos);
2737            match <crate::r#type::instruction::tcgen05_mma::section_5::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1 as PtxParser>::parse(stream) {
2738                Ok(inst) => return Ok(Inst::Tcgen05MmaCtaGroupKindI8AshiftCollectorUsage1(inst)),
2739                Err(_) => {}
2740            }
2741            stream.set_position(start_pos);
2742            match <crate::r#type::instruction::tcgen05_mma_ws_sp::section_0::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage as PtxParser>::parse(stream) {
2743                Ok(inst) => return Ok(Inst::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage(inst)),
2744                Err(_) => {}
2745            }
2746            stream.set_position(start_pos);
2747            match <crate::r#type::instruction::tcgen05_mma_ws_sp::section_0::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage1 as PtxParser>::parse(stream) {
2748                Ok(inst) => return Ok(Inst::Tcgen05MmaWsSpCtaGroup1KindCollectorUsage1(inst)),
2749                Err(_) => {}
2750            }
2751            stream.set_position(start_pos);
2752            match <crate::r#type::instruction::tcgen05_mma_ws_sp::section_1::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage as PtxParser>::parse(stream) {
2753                Ok(inst) => return Ok(Inst::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage(inst)),
2754                Err(_) => {}
2755            }
2756            stream.set_position(start_pos);
2757            match <crate::r#type::instruction::tcgen05_mma_ws_sp::section_1::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage1 as PtxParser>::parse(stream) {
2758                Ok(inst) => return Ok(Inst::Tcgen05MmaWsSpCtaGroup1KindI8CollectorUsage1(inst)),
2759                Err(_) => {}
2760            }
2761            stream.set_position(start_pos);
2762            match <crate::r#type::instruction::tcgen05_mma_ws::section_0::Tcgen05MmaWsCtaGroup1KindCollectorUsage as PtxParser>::parse(stream) {
2763                Ok(inst) => return Ok(Inst::Tcgen05MmaWsCtaGroup1KindCollectorUsage(inst)),
2764                Err(_) => {}
2765            }
2766            stream.set_position(start_pos);
2767            match <crate::r#type::instruction::tcgen05_mma_ws::section_0::Tcgen05MmaWsCtaGroup1KindCollectorUsage1 as PtxParser>::parse(stream) {
2768                Ok(inst) => return Ok(Inst::Tcgen05MmaWsCtaGroup1KindCollectorUsage1(inst)),
2769                Err(_) => {}
2770            }
2771            stream.set_position(start_pos);
2772            match <crate::r#type::instruction::tcgen05_mma_ws::section_1::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage as PtxParser>::parse(stream) {
2773                Ok(inst) => return Ok(Inst::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage(inst)),
2774                Err(_) => {}
2775            }
2776            stream.set_position(start_pos);
2777            match <crate::r#type::instruction::tcgen05_mma_ws::section_1::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1 as PtxParser>::parse(stream) {
2778                Ok(inst) => return Ok(Inst::Tcgen05MmaWsCtaGroup1KindI8CollectorUsage1(inst)),
2779                Err(_) => {}
2780            }
2781            stream.set_position(start_pos);
2782            match <crate::r#type::instruction::tcgen05_shift::section_0::Tcgen05ShiftCtaGroupDown as PtxParser>::parse(stream) {
2783                Ok(inst) => return Ok(Inst::Tcgen05ShiftCtaGroupDown(inst)),
2784                Err(_) => {}
2785            }
2786            stream.set_position(start_pos);
2787            match <crate::r#type::instruction::tcgen05_st::section_0::Tcgen05StSyncAlignedShape1NumUnpackB32 as PtxParser>::parse(stream) {
2788                Ok(inst) => return Ok(Inst::Tcgen05StSyncAlignedShape1NumUnpackB32(inst)),
2789                Err(_) => {}
2790            }
2791            stream.set_position(start_pos);
2792            match <crate::r#type::instruction::tcgen05_st::section_0::Tcgen05StSyncAlignedShape2NumUnpackB32 as PtxParser>::parse(stream) {
2793                Ok(inst) => return Ok(Inst::Tcgen05StSyncAlignedShape2NumUnpackB32(inst)),
2794                Err(_) => {}
2795            }
2796            stream.set_position(start_pos);
2797            match <crate::r#type::instruction::tcgen05_wait::section_0::Tcgen05WaitOperationSyncAligned as PtxParser>::parse(stream) {
2798                Ok(inst) => return Ok(Inst::Tcgen05WaitOperationSyncAligned(inst)),
2799                Err(_) => {}
2800            }
2801        }
2802        "tensormap" => {
2803            stream.set_position(start_pos);
2804            match <crate::r#type::instruction::tensormap_cp_fenceproxy::section_0::TensormapCpFenceproxyCpQualifiersFenceQualifiersSyncAligned as PtxParser>::parse(stream) {
2805                Ok(inst) => return Ok(Inst::TensormapCpFenceproxyCpQualifiersFenceQualifiersSyncAligned(inst)),
2806                Err(_) => {}
2807            }
2808            stream.set_position(start_pos);
2809            match <crate::r#type::instruction::tensormap_replace::section_0::TensormapReplaceModeField1SsB1024Type as PtxParser>::parse(stream) {
2810                Ok(inst) => return Ok(Inst::TensormapReplaceModeField1SsB1024Type(inst)),
2811                Err(_) => {}
2812            }
2813            stream.set_position(start_pos);
2814            match <crate::r#type::instruction::tensormap_replace::section_0::TensormapReplaceModeField2SsB1024Type as PtxParser>::parse(stream) {
2815                Ok(inst) => return Ok(Inst::TensormapReplaceModeField2SsB1024Type(inst)),
2816                Err(_) => {}
2817            }
2818            stream.set_position(start_pos);
2819            match <crate::r#type::instruction::tensormap_replace::section_0::TensormapReplaceModeField3SsB1024Type as PtxParser>::parse(stream) {
2820                Ok(inst) => return Ok(Inst::TensormapReplaceModeField3SsB1024Type(inst)),
2821                Err(_) => {}
2822            }
2823        }
2824        "testp" => {
2825            stream.set_position(start_pos);
2826            match <crate::r#type::instruction::testp::section_0::TestpOpType as PtxParser>::parse(stream) {
2827                Ok(inst) => return Ok(Inst::TestpOpType(inst)),
2828                Err(_) => {}
2829            }
2830        }
2831        "tex" => {
2832            stream.set_position(start_pos);
2833            match <crate::r#type::instruction::tex::section_0::TexGeomV4DtypeCtype as PtxParser>::parse(stream) {
2834                Ok(inst) => return Ok(Inst::TexGeomV4DtypeCtype(inst)),
2835                Err(_) => {}
2836            }
2837            stream.set_position(start_pos);
2838            match <crate::r#type::instruction::tex::section_0::TexGeomV4DtypeCtype1 as PtxParser>::parse(stream) {
2839                Ok(inst) => return Ok(Inst::TexGeomV4DtypeCtype1(inst)),
2840                Err(_) => {}
2841            }
2842            stream.set_position(start_pos);
2843            match <crate::r#type::instruction::tex::section_0::TexGeomV2F16x2Ctype as PtxParser>::parse(stream) {
2844                Ok(inst) => return Ok(Inst::TexGeomV2F16x2Ctype(inst)),
2845                Err(_) => {}
2846            }
2847            stream.set_position(start_pos);
2848            match <crate::r#type::instruction::tex::section_0::TexGeomV2F16x2Ctype1 as PtxParser>::parse(stream) {
2849                Ok(inst) => return Ok(Inst::TexGeomV2F16x2Ctype1(inst)),
2850                Err(_) => {}
2851            }
2852            stream.set_position(start_pos);
2853            match <crate::r#type::instruction::tex::section_0::TexBaseGeomV4DtypeCtype as PtxParser>::parse(stream) {
2854                Ok(inst) => return Ok(Inst::TexBaseGeomV4DtypeCtype(inst)),
2855                Err(_) => {}
2856            }
2857            stream.set_position(start_pos);
2858            match <crate::r#type::instruction::tex::section_0::TexLevelGeomV4DtypeCtype as PtxParser>::parse(stream) {
2859                Ok(inst) => return Ok(Inst::TexLevelGeomV4DtypeCtype(inst)),
2860                Err(_) => {}
2861            }
2862            stream.set_position(start_pos);
2863            match <crate::r#type::instruction::tex::section_0::TexGradGeomV4DtypeCtype as PtxParser>::parse(stream) {
2864                Ok(inst) => return Ok(Inst::TexGradGeomV4DtypeCtype(inst)),
2865                Err(_) => {}
2866            }
2867            stream.set_position(start_pos);
2868            match <crate::r#type::instruction::tex::section_0::TexBaseGeomV2F16x2Ctype as PtxParser>::parse(stream) {
2869                Ok(inst) => return Ok(Inst::TexBaseGeomV2F16x2Ctype(inst)),
2870                Err(_) => {}
2871            }
2872            stream.set_position(start_pos);
2873            match <crate::r#type::instruction::tex::section_0::TexLevelGeomV2F16x2Ctype as PtxParser>::parse(stream) {
2874                Ok(inst) => return Ok(Inst::TexLevelGeomV2F16x2Ctype(inst)),
2875                Err(_) => {}
2876            }
2877            stream.set_position(start_pos);
2878            match <crate::r#type::instruction::tex::section_0::TexGradGeomV2F16x2Ctype as PtxParser>::parse(stream) {
2879                Ok(inst) => return Ok(Inst::TexGradGeomV2F16x2Ctype(inst)),
2880                Err(_) => {}
2881            }
2882        }
2883        "tld4" => {
2884            stream.set_position(start_pos);
2885            match <crate::r#type::instruction::tld4::section_0::Tld4Comp2dV4DtypeF32 as PtxParser>::parse(stream) {
2886                Ok(inst) => return Ok(Inst::Tld4Comp2dV4DtypeF32(inst)),
2887                Err(_) => {}
2888            }
2889            stream.set_position(start_pos);
2890            match <crate::r#type::instruction::tld4::section_0::Tld4CompGeomV4DtypeF32 as PtxParser>::parse(stream) {
2891                Ok(inst) => return Ok(Inst::Tld4CompGeomV4DtypeF32(inst)),
2892                Err(_) => {}
2893            }
2894        }
2895        "trap" => {
2896            stream.set_position(start_pos);
2897            match <crate::r#type::instruction::trap::section_0::Trap as PtxParser>::parse(stream) {
2898                Ok(inst) => return Ok(Inst::Trap(inst)),
2899                Err(_) => {}
2900            }
2901        }
2902        "txq" => {
2903            stream.set_position(start_pos);
2904            match <crate::r#type::instruction::txq::section_0::TxqTqueryB32 as PtxParser>::parse(stream) {
2905                Ok(inst) => return Ok(Inst::TxqTqueryB32(inst)),
2906                Err(_) => {}
2907            }
2908            stream.set_position(start_pos);
2909            match <crate::r#type::instruction::txq::section_0::TxqLevelTlqueryB32 as PtxParser>::parse(stream) {
2910                Ok(inst) => return Ok(Inst::TxqLevelTlqueryB32(inst)),
2911                Err(_) => {}
2912            }
2913            stream.set_position(start_pos);
2914            match <crate::r#type::instruction::txq::section_0::TxqSqueryB32 as PtxParser>::parse(stream) {
2915                Ok(inst) => return Ok(Inst::TxqSqueryB32(inst)),
2916                Err(_) => {}
2917            }
2918        }
2919        "vabsdiff" => {
2920            stream.set_position(start_pos);
2921            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2922                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat(inst)),
2923                Err(_) => {}
2924            }
2925            stream.set_position(start_pos);
2926            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2927                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat(inst)),
2928                Err(_) => {}
2929            }
2930            stream.set_position(start_pos);
2931            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2932                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat(inst)),
2933                Err(_) => {}
2934            }
2935            stream.set_position(start_pos);
2936            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2937                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat(inst)),
2938                Err(_) => {}
2939            }
2940            stream.set_position(start_pos);
2941            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2942                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat(inst)),
2943                Err(_) => {}
2944            }
2945            stream.set_position(start_pos);
2946            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
2947                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSatOp2(inst)),
2948                Err(_) => {}
2949            }
2950            stream.set_position(start_pos);
2951            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
2952                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSatOp2(inst)),
2953                Err(_) => {}
2954            }
2955            stream.set_position(start_pos);
2956            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
2957                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSatOp2(inst)),
2958                Err(_) => {}
2959            }
2960            stream.set_position(start_pos);
2961            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
2962                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSatOp2(inst)),
2963                Err(_) => {}
2964            }
2965            stream.set_position(start_pos);
2966            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
2967                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSatOp2(inst)),
2968                Err(_) => {}
2969            }
2970            stream.set_position(start_pos);
2971            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
2972                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat1(inst)),
2973                Err(_) => {}
2974            }
2975            stream.set_position(start_pos);
2976            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
2977                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat1(inst)),
2978                Err(_) => {}
2979            }
2980            stream.set_position(start_pos);
2981            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
2982                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat1(inst)),
2983                Err(_) => {}
2984            }
2985            stream.set_position(start_pos);
2986            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
2987                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat1(inst)),
2988                Err(_) => {}
2989            }
2990            stream.set_position(start_pos);
2991            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
2992                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat1(inst)),
2993                Err(_) => {}
2994            }
2995        }
2996        "vabsdiff2" => {
2997            stream.set_position(start_pos);
2998            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
2999                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
3000                Err(_) => {}
3001            }
3002            stream.set_position(start_pos);
3003            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3004                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
3005                Err(_) => {}
3006            }
3007            stream.set_position(start_pos);
3008            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3009                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
3010                Err(_) => {}
3011            }
3012            stream.set_position(start_pos);
3013            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3014                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
3015                Err(_) => {}
3016            }
3017            stream.set_position(start_pos);
3018            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3019                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
3020                Err(_) => {}
3021            }
3022            stream.set_position(start_pos);
3023            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3024                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
3025                Err(_) => {}
3026            }
3027            stream.set_position(start_pos);
3028            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3029                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
3030                Err(_) => {}
3031            }
3032            stream.set_position(start_pos);
3033            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3034                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
3035                Err(_) => {}
3036            }
3037            stream.set_position(start_pos);
3038            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3039                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
3040                Err(_) => {}
3041            }
3042            stream.set_position(start_pos);
3043            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3044                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
3045                Err(_) => {}
3046            }
3047            stream.set_position(start_pos);
3048            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3049                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
3050                Err(_) => {}
3051            }
3052            stream.set_position(start_pos);
3053            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3054                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
3055                Err(_) => {}
3056            }
3057        }
3058        "vabsdiff4" => {
3059            stream.set_position(start_pos);
3060            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3061                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
3062                Err(_) => {}
3063            }
3064            stream.set_position(start_pos);
3065            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3066                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
3067                Err(_) => {}
3068            }
3069            stream.set_position(start_pos);
3070            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3071                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
3072                Err(_) => {}
3073            }
3074            stream.set_position(start_pos);
3075            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3076                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
3077                Err(_) => {}
3078            }
3079            stream.set_position(start_pos);
3080            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3081                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
3082                Err(_) => {}
3083            }
3084            stream.set_position(start_pos);
3085            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3086                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
3087                Err(_) => {}
3088            }
3089            stream.set_position(start_pos);
3090            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3091                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
3092                Err(_) => {}
3093            }
3094            stream.set_position(start_pos);
3095            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3096                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
3097                Err(_) => {}
3098            }
3099            stream.set_position(start_pos);
3100            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3101                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
3102                Err(_) => {}
3103            }
3104            stream.set_position(start_pos);
3105            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3106                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
3107                Err(_) => {}
3108            }
3109            stream.set_position(start_pos);
3110            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3111                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
3112                Err(_) => {}
3113            }
3114            stream.set_position(start_pos);
3115            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3116                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
3117                Err(_) => {}
3118            }
3119        }
3120        "vadd" => {
3121            stream.set_position(start_pos);
3122            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3123                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat(inst)),
3124                Err(_) => {}
3125            }
3126            stream.set_position(start_pos);
3127            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3128                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat(inst)),
3129                Err(_) => {}
3130            }
3131            stream.set_position(start_pos);
3132            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3133                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat(inst)),
3134                Err(_) => {}
3135            }
3136            stream.set_position(start_pos);
3137            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3138                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat(inst)),
3139                Err(_) => {}
3140            }
3141            stream.set_position(start_pos);
3142            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3143                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat(inst)),
3144                Err(_) => {}
3145            }
3146            stream.set_position(start_pos);
3147            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3148                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSatOp2(inst)),
3149                Err(_) => {}
3150            }
3151            stream.set_position(start_pos);
3152            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3153                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSatOp2(inst)),
3154                Err(_) => {}
3155            }
3156            stream.set_position(start_pos);
3157            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3158                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSatOp2(inst)),
3159                Err(_) => {}
3160            }
3161            stream.set_position(start_pos);
3162            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3163                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSatOp2(inst)),
3164                Err(_) => {}
3165            }
3166            stream.set_position(start_pos);
3167            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3168                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSatOp2(inst)),
3169                Err(_) => {}
3170            }
3171            stream.set_position(start_pos);
3172            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3173                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat1(inst)),
3174                Err(_) => {}
3175            }
3176            stream.set_position(start_pos);
3177            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3178                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat1(inst)),
3179                Err(_) => {}
3180            }
3181            stream.set_position(start_pos);
3182            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3183                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat1(inst)),
3184                Err(_) => {}
3185            }
3186            stream.set_position(start_pos);
3187            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3188                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat1(inst)),
3189                Err(_) => {}
3190            }
3191            stream.set_position(start_pos);
3192            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3193                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat1(inst)),
3194                Err(_) => {}
3195            }
3196        }
3197        "vadd2" => {
3198            stream.set_position(start_pos);
3199            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3200                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
3201                Err(_) => {}
3202            }
3203            stream.set_position(start_pos);
3204            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3205                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
3206                Err(_) => {}
3207            }
3208            stream.set_position(start_pos);
3209            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3210                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
3211                Err(_) => {}
3212            }
3213            stream.set_position(start_pos);
3214            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3215                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
3216                Err(_) => {}
3217            }
3218            stream.set_position(start_pos);
3219            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3220                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
3221                Err(_) => {}
3222            }
3223            stream.set_position(start_pos);
3224            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3225                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
3226                Err(_) => {}
3227            }
3228            stream.set_position(start_pos);
3229            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3230                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
3231                Err(_) => {}
3232            }
3233            stream.set_position(start_pos);
3234            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3235                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
3236                Err(_) => {}
3237            }
3238            stream.set_position(start_pos);
3239            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3240                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
3241                Err(_) => {}
3242            }
3243            stream.set_position(start_pos);
3244            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3245                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
3246                Err(_) => {}
3247            }
3248            stream.set_position(start_pos);
3249            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3250                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
3251                Err(_) => {}
3252            }
3253            stream.set_position(start_pos);
3254            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3255                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
3256                Err(_) => {}
3257            }
3258        }
3259        "vadd4" => {
3260            stream.set_position(start_pos);
3261            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3262                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
3263                Err(_) => {}
3264            }
3265            stream.set_position(start_pos);
3266            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3267                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
3268                Err(_) => {}
3269            }
3270            stream.set_position(start_pos);
3271            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3272                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
3273                Err(_) => {}
3274            }
3275            stream.set_position(start_pos);
3276            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3277                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
3278                Err(_) => {}
3279            }
3280            stream.set_position(start_pos);
3281            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3282                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
3283                Err(_) => {}
3284            }
3285            stream.set_position(start_pos);
3286            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3287                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
3288                Err(_) => {}
3289            }
3290            stream.set_position(start_pos);
3291            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3292                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
3293                Err(_) => {}
3294            }
3295            stream.set_position(start_pos);
3296            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3297                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
3298                Err(_) => {}
3299            }
3300            stream.set_position(start_pos);
3301            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3302                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
3303                Err(_) => {}
3304            }
3305            stream.set_position(start_pos);
3306            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3307                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
3308                Err(_) => {}
3309            }
3310            stream.set_position(start_pos);
3311            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3312                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
3313                Err(_) => {}
3314            }
3315            stream.set_position(start_pos);
3316            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3317                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
3318                Err(_) => {}
3319            }
3320        }
3321        "vavrg2" => {
3322            stream.set_position(start_pos);
3323            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3324                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
3325                Err(_) => {}
3326            }
3327            stream.set_position(start_pos);
3328            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3329                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
3330                Err(_) => {}
3331            }
3332            stream.set_position(start_pos);
3333            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3334                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
3335                Err(_) => {}
3336            }
3337            stream.set_position(start_pos);
3338            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3339                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
3340                Err(_) => {}
3341            }
3342            stream.set_position(start_pos);
3343            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3344                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
3345                Err(_) => {}
3346            }
3347            stream.set_position(start_pos);
3348            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3349                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
3350                Err(_) => {}
3351            }
3352            stream.set_position(start_pos);
3353            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3354                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
3355                Err(_) => {}
3356            }
3357            stream.set_position(start_pos);
3358            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3359                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
3360                Err(_) => {}
3361            }
3362            stream.set_position(start_pos);
3363            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3364                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
3365                Err(_) => {}
3366            }
3367            stream.set_position(start_pos);
3368            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3369                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
3370                Err(_) => {}
3371            }
3372            stream.set_position(start_pos);
3373            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3374                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
3375                Err(_) => {}
3376            }
3377            stream.set_position(start_pos);
3378            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3379                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
3380                Err(_) => {}
3381            }
3382        }
3383        "vavrg4" => {
3384            stream.set_position(start_pos);
3385            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3386                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
3387                Err(_) => {}
3388            }
3389            stream.set_position(start_pos);
3390            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3391                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
3392                Err(_) => {}
3393            }
3394            stream.set_position(start_pos);
3395            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3396                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
3397                Err(_) => {}
3398            }
3399            stream.set_position(start_pos);
3400            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3401                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
3402                Err(_) => {}
3403            }
3404            stream.set_position(start_pos);
3405            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3406                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
3407                Err(_) => {}
3408            }
3409            stream.set_position(start_pos);
3410            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3411                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
3412                Err(_) => {}
3413            }
3414            stream.set_position(start_pos);
3415            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3416                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
3417                Err(_) => {}
3418            }
3419            stream.set_position(start_pos);
3420            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3421                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
3422                Err(_) => {}
3423            }
3424            stream.set_position(start_pos);
3425            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3426                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
3427                Err(_) => {}
3428            }
3429            stream.set_position(start_pos);
3430            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3431                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
3432                Err(_) => {}
3433            }
3434            stream.set_position(start_pos);
3435            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3436                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
3437                Err(_) => {}
3438            }
3439            stream.set_position(start_pos);
3440            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3441                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
3442                Err(_) => {}
3443            }
3444        }
3445        "vmad" => {
3446            stream.set_position(start_pos);
3447            match <crate::r#type::instruction::vmad::section_0::VmadDtypeAtypeBtypeSatScale as PtxParser>::parse(stream) {
3448                Ok(inst) => return Ok(Inst::VmadDtypeAtypeBtypeSatScale(inst)),
3449                Err(_) => {}
3450            }
3451            stream.set_position(start_pos);
3452            match <crate::r#type::instruction::vmad::section_0::VmadDtypeAtypeBtypePoSatScale as PtxParser>::parse(stream) {
3453                Ok(inst) => return Ok(Inst::VmadDtypeAtypeBtypePoSatScale(inst)),
3454                Err(_) => {}
3455            }
3456        }
3457        "vmax" => {
3458            stream.set_position(start_pos);
3459            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3460                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat(inst)),
3461                Err(_) => {}
3462            }
3463            stream.set_position(start_pos);
3464            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3465                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat(inst)),
3466                Err(_) => {}
3467            }
3468            stream.set_position(start_pos);
3469            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3470                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat(inst)),
3471                Err(_) => {}
3472            }
3473            stream.set_position(start_pos);
3474            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3475                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat(inst)),
3476                Err(_) => {}
3477            }
3478            stream.set_position(start_pos);
3479            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3480                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat(inst)),
3481                Err(_) => {}
3482            }
3483            stream.set_position(start_pos);
3484            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3485                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSatOp2(inst)),
3486                Err(_) => {}
3487            }
3488            stream.set_position(start_pos);
3489            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3490                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSatOp2(inst)),
3491                Err(_) => {}
3492            }
3493            stream.set_position(start_pos);
3494            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3495                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSatOp2(inst)),
3496                Err(_) => {}
3497            }
3498            stream.set_position(start_pos);
3499            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3500                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSatOp2(inst)),
3501                Err(_) => {}
3502            }
3503            stream.set_position(start_pos);
3504            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3505                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSatOp2(inst)),
3506                Err(_) => {}
3507            }
3508            stream.set_position(start_pos);
3509            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3510                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat1(inst)),
3511                Err(_) => {}
3512            }
3513            stream.set_position(start_pos);
3514            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3515                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat1(inst)),
3516                Err(_) => {}
3517            }
3518            stream.set_position(start_pos);
3519            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3520                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat1(inst)),
3521                Err(_) => {}
3522            }
3523            stream.set_position(start_pos);
3524            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3525                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat1(inst)),
3526                Err(_) => {}
3527            }
3528            stream.set_position(start_pos);
3529            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3530                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat1(inst)),
3531                Err(_) => {}
3532            }
3533        }
3534        "vmax2" => {
3535            stream.set_position(start_pos);
3536            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3537                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
3538                Err(_) => {}
3539            }
3540            stream.set_position(start_pos);
3541            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3542                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
3543                Err(_) => {}
3544            }
3545            stream.set_position(start_pos);
3546            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3547                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
3548                Err(_) => {}
3549            }
3550            stream.set_position(start_pos);
3551            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3552                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
3553                Err(_) => {}
3554            }
3555            stream.set_position(start_pos);
3556            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3557                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
3558                Err(_) => {}
3559            }
3560            stream.set_position(start_pos);
3561            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3562                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
3563                Err(_) => {}
3564            }
3565            stream.set_position(start_pos);
3566            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3567                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
3568                Err(_) => {}
3569            }
3570            stream.set_position(start_pos);
3571            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3572                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
3573                Err(_) => {}
3574            }
3575            stream.set_position(start_pos);
3576            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3577                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
3578                Err(_) => {}
3579            }
3580            stream.set_position(start_pos);
3581            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3582                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
3583                Err(_) => {}
3584            }
3585            stream.set_position(start_pos);
3586            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3587                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
3588                Err(_) => {}
3589            }
3590            stream.set_position(start_pos);
3591            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3592                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
3593                Err(_) => {}
3594            }
3595        }
3596        "vmax4" => {
3597            stream.set_position(start_pos);
3598            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3599                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
3600                Err(_) => {}
3601            }
3602            stream.set_position(start_pos);
3603            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3604                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
3605                Err(_) => {}
3606            }
3607            stream.set_position(start_pos);
3608            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3609                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
3610                Err(_) => {}
3611            }
3612            stream.set_position(start_pos);
3613            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3614                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
3615                Err(_) => {}
3616            }
3617            stream.set_position(start_pos);
3618            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3619                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
3620                Err(_) => {}
3621            }
3622            stream.set_position(start_pos);
3623            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3624                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
3625                Err(_) => {}
3626            }
3627            stream.set_position(start_pos);
3628            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3629                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
3630                Err(_) => {}
3631            }
3632            stream.set_position(start_pos);
3633            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3634                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
3635                Err(_) => {}
3636            }
3637            stream.set_position(start_pos);
3638            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3639                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
3640                Err(_) => {}
3641            }
3642            stream.set_position(start_pos);
3643            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3644                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
3645                Err(_) => {}
3646            }
3647            stream.set_position(start_pos);
3648            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3649                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
3650                Err(_) => {}
3651            }
3652            stream.set_position(start_pos);
3653            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3654                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
3655                Err(_) => {}
3656            }
3657        }
3658        "vmin" => {
3659            stream.set_position(start_pos);
3660            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3661                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat(inst)),
3662                Err(_) => {}
3663            }
3664            stream.set_position(start_pos);
3665            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3666                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat(inst)),
3667                Err(_) => {}
3668            }
3669            stream.set_position(start_pos);
3670            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3671                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat(inst)),
3672                Err(_) => {}
3673            }
3674            stream.set_position(start_pos);
3675            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3676                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat(inst)),
3677                Err(_) => {}
3678            }
3679            stream.set_position(start_pos);
3680            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3681                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat(inst)),
3682                Err(_) => {}
3683            }
3684            stream.set_position(start_pos);
3685            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3686                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSatOp2(inst)),
3687                Err(_) => {}
3688            }
3689            stream.set_position(start_pos);
3690            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3691                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSatOp2(inst)),
3692                Err(_) => {}
3693            }
3694            stream.set_position(start_pos);
3695            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3696                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSatOp2(inst)),
3697                Err(_) => {}
3698            }
3699            stream.set_position(start_pos);
3700            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3701                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSatOp2(inst)),
3702                Err(_) => {}
3703            }
3704            stream.set_position(start_pos);
3705            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
3706                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSatOp2(inst)),
3707                Err(_) => {}
3708            }
3709            stream.set_position(start_pos);
3710            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3711                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat1(inst)),
3712                Err(_) => {}
3713            }
3714            stream.set_position(start_pos);
3715            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3716                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat1(inst)),
3717                Err(_) => {}
3718            }
3719            stream.set_position(start_pos);
3720            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3721                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat1(inst)),
3722                Err(_) => {}
3723            }
3724            stream.set_position(start_pos);
3725            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3726                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat1(inst)),
3727                Err(_) => {}
3728            }
3729            stream.set_position(start_pos);
3730            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
3731                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat1(inst)),
3732                Err(_) => {}
3733            }
3734        }
3735        "vmin2" => {
3736            stream.set_position(start_pos);
3737            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3738                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
3739                Err(_) => {}
3740            }
3741            stream.set_position(start_pos);
3742            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3743                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
3744                Err(_) => {}
3745            }
3746            stream.set_position(start_pos);
3747            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3748                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
3749                Err(_) => {}
3750            }
3751            stream.set_position(start_pos);
3752            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3753                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
3754                Err(_) => {}
3755            }
3756            stream.set_position(start_pos);
3757            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3758                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
3759                Err(_) => {}
3760            }
3761            stream.set_position(start_pos);
3762            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3763                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
3764                Err(_) => {}
3765            }
3766            stream.set_position(start_pos);
3767            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3768                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
3769                Err(_) => {}
3770            }
3771            stream.set_position(start_pos);
3772            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3773                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
3774                Err(_) => {}
3775            }
3776            stream.set_position(start_pos);
3777            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3778                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
3779                Err(_) => {}
3780            }
3781            stream.set_position(start_pos);
3782            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3783                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
3784                Err(_) => {}
3785            }
3786            stream.set_position(start_pos);
3787            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3788                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
3789                Err(_) => {}
3790            }
3791            stream.set_position(start_pos);
3792            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3793                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
3794                Err(_) => {}
3795            }
3796        }
3797        "vmin4" => {
3798            stream.set_position(start_pos);
3799            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3800                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
3801                Err(_) => {}
3802            }
3803            stream.set_position(start_pos);
3804            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3805                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
3806                Err(_) => {}
3807            }
3808            stream.set_position(start_pos);
3809            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3810                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
3811                Err(_) => {}
3812            }
3813            stream.set_position(start_pos);
3814            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3815                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
3816                Err(_) => {}
3817            }
3818            stream.set_position(start_pos);
3819            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3820                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
3821                Err(_) => {}
3822            }
3823            stream.set_position(start_pos);
3824            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3825                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
3826                Err(_) => {}
3827            }
3828            stream.set_position(start_pos);
3829            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3830                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
3831                Err(_) => {}
3832            }
3833            stream.set_position(start_pos);
3834            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3835                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
3836                Err(_) => {}
3837            }
3838            stream.set_position(start_pos);
3839            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3840                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
3841                Err(_) => {}
3842            }
3843            stream.set_position(start_pos);
3844            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3845                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
3846                Err(_) => {}
3847            }
3848            stream.set_position(start_pos);
3849            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3850                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
3851                Err(_) => {}
3852            }
3853            stream.set_position(start_pos);
3854            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
3855                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
3856                Err(_) => {}
3857            }
3858        }
3859        "vote" => {
3860            stream.set_position(start_pos);
3861            match <crate::r#type::instruction::vote_sync::section_0::VoteSyncModePred as PtxParser>::parse(stream) {
3862                Ok(inst) => return Ok(Inst::VoteSyncModePred(inst)),
3863                Err(_) => {}
3864            }
3865            stream.set_position(start_pos);
3866            match <crate::r#type::instruction::vote_sync::section_0::VoteSyncBallotB32 as PtxParser>::parse(stream) {
3867                Ok(inst) => return Ok(Inst::VoteSyncBallotB32(inst)),
3868                Err(_) => {}
3869            }
3870            stream.set_position(start_pos);
3871            match <crate::r#type::instruction::vote::section_0::VoteModePred as PtxParser>::parse(stream) {
3872                Ok(inst) => return Ok(Inst::VoteModePred(inst)),
3873                Err(_) => {}
3874            }
3875            stream.set_position(start_pos);
3876            match <crate::r#type::instruction::vote::section_0::VoteBallotB32 as PtxParser>::parse(stream) {
3877                Ok(inst) => return Ok(Inst::VoteBallotB32(inst)),
3878                Err(_) => {}
3879            }
3880        }
3881        "vset" => {
3882            stream.set_position(start_pos);
3883            match <crate::r#type::instruction::vset::section_0::VsetAtypeBtypeCmp as PtxParser>::parse(stream) {
3884                Ok(inst) => return Ok(Inst::VsetAtypeBtypeCmp(inst)),
3885                Err(_) => {}
3886            }
3887            stream.set_position(start_pos);
3888            match <crate::r#type::instruction::vset::section_0::VsetAtypeBtypeCmpOp2 as PtxParser>::parse(stream) {
3889                Ok(inst) => return Ok(Inst::VsetAtypeBtypeCmpOp2(inst)),
3890                Err(_) => {}
3891            }
3892            stream.set_position(start_pos);
3893            match <crate::r#type::instruction::vset::section_0::VsetAtypeBtypeCmp1 as PtxParser>::parse(stream) {
3894                Ok(inst) => return Ok(Inst::VsetAtypeBtypeCmp1(inst)),
3895                Err(_) => {}
3896            }
3897        }
3898        "vset2" => {
3899            stream.set_position(start_pos);
3900            match <crate::r#type::instruction::vset2::section_0::Vset2AtypeBtypeCmp as PtxParser>::parse(stream) {
3901                Ok(inst) => return Ok(Inst::Vset2AtypeBtypeCmp(inst)),
3902                Err(_) => {}
3903            }
3904            stream.set_position(start_pos);
3905            match <crate::r#type::instruction::vset2::section_0::Vset2AtypeBtypeCmpAdd as PtxParser>::parse(stream) {
3906                Ok(inst) => return Ok(Inst::Vset2AtypeBtypeCmpAdd(inst)),
3907                Err(_) => {}
3908            }
3909        }
3910        "vset4" => {
3911            stream.set_position(start_pos);
3912            match <crate::r#type::instruction::vset4::section_0::Vset4AtypeBtypeCmp as PtxParser>::parse(stream) {
3913                Ok(inst) => return Ok(Inst::Vset4AtypeBtypeCmp(inst)),
3914                Err(_) => {}
3915            }
3916            stream.set_position(start_pos);
3917            match <crate::r#type::instruction::vset4::section_0::Vset4AtypeBtypeCmpAdd as PtxParser>::parse(stream) {
3918                Ok(inst) => return Ok(Inst::Vset4AtypeBtypeCmpAdd(inst)),
3919                Err(_) => {}
3920            }
3921        }
3922        "vshl" => {
3923            stream.set_position(start_pos);
3924            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatMode as PtxParser>::parse(stream) {
3925                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatMode(inst)),
3926                Err(_) => {}
3927            }
3928            stream.set_position(start_pos);
3929            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatMode as PtxParser>::parse(stream) {
3930                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatMode(inst)),
3931                Err(_) => {}
3932            }
3933            stream.set_position(start_pos);
3934            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatModeOp2 as PtxParser>::parse(stream) {
3935                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatModeOp2(inst)),
3936                Err(_) => {}
3937            }
3938            stream.set_position(start_pos);
3939            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatModeOp2 as PtxParser>::parse(stream) {
3940                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatModeOp2(inst)),
3941                Err(_) => {}
3942            }
3943            stream.set_position(start_pos);
3944            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatMode1 as PtxParser>::parse(stream) {
3945                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatMode1(inst)),
3946                Err(_) => {}
3947            }
3948            stream.set_position(start_pos);
3949            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatMode1 as PtxParser>::parse(stream) {
3950                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatMode1(inst)),
3951                Err(_) => {}
3952            }
3953        }
3954        "vshr" => {
3955            stream.set_position(start_pos);
3956            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatMode as PtxParser>::parse(stream) {
3957                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatMode(inst)),
3958                Err(_) => {}
3959            }
3960            stream.set_position(start_pos);
3961            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatMode as PtxParser>::parse(stream) {
3962                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatMode(inst)),
3963                Err(_) => {}
3964            }
3965            stream.set_position(start_pos);
3966            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatModeOp2 as PtxParser>::parse(stream) {
3967                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatModeOp2(inst)),
3968                Err(_) => {}
3969            }
3970            stream.set_position(start_pos);
3971            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatModeOp2 as PtxParser>::parse(stream) {
3972                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatModeOp2(inst)),
3973                Err(_) => {}
3974            }
3975            stream.set_position(start_pos);
3976            match <crate::r#type::instruction::vsh::section_0::VshlDtypeAtypeU32SatMode1 as PtxParser>::parse(stream) {
3977                Ok(inst) => return Ok(Inst::VshlDtypeAtypeU32SatMode1(inst)),
3978                Err(_) => {}
3979            }
3980            stream.set_position(start_pos);
3981            match <crate::r#type::instruction::vsh::section_0::VshrDtypeAtypeU32SatMode1 as PtxParser>::parse(stream) {
3982                Ok(inst) => return Ok(Inst::VshrDtypeAtypeU32SatMode1(inst)),
3983                Err(_) => {}
3984            }
3985        }
3986        "vsub" => {
3987            stream.set_position(start_pos);
3988            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3989                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat(inst)),
3990                Err(_) => {}
3991            }
3992            stream.set_position(start_pos);
3993            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3994                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat(inst)),
3995                Err(_) => {}
3996            }
3997            stream.set_position(start_pos);
3998            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
3999                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat(inst)),
4000                Err(_) => {}
4001            }
4002            stream.set_position(start_pos);
4003            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4004                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat(inst)),
4005                Err(_) => {}
4006            }
4007            stream.set_position(start_pos);
4008            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4009                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat(inst)),
4010                Err(_) => {}
4011            }
4012            stream.set_position(start_pos);
4013            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
4014                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSatOp2(inst)),
4015                Err(_) => {}
4016            }
4017            stream.set_position(start_pos);
4018            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
4019                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSatOp2(inst)),
4020                Err(_) => {}
4021            }
4022            stream.set_position(start_pos);
4023            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
4024                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSatOp2(inst)),
4025                Err(_) => {}
4026            }
4027            stream.set_position(start_pos);
4028            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
4029                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSatOp2(inst)),
4030                Err(_) => {}
4031            }
4032            stream.set_position(start_pos);
4033            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSatOp2 as PtxParser>::parse(stream) {
4034                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSatOp2(inst)),
4035                Err(_) => {}
4036            }
4037            stream.set_position(start_pos);
4038            match <crate::r#type::instruction::vop::section_0::VaddDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
4039                Ok(inst) => return Ok(Inst::VaddDtypeAtypeBtypeSat1(inst)),
4040                Err(_) => {}
4041            }
4042            stream.set_position(start_pos);
4043            match <crate::r#type::instruction::vop::section_0::VsubDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
4044                Ok(inst) => return Ok(Inst::VsubDtypeAtypeBtypeSat1(inst)),
4045                Err(_) => {}
4046            }
4047            stream.set_position(start_pos);
4048            match <crate::r#type::instruction::vop::section_0::VabsdiffDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
4049                Ok(inst) => return Ok(Inst::VabsdiffDtypeAtypeBtypeSat1(inst)),
4050                Err(_) => {}
4051            }
4052            stream.set_position(start_pos);
4053            match <crate::r#type::instruction::vop::section_0::VminDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
4054                Ok(inst) => return Ok(Inst::VminDtypeAtypeBtypeSat1(inst)),
4055                Err(_) => {}
4056            }
4057            stream.set_position(start_pos);
4058            match <crate::r#type::instruction::vop::section_0::VmaxDtypeAtypeBtypeSat1 as PtxParser>::parse(stream) {
4059                Ok(inst) => return Ok(Inst::VmaxDtypeAtypeBtypeSat1(inst)),
4060                Err(_) => {}
4061            }
4062        }
4063        "vsub2" => {
4064            stream.set_position(start_pos);
4065            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4066                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeSat(inst)),
4067                Err(_) => {}
4068            }
4069            stream.set_position(start_pos);
4070            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4071                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeSat(inst)),
4072                Err(_) => {}
4073            }
4074            stream.set_position(start_pos);
4075            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4076                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeSat(inst)),
4077                Err(_) => {}
4078            }
4079            stream.set_position(start_pos);
4080            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4081                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeSat(inst)),
4082                Err(_) => {}
4083            }
4084            stream.set_position(start_pos);
4085            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4086                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeSat(inst)),
4087                Err(_) => {}
4088            }
4089            stream.set_position(start_pos);
4090            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4091                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeSat(inst)),
4092                Err(_) => {}
4093            }
4094            stream.set_position(start_pos);
4095            match <crate::r#type::instruction::vop2::section_0::Vadd2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4096                Ok(inst) => return Ok(Inst::Vadd2DtypeAtypeBtypeAdd(inst)),
4097                Err(_) => {}
4098            }
4099            stream.set_position(start_pos);
4100            match <crate::r#type::instruction::vop2::section_0::Vsub2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4101                Ok(inst) => return Ok(Inst::Vsub2DtypeAtypeBtypeAdd(inst)),
4102                Err(_) => {}
4103            }
4104            stream.set_position(start_pos);
4105            match <crate::r#type::instruction::vop2::section_0::Vavrg2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4106                Ok(inst) => return Ok(Inst::Vavrg2DtypeAtypeBtypeAdd(inst)),
4107                Err(_) => {}
4108            }
4109            stream.set_position(start_pos);
4110            match <crate::r#type::instruction::vop2::section_0::Vabsdiff2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4111                Ok(inst) => return Ok(Inst::Vabsdiff2DtypeAtypeBtypeAdd(inst)),
4112                Err(_) => {}
4113            }
4114            stream.set_position(start_pos);
4115            match <crate::r#type::instruction::vop2::section_0::Vmin2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4116                Ok(inst) => return Ok(Inst::Vmin2DtypeAtypeBtypeAdd(inst)),
4117                Err(_) => {}
4118            }
4119            stream.set_position(start_pos);
4120            match <crate::r#type::instruction::vop2::section_0::Vmax2DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4121                Ok(inst) => return Ok(Inst::Vmax2DtypeAtypeBtypeAdd(inst)),
4122                Err(_) => {}
4123            }
4124        }
4125        "vsub4" => {
4126            stream.set_position(start_pos);
4127            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4128                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeSat(inst)),
4129                Err(_) => {}
4130            }
4131            stream.set_position(start_pos);
4132            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4133                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeSat(inst)),
4134                Err(_) => {}
4135            }
4136            stream.set_position(start_pos);
4137            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4138                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeSat(inst)),
4139                Err(_) => {}
4140            }
4141            stream.set_position(start_pos);
4142            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4143                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeSat(inst)),
4144                Err(_) => {}
4145            }
4146            stream.set_position(start_pos);
4147            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4148                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeSat(inst)),
4149                Err(_) => {}
4150            }
4151            stream.set_position(start_pos);
4152            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeSat as PtxParser>::parse(stream) {
4153                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeSat(inst)),
4154                Err(_) => {}
4155            }
4156            stream.set_position(start_pos);
4157            match <crate::r#type::instruction::vop4::section_0::Vadd4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4158                Ok(inst) => return Ok(Inst::Vadd4DtypeAtypeBtypeAdd(inst)),
4159                Err(_) => {}
4160            }
4161            stream.set_position(start_pos);
4162            match <crate::r#type::instruction::vop4::section_0::Vsub4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4163                Ok(inst) => return Ok(Inst::Vsub4DtypeAtypeBtypeAdd(inst)),
4164                Err(_) => {}
4165            }
4166            stream.set_position(start_pos);
4167            match <crate::r#type::instruction::vop4::section_0::Vavrg4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4168                Ok(inst) => return Ok(Inst::Vavrg4DtypeAtypeBtypeAdd(inst)),
4169                Err(_) => {}
4170            }
4171            stream.set_position(start_pos);
4172            match <crate::r#type::instruction::vop4::section_0::Vabsdiff4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4173                Ok(inst) => return Ok(Inst::Vabsdiff4DtypeAtypeBtypeAdd(inst)),
4174                Err(_) => {}
4175            }
4176            stream.set_position(start_pos);
4177            match <crate::r#type::instruction::vop4::section_0::Vmin4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4178                Ok(inst) => return Ok(Inst::Vmin4DtypeAtypeBtypeAdd(inst)),
4179                Err(_) => {}
4180            }
4181            stream.set_position(start_pos);
4182            match <crate::r#type::instruction::vop4::section_0::Vmax4DtypeAtypeBtypeAdd as PtxParser>::parse(stream) {
4183                Ok(inst) => return Ok(Inst::Vmax4DtypeAtypeBtypeAdd(inst)),
4184                Err(_) => {}
4185            }
4186        }
4187        "wgmma" => {
4188            stream.set_position(start_pos);
4189            match <crate::r#type::instruction::wgmma_commit_group::section_0::WgmmaCommitGroupSyncAligned as PtxParser>::parse(stream) {
4190                Ok(inst) => return Ok(Inst::WgmmaCommitGroupSyncAligned(inst)),
4191                Err(_) => {}
4192            }
4193            stream.set_position(start_pos);
4194            match <crate::r#type::instruction::wgmma_fence::section_0::WgmmaFenceSyncAligned as PtxParser>::parse(stream) {
4195                Ok(inst) => return Ok(Inst::WgmmaFenceSyncAligned(inst)),
4196                Err(_) => {}
4197            }
4198            stream.set_position(start_pos);
4199            match <crate::r#type::instruction::wgmma_mma_async_sp::section_0::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16 as PtxParser>::parse(stream) {
4200                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16(inst)),
4201                Err(_) => {}
4202            }
4203            stream.set_position(start_pos);
4204            match <crate::r#type::instruction::wgmma_mma_async_sp::section_0::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161 as PtxParser>::parse(stream) {
4205                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161(inst)),
4206                Err(_) => {}
4207            }
4208            stream.set_position(start_pos);
4209            match <crate::r#type::instruction::wgmma_mma_async_sp::section_1::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16 as PtxParser>::parse(stream) {
4210                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16(inst)),
4211                Err(_) => {}
4212            }
4213            stream.set_position(start_pos);
4214            match <crate::r#type::instruction::wgmma_mma_async_sp::section_1::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161 as PtxParser>::parse(stream) {
4215                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161(inst)),
4216                Err(_) => {}
4217            }
4218            stream.set_position(start_pos);
4219            match <crate::r#type::instruction::wgmma_mma_async_sp::section_2::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32 as PtxParser>::parse(stream) {
4220                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32(inst)),
4221                Err(_) => {}
4222            }
4223            stream.set_position(start_pos);
4224            match <crate::r#type::instruction::wgmma_mma_async_sp::section_2::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321 as PtxParser>::parse(stream) {
4225                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321(inst)),
4226                Err(_) => {}
4227            }
4228            stream.set_position(start_pos);
4229            match <crate::r#type::instruction::wgmma_mma_async_sp::section_3::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype as PtxParser>::parse(stream) {
4230                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype(inst)),
4231                Err(_) => {}
4232            }
4233            stream.set_position(start_pos);
4234            match <crate::r#type::instruction::wgmma_mma_async_sp::section_3::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1 as PtxParser>::parse(stream) {
4235                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1(inst)),
4236                Err(_) => {}
4237            }
4238            stream.set_position(start_pos);
4239            match <crate::r#type::instruction::wgmma_mma_async_sp::section_4::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype as PtxParser>::parse(stream) {
4240                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype(inst)),
4241                Err(_) => {}
4242            }
4243            stream.set_position(start_pos);
4244            match <crate::r#type::instruction::wgmma_mma_async_sp::section_4::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1 as PtxParser>::parse(stream) {
4245                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1(inst)),
4246                Err(_) => {}
4247            }
4248            stream.set_position(start_pos);
4249            match <crate::r#type::instruction::wgmma_mma_async::section_0::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16 as PtxParser>::parse(stream) {
4250                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F16(inst)),
4251                Err(_) => {}
4252            }
4253            stream.set_position(start_pos);
4254            match <crate::r#type::instruction::wgmma_mma_async::section_0::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161 as PtxParser>::parse(stream) {
4255                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeF16F161(inst)),
4256                Err(_) => {}
4257            }
4258            stream.set_position(start_pos);
4259            match <crate::r#type::instruction::wgmma_mma_async::section_1::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16 as PtxParser>::parse(stream) {
4260                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf16(inst)),
4261                Err(_) => {}
4262            }
4263            stream.set_position(start_pos);
4264            match <crate::r#type::instruction::wgmma_mma_async::section_1::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161 as PtxParser>::parse(stream) {
4265                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeBf16Bf161(inst)),
4266                Err(_) => {}
4267            }
4268            stream.set_position(start_pos);
4269            match <crate::r#type::instruction::wgmma_mma_async::section_2::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32 as PtxParser>::parse(stream) {
4270                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf32(inst)),
4271                Err(_) => {}
4272            }
4273            stream.set_position(start_pos);
4274            match <crate::r#type::instruction::wgmma_mma_async::section_2::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321 as PtxParser>::parse(stream) {
4275                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeTf32Tf321(inst)),
4276                Err(_) => {}
4277            }
4278            stream.set_position(start_pos);
4279            match <crate::r#type::instruction::wgmma_mma_async::section_3::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype as PtxParser>::parse(stream) {
4280                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype(inst)),
4281                Err(_) => {}
4282            }
4283            stream.set_position(start_pos);
4284            match <crate::r#type::instruction::wgmma_mma_async::section_3::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1 as PtxParser>::parse(stream) {
4285                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeDtypeAtypeBtype1(inst)),
4286                Err(_) => {}
4287            }
4288            stream.set_position(start_pos);
4289            match <crate::r#type::instruction::wgmma_mma_async::section_4::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype as PtxParser>::parse(stream) {
4290                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype(inst)),
4291                Err(_) => {}
4292            }
4293            stream.set_position(start_pos);
4294            match <crate::r#type::instruction::wgmma_mma_async::section_4::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1 as PtxParser>::parse(stream) {
4295                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeSatfiniteS32AtypeBtype1(inst)),
4296                Err(_) => {}
4297            }
4298            stream.set_position(start_pos);
4299            match <crate::r#type::instruction::wgmma_mma_async::section_5::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc as PtxParser>::parse(stream) {
4300                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc(inst)),
4301                Err(_) => {}
4302            }
4303            stream.set_position(start_pos);
4304            match <crate::r#type::instruction::wgmma_mma_async::section_5::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1 as PtxParser>::parse(stream) {
4305                Ok(inst) => return Ok(Inst::WgmmaMmaAsyncSyncAlignedShapeS32B1B1OpPopc1(inst)),
4306                Err(_) => {}
4307            }
4308            stream.set_position(start_pos);
4309            match <crate::r#type::instruction::wgmma_wait_group::section_0::WgmmaWaitGroupSyncAligned as PtxParser>::parse(stream) {
4310                Ok(inst) => return Ok(Inst::WgmmaWaitGroupSyncAligned(inst)),
4311                Err(_) => {}
4312            }
4313        }
4314        "wmma" => {
4315            stream.set_position(start_pos);
4316            match <crate::r#type::instruction::wmma_load::section_0::WmmaLoadASyncAlignedLayoutShapeSsAtype as PtxParser>::parse(stream) {
4317                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedLayoutShapeSsAtype(inst)),
4318                Err(_) => {}
4319            }
4320            stream.set_position(start_pos);
4321            match <crate::r#type::instruction::wmma_load::section_0::WmmaLoadBSyncAlignedLayoutShapeSsBtype as PtxParser>::parse(stream) {
4322                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedLayoutShapeSsBtype(inst)),
4323                Err(_) => {}
4324            }
4325            stream.set_position(start_pos);
4326            match <crate::r#type::instruction::wmma_load::section_0::WmmaLoadCSyncAlignedLayoutShapeSsCtype as PtxParser>::parse(stream) {
4327                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype(inst)),
4328                Err(_) => {}
4329            }
4330            stream.set_position(start_pos);
4331            match <crate::r#type::instruction::wmma_load::section_1::WmmaLoadASyncAlignedLayoutShapeSsAtype1 as PtxParser>::parse(stream) {
4332                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedLayoutShapeSsAtype1(inst)),
4333                Err(_) => {}
4334            }
4335            stream.set_position(start_pos);
4336            match <crate::r#type::instruction::wmma_load::section_1::WmmaLoadBSyncAlignedLayoutShapeSsBtype1 as PtxParser>::parse(stream) {
4337                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedLayoutShapeSsBtype1(inst)),
4338                Err(_) => {}
4339            }
4340            stream.set_position(start_pos);
4341            match <crate::r#type::instruction::wmma_load::section_1::WmmaLoadCSyncAlignedLayoutShapeSsCtype1 as PtxParser>::parse(stream) {
4342                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype1(inst)),
4343                Err(_) => {}
4344            }
4345            stream.set_position(start_pos);
4346            match <crate::r#type::instruction::wmma_load::section_2::WmmaLoadASyncAlignedLayoutShapeSsAtype2 as PtxParser>::parse(stream) {
4347                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedLayoutShapeSsAtype2(inst)),
4348                Err(_) => {}
4349            }
4350            stream.set_position(start_pos);
4351            match <crate::r#type::instruction::wmma_load::section_2::WmmaLoadBSyncAlignedLayoutShapeSsBtype2 as PtxParser>::parse(stream) {
4352                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedLayoutShapeSsBtype2(inst)),
4353                Err(_) => {}
4354            }
4355            stream.set_position(start_pos);
4356            match <crate::r#type::instruction::wmma_load::section_2::WmmaLoadCSyncAlignedLayoutShapeSsCtype2 as PtxParser>::parse(stream) {
4357                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype2(inst)),
4358                Err(_) => {}
4359            }
4360            stream.set_position(start_pos);
4361            match <crate::r#type::instruction::wmma_load::section_3::WmmaLoadASyncAlignedLayoutShapeSsAtype3 as PtxParser>::parse(stream) {
4362                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedLayoutShapeSsAtype3(inst)),
4363                Err(_) => {}
4364            }
4365            stream.set_position(start_pos);
4366            match <crate::r#type::instruction::wmma_load::section_3::WmmaLoadBSyncAlignedLayoutShapeSsBtype3 as PtxParser>::parse(stream) {
4367                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedLayoutShapeSsBtype3(inst)),
4368                Err(_) => {}
4369            }
4370            stream.set_position(start_pos);
4371            match <crate::r#type::instruction::wmma_load::section_3::WmmaLoadCSyncAlignedLayoutShapeSsCtype3 as PtxParser>::parse(stream) {
4372                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype3(inst)),
4373                Err(_) => {}
4374            }
4375            stream.set_position(start_pos);
4376            match <crate::r#type::instruction::wmma_load::section_4::WmmaLoadASyncAlignedRowShapeSsAtype as PtxParser>::parse(stream) {
4377                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedRowShapeSsAtype(inst)),
4378                Err(_) => {}
4379            }
4380            stream.set_position(start_pos);
4381            match <crate::r#type::instruction::wmma_load::section_4::WmmaLoadBSyncAlignedColShapeSsBtype as PtxParser>::parse(stream) {
4382                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedColShapeSsBtype(inst)),
4383                Err(_) => {}
4384            }
4385            stream.set_position(start_pos);
4386            match <crate::r#type::instruction::wmma_load::section_4::WmmaLoadCSyncAlignedLayoutShapeSsCtype4 as PtxParser>::parse(stream) {
4387                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype4(inst)),
4388                Err(_) => {}
4389            }
4390            stream.set_position(start_pos);
4391            match <crate::r#type::instruction::wmma_load::section_5::WmmaLoadASyncAlignedRowShapeSsAtype1 as PtxParser>::parse(stream) {
4392                Ok(inst) => return Ok(Inst::WmmaLoadASyncAlignedRowShapeSsAtype1(inst)),
4393                Err(_) => {}
4394            }
4395            stream.set_position(start_pos);
4396            match <crate::r#type::instruction::wmma_load::section_5::WmmaLoadBSyncAlignedColShapeSsBtype1 as PtxParser>::parse(stream) {
4397                Ok(inst) => return Ok(Inst::WmmaLoadBSyncAlignedColShapeSsBtype1(inst)),
4398                Err(_) => {}
4399            }
4400            stream.set_position(start_pos);
4401            match <crate::r#type::instruction::wmma_load::section_5::WmmaLoadCSyncAlignedLayoutShapeSsCtype5 as PtxParser>::parse(stream) {
4402                Ok(inst) => return Ok(Inst::WmmaLoadCSyncAlignedLayoutShapeSsCtype5(inst)),
4403                Err(_) => {}
4404            }
4405            stream.set_position(start_pos);
4406            match <crate::r#type::instruction::wmma_mma::section_0::WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype as PtxParser>::parse(stream) {
4407                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedAlayoutBlayoutShapeDtypeCtype(inst)),
4408                Err(_) => {}
4409            }
4410            stream.set_position(start_pos);
4411            match <crate::r#type::instruction::wmma_mma::section_1::WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite as PtxParser>::parse(stream) {
4412                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedAlayoutBlayoutShapeS32AtypeBtypeS32Satfinite(inst)),
4413                Err(_) => {}
4414            }
4415            stream.set_position(start_pos);
4416            match <crate::r#type::instruction::wmma_mma::section_2::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32 as PtxParser>::parse(stream) {
4417                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF32(inst)),
4418                Err(_) => {}
4419            }
4420            stream.set_position(start_pos);
4421            match <crate::r#type::instruction::wmma_mma::section_3::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321 as PtxParser>::parse(stream) {
4422                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedAlayoutBlayoutShapeF32AtypeBtypeF321(inst)),
4423                Err(_) => {}
4424            }
4425            stream.set_position(start_pos);
4426            match <crate::r#type::instruction::wmma_mma::section_4::WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64 as PtxParser>::parse(stream) {
4427                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedAlayoutBlayoutShapeRndF64F64F64F64(inst)),
4428                Err(_) => {}
4429            }
4430            stream.set_position(start_pos);
4431            match <crate::r#type::instruction::wmma_mma::section_5::WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite as PtxParser>::parse(stream) {
4432                Ok(inst) => return Ok(Inst::WmmaMmaSyncAlignedRowColShapeS32AtypeBtypeS32Satfinite(inst)),
4433                Err(_) => {}
4434            }
4435            stream.set_position(start_pos);
4436            match <crate::r#type::instruction::wmma_mma::section_6::WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32 as PtxParser>::parse(stream) {
4437                Ok(inst) => return Ok(Inst::WmmaMmaOpPopcSyncAlignedRowColShapeS32AtypeBtypeS32(inst)),
4438                Err(_) => {}
4439            }
4440            stream.set_position(start_pos);
4441            match <crate::r#type::instruction::wmma_store::section_0::WmmaStoreDSyncAlignedLayoutShapeSsType as PtxParser>::parse(stream) {
4442                Ok(inst) => return Ok(Inst::WmmaStoreDSyncAlignedLayoutShapeSsType(inst)),
4443                Err(_) => {}
4444            }
4445            stream.set_position(start_pos);
4446            match <crate::r#type::instruction::wmma_store::section_1::WmmaStoreDSyncAlignedLayoutShapeSsType1 as PtxParser>::parse(stream) {
4447                Ok(inst) => return Ok(Inst::WmmaStoreDSyncAlignedLayoutShapeSsType1(inst)),
4448                Err(_) => {}
4449            }
4450            stream.set_position(start_pos);
4451            match <crate::r#type::instruction::wmma_store::section_2::WmmaStoreDSyncAlignedLayoutShapeSsType2 as PtxParser>::parse(stream) {
4452                Ok(inst) => return Ok(Inst::WmmaStoreDSyncAlignedLayoutShapeSsType2(inst)),
4453                Err(_) => {}
4454            }
4455            stream.set_position(start_pos);
4456            match <crate::r#type::instruction::wmma_store::section_3::WmmaStoreDSyncAlignedLayoutShapeSsType3 as PtxParser>::parse(stream) {
4457                Ok(inst) => return Ok(Inst::WmmaStoreDSyncAlignedLayoutShapeSsType3(inst)),
4458                Err(_) => {}
4459            }
4460        }
4461        "xor" => {
4462            stream.set_position(start_pos);
4463            match <crate::r#type::instruction::xor::section_0::XorType as PtxParser>::parse(stream) {
4464                Ok(inst) => return Ok(Inst::XorType(inst)),
4465                Err(_) => {}
4466            }
4467        }
4468        _ => {}
4469    }
4470    
4471    // If no parser matched, return error
4472    let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(0..0);
4473    Err(crate::parser::unexpected_value(span, &["valid PTX instruction"], "no matching instruction format"))
4474}
4475