mirrors/ZLUDA
1
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-04-19 16:04:44 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Recover from and report unknown instructions and directives (#295)

Browse source
This commit is contained in:
Andrzej Janik 2024-11-02 15:57:57 +01:00 • committed by GitHub
parent 3870a96592
commit b4cb3ade63
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
5 changed files with 314 additions and 92 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
llvm_zluda/build.rs
View file

@ -16,6 +16,8 @@ fn main() {
let mut cmake = Config::new(r"../ext/llvm-project/llvm"); let mut cmake = Config::new(r"../ext/llvm-project/llvm");
try_use_ninja(&mut cmake); try_use_ninja(&mut cmake);
cmake cmake
// It's not like we can do anything about the warnings
.define("LLVM_ENABLE_WARNINGS", "OFF")
// For some reason Rust always links to release MSVCRT // For some reason Rust always links to release MSVCRT
.define("CMAKE_MSVC_RUNTIME_LIBRARY", "MultiThreadedDLL") .define("CMAKE_MSVC_RUNTIME_LIBRARY", "MultiThreadedDLL")
.define("LLVM_ENABLE_TERMINFO", "OFF") .define("LLVM_ENABLE_TERMINFO", "OFF")

6
llvm_zluda/src/lib.cpp
View file

@ -1,7 +1,7 @@
#include <llvm-c/Core.h> #include <llvm-c/Core.h>
#include "llvm/IR/IRBuilder.h" #include <llvm/IR/IRBuilder.h>
#include "llvm/IR/Type.h" #include <llvm/IR/Type.h>
#include "llvm/IR/Instructions.h" #include <llvm/IR/Instructions.h>
using namespace llvm; using namespace llvm;

4
ptx_parser/src/ast.rs
View file

@ -1264,9 +1264,9 @@ pub enum SetpCompareFloat {
} }
impl TryFrom<(RawSetpCompareOp, ScalarKind)> for SetpCompareInt { impl TryFrom<(RawSetpCompareOp, ScalarKind)> for SetpCompareInt {
type Error = PtxError; type Error = PtxError<'static>;
fn try_from((value, kind): (RawSetpCompareOp, ScalarKind)) -> Result<Self, PtxError> { fn try_from((value, kind): (RawSetpCompareOp, ScalarKind)) -> Result<Self, PtxError<'static>> {
match (value, kind) { match (value, kind) {
(RawSetpCompareOp::Eq, _) => Ok(SetpCompareInt::Eq), (RawSetpCompareOp::Eq, _) => Ok(SetpCompareInt::Eq),
(RawSetpCompareOp::Ne, _) => Ok(SetpCompareInt::NotEq), (RawSetpCompareOp::Ne, _) => Ok(SetpCompareInt::NotEq),

362
ptx_parser/src/lib.rs
View file

@ -9,7 +9,7 @@ use winnow::ascii::dec_uint;
use winnow::combinator::*; use winnow::combinator::*;
use winnow::error::{ErrMode, ErrorKind}; use winnow::error::{ErrMode, ErrorKind};
use winnow::stream::Accumulate; use winnow::stream::Accumulate;
use winnow::token::any; use winnow::token::{any, take_till};
use winnow::{ use winnow::{
error::{ContextError, ParserError}, error::{ContextError, ParserError},
stream::{Offset, Stream, StreamIsPartial}, stream::{Offset, Stream, StreamIsPartial},
@ -86,14 +86,16 @@ impl VectorPrefix {
} }
struct PtxParserState<'a, 'input> { struct PtxParserState<'a, 'input> {
errors: &'a mut Vec<PtxError>, text: &'input str,
errors: &'a mut Vec<PtxError<'input>>,
function_declarations: function_declarations:
FxHashMap<&'input str, (Vec<(ast::Type, StateSpace)>, Vec<(ast::Type, StateSpace)>)>, FxHashMap<&'input str, (Vec<(ast::Type, StateSpace)>, Vec<(ast::Type, StateSpace)>)>,
} }
impl<'a, 'input> PtxParserState<'a, 'input> { impl<'a, 'input> PtxParserState<'a, 'input> {
fn new(errors: &'a mut Vec<PtxError>) -> Self { fn new(text: &'input str, errors: &'a mut Vec<PtxError<'input>>) -> Self {
Self { Self {
text,
errors, errors,
function_declarations: FxHashMap::default(), function_declarations: FxHashMap::default(),
} }
@ -127,10 +129,11 @@ impl<'a, 'input> Debug for PtxParserState<'a, 'input> {
} }
} }
type PtxParser<'a, 'input> = Stateful<&'a [Token<'input>], PtxParserState<'a, 'input>>; type PtxParser<'a, 'input> =
Stateful<&'a [(Token<'input>, logos::Span)], PtxParserState<'a, 'input>>;
fn ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input str> { fn ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input str> {
any.verify_map(|t| { any.verify_map(|(t, _)| {
if let Token::Ident(text) = t { if let Token::Ident(text) = t {
Some(text) Some(text)
} else if let Some(text) = t.opcode_text() { } else if let Some(text) = t.opcode_text() {
@ -143,7 +146,7 @@ fn ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input str>
} }
fn dot_ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input str> { fn dot_ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input str> {
any.verify_map(|t| { any.verify_map(|(t, _)| {
if let Token::DotIdent(text) = t { if let Token::DotIdent(text) = t {
Some(text) Some(text)
} else { } else {
@ -154,7 +157,7 @@ fn dot_ident<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<&'input
} }
fn num<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<(&'input str, u32, bool)> { fn num<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<(&'input str, u32, bool)> {
any.verify_map(|t| { any.verify_map(|(t, _)| {
Some(match t { Some(match t {
Token::Hex(s) => { Token::Hex(s) => {
if s.ends_with('U') { if s.ends_with('U') {
@ -178,7 +181,7 @@ fn num<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<(&'input str,
} }
fn take_error<'a, 'input: 'a, O, E>( fn take_error<'a, 'input: 'a, O, E>(
mut parser: impl Parser<PtxParser<'a, 'input>, Result<O, (O, PtxError)>, E>, mut parser: impl Parser<PtxParser<'a, 'input>, Result<O, (O, PtxError<'input>)>, E>,
) -> impl Parser<PtxParser<'a, 'input>, O, E> { ) -> impl Parser<PtxParser<'a, 'input>, O, E> {
move |input: &mut PtxParser<'a, 'input>| { move |input: &mut PtxParser<'a, 'input>| {
Ok(match parser.parse_next(input)? { Ok(match parser.parse_next(input)? {
@ -218,7 +221,7 @@ fn int_immediate<'a, 'input>(input: &mut PtxParser<'a, 'input>) -> PResult<ast::
} }
fn f32<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<f32> { fn f32<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<f32> {
take_error(any.verify_map(|t| match t { take_error(any.verify_map(|(t, _)| match t {
Token::F32(f) => Some(match u32::from_str_radix(&f[2..], 16) { Token::F32(f) => Some(match u32::from_str_radix(&f[2..], 16) {
Ok(x) => Ok(f32::from_bits(x)), Ok(x) => Ok(f32::from_bits(x)),
Err(err) => Err((0.0, PtxError::from(err))), Err(err) => Err((0.0, PtxError::from(err))),
@ -229,7 +232,7 @@ fn f32<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<f32> {
} }
fn f64<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<f64> { fn f64<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<f64> {
take_error(any.verify_map(|t| match t { take_error(any.verify_map(|(t, _)| match t {
Token::F64(f) => Some(match u64::from_str_radix(&f[2..], 16) { Token::F64(f) => Some(match u64::from_str_radix(&f[2..], 16) {
Ok(x) => Ok(f64::from_bits(x)), Ok(x) => Ok(f64::from_bits(x)),
Err(err) => Err((0.0, PtxError::from(err))), Err(err) => Err((0.0, PtxError::from(err))),
@ -282,10 +285,9 @@ fn immediate_value<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<as
} }
pub fn parse_module_unchecked<'input>(text: &'input str) -> Option<ast::Module<'input>> { pub fn parse_module_unchecked<'input>(text: &'input str) -> Option<ast::Module<'input>> {
let lexer = Token::lexer(text); let input = lex_with_span(text).ok()?;
let input = lexer.collect::<Result<Vec<_>, _>>().ok()?;
let mut errors = Vec::new(); let mut errors = Vec::new();
let state = PtxParserState::new(&mut errors); let state = PtxParserState::new(text, &mut errors);
let parser = PtxParser { let parser = PtxParser {
state, state,
input: &input[..], input: &input[..],
@ -310,7 +312,7 @@ pub fn parse_module_checked<'input>(
None => break, None => break,
}; };
match maybe_token { match maybe_token {
Ok(token) => tokens.push(token), Ok(token) => tokens.push((token, lexer.span())),
Err(mut err) => { Err(mut err) => {
err.0 = lexer.span(); err.0 = lexer.span();
errors.push(PtxError::from(err)) errors.push(PtxError::from(err))
@ -321,7 +323,7 @@ pub fn parse_module_checked<'input>(
return Err(errors); return Err(errors);
} }
let parse_result = { let parse_result = {
let state = PtxParserState::new(&mut errors); let state = PtxParserState::new(text, &mut errors);
let parser = PtxParser { let parser = PtxParser {
state, state,
input: &tokens[..], input: &tokens[..],
@ -340,6 +342,17 @@ pub fn parse_module_checked<'input>(
} }
} }
fn lex_with_span<'input>(
text: &'input str,
) -> Result<Vec<(Token<'input>, logos::Span)>, TokenError> {
let lexer = Token::lexer(text);
let mut result = Vec::new();
for (token, span) in lexer.spanned() {
result.push((token?, span));
}
Ok(result)
}
fn module<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ast::Module<'input>> { fn module<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ast::Module<'input>> {
( (
version, version,
@ -385,13 +398,29 @@ fn shader_model<'a>(stream: &mut &str) -> PResult<(u32, Option<char>)> {
fn directive<'a, 'input>( fn directive<'a, 'input>(
stream: &mut PtxParser<'a, 'input>, stream: &mut PtxParser<'a, 'input>,
) -> PResult<Option<ast::Directive<'input, ast::ParsedOperand<&'input str>>>> { ) -> PResult<Option<ast::Directive<'input, ast::ParsedOperand<&'input str>>>> {
alt(( with_recovery(
function.map(|(linking, func)| Some(ast::Directive::Method(linking, func))), alt((
file.map(|_| None), // When adding a new variant here remember to add its first token into recovery parser down below
section.map(|_| None), function.map(|(linking, func)| Some(ast::Directive::Method(linking, func))),
(module_variable, Token::Semicolon) file.map(|_| None),
.map(|((linking, var), _)| Some(ast::Directive::Variable(linking, var))), section.map(|_| None),
)) (module_variable, Token::Semicolon)
.map(|((linking, var), _)| Some(ast::Directive::Variable(linking, var))),
)),
take_till(1.., |(token, _)| match token {
// visibility
Token::DotExtern | Token::DotVisible | Token::DotWeak
// methods
| Token::DotFunc | Token::DotEntry
// module variables
| Token::DotGlobal | Token::DotConst | Token::DotShared
// other sections
| Token::DotFile | Token::DotSection => true,
_ => false,
}),
PtxError::UnrecognizedDirective,
)
.map(Option::flatten)
.parse_next(stream) .parse_next(stream)
} }
@ -487,9 +516,9 @@ fn linking_directives<'a, 'input>(
repeat( repeat(
0.., 0..,
dispatch! { any; dispatch! { any;
Token::DotExtern => empty.value(ast::LinkingDirective::EXTERN), (Token::DotExtern, _) => empty.value(ast::LinkingDirective::EXTERN),
Token::DotVisible => empty.value(ast::LinkingDirective::VISIBLE), (Token::DotVisible, _) => empty.value(ast::LinkingDirective::VISIBLE),
Token::DotWeak => empty.value(ast::LinkingDirective::WEAK), (Token::DotWeak, _) => empty.value(ast::LinkingDirective::WEAK),
_ => fail _ => fail
}, },
) )
@ -501,10 +530,10 @@ fn tuning_directive<'a, 'input>(
stream: &mut PtxParser<'a, 'input>, stream: &mut PtxParser<'a, 'input>,
) -> PResult<ast::TuningDirective> { ) -> PResult<ast::TuningDirective> {
dispatch! {any; dispatch! {any;
Token::DotMaxnreg => u32.map(ast::TuningDirective::MaxNReg), (Token::DotMaxnreg, _) => u32.map(ast::TuningDirective::MaxNReg),
Token::DotMaxntid => tuple1to3_u32.map(|(nx, ny, nz)| ast::TuningDirective::MaxNtid(nx, ny, nz)), (Token::DotMaxntid, _) => tuple1to3_u32.map(|(nx, ny, nz)| ast::TuningDirective::MaxNtid(nx, ny, nz)),
Token::DotReqntid => tuple1to3_u32.map(|(nx, ny, nz)| ast::TuningDirective::ReqNtid(nx, ny, nz)), (Token::DotReqntid, _) => tuple1to3_u32.map(|(nx, ny, nz)| ast::TuningDirective::ReqNtid(nx, ny, nz)),
Token::DotMinnctapersm => u32.map(ast::TuningDirective::MinNCtaPerSm), (Token::DotMinnctapersm, _) => u32.map(ast::TuningDirective::MinNCtaPerSm),
_ => fail _ => fail
} }
.parse_next(stream) .parse_next(stream)
@ -514,10 +543,10 @@ fn method_declaration<'a, 'input>(
stream: &mut PtxParser<'a, 'input>, stream: &mut PtxParser<'a, 'input>,
) -> PResult<ast::MethodDeclaration<'input, &'input str>> { ) -> PResult<ast::MethodDeclaration<'input, &'input str>> {
dispatch! {any; dispatch! {any;
Token::DotEntry => (ident, kernel_arguments).map(|(name, input_arguments)| ast::MethodDeclaration{ (Token::DotEntry, _) => (ident, kernel_arguments).map(|(name, input_arguments)| ast::MethodDeclaration{
return_arguments: Vec::new(), name: ast::MethodName::Kernel(name), input_arguments, shared_mem: None return_arguments: Vec::new(), name: ast::MethodName::Kernel(name), input_arguments, shared_mem: None
}), }),
Token::DotFunc => (opt(fn_arguments), ident, fn_arguments).map(|(return_arguments, name,input_arguments)| { (Token::DotFunc, _) => (opt(fn_arguments), ident, fn_arguments).map(|(return_arguments, name,input_arguments)| {
let return_arguments = return_arguments.unwrap_or_else(|| Vec::new()); let return_arguments = return_arguments.unwrap_or_else(|| Vec::new());
let name = ast::MethodName::Func(name); let name = ast::MethodName::Func(name);
ast::MethodDeclaration{ return_arguments, name, input_arguments, shared_mem: None } ast::MethodDeclaration{ return_arguments, name, input_arguments, shared_mem: None }
@ -557,8 +586,8 @@ fn kernel_input<'a, 'input>(
fn fn_input<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ast::Variable<&'input str>> { fn fn_input<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ast::Variable<&'input str>> {
dispatch! { any; dispatch! { any;
Token::DotParam => method_parameter(StateSpace::Param), (Token::DotParam, _) => method_parameter(StateSpace::Param),
Token::DotReg => method_parameter(StateSpace::Reg), (Token::DotReg, _) => method_parameter(StateSpace::Reg),
_ => fail _ => fail
} }
.parse_next(stream) .parse_next(stream)
@ -606,8 +635,8 @@ fn function_body<'a, 'input>(
stream: &mut PtxParser<'a, 'input>, stream: &mut PtxParser<'a, 'input>,
) -> PResult<Option<Vec<ast::Statement<ParsedOperandStr<'input>>>>> { ) -> PResult<Option<Vec<ast::Statement<ParsedOperandStr<'input>>>>> {
dispatch! {any; dispatch! {any;
Token::LBrace => terminated(repeat_without_none(statement), Token::RBrace).map(Some), (Token::LBrace, _) => terminated(repeat_without_none(statement), Token::RBrace).map(Some),
Token::Semicolon => empty.map(|_| None), (Token::Semicolon, _) => empty.map(|_| None),
_ => fail _ => fail
} }
.parse_next(stream) .parse_next(stream)
@ -616,22 +645,122 @@ fn function_body<'a, 'input>(
fn statement<'a, 'input>( fn statement<'a, 'input>(
stream: &mut PtxParser<'a, 'input>, stream: &mut PtxParser<'a, 'input>,
) -> PResult<Option<Statement<ParsedOperandStr<'input>>>> { ) -> PResult<Option<Statement<ParsedOperandStr<'input>>>> {
alt(( with_recovery(
label.map(Some), alt((
debug_directive.map(|_| None), label.map(Some),
terminated( debug_directive.map(|_| None),
method_space terminated(
.flat_map(|space| multi_variable(false, space)) method_space
.map(|var| Some(Statement::Variable(var))), .flat_map(|space| multi_variable(false, space))
Token::Semicolon, .map(|var| Some(Statement::Variable(var))),
Token::Semicolon,
),
predicated_instruction.map(Some),
pragma.map(|_| None),
block_statement.map(Some),
)),
take_till_inclusive(
|(t, _)| *t == Token::RBrace,
|(t, _)| match t {
Token::Semicolon | Token::Colon => true,
_ => false,
},
), ),
predicated_instruction.map(Some), PtxError::UnrecognizedStatement,
pragma.map(|_| None), )
block_statement.map(Some), .map(Option::flatten)
))
.parse_next(stream) .parse_next(stream)
} }
fn take_till_inclusive<I: Stream, E: ParserError<I>>(
backtrack_token: impl Fn(&I::Token) -> bool,
end_token: impl Fn(&I::Token) -> bool,
) -> impl Parser<I, <I as Stream>::Slice, E> {
fn get_offset<I: Stream>(
input: &mut I,
backtrack_token: &impl Fn(&I::Token) -> bool,
end_token: &impl Fn(&I::Token) -> bool,
should_backtrack: &mut bool,
) -> usize {
*should_backtrack = false;
let mut hit = false;
for (offset, token) in input.iter_offsets() {
if hit {
return offset;
} else {
if backtrack_token(&token) {
*should_backtrack = true;
return offset;
}
if end_token(&token) {
hit = true;
}
}
}
input.eof_offset()
}
move |stream: &mut I| {
let mut should_backtrack = false;
let offset = get_offset(stream, &backtrack_token, &end_token, &mut should_backtrack);
let result = stream.next_slice(offset);
if should_backtrack {
Err(ErrMode::from_error_kind(
stream,
winnow::error::ErrorKind::Token,
))
} else {
Ok(result)
}
}
}
/*
pub fn take_till_or_backtrack_eof<Set, Input, Error>(
set: Set,
) -> impl Parser<Input, <Input as Stream>::Slice, Error>
where
Input: StreamIsPartial + Stream,
Set: winnow::stream::ContainsToken<<Input as Stream>::Token>,
Error: ParserError<Input>,
{
move |stream: &mut Input| {
if stream.eof_offset() == 0 {
return ;
}
take_till(0.., set)
}
}
*/
fn with_recovery<'a, 'input: 'a, T>(
mut parser: impl Parser<PtxParser<'a, 'input>, T, ContextError>,
mut recovery: impl Parser<PtxParser<'a, 'input>, &'a [(Token<'input>, logos::Span)], ContextError>,
mut error: impl FnMut(Option<&'input str>) -> PtxError<'input>,
) -> impl Parser<PtxParser<'a, 'input>, Option<T>, ContextError> {
move |stream: &mut PtxParser<'a, 'input>| {
let input_start = stream.input.first().map(|(_, s)| s).cloned();
let stream_start = stream.checkpoint();
match parser.parse_next(stream) {
Ok(value) => Ok(Some(value)),
Err(ErrMode::Backtrack(_)) => {
stream.reset(&stream_start);
let tokens = recovery.parse_next(stream)?;
let range = match input_start {
Some(start) => {
Some(&stream.state.text[start.start..tokens.last().unwrap().1.end])
}
// We could handle `(Some(start), None)``, but this whole error recovery is to
// recover from unknown instructions, so we don't care about early end of stream
_ => None,
};
stream.state.errors.push(error(range));
Ok(None)
}
Err(err) => Err(err),
}
}
}
fn pragma<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<()> { fn pragma<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<()> {
(Token::DotPragma, Token::String, Token::Semicolon) (Token::DotPragma, Token::String, Token::Semicolon)
.void() .void()
@ -746,7 +875,7 @@ fn array_initializer<'a, 'input: 'a>(
} }
delimited( delimited(
Token::LBrace, Token::LBrace,
separated( separated::<_, (), (), _, _, _, _>(
0..=array_dimensions[0] as usize, 0..=array_dimensions[0] as usize,
single_value_append(&mut result, type_), single_value_append(&mut result, type_),
Token::Comma, Token::Comma,
@ -935,7 +1064,7 @@ fn vector_prefix<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<Opti
} }
fn scalar_type<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ScalarType> { fn scalar_type<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<ScalarType> {
any.verify_map(|t| { any.verify_map(|(t, _)| {
Some(match t { Some(match t {
Token::DotS8 => ScalarType::S8, Token::DotS8 => ScalarType::S8,
Token::DotS16 => ScalarType::S16, Token::DotS16 => ScalarType::S16,
@ -1001,8 +1130,8 @@ fn debug_directive<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<()
ident_literal("function_name"), ident_literal("function_name"),
ident, ident,
dispatch! { any; dispatch! { any;
Token::Comma => (ident_literal("inlined_at"), u32, u32, u32).void(), (Token::Comma, _) => (ident_literal("inlined_at"), u32, u32, u32).void(),
Token::Plus => (u32, Token::Comma, ident_literal("inlined_at"), u32, u32, u32).void(), (Token::Plus, _) => (u32, Token::Comma, ident_literal("inlined_at"), u32, u32, u32).void(),
_ => fail _ => fail
}, },
)), )),
@ -1033,13 +1162,14 @@ fn repeat_without_none<Input: Stream, Output, Error: ParserError<Input>>(
fn ident_literal< fn ident_literal<
'a, 'a,
'input, 'input,
I: Stream<Token = Token<'input>> + StreamIsPartial, X,
I: Stream<Token = (Token<'input>, X)> + StreamIsPartial,
E: ParserError<I>, E: ParserError<I>,
>( >(
s: &'input str, s: &'input str,
) -> impl Parser<I, (), E> + 'input { ) -> impl Parser<I, (), E> + 'input {
move |stream: &mut I| { move |stream: &mut I| {
any.verify(|t| matches!(t, Token::Ident(text) if *text == s)) any.verify(|(t, _)| matches!(t, Token::Ident(text) if *text == s))
.void() .void()
.parse_next(stream) .parse_next(stream)
} }
@ -1086,8 +1216,8 @@ impl<Ident> ast::ParsedOperand<Ident> {
let (_, r1, _, r2) = (Token::LBrace, ident, Token::Comma, ident).parse_next(stream)?; let (_, r1, _, r2) = (Token::LBrace, ident, Token::Comma, ident).parse_next(stream)?;
// TODO: parse .v8 literals // TODO: parse .v8 literals
dispatch! {any; dispatch! {any;
Token::RBrace => empty.map(|_| vec![r1, r2]), (Token::RBrace, _) => empty.map(|_| vec![r1, r2]),
Token::Comma => (ident, Token::Comma, ident, Token::RBrace).map(|(r3, _, r4, _)| vec![r1, r2, r3, r4]), (Token::Comma, _) => (ident, Token::Comma, ident, Token::RBrace).map(|(r3, _, r4, _)| vec![r1, r2, r3, r4]),
_ => fail _ => fail
} }
.parse_next(stream) .parse_next(stream)
@ -1102,7 +1232,7 @@ impl<Ident> ast::ParsedOperand<Ident> {
} }
#[derive(Debug, thiserror::Error)] #[derive(Debug, thiserror::Error)]
pub enum PtxError { pub enum PtxError<'input> {
#[error("{source}")] #[error("{source}")]
ParseInt { ParseInt {
#[from] #[from]
@ -1146,10 +1276,10 @@ pub enum PtxError {
ArrayInitalizer, ArrayInitalizer,
#[error("")] #[error("")]
NonExternPointer, NonExternPointer,
#[error("{start}:{end}")] #[error("{0:?}")]
UnrecognizedStatement { start: usize, end: usize }, UnrecognizedStatement(Option<&'input str>),
#[error("{start}:{end}")] #[error("{0:?}")]
UnrecognizedDirective { start: usize, end: usize }, UnrecognizedDirective(Option<&'input str>),
} }
#[derive(Debug)] #[derive(Debug)]
@ -1244,11 +1374,11 @@ impl<'a, T> StreamIsPartial for ReverseStream<'a, T> {
} }
} }
impl<'input, I: Stream<Token = Self> + StreamIsPartial, E: ParserError<I>> Parser<I, Self, E> impl<'input, X, I: Stream<Token = (Self, X)> + StreamIsPartial, E: ParserError<I>>
for Token<'input> Parser<I, (Self, X), E> for Token<'input>
{ {
fn parse_next(&mut self, input: &mut I) -> PResult<Self, E> { fn parse_next(&mut self, input: &mut I) -> PResult<(Self, X), E> {
any.verify(|t| t == self).parse_next(input) any.verify(|(t, _)| t == self).parse_next(input)
} }
} }
@ -1257,7 +1387,7 @@ fn bra<'a, 'input>(
) -> PResult<ast::Instruction<ParsedOperandStr<'input>>> { ) -> PResult<ast::Instruction<ParsedOperandStr<'input>>> {
preceded( preceded(
opt(Token::DotUni), opt(Token::DotUni),
any.verify_map(|t| match t { any.verify_map(|(t, _)| match t {
Token::Ident(ident) => Some(ast::Instruction::Bra { Token::Ident(ident) => Some(ast::Instruction::Bra {
arguments: BraArgs { src: ident }, arguments: BraArgs { src: ident },
}), }),
@ -3224,21 +3354,27 @@ derive_parser!(
#[cfg(test)] #[cfg(test)]
mod tests { mod tests {
use crate::parse_module_checked;
use crate::PtxError;
use super::target; use super::target;
use super::PtxParserState; use super::PtxParserState;
use super::Token; use super::Token;
use logos::Logos; use logos::Logos;
use logos::Span;
use winnow::prelude::*; use winnow::prelude::*;
#[test] #[test]
fn sm_11() { fn sm_11() {
let tokens = Token::lexer(".target sm_11") let text = ".target sm_11";
let tokens = Token::lexer(text)
.map(|t| t.map(|t| (t, Span::default())))
.collect::<Result<Vec<_>, _>>() .collect::<Result<Vec<_>, _>>()
.unwrap(); .unwrap();
let mut errors = Vec::new(); let mut errors = Vec::new();
let stream = super::PtxParser { let stream = super::PtxParser {
input: &tokens[..], input: &tokens[..],
state: PtxParserState::new(&mut errors), state: PtxParserState::new(text, &mut errors),
}; };
assert_eq!(target.parse(stream).unwrap(), (11, None)); assert_eq!(target.parse(stream).unwrap(), (11, None));
assert_eq!(errors.len(), 0); assert_eq!(errors.len(), 0);
@ -3246,13 +3382,15 @@ mod tests {
#[test] #[test]
fn sm_90a() { fn sm_90a() {
let tokens = Token::lexer(".target sm_90a") let text = ".target sm_90a";
let tokens = Token::lexer(text)
.map(|t| t.map(|t| (t, Span::default())))
.collect::<Result<Vec<_>, _>>() .collect::<Result<Vec<_>, _>>()
.unwrap(); .unwrap();
let mut errors = Vec::new(); let mut errors = Vec::new();
let stream = super::PtxParser { let stream = super::PtxParser {
input: &tokens[..], input: &tokens[..],
state: PtxParserState::new(&mut errors), state: PtxParserState::new(text, &mut errors),
}; };
assert_eq!(target.parse(stream).unwrap(), (90, Some('a'))); assert_eq!(target.parse(stream).unwrap(), (90, Some('a')));
assert_eq!(errors.len(), 0); assert_eq!(errors.len(), 0);
@ -3260,15 +3398,97 @@ mod tests {
#[test] #[test]
fn sm_90ab() { fn sm_90ab() {
let tokens = Token::lexer(".target sm_90ab") let text = ".target sm_90ab";
let tokens = Token::lexer(text)
.map(|t| t.map(|t| (t, Span::default())))
.collect::<Result<Vec<_>, _>>() .collect::<Result<Vec<_>, _>>()
.unwrap(); .unwrap();
let mut errors = Vec::new(); let mut errors = Vec::new();
let stream = super::PtxParser { let stream = super::PtxParser {
input: &tokens[..], input: &tokens[..],
state: PtxParserState::new(&mut errors), state: PtxParserState::new(text, &mut errors),
}; };
assert!(target.parse(stream).is_err()); assert!(target.parse(stream).is_err());
assert_eq!(errors.len(), 0); assert_eq!(errors.len(), 0);
} }
#[test]
fn report_unknown_intruction() {
let text = "
.version 6.5
.target sm_30
.address_size 64
.visible .entry add(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .u64 temp;
.reg .u64 temp2;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.u64 temp, [in_addr];
unknown_op1.asdf foobar;
add.u64 temp2, temp, 1;
unknown_op2 temp2, temp;
st.u64 [out_addr], temp2;
ret;
}";
let errors = parse_module_checked(text).err().unwrap();
assert_eq!(errors.len(), 2);
assert!(matches!(
errors[0],
PtxError::UnrecognizedStatement(Some("unknown_op1.asdf foobar;"))
));
assert!(matches!(
errors[1],
PtxError::UnrecognizedStatement(Some("unknown_op2 temp2, temp;"))
));
}
#[test]
fn report_unknown_directive() {
let text = "
.version 6.5
.target sm_30
.address_size 64
.broken_directive_fail; 34; {
.visible .entry add(
.param .u64 input,
.param .u64 output
)
{
.reg .u64 in_addr;
.reg .u64 out_addr;
.reg .u64 temp;
.reg .u64 temp2;
ld.param.u64 in_addr, [input];
ld.param.u64 out_addr, [output];
ld.u64 temp, [in_addr];
add.u64 temp2, temp, 1;
st.u64 [out_addr], temp2;
ret;
}
section foobar }";
let errors = parse_module_checked(text).err().unwrap();
assert_eq!(errors.len(), 2);
assert!(matches!(
errors[0],
PtxError::UnrecognizedDirective(Some(".broken_directive_fail; 34; {"))
));
assert!(matches!(
errors[1],
PtxError::UnrecognizedDirective(Some("section foobar }"))
));
}
} }

32
ptx_parser_macros/src/lib.rs
View file

@ -479,7 +479,7 @@ fn emit_parse_function(
Some(selection_keys) => { Some(selection_keys) => {
let selection_keys = selection_keys.iter().map(|k| k.dot_capitalized()); let selection_keys = selection_keys.iter().map(|k| k.dot_capitalized());
quote! { quote! {
else if false #(|| modifiers.contains(& #type_name :: #selection_keys))* { else if false #(|| modifiers.iter().any(|(t, _)| *t == #type_name :: #selection_keys))* {
#def_parser #def_parser
} }
} }
@ -508,7 +508,7 @@ fn emit_parse_function(
quote! { quote! {
#opcode_variant => { #opcode_variant => {
let modifers_start = stream.checkpoint(); let modifers_start = stream.checkpoint();
let modifiers = take_while(0.., Token::modifier).parse_next(stream)?; let modifiers = take_while(0.., |(t,_)| Token::modifier(t)).parse_next(stream)?;
#selectors #selectors
} }
} }
@ -556,7 +556,7 @@ fn emit_parse_function(
use winnow::Parser; use winnow::Parser;
use winnow::token::*; use winnow::token::*;
use winnow::combinator::*; use winnow::combinator::*;
let opcode = any.parse_next(stream)?; let opcode = any.parse_next(stream)?.0;
let modifiers_start = stream.checkpoint(); let modifiers_start = stream.checkpoint();
Ok(match opcode { Ok(match opcode {
#( #(
@ -594,11 +594,11 @@ fn emit_definition_parser(
let variant = value.dot_capitalized(); let variant = value.dot_capitalized();
if *optional { if *optional {
quote! { quote! {
#arg_name = opt(any.verify(|t| *t == #token_type :: #variant)).parse_next(&mut stream)?.is_some(); #arg_name = opt(any.verify(|(t, _)| *t == #token_type :: #variant)).parse_next(&mut stream)?.is_some();
} }
} else { } else {
quote! { quote! {
any.verify(|t| *t == #token_type :: #variant).parse_next(&mut stream)?; any.verify(|(t, _)| *t == #token_type :: #variant).parse_next(&mut stream)?;
} }
} }
} }
@ -607,7 +607,7 @@ fn emit_definition_parser(
let variant = value.dot_capitalized(); let variant = value.dot_capitalized();
let parsed_variant = value.variant_capitalized(); let parsed_variant = value.variant_capitalized();
quote! { quote! {
any.verify(|t| *t == #token_type :: #variant).parse_next(&mut stream)?; any.verify(|(t, _)| *t == #token_type :: #variant).parse_next(&mut stream)?;
#variable = #type_ :: #parsed_variant; #variable = #type_ :: #parsed_variant;
} }
} }
@ -623,7 +623,7 @@ fn emit_definition_parser(
}); });
if *optional { if *optional {
quote! { quote! {
#arg_name = opt(any.verify_map(|tok| { #arg_name = opt(any.verify_map(|(tok, _)| {
Some(match tok { Some(match tok {
#(#variants)* #(#variants)*
_ => return None _ => return None
@ -632,7 +632,7 @@ fn emit_definition_parser(
} }
} else { } else {
quote! { quote! {
#arg_name = any.verify_map(|tok| { #arg_name = any.verify_map(|(tok, _)| {
Some(match tok { Some(match tok {
#(#variants)* #(#variants)*
_ => return None _ => return None
@ -761,11 +761,11 @@ fn emit_definition_parser(
let comma = if idx == 0 || arg.pre_pipe { let comma = if idx == 0 || arg.pre_pipe {
quote! { empty } quote! { empty }
} else { } else {
quote! { any.verify(|t| *t == #token_type::Comma).void() } quote! { any.verify(|(t, _)| *t == #token_type::Comma).void() }
}; };
let pre_bracket = if arg.pre_bracket { let pre_bracket = if arg.pre_bracket {
quote! { quote! {
any.verify(|t| *t == #token_type::LBracket).void() any.verify(|(t, _)| *t == #token_type::LBracket).void()
} }
} else { } else {
quote! { quote! {
@ -774,7 +774,7 @@ fn emit_definition_parser(
}; };
let pre_pipe = if arg.pre_pipe { let pre_pipe = if arg.pre_pipe {
quote! { quote! {
any.verify(|t| *t == #token_type::Pipe).void() any.verify(|(t, _)| *t == #token_type::Pipe).void()
} }
} else { } else {
quote! { quote! {
@ -783,7 +783,7 @@ fn emit_definition_parser(
}; };
let can_be_negated = if arg.can_be_negated { let can_be_negated = if arg.can_be_negated {
quote! { quote! {
opt(any.verify(|t| *t == #token_type::Not)).map(|o| o.is_some()) opt(any.verify(|(t, _)| *t == #token_type::Not)).map(|o| o.is_some())
} }
} else { } else {
quote! { quote! {
@ -797,7 +797,7 @@ fn emit_definition_parser(
}; };
let post_bracket = if arg.post_bracket { let post_bracket = if arg.post_bracket {
quote! { quote! {
any.verify(|t| *t == #token_type::RBracket).void() any.verify(|(t, _)| *t == #token_type::RBracket).void()
} }
} else { } else {
quote! { quote! {
@ -806,7 +806,7 @@ fn emit_definition_parser(
}; };
let unified = if arg.unified { let unified = if arg.unified {
quote! { quote! {
opt(any.verify(|t| *t == #token_type::DotUnified).void()).map(|u| u.is_some()) opt(any.verify(|(t, _)| *t == #token_type::DotUnified).void()).map(|u| u.is_some())
} }
} else { } else {
quote! { quote! {
@ -854,8 +854,8 @@ fn emit_definition_parser(
{ {
let mut stream = ReverseStream(modifiers); let mut stream = ReverseStream(modifiers);
#(#ordered_parse)* #(#ordered_parse)*
let mut stream: &[#token_type] = stream.0; let mut stream: &[_] = stream.0;
for token in stream.iter().copied() { for (token, _) in stream.iter().cloned() {
match token { match token {
#(#unordered_parse)* #(#unordered_parse)*
_ => #return_error_ref _ => #return_error_ref