ptx_parser/parser/instruction/
tcgen05_fence.rs

1//! Original PTX specification:
2//!
3//! tcgen05.fence::before_thread_sync ;
4//! tcgen05.fence::after_thread_sync  ;
5
6#![allow(unused)]
7
8use crate::parser::{
9    PtxParseError, PtxParser, PtxTokenStream, Span,
10    util::{
11        between, comma_p, directive_p, exclamation_p, lbracket_p, lparen_p, map, minus_p, optional,
12        pipe_p, rbracket_p, rparen_p, semicolon_p, sep_by, string_p, try_map,
13    },
14};
15use crate::r#type::common::*;
16use crate::{alt, ok, seq_n};
17
18pub mod section_0 {
19    use super::*;
20    use crate::r#type::instruction::tcgen05_fence::section_0::*;
21
22    impl PtxParser for Tcgen05FenceBeforeThreadSync {
23        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
24            try_map(
25                seq_n!(
26                    string_p("tcgen05"),
27                    string_p(".fence::before_thread_sync"),
28                    semicolon_p()
29                ),
30                |(_, fence_before_thread_sync, _), span| {
31                    ok!(Tcgen05FenceBeforeThreadSync {
32                        fence_before_thread_sync = fence_before_thread_sync,
33
34                    })
35                },
36            )
37        }
38    }
39
40    impl PtxParser for Tcgen05FenceAfterThreadSync {
41        fn parse() -> impl Fn(&mut PtxTokenStream) -> Result<(Self, Span), PtxParseError> {
42            try_map(
43                seq_n!(
44                    string_p("tcgen05"),
45                    string_p(".fence::after_thread_sync"),
46                    semicolon_p()
47                ),
48                |(_, fence_after_thread_sync, _), span| {
49                    ok!(Tcgen05FenceAfterThreadSync {
50                        fence_after_thread_sync = fence_after_thread_sync,
51
52                    })
53                },
54            )
55        }
56    }
57}
ptx_parser/parser/instruction/tcgen05_fence.rs

ptx_parser/parser/instruction/
tcgen05_fence.rs
