1#![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
179pub(crate) fn parse_instruction_inner(stream: &mut PtxTokenStream) -> Result<Inst, PtxParseError> {
181 let start_pos = stream.position();
182
183 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 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 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