diff --git a/.cargo/config.toml b/.cargo/config.toml index e9517f0e..0cb4dc23 100644 --- a/.cargo/config.toml +++ b/.cargo/config.toml @@ -12,8 +12,8 @@ accelsim-sim = "run --release --package accelsim-sim --" playground-sim = "run --release --package playground-sim --" pchase = "run --release --package gpucachesim-benchmarks --bin pchase --" -t = "test -- --test-threads=1" -tt = "test --release -- --test-threads=1" +t = "test --features local-data -- --test-threads=1" +tt = "test --release --features local-data -- --test-threads=1" -tv = "test -- --nocapture --test-threads=1" -ttv = "test --release -- --nocapture --test-threads=1" +tv = "test --features local-data -- --nocapture --test-threads=1" +ttv = "test --release --features local-data -- --nocapture --test-threads=1" diff --git a/.gitignore b/.gitignore index fbc0dd0d..a6f6ac85 100644 --- a/.gitignore +++ b/.gitignore @@ -22,6 +22,7 @@ /debug.* /debug*.* +/temp-traces *debug-trace .ipynb_checkpoints Pipfile.lock diff --git a/Cargo.lock b/Cargo.lock index da87a330..cab3f97a 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -689,16 +689,6 @@ version = "1.0.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "acbf1af155f9b9ef647e42cdc158db4b64a1b61f743629225fde6f3e0be2a7c7" -[[package]] -name = "colored" -version = "2.1.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cbf2150cce219b664a8a70df7a1f933836724b503f8a413af9365b4dcc4d90b8" -dependencies = [ - "lazy_static", - "windows-sys 0.48.0", -] - [[package]] name = "colorsys" version = "0.6.7" @@ -2546,19 +2536,6 @@ dependencies = [ "syn 2.0.52", ] -[[package]] -name = "pest-test" -version = "0.1.6" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "a426ffa96d4df49eb2a43b010059dc085c4a0b655fd229de11b49340d1deaffe" -dependencies = [ - "colored", - "pest", - "pest_derive", - "snailquote", - "thiserror", -] - [[package]] name = "pest_derive" version = "2.7.8" @@ -2852,7 +2829,6 @@ dependencies = [ "once_cell", "pest", "pest-ast", - "pest-test", "pest_derive", "snailquote", "thiserror", diff --git a/accelsim/.gitignore b/accelsim/.gitignore index fe8b37c0..703fbc1b 100644 --- a/accelsim/.gitignore +++ b/accelsim/.gitignore @@ -1,5 +1,11 @@ .DS_Store +# make sure, this should also be excluded in the simulator directory +accel-sim-framework-dev/gpu-simulator/build +accel-sim-framework-dev/gpu-simulator/bin +accel-sim-framework-dev/gpu-simulator/gpgpu-sim/build +accel-sim-framework-dev/gpu-simulator/gpgpu-sim/bin + # build logs build.log* accel-sim-framework-dev-*/ diff --git a/accelsim/.ignore b/accelsim/.ignore index b9796036..8bacde92 100644 --- a/accelsim/.ignore +++ b/accelsim/.ignore @@ -1 +1,5 @@ accel-sim-framework-dev/util/ +accel-sim-framework-dev/gpu-simulator/build +accel-sim-framework-dev/gpu-simulator/bin +accel-sim-framework-dev/gpu-simulator/gpgpu-sim/build +accel-sim-framework-dev/gpu-simulator/gpgpu-sim/bin diff --git a/exec/src/cfg.rs b/exec/src/cfg.rs index 9accef9a..a1f6f632 100644 --- a/exec/src/cfg.rs +++ b/exec/src/cfg.rs @@ -571,9 +571,11 @@ pub mod visit { self.graph[*last_branch_node_idx], self.graph[node_idx], ); - // continue from this reconvergence point + + // all paths from last branch node have been traversed self.dominator_stack.pop(); + // continue from this reconvergence point for (outgoing_edge_idx, next_node_idx) in self.graph.outgoing_neigbors(node_idx) { diff --git a/exec/src/tracegen.rs b/exec/src/tracegen.rs index 515ca231..d26e4a02 100644 --- a/exec/src/tracegen.rs +++ b/exec/src/tracegen.rs @@ -677,14 +677,14 @@ impl TraceGenerator for Tracer { let is_load = access.kind == model::MemAccessKind::Load; let is_store = access.kind == model::MemAccessKind::Store; let mut instr_opcode = match access.mem_space { - model::MemorySpace::Local if is_load => "LDL.E".to_string(), - model::MemorySpace::Global if is_load => "LDG.E".to_string(), - model::MemorySpace::Shared if is_load => "LDS.E".to_string(), + model::MemorySpace::Local if is_load => "LDL".to_string(), + model::MemorySpace::Global if is_load => "LDG".to_string(), + model::MemorySpace::Shared if is_load => "LDS".to_string(), // MemorySpace::Texture if is_load => "LDG".to_string(), - model::MemorySpace::Constant if is_load => "LDC.E".to_string(), - model::MemorySpace::Local if is_store => "STL.E".to_string(), - model::MemorySpace::Global if is_store => "STG.E".to_string(), - model::MemorySpace::Shared if is_store => "STS.E".to_string(), + model::MemorySpace::Constant if is_load => "LDC".to_string(), + model::MemorySpace::Local if is_store => "STL".to_string(), + model::MemorySpace::Global if is_store => "STG".to_string(), + model::MemorySpace::Shared if is_store => "STS".to_string(), // MemorySpace::Texture if is_store => "LDG".to_string(), model::MemorySpace::Constant if is_store => { todo!("constant store") @@ -692,6 +692,10 @@ impl TraceGenerator for Tracer { other => panic!("unknown memory space {other:?}"), }; + if access.size >= 64 { + instr_opcode += ".E"; + } + if access.bypass_l1 { instr_opcode += ".CG"; } diff --git a/ptx/Cargo.toml b/ptx/Cargo.toml index 4946d505..cacc90f4 100644 --- a/ptx/Cargo.toml +++ b/ptx/Cargo.toml @@ -29,5 +29,5 @@ bytes = "1" [dev-dependencies] once_cell = "1" snailquote = "0" -pest-test = "0" +# pest-test = "0" diff = { path = "../diff" } diff --git a/ptx/src/parser.rs b/ptx/src/parser.rs index ed44761e..b9166436 100644 --- a/ptx/src/parser.rs +++ b/ptx/src/parser.rs @@ -6,56 +6,16 @@ pub struct Parser; mod tests { use super::{Parser as PTXParser, Rule}; use color_eyre::eyre; - use expression::{Expression as E, Rule as R}; + use expression::{Expression as E, Rule as RL}; use pest::Parser; - // use pest_test::model::Expression; - // use pest_test::{ - // model::{Expression, TestCase}, - // // parser::ParserError, - // TestError, - // }; - - // pub static PTX_PARSER: once_cell::sync::Lazy = - // once_cell::sync::Lazy::new(|| PTXParser::de()); pub mod expression { - // use crate::parser::Rule; - // use colored::{Color, Colorize}; use pest::{ iterators::{Pair, Pairs}, - // iterators::Pair, - Parser, - RuleType, + Parser, RuleType, }; use snailquote::unescape; use std::collections::HashSet; - // fmt::{Display, Result as std::fmt::Result, Write}, - // use thiserror::Error; - - // pub mod parser { - // use pest::{error::Error as PestError, iterators::Pair, Parser, RuleType}; - // use std::marker::PhantomData; - // use thiserror::Error; - - // #[derive(Error, Debug)] - // pub enum ParserError { - // #[error("Error parsing source text")] - // Pest { source: Box> }, - // #[error("Empty parse tree")] - // Empty, - // } - - // pub fn parse>( - // text: &str, - // rule: R, - // _: PhantomData

