1#![allow(unused)]
70
71use crate::lexer::PtxToken;
72use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
73use crate::r#type::common::*;
74
75pub mod section_0 {
76 use super::*;
77 use crate::r#type::instruction::wgmma_mma_async_sp::section_0::*;
78
79 impl PtxParser for Dtype {
84 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
85 {
87 let saved_pos = stream.position();
88 if stream.expect_string(".f16").is_ok() {
89 return Ok(Dtype::F16);
90 }
91 stream.set_position(saved_pos);
92 }
93 let saved_pos = stream.position();
94 {
96 let saved_pos = stream.position();
97 if stream.expect_string(".f32").is_ok() {
98 return Ok(Dtype::F32);
99 }
100 stream.set_position(saved_pos);
101 }
102 stream.set_position(saved_pos);
103 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
104 let expected = &[".f16", ".f32"];
105 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
106 Err(crate::parser::unexpected_value(span, expected, found))
107 }
108 }
109
110 impl PtxParser for Shape {
111 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
112 {
114 let saved_pos = stream.position();
115 if stream.expect_string(".m64n104k32").is_ok() {
116 return Ok(Shape::M64n104k32);
117 }
118 stream.set_position(saved_pos);
119 }
120 let saved_pos = stream.position();
121 {
123 let saved_pos = stream.position();
124 if stream.expect_string(".m64n112k32").is_ok() {
125 return Ok(Shape::M64n112k32);
126 }
127 stream.set_position(saved_pos);
128 }
129 stream.set_position(saved_pos);
130 let saved_pos = stream.position();
131 {
133 let saved_pos = stream.position();
134 if stream.expect_string(".m64n120k32").is_ok() {
135 return Ok(Shape::M64n120k32);
136 }
137 stream.set_position(saved_pos);
138 }
139 stream.set_position(saved_pos);
140 let saved_pos = stream.position();
141 {
143 let saved_pos = stream.position();
144 if stream.expect_string(".m64n128k32").is_ok() {
145 return Ok(Shape::M64n128k32);
146 }
147 stream.set_position(saved_pos);
148 }
149 stream.set_position(saved_pos);
150 let saved_pos = stream.position();
151 {
153 let saved_pos = stream.position();
154 if stream.expect_string(".m64n136k32").is_ok() {
155 return Ok(Shape::M64n136k32);
156 }
157 stream.set_position(saved_pos);
158 }
159 stream.set_position(saved_pos);
160 let saved_pos = stream.position();
161 {
163 let saved_pos = stream.position();
164 if stream.expect_string(".m64n144k32").is_ok() {
165 return Ok(Shape::M64n144k32);
166 }
167 stream.set_position(saved_pos);
168 }
169 stream.set_position(saved_pos);
170 let saved_pos = stream.position();
171 {
173 let saved_pos = stream.position();
174 if stream.expect_string(".m64n152k32").is_ok() {
175 return Ok(Shape::M64n152k32);
176 }
177 stream.set_position(saved_pos);
178 }
179 stream.set_position(saved_pos);
180 let saved_pos = stream.position();
181 {
183 let saved_pos = stream.position();
184 if stream.expect_string(".m64n160k32").is_ok() {
185 return Ok(Shape::M64n160k32);
186 }
187 stream.set_position(saved_pos);
188 }
189 stream.set_position(saved_pos);
190 let saved_pos = stream.position();
191 {
193 let saved_pos = stream.position();
194 if stream.expect_string(".m64n168k32").is_ok() {
195 return Ok(Shape::M64n168k32);
196 }
197 stream.set_position(saved_pos);
198 }
199 stream.set_position(saved_pos);
200 let saved_pos = stream.position();
201 {
203 let saved_pos = stream.position();
204 if stream.expect_string(".m64n176k32").is_ok() {
205 return Ok(Shape::M64n176k32);
206 }
207 stream.set_position(saved_pos);
208 }
209 stream.set_position(saved_pos);
210 let saved_pos = stream.position();
211 {
213 let saved_pos = stream.position();
214 if stream.expect_string(".m64n184k32").is_ok() {
215 return Ok(Shape::M64n184k32);
216 }
217 stream.set_position(saved_pos);
218 }
219 stream.set_position(saved_pos);
220 let saved_pos = stream.position();
221 {
223 let saved_pos = stream.position();
224 if stream.expect_string(".m64n192k32").is_ok() {
225 return Ok(Shape::M64n192k32);
226 }
227 stream.set_position(saved_pos);
228 }
229 stream.set_position(saved_pos);
230 let saved_pos = stream.position();
231 {
233 let saved_pos = stream.position();
234 if stream.expect_string(".m64n200k32").is_ok() {
235 return Ok(Shape::M64n200k32);
236 }
237 stream.set_position(saved_pos);
238 }
239 stream.set_position(saved_pos);
240 let saved_pos = stream.position();
241 {
243 let saved_pos = stream.position();
244 if stream.expect_string(".m64n208k32").is_ok() {
245 return Ok(Shape::M64n208k32);
246 }
247 stream.set_position(saved_pos);
248 }
249 stream.set_position(saved_pos);
250 let saved_pos = stream.position();
251 {
253 let saved_pos = stream.position();
254 if stream.expect_string(".m64n216k32").is_ok() {
255 return Ok(Shape::M64n216k32);
256 }
257 stream.set_position(saved_pos);
258 }
259 stream.set_position(saved_pos);
260 let saved_pos = stream.position();
261 {
263 let saved_pos = stream.position();
264 if stream.expect_string(".m64n224k32").is_ok() {
265 return Ok(Shape::M64n224k32);
266 }
267 stream.set_position(saved_pos);
268 }
269 stream.set_position(saved_pos);
270 let saved_pos = stream.position();
271 {
273 let saved_pos = stream.position();
274 if stream.expect_string(".m64n232k32").is_ok() {
275 return Ok(Shape::M64n232k32);
276 }
277 stream.set_position(saved_pos);
278 }
279 stream.set_position(saved_pos);
280 let saved_pos = stream.position();
281 {
283 let saved_pos = stream.position();
284 if stream.expect_string(".m64n240k32").is_ok() {
285 return Ok(Shape::M64n240k32);
286 }
287 stream.set_position(saved_pos);
288 }
289 stream.set_position(saved_pos);
290 let saved_pos = stream.position();
291 {
293 let saved_pos = stream.position();
294 if stream.expect_string(".m64n248k32").is_ok() {
295 return Ok(Shape::M64n248k32);
296 }
297 stream.set_position(saved_pos);
298 }
299 stream.set_position(saved_pos);
300 let saved_pos = stream.position();
301 {
303 let saved_pos = stream.position();
304 if stream.expect_string(".m64n256k32").is_ok() {
305 return Ok(Shape::M64n256k32);
306 }
307 stream.set_position(saved_pos);
308 }
309 stream.set_position(saved_pos);
310 let saved_pos = stream.position();
311 {
313 let saved_pos = stream.position();
314 if stream.expect_string(".m64n16k32").is_ok() {
315 return Ok(Shape::M64n16k32);
316 }
317 stream.set_position(saved_pos);
318 }
319 stream.set_position(saved_pos);
320 let saved_pos = stream.position();
321 {
323 let saved_pos = stream.position();
324 if stream.expect_string(".m64n24k32").is_ok() {
325 return Ok(Shape::M64n24k32);
326 }
327 stream.set_position(saved_pos);
328 }
329 stream.set_position(saved_pos);
330 let saved_pos = stream.position();
331 {
333 let saved_pos = stream.position();
334 if stream.expect_string(".m64n32k32").is_ok() {
335 return Ok(Shape::M64n32k32);
336 }
337 stream.set_position(saved_pos);
338 }
339 stream.set_position(saved_pos);
340 let saved_pos = stream.position();
341 {
343 let saved_pos = stream.position();
344 if stream.expect_string(".m64n40k32").is_ok() {
345 return Ok(Shape::M64n40k32);
346 }
347 stream.set_position(saved_pos);
348 }
349 stream.set_position(saved_pos);
350 let saved_pos = stream.position();
351 {
353 let saved_pos = stream.position();
354 if stream.expect_string(".m64n48k32").is_ok() {
355 return Ok(Shape::M64n48k32);
356 }
357 stream.set_position(saved_pos);
358 }
359 stream.set_position(saved_pos);
360 let saved_pos = stream.position();
361 {
363 let saved_pos = stream.position();
364 if stream.expect_string(".m64n56k32").is_ok() {
365 return Ok(Shape::M64n56k32);
366 }
367 stream.set_position(saved_pos);
368 }
369 stream.set_position(saved_pos);
370 let saved_pos = stream.position();
371 {
373 let saved_pos = stream.position();
374 if stream.expect_string(".m64n64k32").is_ok() {
375 return Ok(Shape::M64n64k32);
376 }
377 stream.set_position(saved_pos);
378 }
379 stream.set_position(saved_pos);
380 let saved_pos = stream.position();
381 {
383 let saved_pos = stream.position();
384 if stream.expect_string(".m64n72k32").is_ok() {
385 return Ok(Shape::M64n72k32);
386 }
387 stream.set_position(saved_pos);
388 }
389 stream.set_position(saved_pos);
390 let saved_pos = stream.position();
391 {
393 let saved_pos = stream.position();
394 if stream.expect_string(".m64n80k32").is_ok() {
395 return Ok(Shape::M64n80k32);
396 }
397 stream.set_position(saved_pos);
398 }
399 stream.set_position(saved_pos);
400 let saved_pos = stream.position();
401 {
403 let saved_pos = stream.position();
404 if stream.expect_string(".m64n88k32").is_ok() {
405 return Ok(Shape::M64n88k32);
406 }
407 stream.set_position(saved_pos);
408 }
409 stream.set_position(saved_pos);
410 let saved_pos = stream.position();
411 {
413 let saved_pos = stream.position();
414 if stream.expect_string(".m64n96k32").is_ok() {
415 return Ok(Shape::M64n96k32);
416 }
417 stream.set_position(saved_pos);
418 }
419 stream.set_position(saved_pos);
420 let saved_pos = stream.position();
421 {
423 let saved_pos = stream.position();
424 if stream.expect_string(".m64n8k32").is_ok() {
425 return Ok(Shape::M64n8k32);
426 }
427 stream.set_position(saved_pos);
428 }
429 stream.set_position(saved_pos);
430 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
431 let expected = &[".m64n104k32", ".m64n112k32", ".m64n120k32", ".m64n128k32", ".m64n136k32", ".m64n144k32", ".m64n152k32", ".m64n160k32", ".m64n168k32", ".m64n176k32", ".m64n184k32", ".m64n192k32", ".m64n200k32", ".m64n208k32", ".m64n216k32", ".m64n224k32", ".m64n232k32", ".m64n240k32", ".m64n248k32", ".m64n256k32", ".m64n16k32", ".m64n24k32", ".m64n32k32", ".m64n40k32", ".m64n48k32", ".m64n56k32", ".m64n64k32", ".m64n72k32", ".m64n80k32", ".m64n88k32", ".m64n96k32", ".m64n8k32"];
432 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
433 Err(crate::parser::unexpected_value(span, expected, found))
434 }
435 }
436
437 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16 {
438 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
439 stream.expect_string("wgmma")?;
440 stream.expect_string(".mma_async")?;
441 let mma_async = ();
442 stream.expect_complete()?;
443 stream.expect_string(".sp")?;
444 let sp = ();
445 stream.expect_complete()?;
446 stream.expect_string(".sync")?;
447 let sync = ();
448 stream.expect_complete()?;
449 stream.expect_string(".aligned")?;
450 let aligned = ();
451 stream.expect_complete()?;
452 let shape = Shape::parse(stream)?;
453 stream.expect_complete()?;
454 let dtype = Dtype::parse(stream)?;
455 stream.expect_complete()?;
456 stream.expect_string(".f16")?;
457 let f16 = ();
458 stream.expect_complete()?;
459 stream.expect_string(".f16")?;
460 let f162 = ();
461 stream.expect_complete()?;
462 let d = GeneralOperand::parse(stream)?;
463 stream.expect_complete()?;
464 stream.expect(&PtxToken::Comma)?;
465 let a_desc = GeneralOperand::parse(stream)?;
466 stream.expect_complete()?;
467 stream.expect(&PtxToken::Comma)?;
468 let b_desc = GeneralOperand::parse(stream)?;
469 stream.expect_complete()?;
470 stream.expect(&PtxToken::Comma)?;
471 let sp_meta = GeneralOperand::parse(stream)?;
472 stream.expect_complete()?;
473 stream.expect(&PtxToken::Comma)?;
474 let sp_sel = GeneralOperand::parse(stream)?;
475 stream.expect_complete()?;
476 stream.expect(&PtxToken::Comma)?;
477 let scale_d = GeneralOperand::parse(stream)?;
478 stream.expect_complete()?;
479 stream.expect(&PtxToken::Comma)?;
480 let imm_scale_a = GeneralOperand::parse(stream)?;
481 stream.expect_complete()?;
482 stream.expect(&PtxToken::Comma)?;
483 let imm_scale_b = GeneralOperand::parse(stream)?;
484 stream.expect_complete()?;
485 stream.expect(&PtxToken::Comma)?;
486 let imm_trans_a = GeneralOperand::parse(stream)?;
487 stream.expect_complete()?;
488 stream.expect(&PtxToken::Comma)?;
489 let imm_trans_b = GeneralOperand::parse(stream)?;
490 stream.expect_complete()?;
491 stream.expect_complete()?;
492 stream.expect(&PtxToken::Semicolon)?;
493 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F16 {
494 mma_async,
495 sp,
496 sync,
497 aligned,
498 shape,
499 dtype,
500 f16,
501 f162,
502 d,
503 a_desc,
504 b_desc,
505 sp_meta,
506 sp_sel,
507 scale_d,
508 imm_scale_a,
509 imm_scale_b,
510 imm_trans_a,
511 imm_trans_b,
512 })
513 }
514 }
515
516
517 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161 {
518 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
519 stream.expect_string("wgmma")?;
520 stream.expect_string(".mma_async")?;
521 let mma_async = ();
522 stream.expect_complete()?;
523 stream.expect_string(".sp")?;
524 let sp = ();
525 stream.expect_complete()?;
526 stream.expect_string(".sync")?;
527 let sync = ();
528 stream.expect_complete()?;
529 stream.expect_string(".aligned")?;
530 let aligned = ();
531 stream.expect_complete()?;
532 let shape = Shape::parse(stream)?;
533 stream.expect_complete()?;
534 let dtype = Dtype::parse(stream)?;
535 stream.expect_complete()?;
536 stream.expect_string(".f16")?;
537 let f16 = ();
538 stream.expect_complete()?;
539 stream.expect_string(".f16")?;
540 let f162 = ();
541 stream.expect_complete()?;
542 let d = GeneralOperand::parse(stream)?;
543 stream.expect_complete()?;
544 stream.expect(&PtxToken::Comma)?;
545 let a = GeneralOperand::parse(stream)?;
546 stream.expect_complete()?;
547 stream.expect(&PtxToken::Comma)?;
548 let b_desc = GeneralOperand::parse(stream)?;
549 stream.expect_complete()?;
550 stream.expect(&PtxToken::Comma)?;
551 let sp_meta = GeneralOperand::parse(stream)?;
552 stream.expect_complete()?;
553 stream.expect(&PtxToken::Comma)?;
554 let sp_sel = GeneralOperand::parse(stream)?;
555 stream.expect_complete()?;
556 stream.expect(&PtxToken::Comma)?;
557 let scale_d = GeneralOperand::parse(stream)?;
558 stream.expect_complete()?;
559 stream.expect(&PtxToken::Comma)?;
560 let imm_scale_a = GeneralOperand::parse(stream)?;
561 stream.expect_complete()?;
562 stream.expect(&PtxToken::Comma)?;
563 let imm_scale_b = GeneralOperand::parse(stream)?;
564 stream.expect_complete()?;
565 stream.expect(&PtxToken::Comma)?;
566 let imm_trans_b = GeneralOperand::parse(stream)?;
567 stream.expect_complete()?;
568 stream.expect_complete()?;
569 stream.expect(&PtxToken::Semicolon)?;
570 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeF16F161 {
571 mma_async,
572 sp,
573 sync,
574 aligned,
575 shape,
576 dtype,
577 f16,
578 f162,
579 d,
580 a,
581 b_desc,
582 sp_meta,
583 sp_sel,
584 scale_d,
585 imm_scale_a,
586 imm_scale_b,
587 imm_trans_b,
588 })
589 }
590 }
591
592
593}
594
595pub mod section_1 {
596 use super::*;
597 use crate::r#type::instruction::wgmma_mma_async_sp::section_1::*;
598
599 impl PtxParser for Dtype {
604 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
605 {
607 let saved_pos = stream.position();
608 if stream.expect_string(".f32").is_ok() {
609 return Ok(Dtype::F32);
610 }
611 stream.set_position(saved_pos);
612 }
613 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
614 let expected = &[".f32"];
615 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
616 Err(crate::parser::unexpected_value(span, expected, found))
617 }
618 }
619
620 impl PtxParser for Shape {
621 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
622 {
624 let saved_pos = stream.position();
625 if stream.expect_string(".m64n104k32").is_ok() {
626 return Ok(Shape::M64n104k32);
627 }
628 stream.set_position(saved_pos);
629 }
630 let saved_pos = stream.position();
631 {
633 let saved_pos = stream.position();
634 if stream.expect_string(".m64n112k32").is_ok() {
635 return Ok(Shape::M64n112k32);
636 }
637 stream.set_position(saved_pos);
638 }
639 stream.set_position(saved_pos);
640 let saved_pos = stream.position();
641 {
643 let saved_pos = stream.position();
644 if stream.expect_string(".m64n120k32").is_ok() {
645 return Ok(Shape::M64n120k32);
646 }
647 stream.set_position(saved_pos);
648 }
649 stream.set_position(saved_pos);
650 let saved_pos = stream.position();
651 {
653 let saved_pos = stream.position();
654 if stream.expect_string(".m64n128k32").is_ok() {
655 return Ok(Shape::M64n128k32);
656 }
657 stream.set_position(saved_pos);
658 }
659 stream.set_position(saved_pos);
660 let saved_pos = stream.position();
661 {
663 let saved_pos = stream.position();
664 if stream.expect_string(".m64n136k32").is_ok() {
665 return Ok(Shape::M64n136k32);
666 }
667 stream.set_position(saved_pos);
668 }
669 stream.set_position(saved_pos);
670 let saved_pos = stream.position();
671 {
673 let saved_pos = stream.position();
674 if stream.expect_string(".m64n144k32").is_ok() {
675 return Ok(Shape::M64n144k32);
676 }
677 stream.set_position(saved_pos);
678 }
679 stream.set_position(saved_pos);
680 let saved_pos = stream.position();
681 {
683 let saved_pos = stream.position();
684 if stream.expect_string(".m64n152k32").is_ok() {
685 return Ok(Shape::M64n152k32);
686 }
687 stream.set_position(saved_pos);
688 }
689 stream.set_position(saved_pos);
690 let saved_pos = stream.position();
691 {
693 let saved_pos = stream.position();
694 if stream.expect_string(".m64n160k32").is_ok() {
695 return Ok(Shape::M64n160k32);
696 }
697 stream.set_position(saved_pos);
698 }
699 stream.set_position(saved_pos);
700 let saved_pos = stream.position();
701 {
703 let saved_pos = stream.position();
704 if stream.expect_string(".m64n168k32").is_ok() {
705 return Ok(Shape::M64n168k32);
706 }
707 stream.set_position(saved_pos);
708 }
709 stream.set_position(saved_pos);
710 let saved_pos = stream.position();
711 {
713 let saved_pos = stream.position();
714 if stream.expect_string(".m64n176k32").is_ok() {
715 return Ok(Shape::M64n176k32);
716 }
717 stream.set_position(saved_pos);
718 }
719 stream.set_position(saved_pos);
720 let saved_pos = stream.position();
721 {
723 let saved_pos = stream.position();
724 if stream.expect_string(".m64n184k32").is_ok() {
725 return Ok(Shape::M64n184k32);
726 }
727 stream.set_position(saved_pos);
728 }
729 stream.set_position(saved_pos);
730 let saved_pos = stream.position();
731 {
733 let saved_pos = stream.position();
734 if stream.expect_string(".m64n192k32").is_ok() {
735 return Ok(Shape::M64n192k32);
736 }
737 stream.set_position(saved_pos);
738 }
739 stream.set_position(saved_pos);
740 let saved_pos = stream.position();
741 {
743 let saved_pos = stream.position();
744 if stream.expect_string(".m64n200k32").is_ok() {
745 return Ok(Shape::M64n200k32);
746 }
747 stream.set_position(saved_pos);
748 }
749 stream.set_position(saved_pos);
750 let saved_pos = stream.position();
751 {
753 let saved_pos = stream.position();
754 if stream.expect_string(".m64n208k32").is_ok() {
755 return Ok(Shape::M64n208k32);
756 }
757 stream.set_position(saved_pos);
758 }
759 stream.set_position(saved_pos);
760 let saved_pos = stream.position();
761 {
763 let saved_pos = stream.position();
764 if stream.expect_string(".m64n216k32").is_ok() {
765 return Ok(Shape::M64n216k32);
766 }
767 stream.set_position(saved_pos);
768 }
769 stream.set_position(saved_pos);
770 let saved_pos = stream.position();
771 {
773 let saved_pos = stream.position();
774 if stream.expect_string(".m64n224k32").is_ok() {
775 return Ok(Shape::M64n224k32);
776 }
777 stream.set_position(saved_pos);
778 }
779 stream.set_position(saved_pos);
780 let saved_pos = stream.position();
781 {
783 let saved_pos = stream.position();
784 if stream.expect_string(".m64n232k32").is_ok() {
785 return Ok(Shape::M64n232k32);
786 }
787 stream.set_position(saved_pos);
788 }
789 stream.set_position(saved_pos);
790 let saved_pos = stream.position();
791 {
793 let saved_pos = stream.position();
794 if stream.expect_string(".m64n240k32").is_ok() {
795 return Ok(Shape::M64n240k32);
796 }
797 stream.set_position(saved_pos);
798 }
799 stream.set_position(saved_pos);
800 let saved_pos = stream.position();
801 {
803 let saved_pos = stream.position();
804 if stream.expect_string(".m64n248k32").is_ok() {
805 return Ok(Shape::M64n248k32);
806 }
807 stream.set_position(saved_pos);
808 }
809 stream.set_position(saved_pos);
810 let saved_pos = stream.position();
811 {
813 let saved_pos = stream.position();
814 if stream.expect_string(".m64n256k32").is_ok() {
815 return Ok(Shape::M64n256k32);
816 }
817 stream.set_position(saved_pos);
818 }
819 stream.set_position(saved_pos);
820 let saved_pos = stream.position();
821 {
823 let saved_pos = stream.position();
824 if stream.expect_string(".m64n16k32").is_ok() {
825 return Ok(Shape::M64n16k32);
826 }
827 stream.set_position(saved_pos);
828 }
829 stream.set_position(saved_pos);
830 let saved_pos = stream.position();
831 {
833 let saved_pos = stream.position();
834 if stream.expect_string(".m64n24k32").is_ok() {
835 return Ok(Shape::M64n24k32);
836 }
837 stream.set_position(saved_pos);
838 }
839 stream.set_position(saved_pos);
840 let saved_pos = stream.position();
841 {
843 let saved_pos = stream.position();
844 if stream.expect_string(".m64n32k32").is_ok() {
845 return Ok(Shape::M64n32k32);
846 }
847 stream.set_position(saved_pos);
848 }
849 stream.set_position(saved_pos);
850 let saved_pos = stream.position();
851 {
853 let saved_pos = stream.position();
854 if stream.expect_string(".m64n40k32").is_ok() {
855 return Ok(Shape::M64n40k32);
856 }
857 stream.set_position(saved_pos);
858 }
859 stream.set_position(saved_pos);
860 let saved_pos = stream.position();
861 {
863 let saved_pos = stream.position();
864 if stream.expect_string(".m64n48k32").is_ok() {
865 return Ok(Shape::M64n48k32);
866 }
867 stream.set_position(saved_pos);
868 }
869 stream.set_position(saved_pos);
870 let saved_pos = stream.position();
871 {
873 let saved_pos = stream.position();
874 if stream.expect_string(".m64n56k32").is_ok() {
875 return Ok(Shape::M64n56k32);
876 }
877 stream.set_position(saved_pos);
878 }
879 stream.set_position(saved_pos);
880 let saved_pos = stream.position();
881 {
883 let saved_pos = stream.position();
884 if stream.expect_string(".m64n64k32").is_ok() {
885 return Ok(Shape::M64n64k32);
886 }
887 stream.set_position(saved_pos);
888 }
889 stream.set_position(saved_pos);
890 let saved_pos = stream.position();
891 {
893 let saved_pos = stream.position();
894 if stream.expect_string(".m64n72k32").is_ok() {
895 return Ok(Shape::M64n72k32);
896 }
897 stream.set_position(saved_pos);
898 }
899 stream.set_position(saved_pos);
900 let saved_pos = stream.position();
901 {
903 let saved_pos = stream.position();
904 if stream.expect_string(".m64n80k32").is_ok() {
905 return Ok(Shape::M64n80k32);
906 }
907 stream.set_position(saved_pos);
908 }
909 stream.set_position(saved_pos);
910 let saved_pos = stream.position();
911 {
913 let saved_pos = stream.position();
914 if stream.expect_string(".m64n88k32").is_ok() {
915 return Ok(Shape::M64n88k32);
916 }
917 stream.set_position(saved_pos);
918 }
919 stream.set_position(saved_pos);
920 let saved_pos = stream.position();
921 {
923 let saved_pos = stream.position();
924 if stream.expect_string(".m64n96k32").is_ok() {
925 return Ok(Shape::M64n96k32);
926 }
927 stream.set_position(saved_pos);
928 }
929 stream.set_position(saved_pos);
930 let saved_pos = stream.position();
931 {
933 let saved_pos = stream.position();
934 if stream.expect_string(".m64n8k32").is_ok() {
935 return Ok(Shape::M64n8k32);
936 }
937 stream.set_position(saved_pos);
938 }
939 stream.set_position(saved_pos);
940 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
941 let expected = &[".m64n104k32", ".m64n112k32", ".m64n120k32", ".m64n128k32", ".m64n136k32", ".m64n144k32", ".m64n152k32", ".m64n160k32", ".m64n168k32", ".m64n176k32", ".m64n184k32", ".m64n192k32", ".m64n200k32", ".m64n208k32", ".m64n216k32", ".m64n224k32", ".m64n232k32", ".m64n240k32", ".m64n248k32", ".m64n256k32", ".m64n16k32", ".m64n24k32", ".m64n32k32", ".m64n40k32", ".m64n48k32", ".m64n56k32", ".m64n64k32", ".m64n72k32", ".m64n80k32", ".m64n88k32", ".m64n96k32", ".m64n8k32"];
942 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
943 Err(crate::parser::unexpected_value(span, expected, found))
944 }
945 }
946
947 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16 {
948 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
949 stream.expect_string("wgmma")?;
950 stream.expect_string(".mma_async")?;
951 let mma_async = ();
952 stream.expect_complete()?;
953 stream.expect_string(".sp")?;
954 let sp = ();
955 stream.expect_complete()?;
956 stream.expect_string(".sync")?;
957 let sync = ();
958 stream.expect_complete()?;
959 stream.expect_string(".aligned")?;
960 let aligned = ();
961 stream.expect_complete()?;
962 let shape = Shape::parse(stream)?;
963 stream.expect_complete()?;
964 let dtype = Dtype::parse(stream)?;
965 stream.expect_complete()?;
966 stream.expect_string(".bf16")?;
967 let bf16 = ();
968 stream.expect_complete()?;
969 stream.expect_string(".bf16")?;
970 let bf162 = ();
971 stream.expect_complete()?;
972 let d = GeneralOperand::parse(stream)?;
973 stream.expect_complete()?;
974 stream.expect(&PtxToken::Comma)?;
975 let a_desc = GeneralOperand::parse(stream)?;
976 stream.expect_complete()?;
977 stream.expect(&PtxToken::Comma)?;
978 let b_desc = GeneralOperand::parse(stream)?;
979 stream.expect_complete()?;
980 stream.expect(&PtxToken::Comma)?;
981 let sp_meta = GeneralOperand::parse(stream)?;
982 stream.expect_complete()?;
983 stream.expect(&PtxToken::Comma)?;
984 let sp_sel = GeneralOperand::parse(stream)?;
985 stream.expect_complete()?;
986 stream.expect(&PtxToken::Comma)?;
987 let scale_d = GeneralOperand::parse(stream)?;
988 stream.expect_complete()?;
989 stream.expect(&PtxToken::Comma)?;
990 let imm_scale_a = GeneralOperand::parse(stream)?;
991 stream.expect_complete()?;
992 stream.expect(&PtxToken::Comma)?;
993 let imm_scale_b = GeneralOperand::parse(stream)?;
994 stream.expect_complete()?;
995 stream.expect(&PtxToken::Comma)?;
996 let imm_trans_a = GeneralOperand::parse(stream)?;
997 stream.expect_complete()?;
998 stream.expect(&PtxToken::Comma)?;
999 let imm_trans_b = GeneralOperand::parse(stream)?;
1000 stream.expect_complete()?;
1001 stream.expect_complete()?;
1002 stream.expect(&PtxToken::Semicolon)?;
1003 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf16 {
1004 mma_async,
1005 sp,
1006 sync,
1007 aligned,
1008 shape,
1009 dtype,
1010 bf16,
1011 bf162,
1012 d,
1013 a_desc,
1014 b_desc,
1015 sp_meta,
1016 sp_sel,
1017 scale_d,
1018 imm_scale_a,
1019 imm_scale_b,
1020 imm_trans_a,
1021 imm_trans_b,
1022 })
1023 }
1024 }
1025
1026
1027 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161 {
1028 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1029 stream.expect_string("wgmma")?;
1030 stream.expect_string(".mma_async")?;
1031 let mma_async = ();
1032 stream.expect_complete()?;
1033 stream.expect_string(".sp")?;
1034 let sp = ();
1035 stream.expect_complete()?;
1036 stream.expect_string(".sync")?;
1037 let sync = ();
1038 stream.expect_complete()?;
1039 stream.expect_string(".aligned")?;
1040 let aligned = ();
1041 stream.expect_complete()?;
1042 let shape = Shape::parse(stream)?;
1043 stream.expect_complete()?;
1044 let dtype = Dtype::parse(stream)?;
1045 stream.expect_complete()?;
1046 stream.expect_string(".bf16")?;
1047 let bf16 = ();
1048 stream.expect_complete()?;
1049 stream.expect_string(".bf16")?;
1050 let bf162 = ();
1051 stream.expect_complete()?;
1052 let d = GeneralOperand::parse(stream)?;
1053 stream.expect_complete()?;
1054 stream.expect(&PtxToken::Comma)?;
1055 let a = GeneralOperand::parse(stream)?;
1056 stream.expect_complete()?;
1057 stream.expect(&PtxToken::Comma)?;
1058 let b_desc = GeneralOperand::parse(stream)?;
1059 stream.expect_complete()?;
1060 stream.expect(&PtxToken::Comma)?;
1061 let sp_meta = GeneralOperand::parse(stream)?;
1062 stream.expect_complete()?;
1063 stream.expect(&PtxToken::Comma)?;
1064 let sp_sel = GeneralOperand::parse(stream)?;
1065 stream.expect_complete()?;
1066 stream.expect(&PtxToken::Comma)?;
1067 let scale_d = GeneralOperand::parse(stream)?;
1068 stream.expect_complete()?;
1069 stream.expect(&PtxToken::Comma)?;
1070 let imm_scale_a = GeneralOperand::parse(stream)?;
1071 stream.expect_complete()?;
1072 stream.expect(&PtxToken::Comma)?;
1073 let imm_scale_b = GeneralOperand::parse(stream)?;
1074 stream.expect_complete()?;
1075 stream.expect(&PtxToken::Comma)?;
1076 let imm_trans_b = GeneralOperand::parse(stream)?;
1077 stream.expect_complete()?;
1078 stream.expect_complete()?;
1079 stream.expect(&PtxToken::Semicolon)?;
1080 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeBf16Bf161 {
1081 mma_async,
1082 sp,
1083 sync,
1084 aligned,
1085 shape,
1086 dtype,
1087 bf16,
1088 bf162,
1089 d,
1090 a,
1091 b_desc,
1092 sp_meta,
1093 sp_sel,
1094 scale_d,
1095 imm_scale_a,
1096 imm_scale_b,
1097 imm_trans_b,
1098 })
1099 }
1100 }
1101
1102
1103}
1104
1105pub mod section_2 {
1106 use super::*;
1107 use crate::r#type::instruction::wgmma_mma_async_sp::section_2::*;
1108
1109 impl PtxParser for Dtype {
1114 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1115 {
1117 let saved_pos = stream.position();
1118 if stream.expect_string(".f32").is_ok() {
1119 return Ok(Dtype::F32);
1120 }
1121 stream.set_position(saved_pos);
1122 }
1123 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1124 let expected = &[".f32"];
1125 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1126 Err(crate::parser::unexpected_value(span, expected, found))
1127 }
1128 }
1129
1130 impl PtxParser for Shape {
1131 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1132 {
1134 let saved_pos = stream.position();
1135 if stream.expect_string(".m64n104k16").is_ok() {
1136 return Ok(Shape::M64n104k16);
1137 }
1138 stream.set_position(saved_pos);
1139 }
1140 let saved_pos = stream.position();
1141 {
1143 let saved_pos = stream.position();
1144 if stream.expect_string(".m64n112k16").is_ok() {
1145 return Ok(Shape::M64n112k16);
1146 }
1147 stream.set_position(saved_pos);
1148 }
1149 stream.set_position(saved_pos);
1150 let saved_pos = stream.position();
1151 {
1153 let saved_pos = stream.position();
1154 if stream.expect_string(".m64n120k16").is_ok() {
1155 return Ok(Shape::M64n120k16);
1156 }
1157 stream.set_position(saved_pos);
1158 }
1159 stream.set_position(saved_pos);
1160 let saved_pos = stream.position();
1161 {
1163 let saved_pos = stream.position();
1164 if stream.expect_string(".m64n128k16").is_ok() {
1165 return Ok(Shape::M64n128k16);
1166 }
1167 stream.set_position(saved_pos);
1168 }
1169 stream.set_position(saved_pos);
1170 let saved_pos = stream.position();
1171 {
1173 let saved_pos = stream.position();
1174 if stream.expect_string(".m64n136k16").is_ok() {
1175 return Ok(Shape::M64n136k16);
1176 }
1177 stream.set_position(saved_pos);
1178 }
1179 stream.set_position(saved_pos);
1180 let saved_pos = stream.position();
1181 {
1183 let saved_pos = stream.position();
1184 if stream.expect_string(".m64n144k16").is_ok() {
1185 return Ok(Shape::M64n144k16);
1186 }
1187 stream.set_position(saved_pos);
1188 }
1189 stream.set_position(saved_pos);
1190 let saved_pos = stream.position();
1191 {
1193 let saved_pos = stream.position();
1194 if stream.expect_string(".m64n152k16").is_ok() {
1195 return Ok(Shape::M64n152k16);
1196 }
1197 stream.set_position(saved_pos);
1198 }
1199 stream.set_position(saved_pos);
1200 let saved_pos = stream.position();
1201 {
1203 let saved_pos = stream.position();
1204 if stream.expect_string(".m64n160k16").is_ok() {
1205 return Ok(Shape::M64n160k16);
1206 }
1207 stream.set_position(saved_pos);
1208 }
1209 stream.set_position(saved_pos);
1210 let saved_pos = stream.position();
1211 {
1213 let saved_pos = stream.position();
1214 if stream.expect_string(".m64n168k16").is_ok() {
1215 return Ok(Shape::M64n168k16);
1216 }
1217 stream.set_position(saved_pos);
1218 }
1219 stream.set_position(saved_pos);
1220 let saved_pos = stream.position();
1221 {
1223 let saved_pos = stream.position();
1224 if stream.expect_string(".m64n176k16").is_ok() {
1225 return Ok(Shape::M64n176k16);
1226 }
1227 stream.set_position(saved_pos);
1228 }
1229 stream.set_position(saved_pos);
1230 let saved_pos = stream.position();
1231 {
1233 let saved_pos = stream.position();
1234 if stream.expect_string(".m64n184k16").is_ok() {
1235 return Ok(Shape::M64n184k16);
1236 }
1237 stream.set_position(saved_pos);
1238 }
1239 stream.set_position(saved_pos);
1240 let saved_pos = stream.position();
1241 {
1243 let saved_pos = stream.position();
1244 if stream.expect_string(".m64n192k16").is_ok() {
1245 return Ok(Shape::M64n192k16);
1246 }
1247 stream.set_position(saved_pos);
1248 }
1249 stream.set_position(saved_pos);
1250 let saved_pos = stream.position();
1251 {
1253 let saved_pos = stream.position();
1254 if stream.expect_string(".m64n200k16").is_ok() {
1255 return Ok(Shape::M64n200k16);
1256 }
1257 stream.set_position(saved_pos);
1258 }
1259 stream.set_position(saved_pos);
1260 let saved_pos = stream.position();
1261 {
1263 let saved_pos = stream.position();
1264 if stream.expect_string(".m64n208k16").is_ok() {
1265 return Ok(Shape::M64n208k16);
1266 }
1267 stream.set_position(saved_pos);
1268 }
1269 stream.set_position(saved_pos);
1270 let saved_pos = stream.position();
1271 {
1273 let saved_pos = stream.position();
1274 if stream.expect_string(".m64n216k16").is_ok() {
1275 return Ok(Shape::M64n216k16);
1276 }
1277 stream.set_position(saved_pos);
1278 }
1279 stream.set_position(saved_pos);
1280 let saved_pos = stream.position();
1281 {
1283 let saved_pos = stream.position();
1284 if stream.expect_string(".m64n224k16").is_ok() {
1285 return Ok(Shape::M64n224k16);
1286 }
1287 stream.set_position(saved_pos);
1288 }
1289 stream.set_position(saved_pos);
1290 let saved_pos = stream.position();
1291 {
1293 let saved_pos = stream.position();
1294 if stream.expect_string(".m64n232k16").is_ok() {
1295 return Ok(Shape::M64n232k16);
1296 }
1297 stream.set_position(saved_pos);
1298 }
1299 stream.set_position(saved_pos);
1300 let saved_pos = stream.position();
1301 {
1303 let saved_pos = stream.position();
1304 if stream.expect_string(".m64n240k16").is_ok() {
1305 return Ok(Shape::M64n240k16);
1306 }
1307 stream.set_position(saved_pos);
1308 }
1309 stream.set_position(saved_pos);
1310 let saved_pos = stream.position();
1311 {
1313 let saved_pos = stream.position();
1314 if stream.expect_string(".m64n248k16").is_ok() {
1315 return Ok(Shape::M64n248k16);
1316 }
1317 stream.set_position(saved_pos);
1318 }
1319 stream.set_position(saved_pos);
1320 let saved_pos = stream.position();
1321 {
1323 let saved_pos = stream.position();
1324 if stream.expect_string(".m64n256k16").is_ok() {
1325 return Ok(Shape::M64n256k16);
1326 }
1327 stream.set_position(saved_pos);
1328 }
1329 stream.set_position(saved_pos);
1330 let saved_pos = stream.position();
1331 {
1333 let saved_pos = stream.position();
1334 if stream.expect_string(".m64n16k16").is_ok() {
1335 return Ok(Shape::M64n16k16);
1336 }
1337 stream.set_position(saved_pos);
1338 }
1339 stream.set_position(saved_pos);
1340 let saved_pos = stream.position();
1341 {
1343 let saved_pos = stream.position();
1344 if stream.expect_string(".m64n24k16").is_ok() {
1345 return Ok(Shape::M64n24k16);
1346 }
1347 stream.set_position(saved_pos);
1348 }
1349 stream.set_position(saved_pos);
1350 let saved_pos = stream.position();
1351 {
1353 let saved_pos = stream.position();
1354 if stream.expect_string(".m64n32k16").is_ok() {
1355 return Ok(Shape::M64n32k16);
1356 }
1357 stream.set_position(saved_pos);
1358 }
1359 stream.set_position(saved_pos);
1360 let saved_pos = stream.position();
1361 {
1363 let saved_pos = stream.position();
1364 if stream.expect_string(".m64n40k16").is_ok() {
1365 return Ok(Shape::M64n40k16);
1366 }
1367 stream.set_position(saved_pos);
1368 }
1369 stream.set_position(saved_pos);
1370 let saved_pos = stream.position();
1371 {
1373 let saved_pos = stream.position();
1374 if stream.expect_string(".m64n48k16").is_ok() {
1375 return Ok(Shape::M64n48k16);
1376 }
1377 stream.set_position(saved_pos);
1378 }
1379 stream.set_position(saved_pos);
1380 let saved_pos = stream.position();
1381 {
1383 let saved_pos = stream.position();
1384 if stream.expect_string(".m64n56k16").is_ok() {
1385 return Ok(Shape::M64n56k16);
1386 }
1387 stream.set_position(saved_pos);
1388 }
1389 stream.set_position(saved_pos);
1390 let saved_pos = stream.position();
1391 {
1393 let saved_pos = stream.position();
1394 if stream.expect_string(".m64n64k16").is_ok() {
1395 return Ok(Shape::M64n64k16);
1396 }
1397 stream.set_position(saved_pos);
1398 }
1399 stream.set_position(saved_pos);
1400 let saved_pos = stream.position();
1401 {
1403 let saved_pos = stream.position();
1404 if stream.expect_string(".m64n72k16").is_ok() {
1405 return Ok(Shape::M64n72k16);
1406 }
1407 stream.set_position(saved_pos);
1408 }
1409 stream.set_position(saved_pos);
1410 let saved_pos = stream.position();
1411 {
1413 let saved_pos = stream.position();
1414 if stream.expect_string(".m64n80k16").is_ok() {
1415 return Ok(Shape::M64n80k16);
1416 }
1417 stream.set_position(saved_pos);
1418 }
1419 stream.set_position(saved_pos);
1420 let saved_pos = stream.position();
1421 {
1423 let saved_pos = stream.position();
1424 if stream.expect_string(".m64n88k16").is_ok() {
1425 return Ok(Shape::M64n88k16);
1426 }
1427 stream.set_position(saved_pos);
1428 }
1429 stream.set_position(saved_pos);
1430 let saved_pos = stream.position();
1431 {
1433 let saved_pos = stream.position();
1434 if stream.expect_string(".m64n96k16").is_ok() {
1435 return Ok(Shape::M64n96k16);
1436 }
1437 stream.set_position(saved_pos);
1438 }
1439 stream.set_position(saved_pos);
1440 let saved_pos = stream.position();
1441 {
1443 let saved_pos = stream.position();
1444 if stream.expect_string(".m64n8k16").is_ok() {
1445 return Ok(Shape::M64n8k16);
1446 }
1447 stream.set_position(saved_pos);
1448 }
1449 stream.set_position(saved_pos);
1450 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1451 let expected = &[".m64n104k16", ".m64n112k16", ".m64n120k16", ".m64n128k16", ".m64n136k16", ".m64n144k16", ".m64n152k16", ".m64n160k16", ".m64n168k16", ".m64n176k16", ".m64n184k16", ".m64n192k16", ".m64n200k16", ".m64n208k16", ".m64n216k16", ".m64n224k16", ".m64n232k16", ".m64n240k16", ".m64n248k16", ".m64n256k16", ".m64n16k16", ".m64n24k16", ".m64n32k16", ".m64n40k16", ".m64n48k16", ".m64n56k16", ".m64n64k16", ".m64n72k16", ".m64n80k16", ".m64n88k16", ".m64n96k16", ".m64n8k16"];
1452 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1453 Err(crate::parser::unexpected_value(span, expected, found))
1454 }
1455 }
1456
1457 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32 {
1458 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1459 stream.expect_string("wgmma")?;
1460 stream.expect_string(".mma_async")?;
1461 let mma_async = ();
1462 stream.expect_complete()?;
1463 stream.expect_string(".sp")?;
1464 let sp = ();
1465 stream.expect_complete()?;
1466 stream.expect_string(".sync")?;
1467 let sync = ();
1468 stream.expect_complete()?;
1469 stream.expect_string(".aligned")?;
1470 let aligned = ();
1471 stream.expect_complete()?;
1472 let shape = Shape::parse(stream)?;
1473 stream.expect_complete()?;
1474 let dtype = Dtype::parse(stream)?;
1475 stream.expect_complete()?;
1476 stream.expect_string(".tf32")?;
1477 let tf32 = ();
1478 stream.expect_complete()?;
1479 stream.expect_string(".tf32")?;
1480 let tf322 = ();
1481 stream.expect_complete()?;
1482 let d = GeneralOperand::parse(stream)?;
1483 stream.expect_complete()?;
1484 stream.expect(&PtxToken::Comma)?;
1485 let a_desc = GeneralOperand::parse(stream)?;
1486 stream.expect_complete()?;
1487 stream.expect(&PtxToken::Comma)?;
1488 let b_desc = GeneralOperand::parse(stream)?;
1489 stream.expect_complete()?;
1490 stream.expect(&PtxToken::Comma)?;
1491 let sp_meta = GeneralOperand::parse(stream)?;
1492 stream.expect_complete()?;
1493 stream.expect(&PtxToken::Comma)?;
1494 let sp_sel = GeneralOperand::parse(stream)?;
1495 stream.expect_complete()?;
1496 stream.expect(&PtxToken::Comma)?;
1497 let scale_d = GeneralOperand::parse(stream)?;
1498 stream.expect_complete()?;
1499 stream.expect(&PtxToken::Comma)?;
1500 let imm_scale_a = GeneralOperand::parse(stream)?;
1501 stream.expect_complete()?;
1502 stream.expect(&PtxToken::Comma)?;
1503 let imm_scale_b = GeneralOperand::parse(stream)?;
1504 stream.expect_complete()?;
1505 stream.expect_complete()?;
1506 stream.expect(&PtxToken::Semicolon)?;
1507 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf32 {
1508 mma_async,
1509 sp,
1510 sync,
1511 aligned,
1512 shape,
1513 dtype,
1514 tf32,
1515 tf322,
1516 d,
1517 a_desc,
1518 b_desc,
1519 sp_meta,
1520 sp_sel,
1521 scale_d,
1522 imm_scale_a,
1523 imm_scale_b,
1524 })
1525 }
1526 }
1527
1528
1529 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321 {
1530 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1531 stream.expect_string("wgmma")?;
1532 stream.expect_string(".mma_async")?;
1533 let mma_async = ();
1534 stream.expect_complete()?;
1535 stream.expect_string(".sp")?;
1536 let sp = ();
1537 stream.expect_complete()?;
1538 stream.expect_string(".sync")?;
1539 let sync = ();
1540 stream.expect_complete()?;
1541 stream.expect_string(".aligned")?;
1542 let aligned = ();
1543 stream.expect_complete()?;
1544 let shape = Shape::parse(stream)?;
1545 stream.expect_complete()?;
1546 let dtype = Dtype::parse(stream)?;
1547 stream.expect_complete()?;
1548 stream.expect_string(".tf32")?;
1549 let tf32 = ();
1550 stream.expect_complete()?;
1551 stream.expect_string(".tf32")?;
1552 let tf322 = ();
1553 stream.expect_complete()?;
1554 let d = GeneralOperand::parse(stream)?;
1555 stream.expect_complete()?;
1556 stream.expect(&PtxToken::Comma)?;
1557 let a = GeneralOperand::parse(stream)?;
1558 stream.expect_complete()?;
1559 stream.expect(&PtxToken::Comma)?;
1560 let b_desc = GeneralOperand::parse(stream)?;
1561 stream.expect_complete()?;
1562 stream.expect(&PtxToken::Comma)?;
1563 let sp_meta = GeneralOperand::parse(stream)?;
1564 stream.expect_complete()?;
1565 stream.expect(&PtxToken::Comma)?;
1566 let sp_sel = GeneralOperand::parse(stream)?;
1567 stream.expect_complete()?;
1568 stream.expect(&PtxToken::Comma)?;
1569 let scale_d = GeneralOperand::parse(stream)?;
1570 stream.expect_complete()?;
1571 stream.expect(&PtxToken::Comma)?;
1572 let imm_scale_a = GeneralOperand::parse(stream)?;
1573 stream.expect_complete()?;
1574 stream.expect(&PtxToken::Comma)?;
1575 let imm_scale_b = GeneralOperand::parse(stream)?;
1576 stream.expect_complete()?;
1577 stream.expect_complete()?;
1578 stream.expect(&PtxToken::Semicolon)?;
1579 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeTf32Tf321 {
1580 mma_async,
1581 sp,
1582 sync,
1583 aligned,
1584 shape,
1585 dtype,
1586 tf32,
1587 tf322,
1588 d,
1589 a,
1590 b_desc,
1591 sp_meta,
1592 sp_sel,
1593 scale_d,
1594 imm_scale_a,
1595 imm_scale_b,
1596 })
1597 }
1598 }
1599
1600
1601}
1602
1603pub mod section_3 {
1604 use super::*;
1605 use crate::r#type::instruction::wgmma_mma_async_sp::section_3::*;
1606
1607 impl PtxParser for Atype {
1612 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1613 {
1615 let saved_pos = stream.position();
1616 if stream.expect_string(".e4m3").is_ok() {
1617 return Ok(Atype::E4m3);
1618 }
1619 stream.set_position(saved_pos);
1620 }
1621 let saved_pos = stream.position();
1622 {
1624 let saved_pos = stream.position();
1625 if stream.expect_string(".e5m2").is_ok() {
1626 return Ok(Atype::E5m2);
1627 }
1628 stream.set_position(saved_pos);
1629 }
1630 stream.set_position(saved_pos);
1631 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1632 let expected = &[".e4m3", ".e5m2"];
1633 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1634 Err(crate::parser::unexpected_value(span, expected, found))
1635 }
1636 }
1637
1638 impl PtxParser for Btype {
1639 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1640 {
1642 let saved_pos = stream.position();
1643 if stream.expect_string(".e4m3").is_ok() {
1644 return Ok(Btype::E4m3);
1645 }
1646 stream.set_position(saved_pos);
1647 }
1648 let saved_pos = stream.position();
1649 {
1651 let saved_pos = stream.position();
1652 if stream.expect_string(".e5m2").is_ok() {
1653 return Ok(Btype::E5m2);
1654 }
1655 stream.set_position(saved_pos);
1656 }
1657 stream.set_position(saved_pos);
1658 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1659 let expected = &[".e4m3", ".e5m2"];
1660 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1661 Err(crate::parser::unexpected_value(span, expected, found))
1662 }
1663 }
1664
1665 impl PtxParser for Dtype {
1666 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1667 {
1669 let saved_pos = stream.position();
1670 if stream.expect_string(".f16").is_ok() {
1671 return Ok(Dtype::F16);
1672 }
1673 stream.set_position(saved_pos);
1674 }
1675 let saved_pos = stream.position();
1676 {
1678 let saved_pos = stream.position();
1679 if stream.expect_string(".f32").is_ok() {
1680 return Ok(Dtype::F32);
1681 }
1682 stream.set_position(saved_pos);
1683 }
1684 stream.set_position(saved_pos);
1685 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
1686 let expected = &[".f16", ".f32"];
1687 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
1688 Err(crate::parser::unexpected_value(span, expected, found))
1689 }
1690 }
1691
1692 impl PtxParser for Shape {
1693 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
1694 {
1696 let saved_pos = stream.position();
1697 if stream.expect_string(".m64n104k64").is_ok() {
1698 return Ok(Shape::M64n104k64);
1699 }
1700 stream.set_position(saved_pos);
1701 }
1702 let saved_pos = stream.position();
1703 {
1705 let saved_pos = stream.position();
1706 if stream.expect_string(".m64n112k64").is_ok() {
1707 return Ok(Shape::M64n112k64);
1708 }
1709 stream.set_position(saved_pos);
1710 }
1711 stream.set_position(saved_pos);
1712 let saved_pos = stream.position();
1713 {
1715 let saved_pos = stream.position();
1716 if stream.expect_string(".m64n120k64").is_ok() {
1717 return Ok(Shape::M64n120k64);
1718 }
1719 stream.set_position(saved_pos);
1720 }
1721 stream.set_position(saved_pos);
1722 let saved_pos = stream.position();
1723 {
1725 let saved_pos = stream.position();
1726 if stream.expect_string(".m64n128k64").is_ok() {
1727 return Ok(Shape::M64n128k64);
1728 }
1729 stream.set_position(saved_pos);
1730 }
1731 stream.set_position(saved_pos);
1732 let saved_pos = stream.position();
1733 {
1735 let saved_pos = stream.position();
1736 if stream.expect_string(".m64n136k64").is_ok() {
1737 return Ok(Shape::M64n136k64);
1738 }
1739 stream.set_position(saved_pos);
1740 }
1741 stream.set_position(saved_pos);
1742 let saved_pos = stream.position();
1743 {
1745 let saved_pos = stream.position();
1746 if stream.expect_string(".m64n144k64").is_ok() {
1747 return Ok(Shape::M64n144k64);
1748 }
1749 stream.set_position(saved_pos);
1750 }
1751 stream.set_position(saved_pos);
1752 let saved_pos = stream.position();
1753 {
1755 let saved_pos = stream.position();
1756 if stream.expect_string(".m64n152k64").is_ok() {
1757 return Ok(Shape::M64n152k64);
1758 }
1759 stream.set_position(saved_pos);
1760 }
1761 stream.set_position(saved_pos);
1762 let saved_pos = stream.position();
1763 {
1765 let saved_pos = stream.position();
1766 if stream.expect_string(".m64n160k64").is_ok() {
1767 return Ok(Shape::M64n160k64);
1768 }
1769 stream.set_position(saved_pos);
1770 }
1771 stream.set_position(saved_pos);
1772 let saved_pos = stream.position();
1773 {
1775 let saved_pos = stream.position();
1776 if stream.expect_string(".m64n168k64").is_ok() {
1777 return Ok(Shape::M64n168k64);
1778 }
1779 stream.set_position(saved_pos);
1780 }
1781 stream.set_position(saved_pos);
1782 let saved_pos = stream.position();
1783 {
1785 let saved_pos = stream.position();
1786 if stream.expect_string(".m64n176k64").is_ok() {
1787 return Ok(Shape::M64n176k64);
1788 }
1789 stream.set_position(saved_pos);
1790 }
1791 stream.set_position(saved_pos);
1792 let saved_pos = stream.position();
1793 {
1795 let saved_pos = stream.position();
1796 if stream.expect_string(".m64n184k64").is_ok() {
1797 return Ok(Shape::M64n184k64);
1798 }
1799 stream.set_position(saved_pos);
1800 }
1801 stream.set_position(saved_pos);
1802 let saved_pos = stream.position();
1803 {
1805 let saved_pos = stream.position();
1806 if stream.expect_string(".m64n192k64").is_ok() {
1807 return Ok(Shape::M64n192k64);
1808 }
1809 stream.set_position(saved_pos);
1810 }
1811 stream.set_position(saved_pos);
1812 let saved_pos = stream.position();
1813 {
1815 let saved_pos = stream.position();
1816 if stream.expect_string(".m64n200k64").is_ok() {
1817 return Ok(Shape::M64n200k64);
1818 }
1819 stream.set_position(saved_pos);
1820 }
1821 stream.set_position(saved_pos);
1822 let saved_pos = stream.position();
1823 {
1825 let saved_pos = stream.position();
1826 if stream.expect_string(".m64n208k64").is_ok() {
1827 return Ok(Shape::M64n208k64);
1828 }
1829 stream.set_position(saved_pos);
1830 }
1831 stream.set_position(saved_pos);
1832 let saved_pos = stream.position();
1833 {
1835 let saved_pos = stream.position();
1836 if stream.expect_string(".m64n216k64").is_ok() {
1837 return Ok(Shape::M64n216k64);
1838 }
1839 stream.set_position(saved_pos);
1840 }
1841 stream.set_position(saved_pos);
1842 let saved_pos = stream.position();
1843 {
1845 let saved_pos = stream.position();
1846 if stream.expect_string(".m64n224k64").is_ok() {
1847 return Ok(Shape::M64n224k64);
1848 }
1849 stream.set_position(saved_pos);
1850 }
1851 stream.set_position(saved_pos);
1852 let saved_pos = stream.position();
1853 {
1855 let saved_pos = stream.position();
1856 if stream.expect_string(".m64n232k64").is_ok() {
1857 return Ok(Shape::M64n232k64);
1858 }
1859 stream.set_position(saved_pos);
1860 }
1861 stream.set_position(saved_pos);
1862 let saved_pos = stream.position();
1863 {
1865 let saved_pos = stream.position();
1866 if stream.expect_string(".m64n240k64").is_ok() {
1867 return Ok(Shape::M64n240k64);
1868 }
1869 stream.set_position(saved_pos);
1870 }
1871 stream.set_position(saved_pos);
1872 let saved_pos = stream.position();
1873 {
1875 let saved_pos = stream.position();
1876 if stream.expect_string(".m64n248k64").is_ok() {
1877 return Ok(Shape::M64n248k64);
1878 }
1879 stream.set_position(saved_pos);
1880 }
1881 stream.set_position(saved_pos);
1882 let saved_pos = stream.position();
1883 {
1885 let saved_pos = stream.position();
1886 if stream.expect_string(".m64n256k64").is_ok() {
1887 return Ok(Shape::M64n256k64);
1888 }
1889 stream.set_position(saved_pos);
1890 }
1891 stream.set_position(saved_pos);
1892 let saved_pos = stream.position();
1893 {
1895 let saved_pos = stream.position();
1896 if stream.expect_string(".m64n16k64").is_ok() {
1897 return Ok(Shape::M64n16k64);
1898 }
1899 stream.set_position(saved_pos);
1900 }
1901 stream.set_position(saved_pos);
1902 let saved_pos = stream.position();
1903 {
1905 let saved_pos = stream.position();
1906 if stream.expect_string(".m64n24k64").is_ok() {
1907 return Ok(Shape::M64n24k64);
1908 }
1909 stream.set_position(saved_pos);
1910 }
1911 stream.set_position(saved_pos);
1912 let saved_pos = stream.position();
1913 {
1915 let saved_pos = stream.position();
1916 if stream.expect_string(".m64n32k64").is_ok() {
1917 return Ok(Shape::M64n32k64);
1918 }
1919 stream.set_position(saved_pos);
1920 }
1921 stream.set_position(saved_pos);
1922 let saved_pos = stream.position();
1923 {
1925 let saved_pos = stream.position();
1926 if stream.expect_string(".m64n40k64").is_ok() {
1927 return Ok(Shape::M64n40k64);
1928 }
1929 stream.set_position(saved_pos);
1930 }
1931 stream.set_position(saved_pos);
1932 let saved_pos = stream.position();
1933 {
1935 let saved_pos = stream.position();
1936 if stream.expect_string(".m64n48k64").is_ok() {
1937 return Ok(Shape::M64n48k64);
1938 }
1939 stream.set_position(saved_pos);
1940 }
1941 stream.set_position(saved_pos);
1942 let saved_pos = stream.position();
1943 {
1945 let saved_pos = stream.position();
1946 if stream.expect_string(".m64n56k64").is_ok() {
1947 return Ok(Shape::M64n56k64);
1948 }
1949 stream.set_position(saved_pos);
1950 }
1951 stream.set_position(saved_pos);
1952 let saved_pos = stream.position();
1953 {
1955 let saved_pos = stream.position();
1956 if stream.expect_string(".m64n64k64").is_ok() {
1957 return Ok(Shape::M64n64k64);
1958 }
1959 stream.set_position(saved_pos);
1960 }
1961 stream.set_position(saved_pos);
1962 let saved_pos = stream.position();
1963 {
1965 let saved_pos = stream.position();
1966 if stream.expect_string(".m64n72k64").is_ok() {
1967 return Ok(Shape::M64n72k64);
1968 }
1969 stream.set_position(saved_pos);
1970 }
1971 stream.set_position(saved_pos);
1972 let saved_pos = stream.position();
1973 {
1975 let saved_pos = stream.position();
1976 if stream.expect_string(".m64n80k64").is_ok() {
1977 return Ok(Shape::M64n80k64);
1978 }
1979 stream.set_position(saved_pos);
1980 }
1981 stream.set_position(saved_pos);
1982 let saved_pos = stream.position();
1983 {
1985 let saved_pos = stream.position();
1986 if stream.expect_string(".m64n88k64").is_ok() {
1987 return Ok(Shape::M64n88k64);
1988 }
1989 stream.set_position(saved_pos);
1990 }
1991 stream.set_position(saved_pos);
1992 let saved_pos = stream.position();
1993 {
1995 let saved_pos = stream.position();
1996 if stream.expect_string(".m64n96k64").is_ok() {
1997 return Ok(Shape::M64n96k64);
1998 }
1999 stream.set_position(saved_pos);
2000 }
2001 stream.set_position(saved_pos);
2002 let saved_pos = stream.position();
2003 {
2005 let saved_pos = stream.position();
2006 if stream.expect_string(".m64n8k64").is_ok() {
2007 return Ok(Shape::M64n8k64);
2008 }
2009 stream.set_position(saved_pos);
2010 }
2011 stream.set_position(saved_pos);
2012 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2013 let expected = &[".m64n104k64", ".m64n112k64", ".m64n120k64", ".m64n128k64", ".m64n136k64", ".m64n144k64", ".m64n152k64", ".m64n160k64", ".m64n168k64", ".m64n176k64", ".m64n184k64", ".m64n192k64", ".m64n200k64", ".m64n208k64", ".m64n216k64", ".m64n224k64", ".m64n232k64", ".m64n240k64", ".m64n248k64", ".m64n256k64", ".m64n16k64", ".m64n24k64", ".m64n32k64", ".m64n40k64", ".m64n48k64", ".m64n56k64", ".m64n64k64", ".m64n72k64", ".m64n80k64", ".m64n88k64", ".m64n96k64", ".m64n8k64"];
2014 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2015 Err(crate::parser::unexpected_value(span, expected, found))
2016 }
2017 }
2018
2019 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype {
2020 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2021 stream.expect_string("wgmma")?;
2022 stream.expect_string(".mma_async")?;
2023 let mma_async = ();
2024 stream.expect_complete()?;
2025 stream.expect_string(".sp")?;
2026 let sp = ();
2027 stream.expect_complete()?;
2028 stream.expect_string(".sync")?;
2029 let sync = ();
2030 stream.expect_complete()?;
2031 stream.expect_string(".aligned")?;
2032 let aligned = ();
2033 stream.expect_complete()?;
2034 let shape = Shape::parse(stream)?;
2035 stream.expect_complete()?;
2036 let dtype = Dtype::parse(stream)?;
2037 stream.expect_complete()?;
2038 let atype = Atype::parse(stream)?;
2039 stream.expect_complete()?;
2040 let btype = Btype::parse(stream)?;
2041 stream.expect_complete()?;
2042 let d = GeneralOperand::parse(stream)?;
2043 stream.expect_complete()?;
2044 stream.expect(&PtxToken::Comma)?;
2045 let a_desc = GeneralOperand::parse(stream)?;
2046 stream.expect_complete()?;
2047 stream.expect(&PtxToken::Comma)?;
2048 let b_desc = GeneralOperand::parse(stream)?;
2049 stream.expect_complete()?;
2050 stream.expect(&PtxToken::Comma)?;
2051 let sp_meta = GeneralOperand::parse(stream)?;
2052 stream.expect_complete()?;
2053 stream.expect(&PtxToken::Comma)?;
2054 let sp_sel = GeneralOperand::parse(stream)?;
2055 stream.expect_complete()?;
2056 stream.expect(&PtxToken::Comma)?;
2057 let scale_d = GeneralOperand::parse(stream)?;
2058 stream.expect_complete()?;
2059 stream.expect(&PtxToken::Comma)?;
2060 let imm_scale_a = GeneralOperand::parse(stream)?;
2061 stream.expect_complete()?;
2062 stream.expect(&PtxToken::Comma)?;
2063 let imm_scale_b = GeneralOperand::parse(stream)?;
2064 stream.expect_complete()?;
2065 stream.expect_complete()?;
2066 stream.expect(&PtxToken::Semicolon)?;
2067 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype {
2068 mma_async,
2069 sp,
2070 sync,
2071 aligned,
2072 shape,
2073 dtype,
2074 atype,
2075 btype,
2076 d,
2077 a_desc,
2078 b_desc,
2079 sp_meta,
2080 sp_sel,
2081 scale_d,
2082 imm_scale_a,
2083 imm_scale_b,
2084 })
2085 }
2086 }
2087
2088
2089 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1 {
2090 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2091 stream.expect_string("wgmma")?;
2092 stream.expect_string(".mma_async")?;
2093 let mma_async = ();
2094 stream.expect_complete()?;
2095 stream.expect_string(".sp")?;
2096 let sp = ();
2097 stream.expect_complete()?;
2098 stream.expect_string(".sync")?;
2099 let sync = ();
2100 stream.expect_complete()?;
2101 stream.expect_string(".aligned")?;
2102 let aligned = ();
2103 stream.expect_complete()?;
2104 let shape = Shape::parse(stream)?;
2105 stream.expect_complete()?;
2106 let dtype = Dtype::parse(stream)?;
2107 stream.expect_complete()?;
2108 let atype = Atype::parse(stream)?;
2109 stream.expect_complete()?;
2110 let btype = Btype::parse(stream)?;
2111 stream.expect_complete()?;
2112 let d = GeneralOperand::parse(stream)?;
2113 stream.expect_complete()?;
2114 stream.expect(&PtxToken::Comma)?;
2115 let a = GeneralOperand::parse(stream)?;
2116 stream.expect_complete()?;
2117 stream.expect(&PtxToken::Comma)?;
2118 let b_desc = GeneralOperand::parse(stream)?;
2119 stream.expect_complete()?;
2120 stream.expect(&PtxToken::Comma)?;
2121 let sp_meta = GeneralOperand::parse(stream)?;
2122 stream.expect_complete()?;
2123 stream.expect(&PtxToken::Comma)?;
2124 let sp_sel = GeneralOperand::parse(stream)?;
2125 stream.expect_complete()?;
2126 stream.expect(&PtxToken::Comma)?;
2127 let scale_d = GeneralOperand::parse(stream)?;
2128 stream.expect_complete()?;
2129 stream.expect(&PtxToken::Comma)?;
2130 let imm_scale_a = GeneralOperand::parse(stream)?;
2131 stream.expect_complete()?;
2132 stream.expect(&PtxToken::Comma)?;
2133 let imm_scale_b = GeneralOperand::parse(stream)?;
2134 stream.expect_complete()?;
2135 stream.expect_complete()?;
2136 stream.expect(&PtxToken::Semicolon)?;
2137 Ok(WgmmaMmaAsyncSpSyncAlignedShapeDtypeAtypeBtype1 {
2138 mma_async,
2139 sp,
2140 sync,
2141 aligned,
2142 shape,
2143 dtype,
2144 atype,
2145 btype,
2146 d,
2147 a,
2148 b_desc,
2149 sp_meta,
2150 sp_sel,
2151 scale_d,
2152 imm_scale_a,
2153 imm_scale_b,
2154 })
2155 }
2156 }
2157
2158
2159}
2160
2161pub mod section_4 {
2162 use super::*;
2163 use crate::r#type::instruction::wgmma_mma_async_sp::section_4::*;
2164
2165 impl PtxParser for Atype {
2170 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2171 {
2173 let saved_pos = stream.position();
2174 if stream.expect_string(".s8").is_ok() {
2175 return Ok(Atype::S8);
2176 }
2177 stream.set_position(saved_pos);
2178 }
2179 let saved_pos = stream.position();
2180 {
2182 let saved_pos = stream.position();
2183 if stream.expect_string(".u8").is_ok() {
2184 return Ok(Atype::U8);
2185 }
2186 stream.set_position(saved_pos);
2187 }
2188 stream.set_position(saved_pos);
2189 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2190 let expected = &[".s8", ".u8"];
2191 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2192 Err(crate::parser::unexpected_value(span, expected, found))
2193 }
2194 }
2195
2196 impl PtxParser for Btype {
2197 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2198 {
2200 let saved_pos = stream.position();
2201 if stream.expect_string(".s8").is_ok() {
2202 return Ok(Btype::S8);
2203 }
2204 stream.set_position(saved_pos);
2205 }
2206 let saved_pos = stream.position();
2207 {
2209 let saved_pos = stream.position();
2210 if stream.expect_string(".u8").is_ok() {
2211 return Ok(Btype::U8);
2212 }
2213 stream.set_position(saved_pos);
2214 }
2215 stream.set_position(saved_pos);
2216 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2217 let expected = &[".s8", ".u8"];
2218 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2219 Err(crate::parser::unexpected_value(span, expected, found))
2220 }
2221 }
2222
2223 impl PtxParser for Shape {
2224 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2225 {
2227 let saved_pos = stream.position();
2228 if stream.expect_string(".m64n112k64").is_ok() {
2229 return Ok(Shape::M64n112k64);
2230 }
2231 stream.set_position(saved_pos);
2232 }
2233 let saved_pos = stream.position();
2234 {
2236 let saved_pos = stream.position();
2237 if stream.expect_string(".m64n128k64").is_ok() {
2238 return Ok(Shape::M64n128k64);
2239 }
2240 stream.set_position(saved_pos);
2241 }
2242 stream.set_position(saved_pos);
2243 let saved_pos = stream.position();
2244 {
2246 let saved_pos = stream.position();
2247 if stream.expect_string(".m64n144k64").is_ok() {
2248 return Ok(Shape::M64n144k64);
2249 }
2250 stream.set_position(saved_pos);
2251 }
2252 stream.set_position(saved_pos);
2253 let saved_pos = stream.position();
2254 {
2256 let saved_pos = stream.position();
2257 if stream.expect_string(".m64n160k64").is_ok() {
2258 return Ok(Shape::M64n160k64);
2259 }
2260 stream.set_position(saved_pos);
2261 }
2262 stream.set_position(saved_pos);
2263 let saved_pos = stream.position();
2264 {
2266 let saved_pos = stream.position();
2267 if stream.expect_string(".m64n176k64").is_ok() {
2268 return Ok(Shape::M64n176k64);
2269 }
2270 stream.set_position(saved_pos);
2271 }
2272 stream.set_position(saved_pos);
2273 let saved_pos = stream.position();
2274 {
2276 let saved_pos = stream.position();
2277 if stream.expect_string(".m64n192k64").is_ok() {
2278 return Ok(Shape::M64n192k64);
2279 }
2280 stream.set_position(saved_pos);
2281 }
2282 stream.set_position(saved_pos);
2283 let saved_pos = stream.position();
2284 {
2286 let saved_pos = stream.position();
2287 if stream.expect_string(".m64n208k64").is_ok() {
2288 return Ok(Shape::M64n208k64);
2289 }
2290 stream.set_position(saved_pos);
2291 }
2292 stream.set_position(saved_pos);
2293 let saved_pos = stream.position();
2294 {
2296 let saved_pos = stream.position();
2297 if stream.expect_string(".m64n224k64").is_ok() {
2298 return Ok(Shape::M64n224k64);
2299 }
2300 stream.set_position(saved_pos);
2301 }
2302 stream.set_position(saved_pos);
2303 let saved_pos = stream.position();
2304 {
2306 let saved_pos = stream.position();
2307 if stream.expect_string(".m64n240k64").is_ok() {
2308 return Ok(Shape::M64n240k64);
2309 }
2310 stream.set_position(saved_pos);
2311 }
2312 stream.set_position(saved_pos);
2313 let saved_pos = stream.position();
2314 {
2316 let saved_pos = stream.position();
2317 if stream.expect_string(".m64n256k64").is_ok() {
2318 return Ok(Shape::M64n256k64);
2319 }
2320 stream.set_position(saved_pos);
2321 }
2322 stream.set_position(saved_pos);
2323 let saved_pos = stream.position();
2324 {
2326 let saved_pos = stream.position();
2327 if stream.expect_string(".m64n16k64").is_ok() {
2328 return Ok(Shape::M64n16k64);
2329 }
2330 stream.set_position(saved_pos);
2331 }
2332 stream.set_position(saved_pos);
2333 let saved_pos = stream.position();
2334 {
2336 let saved_pos = stream.position();
2337 if stream.expect_string(".m64n24k64").is_ok() {
2338 return Ok(Shape::M64n24k64);
2339 }
2340 stream.set_position(saved_pos);
2341 }
2342 stream.set_position(saved_pos);
2343 let saved_pos = stream.position();
2344 {
2346 let saved_pos = stream.position();
2347 if stream.expect_string(".m64n32k64").is_ok() {
2348 return Ok(Shape::M64n32k64);
2349 }
2350 stream.set_position(saved_pos);
2351 }
2352 stream.set_position(saved_pos);
2353 let saved_pos = stream.position();
2354 {
2356 let saved_pos = stream.position();
2357 if stream.expect_string(".m64n48k64").is_ok() {
2358 return Ok(Shape::M64n48k64);
2359 }
2360 stream.set_position(saved_pos);
2361 }
2362 stream.set_position(saved_pos);
2363 let saved_pos = stream.position();
2364 {
2366 let saved_pos = stream.position();
2367 if stream.expect_string(".m64n64k64").is_ok() {
2368 return Ok(Shape::M64n64k64);
2369 }
2370 stream.set_position(saved_pos);
2371 }
2372 stream.set_position(saved_pos);
2373 let saved_pos = stream.position();
2374 {
2376 let saved_pos = stream.position();
2377 if stream.expect_string(".m64n80k64").is_ok() {
2378 return Ok(Shape::M64n80k64);
2379 }
2380 stream.set_position(saved_pos);
2381 }
2382 stream.set_position(saved_pos);
2383 let saved_pos = stream.position();
2384 {
2386 let saved_pos = stream.position();
2387 if stream.expect_string(".m64n96k64").is_ok() {
2388 return Ok(Shape::M64n96k64);
2389 }
2390 stream.set_position(saved_pos);
2391 }
2392 stream.set_position(saved_pos);
2393 let saved_pos = stream.position();
2394 {
2396 let saved_pos = stream.position();
2397 if stream.expect_string(".m64n8k64").is_ok() {
2398 return Ok(Shape::M64n8k64);
2399 }
2400 stream.set_position(saved_pos);
2401 }
2402 stream.set_position(saved_pos);
2403 let span = stream.peek().map(|(_, s)| s.clone()).unwrap_or(Span { start: 0, end: 0 });
2404 let expected = &[".m64n112k64", ".m64n128k64", ".m64n144k64", ".m64n160k64", ".m64n176k64", ".m64n192k64", ".m64n208k64", ".m64n224k64", ".m64n240k64", ".m64n256k64", ".m64n16k64", ".m64n24k64", ".m64n32k64", ".m64n48k64", ".m64n64k64", ".m64n80k64", ".m64n96k64", ".m64n8k64"];
2405 let found = stream.peek().map(|(t, _)| format!("{:?}", t)).unwrap_or_else(|_| "<end of input>".to_string());
2406 Err(crate::parser::unexpected_value(span, expected, found))
2407 }
2408 }
2409
2410 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype {
2411 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2412 stream.expect_string("wgmma")?;
2413 stream.expect_string(".mma_async")?;
2414 let mma_async = ();
2415 stream.expect_complete()?;
2416 stream.expect_string(".sp")?;
2417 let sp = ();
2418 stream.expect_complete()?;
2419 stream.expect_string(".sync")?;
2420 let sync = ();
2421 stream.expect_complete()?;
2422 stream.expect_string(".aligned")?;
2423 let aligned = ();
2424 stream.expect_complete()?;
2425 let shape = Shape::parse(stream)?;
2426 stream.expect_complete()?;
2427 let saved_pos = stream.position();
2428 let satfinite = stream.expect_string(".satfinite").is_ok();
2429 if !satfinite {
2430 stream.set_position(saved_pos);
2431 }
2432 stream.expect_complete()?;
2433 stream.expect_string(".s32")?;
2434 let s32 = ();
2435 stream.expect_complete()?;
2436 let atype = Atype::parse(stream)?;
2437 stream.expect_complete()?;
2438 let btype = Btype::parse(stream)?;
2439 stream.expect_complete()?;
2440 let d = GeneralOperand::parse(stream)?;
2441 stream.expect_complete()?;
2442 stream.expect(&PtxToken::Comma)?;
2443 let a_desc = GeneralOperand::parse(stream)?;
2444 stream.expect_complete()?;
2445 stream.expect(&PtxToken::Comma)?;
2446 let b_desc = GeneralOperand::parse(stream)?;
2447 stream.expect_complete()?;
2448 stream.expect(&PtxToken::Comma)?;
2449 let sp_meta = GeneralOperand::parse(stream)?;
2450 stream.expect_complete()?;
2451 stream.expect(&PtxToken::Comma)?;
2452 let sp_sel = GeneralOperand::parse(stream)?;
2453 stream.expect_complete()?;
2454 stream.expect(&PtxToken::Comma)?;
2455 let scale_d = GeneralOperand::parse(stream)?;
2456 stream.expect_complete()?;
2457 stream.expect_complete()?;
2458 stream.expect(&PtxToken::Semicolon)?;
2459 Ok(WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype {
2460 mma_async,
2461 sp,
2462 sync,
2463 aligned,
2464 shape,
2465 satfinite,
2466 s32,
2467 atype,
2468 btype,
2469 d,
2470 a_desc,
2471 b_desc,
2472 sp_meta,
2473 sp_sel,
2474 scale_d,
2475 })
2476 }
2477 }
2478
2479
2480 impl PtxParser for WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1 {
2481 fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
2482 stream.expect_string("wgmma")?;
2483 stream.expect_string(".mma_async")?;
2484 let mma_async = ();
2485 stream.expect_complete()?;
2486 stream.expect_string(".sp")?;
2487 let sp = ();
2488 stream.expect_complete()?;
2489 stream.expect_string(".sync")?;
2490 let sync = ();
2491 stream.expect_complete()?;
2492 stream.expect_string(".aligned")?;
2493 let aligned = ();
2494 stream.expect_complete()?;
2495 let shape = Shape::parse(stream)?;
2496 stream.expect_complete()?;
2497 let saved_pos = stream.position();
2498 let satfinite = stream.expect_string(".satfinite").is_ok();
2499 if !satfinite {
2500 stream.set_position(saved_pos);
2501 }
2502 stream.expect_complete()?;
2503 stream.expect_string(".s32")?;
2504 let s32 = ();
2505 stream.expect_complete()?;
2506 let atype = Atype::parse(stream)?;
2507 stream.expect_complete()?;
2508 let btype = Btype::parse(stream)?;
2509 stream.expect_complete()?;
2510 let d = GeneralOperand::parse(stream)?;
2511 stream.expect_complete()?;
2512 stream.expect(&PtxToken::Comma)?;
2513 let a = GeneralOperand::parse(stream)?;
2514 stream.expect_complete()?;
2515 stream.expect(&PtxToken::Comma)?;
2516 let b_desc = GeneralOperand::parse(stream)?;
2517 stream.expect_complete()?;
2518 stream.expect(&PtxToken::Comma)?;
2519 let sp_meta = GeneralOperand::parse(stream)?;
2520 stream.expect_complete()?;
2521 stream.expect(&PtxToken::Comma)?;
2522 let sp_sel = GeneralOperand::parse(stream)?;
2523 stream.expect_complete()?;
2524 stream.expect(&PtxToken::Comma)?;
2525 let scale_d = GeneralOperand::parse(stream)?;
2526 stream.expect_complete()?;
2527 stream.expect_complete()?;
2528 stream.expect(&PtxToken::Semicolon)?;
2529 Ok(WgmmaMmaAsyncSpSyncAlignedShapeSatfiniteS32AtypeBtype1 {
2530 mma_async,
2531 sp,
2532 sync,
2533 aligned,
2534 shape,
2535 satfinite,
2536 s32,
2537 atype,
2538 btype,
2539 d,
2540 a,
2541 b_desc,
2542 sp_meta,
2543 sp_sel,
2544 scale_d,
2545 })
2546 }
2547 }
2548
2549
2550}
2551