diff options
author | Andrzej Janik <[email protected]> | 2024-09-04 15:32:12 +0200 |
---|---|---|
committer | Andrzej Janik <[email protected]> | 2024-09-04 15:32:12 +0200 |
commit | 061312cf8f3762a2ca07b938748a32aa8a1159b1 (patch) | |
tree | 42ed1179594a362f417266d08842aa186028204a | |
parent | aa98ab9e03c37094d745429a8114ee071676f7a7 (diff) | |
download | ZLUDA-parser_rewrite.tar.gz ZLUDA-parser_rewrite.zip |
Document wtf is going on with parsing macrosparser_rewrite
-rw-r--r-- | ptx_parser/src/ast.rs | 15 | ||||
-rw-r--r-- | ptx_parser/src/lib.rs | 69 | ||||
-rw-r--r-- | ptx_parser_macros/src/lib.rs | 2 |
3 files changed, 77 insertions, 9 deletions
diff --git a/ptx_parser/src/ast.rs b/ptx_parser/src/ast.rs index ad44ab7..d0dc303 100644 --- a/ptx_parser/src/ast.rs +++ b/ptx_parser/src/ast.rs @@ -13,6 +13,21 @@ pub enum Statement<P: Operand> { Block(Vec<Statement<P>>),
}
+// We define the instruction enum through the macro instead of normally, because we have some of how
+// we use this type in the compilee. Each instruction can be logically split into two parts:
+// properties that define instruction semantics (e.g. is memory load volatile?) that don't change
+// during compilation and arguments (e.g. memory load source and destination) that evolve during
+// compilation. To support compilation passes we need to be able to visit (and change) every
+// argument in a generic way. This macro has visibility over all the fields. Consequently, we use it
+// to generate visitor functions. There re three functions to support three different semantics:
+// visit-by-ref, visit-by-mutable-ref, visit-and-map. In a previous version of the compiler it was
+// done by hand and was very limiting (we supported only visit-and-map).
+// The visitor must implement appropriate visitor trait defined below this macro. For convenience,
+// we implemented visitors for some corresponding FnMut(...) types.
+// Properties in this macro are used to encode information about the instruction arguments (what
+// Rust type is used for it post-parsing, what PTX type does it expect, what PTX address space does
+// it expect, etc.).
+// This information is then available to a visitor.
ptx_parser_macros::generate_instruction_type!(
pub enum Instruction<T: Operand> {
Mov {
diff --git a/ptx_parser/src/lib.rs b/ptx_parser/src/lib.rs index ed2cf2a..f842ace 100644 --- a/ptx_parser/src/lib.rs +++ b/ptx_parser/src/lib.rs @@ -1335,14 +1335,6 @@ fn empty_call<'input>( } } -// Modifiers are turned into arguments to the blocks, with type: -// * If it is an alternative: -// * If it is mandatory then its type is Foo (as defined by the relevant rule) -// * If it is optional then its type is Option<Foo> -// * Otherwise: -// * If it is mandatory then it is skipped -// * If it is optional then its type is `bool` - type ParsedOperandStr<'input> = ast::ParsedOperand<&'input str>; #[derive(Clone, PartialEq, Default, Debug, Display)] @@ -1351,6 +1343,67 @@ pub struct TokenError(std::ops::Range<usize>); impl std::error::Error for TokenError {} +// This macro is responsible for generating parser code for instruction parser. +// Instruction parsing is by far the most complex part of parsing PTX code: +// * There are tens of instruction kinds, each with slightly different parsing rules +// * After parsing, each instruction needs to do some early validation and generate a specific, +// strongly-typed object. We want strong-typing because we have a single PTX parser frontend, but +// there can be multiple different code emitter backends +// * Most importantly, instruction modifiers can come in aby order, so e.g. both +// `ld.relaxed.global.u32 a, b` and `ld.global.relaxed.u32 a, b` are equally valid. This makes +// classic parsing generators fail: if we tried to generate parsing rules that cover every possible +// ordering we'd need thousands of rules. This is not a purely theoretical problem. NVCC and Clang +// will always emit modifiers in the correct order, but people who write inline assembly usually +// get it wrong (even first party developers) +// +// This macro exists purely to generate repetitive code for parsing each instruction. It is +// _not_ self-contained and is _not_ general-purpose: it relies on certain types and functions from +// the enclosing module +// +// derive_parser!(...) input is split into three parts: +// * Token type definition +// * Partial enums +// * Parsing definitions +// +// Token type definition: +// This is the enum type that will be usesby the instruction parser. For every instruction and +// modifier, derive_parser!(...) will add appropriate variant into this type. So e.g. if there is a +// rule for for `bar.sync` then those two variants wil be appended to the Token enum: +// #[token("bar")] Bar, +// #[token(".sync")] DotSync, +// +// Partial enums: +// With proper annotations, derive_parser!(...) parsing definitions are able to interpret +// instruction modifiers as variants of a single enum type. So e.g. for definitions `ld.u32` and +// `ld.u64` the macro can generate `enum ScalarType { U32, U64 }`. The problem is that for some +// (but not all) of those generated enum types we want to add some attributes and additional +// variants. In order to do so, you need to define this enum and derive_parser!(...) will append to +// the type instead of creating a new type. This is sort of replacement for partial classes known +// from C# +// +// Parsing definitions: +// Parsing definitions consist of a list of patterns and rules: +// * Pattern consists of: +// * Opcode: `ld` +// * Modifiers, always start with a dot: `.global`, `.relaxed`. Optionals are enclosed in braces +// * Arguments: `a`, `b`. Optionals are enclosed in braces +// * Code block: => { <code expression> }. Code blocks implictly take all modifiers ansd arguments +// as parameters. All modifiers and arguments are passed to the code block: +// * If it is an alternative (as defined in rules list later): +// * If it is mandatory then its type is Foo (as defined by the relevant rule) +// * If it is optional then its type is Option<Foo> +// * Otherwise: +// * If it is mandatory then it is skipped +// * If it is optional then its type is `bool` +// * List of rules. They are associated with the preceding patterns (until different opcode or +// different rules). Rules are used to resolve modifiers. There are two types of rules: +// * Normal rule: `.foobar: FoobarEnum => { .a, .b, .c }`. This means that instead of `.foobar` we +// expecte one of `.a`, `.b`, `.c` and will emit value FoobarEnum::DotA, FoobarEnum::DotB, +// FoobarEnum::DotC appropriately +// * Type-only rule: `FoobarEnum => { .a, .b, .c }` this means that all the occurences of `.a` will +// emit FoobarEnum::DotA to the code block. This helps to avoid copy-paste errors +// Additionally, you can opt out from the usual parsing rule generation with a special `<=` pattern. +// See `call` instruction to see it in action derive_parser!( #[derive(Logos, PartialEq, Eq, Debug, Clone, Copy)] #[logos(skip r"(?:\s+)|(?://[^\n\r]*[\n\r]*)|(?:/\*[^*]*\*+(?:[^/*][^*]*\*+)*/)")] diff --git a/ptx_parser_macros/src/lib.rs b/ptx_parser_macros/src/lib.rs index a2f8396..5f47fac 100644 --- a/ptx_parser_macros/src/lib.rs +++ b/ptx_parser_macros/src/lib.rs @@ -16,7 +16,7 @@ use syn::{ // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#packed-integer-data-types
#[rustfmt::skip]
static POSTFIX_MODIFIERS: &[&str] = &[
- ".v2", ".v4",
+ ".v2", ".v4", ".v8",
".s8", ".s16", ".s16x2", ".s32", ".s64",
".u8", ".u16", ".u16x2", ".u32", ".u64",
".f16", ".f16x2", ".f32", ".f64",
|