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::lexer::PtxToken;
9use crate::parser::{PtxParseError, PtxParser, PtxTokenStream, Span};
10use crate::r#type::common::*;
11
12pub mod section_0 {
13    use super::*;
14    use crate::r#type::instruction::tcgen05_fence::section_0::*;
15
16    impl PtxParser for Tcgen05FenceBeforeThreadSync {
17        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
18            stream.expect_string("tcgen05")?;
19            stream.expect_string(".fence::before_thread_sync")?;
20            let fence_before_thread_sync = ();
21            stream.expect_complete()?;
22            stream.expect_complete()?;
23            stream.expect(&PtxToken::Semicolon)?;
24            Ok(Tcgen05FenceBeforeThreadSync {
25                fence_before_thread_sync,
26            })
27        }
28    }
29
30    impl PtxParser for Tcgen05FenceAfterThreadSync {
31        fn parse(stream: &mut PtxTokenStream) -> Result<Self, PtxParseError> {
32            stream.expect_string("tcgen05")?;
33            stream.expect_string(".fence::after_thread_sync")?;
34            let fence_after_thread_sync = ();
35            stream.expect_complete()?;
36            stream.expect_complete()?;
37            stream.expect(&PtxToken::Semicolon)?;
38            Ok(Tcgen05FenceAfterThreadSync {
39                fence_after_thread_sync,
40            })
41        }
42    }
43}