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