diff --git a/src/Lexers.jl b/src/Lexers.jl index 5dfbc63..48ca0c4 100644 --- a/src/Lexers.jl +++ b/src/Lexers.jl @@ -91,8 +91,9 @@ export JuliaLexer, JuliaConsoleLexer, MatlabLexer, + PTXLexer, RLexer, - TOMLLexer + TOMLLexer, "A FORTRAN 90 source code lexer." abstract type FortranLexer <: AbstractLexer end @@ -102,6 +103,8 @@ abstract type JuliaLexer <: AbstractLexer end abstract type JuliaConsoleLexer <: AbstractLexer end "A lexer for MATLAB source code." abstract type MatlabLexer <: AbstractLexer end +"PTX ISA(NVIDIA GPU) Lexer" +abstract type PTXLexer <: AbstractLexer end "A lexer for the R language." abstract type RLexer <: AbstractLexer end "TOML (Tom's Obvious, Minimal Language) lexer." @@ -115,6 +118,7 @@ using ..Highlights.Tokens include("lexers/fortran.jl") include("lexers/julia.jl") include("lexers/matlab.jl") +include("lexers/ptx.jl") include("lexers/r.jl") include("lexers/toml.jl") diff --git a/src/lexers/ptx-constants.jl b/src/lexers/ptx-constants.jl new file mode 100644 index 0000000..2019c0a --- /dev/null +++ b/src/lexers/ptx-constants.jl @@ -0,0 +1,250 @@ +const types = [ + "s8", + "s16", + "s32", + "s64", + "u8", + "u16,", + "u32", + "u64", + "f16", + "f16x2", + "f32", + "f64", + "b8", + "b16", + "b32", + "b64", + "pred", +] + +const directives = [ + "address_size", + "align", + "branchtargets", + "callprototype", + "calltargets", + "const", + "entry", + "extern", + "file", + "func", + "global", + "loc", + "local", + "maxnctapersm", + "maxnreg", + "maxntid", + "minnctapersm", + "param", + "pragma", + "reg", + "reqntid", + "section", + "shared", + "sreg", + "target", + "tex", + "version", + "visible", + "weak", +] + +const predefined_identifiers = [ + "%clock", + "%clock64", + "%ctaid", + "%envreg<32>", + "%gridid", + "%laneid", + "%lanemask_eq", + "%lanemask_ge", + "%lanemask_gt", + "%lanemask_le", + "%lanemask_lt", + "%nctaid", + "%nsmid", + "%ntid", + "%nwarpid", + "%pm0", + "%pm1", + "%pm2", + "%pm3", + "%pm4", + "%pm5", + "%pm6", + "%pm7", + "%smid", + "%tid", + "%warpid", + "WARP_SZ", +] + +const ptx_instructions = [ + "abs", + "activemask", + "add", + "addc", + "alloca", + "and", + "applypriority", + "atom", + "bar", + "barrier", + "bfe", + "bfi", + "bfind", + "bmsk", + "bra", + "brev", + "brkpt", + "brx", + "call", + "clz", + "cnot", + "copysign", + "cos", + "cp", + "createpolicy", + "cvt", + "cvta", + "discard", + "div", + "dp2a", + "dp4a", + "ex2", + "exit", + "fence", + "fma", + "fns", + "isspacep", + "istypep", + "ld", + "ldmatrix", + "ldu", + "lg2", + "lop3", + "mad", + "mad24", + "madc", + "match", + "max", + "mbarrier", + "membar", + "min", + "mma", + "mov", + "mul", + "mul24", + "nanosleep", + "neg", + "not", + "or", + "pmevent", + "popc", + "prefetch", + "prefetchu", + "prmt", + "rcp", + "red", + "redux", + "rem", + "ret", + "rsqrt", + "sad", + "selp", + "set", + "setp", + "shf", + "shfl", + "shl", + "shr", + "sin", + "slct", + "sqrt", + "st", + "stackrestore", + "stacksave", + "sub", + "subc", + "suld", + "suq", + "sured", + "sust", + "szext", + "tanh", + "testp", + "tex", + "tld4", + "trap", + "txq", + "vabsdiff", + "vabsdiff2", + "vabsdiff4", + "vadd", + "vadd2", + "vadd4", + "vavrg2", + "vavrg4", + "vmad", + "vmax", + "vmax2", + "vmax4", + "vmin", + "vmin2", + "vmin4", + "vote", + "vset", + "vset2", + "vset4", + "vshl", + "vshr", + "vsub", + "vsub2", + "vsub4", + "wmma", + "xor", + "uni" +] + +operators_comparison_sint = ["eq", "ne", "lt", "le", "gt", "ge"] +operators_comparison_uint = ["eq", "ne", "lo", "ls", "hi", "hs"] +operators_comparison_bit = ["eq", "ne"] + +operators_comparison_float = ["eq", "ne", "lt", "le", "gt", "ge"] +operators_comparison_nanfloat = ["equ", "neu", "ltu", "leu", "gtu", "geu"] +operators_comparison_nan = ["num", "nan"] + +modifiers_int = ["rni", "rzi", "rmi", "rpi"] +modifiers_float = ["rn", "rna", "rz", "rm", "rp"] +modifiers = sort(unique([modifiers_int..., modifiers_float...])) + +state_spaces = ["reg", "sreg", "const", "global", "local", "param", "shared", "tex"] + +const operators = sort( + unique([ + operators_comparison_sint..., + operators_comparison_uint..., + operators_comparison_bit..., + operators_comparison_float..., + operators_comparison_nanfloat..., + operators_comparison_nan..., + modifiers..., + state_spaces..., + types..., + ]), + +r_followsym = "[a-zA-Z0-9_\$]" +r_identifier = "(?:[a-zA-Z]$r_followsym*)|(?:[_\$%]$r_followsym+)" + +r_hex = "0[xX][A-F]+U?" +r_octal = "0[0-8]+U?" +r_binary = "0[bB][01]+U?" +r_decimal = "[0-9]+(?:\\.[0-9]+)?U?" +r_float = "0[fF]{hexdigit}{8}" +r_double = "0[dD]{hexdigit}{16}" + +r_number = join( + map(x -> "(?:" * x * ")", [r_hex, r_octal, r_binary, r_decimal, r_float, r_double]), + "|", +) diff --git a/src/lexers/ptx.jl b/src/lexers/ptx.jl new file mode 100644 index 0000000..4c29ff2 --- /dev/null +++ b/src/lexers/ptx.jl @@ -0,0 +1,46 @@ +include("./ptx-constants.jl") + +@lexer PTXLexer Dict( + :name => "PTX Lexer", + :description => "A PTX(Parallel Thread Execution) lexer. Virtual ISA for NVIDIA GPU's", + :tokens => Dict( + :root => [ + (r"//[^\n]*\n", COMMENT_SINGLE), + (r"[ \t\n]+", WHITESPACE), + (r"\w+:", NAME_LABEL), + (r"bra[^\n]+\n", :branch), + (Regex("(" * r_number * ")") * r"([ \t()\n,;\]])", (NUMBER, TEXT)), + (r"(\.\w+)([ \t,\.])", (:dot_token, TEXT)), + (r"(\w+)([ \t,\.\n;\[()\+\]])", (:identifier, TEXT)), + (r"@?%\w+(?:<\d+>)?(\.x)?", :register), + (r"/\*", COMMENT_MULTILINE, :multiline_comments), + ], + :branch => [ + (r"bra", KEYWORD), + (r"[ \t]+", WHITESPACE), + (r"(\.\w+)([ \t,\.])", (:dot_token, TEXT)), + (r"(\w+)([^ \t\.])", (NAME_LABEL, TEXT)), + (r";", TEXT, :__pop__) + ], + :dot_token => [ + (r"\.func[^\w]", NAME_FUNCTION, :__pop__), + (r"\." * Lexers.words(vcat(ptx_instructions, directives, operators)), KEYWORD, :__pop__), + (r"\.\w+[^\w]", TEXT, :__pop__), + (r"\w+[^\w]", TEXT, :__pop__) + ], + :register => [ + (r"@?%\w+(?:<\d+>)?(\.x)?", LITERAL, :__pop__) + ], + :identifier => [ + (Lexers.words(predefined_identifiers), KEYWORD, :__pop__), + (Lexers.words(ptx_instructions), KEYWORD, :__pop__), + (Lexers.words(operators), KEYWORD, :__pop__), + (Regex(r_identifier), TEXT, :__pop__), + ], + :multiline_comments => [ + (r"/\*", COMMENT_MULTILINE, :__push__), + (r"\*/", COMMENT_MULTILINE, :__pop__), + (r"[^/\*]+", COMMENT_MULTILINE), + ], + ) +)