, - // ) -> Result, ParserError> { - // P::parse(rule, text) - // .map_err(|source| ParserError::Pest { - // source: Box::new(source), - // }) - // .and_then(|mut code_pairs| code_pairs.next().ok_or(ParserError::Empty)) - // } pub mod parser { #[derive(pest_derive::Parser)] @@ -63,7 +23,7 @@ mod tests { pub struct Parser; #[derive(thiserror::Error, Debug)] - pub enum ParseError { + pub enum Error { #[error(transparent)] Parse(#[from] pest::error::Error), #[error("Missing skip depth")] @@ -81,18 +41,6 @@ mod tests { } } - // fn assert_rule(pair: Pair<'_, Rule>, rule: Rule) -> Result, ParseError> { - // if pair.as_rule() == rule { - // Ok(pair) - // } else { - // Err(ParseError::UnexpectedPair( - // pair.as_str()[..20].to_string(), - // pair.as_rule(), - // rule, - // )) - // } - // } - /// Options for building expressions #[derive(Debug, PartialEq, Eq)] pub struct Options @@ -113,130 +61,83 @@ mod tests { } } - // #[derive(Clone, Hash, PartialEq, Eq, PartialOrd, Ord)] - // pub enum TypedExpression<'a, R> { - // /// Terminal expression - // T(R, Option<&'a str>), - // // T(&'a str, Option<&'a str>), - // /// Nonterminal expression - // NT(R, Vec>), - // // NT(&'a str, Vec>), - // Skip { - // depth: usize, - // next: Box>, - // }, - // /// Did not match the rule - // Empty, - // } - + /// An expression is a high-level representation of parsed rules. + /// + /// There exist + /// - `Empty` expressions. + /// - `Skip` expressions. + /// - Terminal (`T`) expressions with rule name and its terminal value. + /// - NonTerminal (`NT`) expressions with rule name and a collection + /// of child expressions. + /// + /// Unfortunately, cannot use &'a str for terminal values because of + /// dequoting. + /// If that could be done in the parser, this could be avoided. #[derive(Clone, Hash, PartialEq, Eq, PartialOrd, Ord)] - pub enum Expression<'a, R> { + // pub enum Expression<'a, R> { + pub enum Expression { /// Terminal expression - T(R, Option<&'a str>), - // T(&'a str, Option<&'a str>), + // T(R, Option<&'a str>), + T(R, Option), /// Nonterminal expression - NT(R, Vec>), - // NT(&'a str, Vec>), + NT(R, Vec>), Skip { depth: usize, - next: Box>, + next: Box>, }, - /// Did not match the rule + /// No tokens parsed Empty, } - // #[derive(Clone, Debug, Hash, PartialEq, Eq, PartialOrd, Ord)] - // pub enum Expression { - // /// Terminal expression - // T(String, Option), - // /// Nonterminal expression - // NT(String, Vec), - // Skip { - // depth: usize, - // next: Box, - // }, - // } - - // #[derive(Clone, Debug, Hash, PartialEq, Eq, PartialOrd, Ord)] - // pub enum Expression { - // Terminal { - // name: String, - // value: Option, - // }, - // NonTerminal { - // name: String, - // children: Vec, - // }, - // Skip { - // depth: usize, - // next: Box, - // }, - // } - - // pest_test::model:: - - // impl Expression { - impl<'a> Expression<'a, String> - // where - // R: RuleType, - { - // pub fn parse(text: &str) -> Result, ParserError> - pub fn parse_expression(pair: Pair<'a, Rule>) -> Result { + // impl<'a> Expression<'a, String> { + impl Expression { + /// Parse an expression from high level syntax + pub fn parse_expression(pair: Pair<'_, parser::Rule>) -> Result { dbg!(&pair.as_rule()); - assert_eq!(pair.as_rule(), Rule::expression); - - // // forward to next expression - // while pair.as_rule() != Rule::expression { - // pair = pair.into_inner().next().unwrap(); - // } + assert_eq!(pair.as_rule(), parser::Rule::expression); let mut inner = pair.into_inner(); - // let skip_depth: usize = - // if inner.peek().map(|pair| pair.as_rule()) == Some(Rule::skip) { - let skip_depth = match inner.peek() { - Some(pair) if pair.as_rule() == Rule::skip => { + Some(pair) if pair.as_rule() == parser::Rule::skip => { let depth_pair = inner .next() .unwrap() .into_inner() .next() - .ok_or_else(|| ParseError::MissingSkipDepth)?; + .ok_or_else(|| parser::Error::MissingSkipDepth)?; // .and_then(|pair| assert_rule(pair, Rule::int))?; depth_pair .as_str() .parse() - .map_err(|source| ParseError::BadSkipDepth { source })? + .map_err(|source| parser::Error::BadSkipDepth { source })? } _ => 0, }; dbg!(&skip_depth); - let pair = inner.next().ok_or_else(|| ParseError::MissingRuleName)?; + let pair = inner.next().ok_or_else(|| parser::Error::MissingRuleName)?; dbg!(&pair); // let pair = assert_rule(pair, Rule::identifier)?; let rule = pair.as_str().to_owned(); - // .(|pair| assert_rule(pair, Rule::identifier)) - // .map(|pair| pair.as_str().to_owned())?; dbg!(&rule); let expr = match inner.next() { None => Self::T(rule, None), Some(pair) => match pair.as_rule() { - Rule::sub_expressions => { + parser::Rule::sub_expressions => { let children: Vec<_> = pair .into_inner() .map(Self::parse_expression) .collect::>()?; Self::NT(rule, children) } - Rule::string => { + parser::Rule::string => { let value = pair.as_str().trim(); - // let value = unescape(s) - // .map_err(|err| ParseError::Unescape(s.to_string(), err))?; + let value = unescape(value) + .map_err(|err| parser::Error::Unescape(value.to_string(), err))?; Self::T(rule, Some(value)) } - Rule::EOI => Expression::Empty, - other => return Err(ParseError::UnexpectedRule(other)), + parser::Rule::EOI => Expression::Empty, + other => return Err(parser::Error::UnexpectedRule(other)), }, }; if skip_depth == 0 { @@ -249,24 +150,21 @@ mod tests { } } - // pub fn parse(text: &'a str) -> Result> { - pub fn parse(text: &'a str) -> Result { - let mut parsed = ExpressionParser::parse(Rule::root, text)?; + pub fn parse(text: &str) -> Result { + let mut parsed = parser::Parser::parse(parser::Rule::root, text)?; if let Some(pair) = parsed.next().and_then(|pair| pair.into_inner().next()) { Self::parse_expression(pair) } else { Ok(Self::Empty) } - // let parsed = ExpressionParser::parse(Rule::root, text)?; - // let options = Options::default(); - // Ok(Self::new(parsed, &options)) } } + /// Wrapper type for rules. #[derive(Clone, Hash, PartialEq, Eq, PartialOrd, Ord)] - pub struct RuleName(R); + pub struct Rule(pub R); - impl std::fmt::Debug for RuleName + impl std::fmt::Debug for Rule where R: RuleType, { @@ -275,7 +173,7 @@ mod tests { } } - impl std::fmt::Display for RuleName + impl std::fmt::Display for Rule where R: RuleType, { @@ -284,11 +182,9 @@ mod tests { } } - impl<'a, R> Expression<'a, RuleName> { - // pub fn try_from_code( - // pub fn new(pair: Pair<'a, R>, options: &Options) -> Result - // pub fn new(pair: Pairs<'a, R>, options: &Options) -> Result - pub fn new(mut pairs: Pairs<'a, R>, options: &Options) -> Self + // impl<'a, R> Expression<'a, Rule> { + impl Expression> { + pub fn new(mut pairs: Pairs<'_, R>, options: &Options) -> Self where R: RuleType, { @@ -299,45 +195,36 @@ mod tests { } } - pub fn from_pair(pair: Pair<'a, R>, options: &Options) -> Self + pub fn from_pair(pair: Pair<'_, R>, options: &Options) -> Self where R: RuleType, { let rule = pair.as_rule(); - // let name = format!("{:?}", pair.as_rule()); - // let name = pair.as_rule().as_str(); // format!("{:?}", pair.as_rule()); let value = pair.as_str(); - // let children: Result, ModelError> = pair let children: Vec<_> = pair .into_inner() .filter(|pair| !options.skip_rules.contains(&pair.as_rule())) .map(|pair| Self::from_pair(pair, options)) .collect(); - // .collect::, _>>()?; - // .collect(); if children.is_empty() { - // Ok(Self::T(name, Some(value))) - Self::T(RuleName(rule), Some(value)) + Self::T(Rule(rule), Some(value.to_string())) } else { - // Ok(Self::NT(name, children)) - Self::NT(RuleName(rule), children) + Self::NT(Rule(rule), children) } + } - // match children { - // Ok(children) if children.is_empty() => Ok(Self::T(name, Some(value))), - // Ok(children) => Ok(Self::NT(name, children)), - // Err(e) => Err(e), - // } + pub fn rule(&self) -> Option + where + R: RuleType, + { + match self { + Self::T(rule, _) => Some(rule.to_string()), + Self::NT(rule, _) => Some(rule.to_string()), + Self::Skip { depth: _, next } => next.rule(), + Self::Empty => None, + } } - // pub fn name(&self) -> &String { - // pub fn name(&self) -> &str { - // match self { - // Self::T(name, _) => name, - // Self::NT(name, _) => name, - // Self::Skip { depth: _, next } => next.name(), - // } - // } // // pub fn skip_depth(&self) -> usize { // match self { @@ -515,7 +402,7 @@ mod tests { } } - impl<'a, R> std::fmt::Display for Expression<'a, R> + impl std::fmt::Display for Expression where R: std::fmt::Display, { @@ -524,7 +411,7 @@ mod tests { } } - impl<'a, R> std::fmt::Debug for Expression<'a, R> + impl std::fmt::Debug for Expression where R: std::fmt::Display, { @@ -534,37 +421,20 @@ mod tests { } } - macro_rules! assert_parses_to { - ($rule:expr, $input:expr, $want:expr) => { - let parsed = PTXParser::parse($rule, $input)?.next().unwrap(); - let skip_rules = std::collections::HashSet::new(); - let have = pest_test::model::Expression::try_from_code(parsed, &skip_rules)?; - dbg!(&have); - let have = have.to_string(); - let want = $want.to_string(); - diff::assert_eq!(have: have, want: want); - }; - } - - fn assert_parses_to_typed( - rule: Rule, - input: &str, - want: expression::Expression>, - ) { - let parsed = PTXParser::parse(rule, input).expect("can parse"); + fn assert_parses_to_typed(rule: Rule, input: &str, want: E>) -> eyre::Result<()> { + let parsed = PTXParser::parse(rule, input)?; let options = expression::Options::default(); let have = E::new(parsed, &options); dbg!(&have); diff::assert_eq!(have: have, want: want); + Ok(()) } - // impl PartialEq> for E<'_, Rule> { - impl PartialEq> for E<'_, expression::RuleName> + impl PartialEq> for E> where R: pest::RuleType, { fn eq(&self, other: &E) -> bool { - dbg!(self, other); match (self, other) { (E::Empty, E::Empty) => true, (E::T(lrule, lvalue), E::T(rrule, rvalue)) => lrule == rrule && lvalue == rvalue, @@ -574,57 +444,26 @@ mod tests { } } - // impl PartialEq for Rule { - impl PartialEq for expression::RuleName + impl PartialEq for RL where R: pest::RuleType, { fn eq(&self, other: &String) -> bool { - dbg!(&self); - dbg!(&other); self.to_string() == other.to_string() - // format!("{:?}", self).eq(other.as_str()) } } - fn assert_parses_to_concise(rule: Rule, input: &str, want: &str) -> eyre::Result<()> { - let parsed = PTXParser::parse(rule, input)?; // .expect("can parse"); + fn assert_parses_to(rule: Rule, input: &str, want: &str) -> eyre::Result<()> { + let parsed = PTXParser::parse(rule, input)?; let options = expression::Options::default(); let have = E::new(parsed, &options); dbg!(&have); - let want = E::parse(want)?; // .expect("can parse expression"); + let want = E::parse(want)?; dbg!(&want); diff::assert_eq!(have: have, want: want); Ok(()) } - // macro_rules! assert_parses_to_typed { - // ($rule:expr, $input:expr, $want:expr) => { - // let parsed = PTXParser::parse($rule, $input)?; // .next().unwrap(); - // // let parsed = parsed.next().unwarp(expression::Expression::Empty)?; - // // .expect(&format!("cannot match rule {:?}", $rule)); - // - // let options = expression::Options::default(); - // let have = E::new(parsed, &options); - // dbg!(&have); - // diff::assert_eq!(have: have, want: $want); - // }; - // } - - // macro_rules! assert_parses_to_new { - // ($rule:expr, $input:expr, $want:expr) => { - // let parsed = PTXParser::parse($rule, $input)?; // .next().unwrap(); - // // let parsed = parsed.next().unwarp(expression::Expression::Empty)?; - // // .expect(&format!("cannot match rule {:?}", $rule)); - // - // let options = expression::Options::default(); - // let have = E::new(parsed, &options); - // let want = E::parse($want).expect("can parse expression"); - // dbg!(&have); - // diff::assert_eq!(have: have, want: $want); - // }; - // } - // #[test] // fn parse_integer_decimal_0() -> eyre::Result<()> { // // let (rule, source, expected) = $value; @@ -676,87 +515,92 @@ mod tests { #[test] fn parse_integer_decimal_7() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "decimal".to_string(), - value: Some("7".to_string()), - }], - }; - assert_parses_to!(Rule::integer, "7", want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "decimal".to_string(), + // value: Some("7".to_string()), + // }], + // }; + let want = r#"(integer (decimal: "7"))"#; + assert_parses_to(Rule::integer, "7", want)?; Ok(()) } #[test] fn parse_integer_decimal_neg_12() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "decimal".to_string(), - value: Some("-12".to_string()), - }], - }; - assert_parses_to!(Rule::integer, "-12", want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "decimal".to_string(), + // value: Some("-12".to_string()), + // }], + // }; + let want = r#"(integer (decimal: "-12"))"#; + assert_parses_to(Rule::integer, "-12", want)?; Ok(()) } #[test] fn parse_integer_decimal_12_u() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "decimal".to_string(), - value: Some("12U".to_string()), - }], - }; - assert_parses_to!(Rule::integer, "12U", want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "decimal".to_string(), + // value: Some("12U".to_string()), + // }], + // }; + let want = r#"(integer (decimal: "12U"))"#; + assert_parses_to(Rule::integer, "12U", want)?; Ok(()) } #[test] fn parse_integer_octal_01110011001() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "octal".to_string(), - value: Some("01110011001".to_string()), - }], - }; + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "octal".to_string(), + // value: Some("01110011001".to_string()), + // }], + // }; - assert_parses_to!(Rule::integer, "01110011001", want); + let want = r#"(integer (octal: "01110011001"))"#; + assert_parses_to(Rule::integer, "01110011001", want)?; Ok(()) } #[test] fn parse_integer_binary_0_b_01110011001() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "binary".to_string(), - value: Some("0b01110011001".to_string()), - }], - }; - - assert_parses_to!(Rule::integer, "0b01110011001", want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "binary".to_string(), + // value: Some("0b01110011001".to_string()), + // }], + // }; + let want = r#"(integer (binary: "0b01110011001"))"#; + assert_parses_to(Rule::integer, "0b01110011001", want)?; Ok(()) } #[test] fn parse_integer_hex_0xaf70d() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "hex".to_string(), - value: Some("0xaf70d".to_string()), - }], - }; + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "hex".to_string(), + // value: Some("0xaf70d".to_string()), + // }], + // }; - assert_parses_to!(Rule::integer, "0xaf70d", want); + let want = r#"(integer (hex: "0xaf70d"))"#; + assert_parses_to(Rule::integer, "0xaf70d", want)?; Ok(()) } @@ -792,89 +636,98 @@ mod tests { Ok(()) } - #[ignore = "old api"] - #[test] - fn parse_instruction_shl_b32_r1_r1_2() -> eyre::Result<()> { - crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "instruction_statement".to_string(), - children: vec![pest_test::model::Expression::NonTerminal { - name: "instruction".to_string(), - children: vec![ - pest_test::model::Expression::NonTerminal { - name: "opcode_spec".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "opcode".to_string(), - value: Some("shl".to_string()), - }], - }, - pest_test::model::Expression::NonTerminal { - name: "operand".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "identifier".to_string(), - value: Some("r1".to_string()), - }], - }, - pest_test::model::Expression::NonTerminal { - name: "operand".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "identifier".to_string(), - value: Some("r1".to_string()), - }], - }, - // Expression::NonTerminal { - // name: "operand".to_string(), - // children: vec![Expression::Terminal { - // name: "literal_operand".to_string(), - // value: Some("r1".to_string()), - // }], - // }, - ], - }], - }; - assert_parses_to!(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want); - Ok(()) - } + // #[ignore = "old api"] + // #[test] + // fn parse_instruction_shl_b32_r1_r1_2_deprecated() -> eyre::Result<()> { + // crate::tests::init_test(); + // let want = pest_test::model::Expression::NonTerminal { + // name: "instruction_statement".to_string(), + // children: vec![pest_test::model::Expression::NonTerminal { + // name: "instruction".to_string(), + // children: vec![ + // pest_test::model::Expression::NonTerminal { + // name: "opcode_spec".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "opcode".to_string(), + // value: Some("shl".to_string()), + // }], + // }, + // pest_test::model::Expression::NonTerminal { + // name: "operand".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "identifier".to_string(), + // value: Some("r1".to_string()), + // }], + // }, + // pest_test::model::Expression::NonTerminal { + // name: "operand".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "identifier".to_string(), + // value: Some("r1".to_string()), + // }], + // }, + // // Expression::NonTerminal { + // // name: "operand".to_string(), + // // children: vec![Expression::Terminal { + // // name: "literal_operand".to_string(), + // // value: Some("r1".to_string()), + // // }], + // // }, + // ], + // }], + // }; + // assert_parses_to!(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want); + // Ok(()) + // } #[test] fn parse_instruction_shl_b32_r1_r1_2_typed() -> eyre::Result<()> { crate::tests::init_test(); - let want = E::NT( - Rule::instruction_statement, + let want: E> = E::NT( + RL(Rule::instruction_statement), vec![E::NT( - Rule::instruction, + RL(Rule::instruction), vec![ E::NT( - Rule::opcode_spec, + RL(Rule::opcode_spec), vec![ - E::T(Rule::opcode, Some("shl")), + E::T(RL(Rule::opcode), Some("shl".to_string())), E::NT( - Rule::option, + RL(Rule::option), vec![E::NT( - Rule::type_spec, - vec![E::T(Rule::scalar_type, Some(".b32"))], + RL(Rule::type_spec), + vec![E::T(RL(Rule::scalar_type), Some(".b32".to_string()))], )], ), ], ), - E::NT(Rule::operand, vec![E::T(Rule::identifier, Some("r1"))]), - E::NT(Rule::operand, vec![E::T(Rule::identifier, Some("r1"))]), E::NT( - Rule::operand, + RL(Rule::operand), + vec![E::T(RL(Rule::identifier), Some("r1".to_string()))], + ), + E::NT( + RL(Rule::operand), + vec![E::T(RL(Rule::identifier), Some("r1".to_string()))], + ), + E::NT( + RL(Rule::operand), vec![E::NT( - Rule::literal_operand, - vec![E::NT(Rule::integer, vec![E::T(Rule::decimal, Some("2"))])], + RL(Rule::literal_operand), + vec![E::NT( + RL(Rule::integer), + vec![E::T(RL(Rule::decimal), Some("2".to_string()))], + )], )], ), ], )], ); - assert_parses_to_typed(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want); + assert_parses_to_typed(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want)?; Ok(()) } #[test] - fn parse_instruction_shl_b32_r1_r1_2_concise() -> eyre::Result<()> { + fn parse_instruction_shl_b32_r1_r1_2() -> eyre::Result<()> { crate::tests::init_test(); let want = r#" (instruction_statement @@ -888,7 +741,51 @@ mod tests { ) ) "#; - assert_parses_to_concise(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want)?; + assert_parses_to(Rule::instruction_statement, "shl.b32 r1, r1, 2;", want)?; + Ok(()) + } + + #[test] + fn parse_instruction_ld_global_b32_r2_array_r1() -> eyre::Result<()> { + crate::tests::init_test(); + let want = r#" +(instruction_statement + (instruction + (opcode_spec + (opcode: "ld") + (option (addressable_spec: ".global")) + (option (type_spec (scalar_type: ".b32"))) + ) + (operand (identifier: "r2")) + (operand (memory_operand + (identifier: "array") + (address_expression (identifier: "r1")) + )) + ) +) + "#; + // this uses a vector operand + // assert_parses_to( + // Rule::array_operand, + // "array[r1]", + // r#"(array_operand + // (identifier: "array") + // (address_expression (identifier: "r1")) + // )"#, + // )?; + // assert_parses_to( + // Rule::memory_operand, + // "array[r1]", + // r#"(memory_operand + // (identifier: "array") + // (address_expression (identifier: "r1")) + // )"#, + // )?; + assert_parses_to( + Rule::instruction_statement, + "ld.global.b32 r2, array[r1];", + want, + )?; Ok(()) } @@ -896,14 +793,14 @@ mod tests { #[test] fn parse_instruction_reg_b32_r1_r2() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "binary".to_string(), - value: Some("0b01110011001".to_string()), - }], - }; - assert_parses_to!(Rule::instruction_statement, ".reg .b32 r1, r2;", want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "binary".to_string(), + // value: Some("0b01110011001".to_string()), + // }], + // }; + // assert_parses_to!(Rule::instruction_statement, ".reg .b32 r1, r2;", want); // assert_parses_to!(Rule::instruction_statement, ".reg .b32 r1, r2;", want); Ok(()) } @@ -921,24 +818,24 @@ mod tests { #[test] fn parse_single_line_comments() -> eyre::Result<()> { crate::tests::init_test(); - let want = pest_test::model::Expression::NonTerminal { - name: "integer".to_string(), - children: vec![pest_test::model::Expression::Terminal { - name: "binary".to_string(), - value: Some("0b01110011001".to_string()), - }], - }; - - let input = " - .reg .b32 r1, r2; - .global .f32 array[N]; - -start: mov.b32 r1, %tid.x; - shl.b32 r1, r1, 2; // shift thread id by 2 bits - ld.global.b32 r2, array[r1]; // thread[tid] gets array[tid] - add.f32 r2, r2, 0.5; // add 1/2"; - - assert_parses_to!(Rule::program, input, want); + // let want = pest_test::model::Expression::NonTerminal { + // name: "integer".to_string(), + // children: vec![pest_test::model::Expression::Terminal { + // name: "binary".to_string(), + // value: Some("0b01110011001".to_string()), + // }], + // }; + // + // let input = " + // .reg .b32 r1, r2; + // .global .f32 array[N]; + // + // start: mov.b32 r1, %tid.x; + // shl.b32 r1, r1, 2; // shift thread id by 2 bits + // ld.global.b32 r2, array[r1]; // thread[tid] gets array[tid] + // add.f32 r2, r2, 0.5; // add 1/2"; + // + // assert_parses_to!(Rule::program, input, want); Ok(()) } diff --git a/ptx/src/ptx.pest b/ptx/src/ptx.pest index a2305693..9473da5d 100644 --- a/ptx/src/ptx.pest +++ b/ptx/src/ptx.pest @@ -58,6 +58,8 @@ predicate = { | ".gtu" | ".neu" | ".cf" | ".sf" | ".nsf" )? } +lohi_option = { ".lo" | ".hi" } + instruction = { opcode_spec ~ operand_list | opcode_spec @@ -71,28 +73,77 @@ operand = { // tex_operand precedes due to "[" ... ~ "," ~ "{" ... "}" "]" // precedes memory_operand, which is just "[" x | x+1 | 1 "]" | tex_operand - // memory_operand precedes tex_operand due to "[" x | x+1 | 1 "]" | memory_operand + // | array_operand // vector precedes due to "{" ... "}" - | "-"? ~ vector_operand + | vector_operand // binary expression of identifiers always precede unary - | "-"? ~ identifier ~ ("+" | "|" | "\\") ~ integer ~ lohi_option? + // | "-"? ~ identifier ~ ("+" | "|" | "\\") ~ integer ~ lohi_option? | ("!" | "-")? ~ identifier } -vector_operand = { - "{" ~ identifier ~ ("," ~ identifier)* ~ "}" + +// Built-in special operand. +builtin_operand = { special_register ~ dimension_modifier | special_register } + + +// Literal operand +// +// Example: +literal_operand = { + // float always precedes because it starts with 0[fF] + // float_exact + // double always precedes integers because it either + // - contains a decimal point (.) + // - starts with 0[dD] + // | double_exact + | integer } + +// Texture operand +// +// Example: +// todo tex_operand = { "[" ~ identifier ~ "," ~ vector_operand ~ "]" } -builtin_operand = { special_register ~ dimension_modifier | special_register } +// Vector operand +// +// Example: +// .reg .v4 .f32 V; +// .reg .f32 a, b, c, d; +// mov.v4.f32 {a,b,c,d}, V; +vector_operand = { + "-"? ~ "{" ~ identifier ~ ("," ~ identifier)* ~ "}" +} + + +// Memory operand +// +// Example: +// ld.shared.u16 r0,[0]; // [immOffset] +// ld.shared.u16 r0,[x]; // [var] +// ld.shared.u16 r0,[x+1]; // [var+immOffset] +// ld.shared.u16 r0,[r1]; // [reg] +// ld.shared.u16 r0,[r1+1]; // [reg+immOffset] +// ld.global.u32 s, a[0]; // var[immOffset] +// ld.global.u32 s, a[N-1]; // var[var+immOffset] +// mov.u32 s, a[1]; // var[immOffset] memory_operand = { - "[" ~ address_expression ~ "]" - // todo: check if the order is brackets is fine - | identifier ~ "[" ~ (twin_operand | address_expression | literal_operand) ~ "]" - | "-" ~ memory_operand + identifier ~ "[" ~ address_expression ~ "]" + | "[" ~ address_expression ~ "]" + // | identifier ~ "[" ~ (twin_operand | address_expression | literal_operand) ~ "]" + // | "-" ~ memory_operand } + +// Array operand +// +// Example: +// ld.global.u32 s, a[0]; +// ld.global.u32 s, a[N-1]; +// mov.u32 s, a[1]; +array_operand = { identifier ~ "[" ~ address_expression ~ "]" } + twin_operand = { // todo: are we missing something here? identifier ~ "+=" ~ identifier ~ lohi_option? @@ -100,10 +151,12 @@ twin_operand = { | identifier ~ "+" ~ identifier ~ lohi_option? } +sign = { "+" | "-" } address_expression = { // precedence because of the plus sign - identifier ~ "+" ~ integer - | identifier ~ lohi_option? + identifier ~ sign ~ integer + // | identifier ~ lohi_option? + | identifier | integer } @@ -355,17 +408,6 @@ scalar_type = { | ".pred" | ".texref" | ".sampleref" | ".surfref" } -lohi_option = { ".lo" | ".hi" } -literal_operand = { - // float always precedes because it starts with 0[fF] - float_exact - // double always precedes integers because it either - // - contains a decimal point (.) - // - starts with 0[dD] - | double_exact - | integer -} - quoted = _{ "\"" ~ string ~ "\"" } string = ${ (!"\"" ~ ANY)* } diff --git a/src/mshr.rs b/src/mshr.rs index 334bb4e0..d040a303 100644 --- a/src/mshr.rs +++ b/src/mshr.rs @@ -125,7 +125,7 @@ impl MSHR for Table { impl Table { #[must_use] pub fn new(num_entries: usize, max_merged: usize) -> Self { - let entries = HashMap::with_capacity(2 * num_entries); + let entries = HashMap::with_capacity(num_entries); Self { num_entries, max_merged, diff --git a/test-apps/test-apps-materialized.yml b/test-apps/test-apps-materialized.yml index 97078f90..75e8aeab 100755 --- a/test-apps/test-apps-materialized.yml +++ b/test-apps/test-apps-materialized.yml @@ -2,7 +2,7 @@ ## ## AUTO GENERATED! DO NOT EDIT ## -## this configuration was materialized from /home/roman/dev/box/test-apps/test-apps.yml on 03/03/2024 16:32:55 +## this configuration was materialized from /home/roman/dev/box/test-apps/test-apps.yml on 22/03/2024 22:07:25 ## config: diff --git a/timings_pie.png b/timings_pie.png new file mode 100644 index 00000000..3d4b0d74 Binary files /dev/null and b/timings_pie.png differ