Download unified-engine/kernel/compiler/parser.js from Snapkitty/repoverse: direct link, hf CLI and curl.
- Browser
- Download file 10.3 kB
-
https://huggingface.co/Snapkitty/repoverse/resolve/main/unified-engine/kernel/compiler/parser.js
- Command line
-
hf download hf://Snapkitty/repoverse/unified-engine/kernel/compiler/parser.js
-
curl -L -o parser.js https://huggingface.co/Snapkitty/repoverse/resolve/main/unified-engine/kernel/compiler/parser.js
10.3 kB
| /** | |
| * Recursive-descent parser for the CUDA-like kernel language. | |
| * Produces the AST in ../ast/ast.js. Every failure raises ParseError with a line. | |
| */ | |
| import { tokenize, TokenType } from './lexer.js'; | |
| export class ParseError extends Error { | |
| constructor(message, line) { | |
| super(`${message} (line ${line})`); | |
| this.name = 'ParseError'; | |
| this.line = line; | |
| } | |
| } | |
| const SCALARS = new Set(['int', 'uint', 'float', 'double', 'void']); | |
| const BUILTINS = new Set([ | |
| 'blockIdx', 'blockDim', 'threadIdx', 'gridDim', | |
| 'global_id', 'local_id', 'block_id', 'global_id_x', 'local_id_x', | |
| ]); | |
| /** Binary operator precedence, higher binds tighter. Follows C. */ | |
| const PRECEDENCE = { | |
| '||': 1, '&&': 2, | |
| '|': 3, '^': 4, '&': 5, | |
| '==': 6, '!=': 6, | |
| '<': 7, '>': 7, '<=': 7, '>=': 7, | |
| '<<': 8, '>>': 8, | |
| '+': 9, '-': 9, | |
| '*': 10, '/': 10, '%': 10, | |
| }; | |
| const ASSIGN_OPS = new Set(['=', '+=', '-=', '*=', '/=', '%=']); | |
| /** Maps CUDA-style builtins onto the backend-neutral IR names. */ | |
| const BUILTIN_MEMBERS = { | |
| blockIdx: { x: 'block_id_x', y: 'block_id_y', z: 'block_id_z' }, | |
| blockDim: { x: 'block_dim_x', y: 'block_dim_y', z: 'block_dim_z' }, | |
| threadIdx: { x: 'local_id_x', y: 'local_id_y', z: 'local_id_z' }, | |
| gridDim: { x: 'grid_dim_x', y: 'grid_dim_y', z: 'grid_dim_z' }, | |
| }; | |
| export function parse(source) { | |
| const tokens = tokenize(source); | |
| let pos = 0; | |
| const peek = (offset = 0) => tokens[Math.min(pos + offset, tokens.length - 1)]; | |
| const at = (value, type = TokenType.PUNCT) => peek().type === type && peek().value === value; | |
| const atKeyword = (value) => peek().type === TokenType.KEYWORD && peek().value === value; | |
| const next = () => tokens[pos++]; | |
| const expect = (value) => { | |
| if (!at(value)) throw new ParseError(`Expected '${value}' but found '${peek().value}'`, peek().line); | |
| return next(); | |
| }; | |
| const atScalar = () => peek().type === TokenType.KEYWORD && SCALARS.has(peek().value) && peek().value !== 'void'; | |
| const expectIdent = () => { | |
| if (peek().type !== TokenType.IDENT) throw new ParseError(`Expected an identifier but found '${peek().value}'`, peek().line); | |
| return next().value; | |
| }; | |
| // ---- parameter type: `const float* A` or `int N` ---- | |
| function parseParamType() { | |
| const isConst = atKeyword('const'); | |
| if (isConst) next(); | |
| const typeToken = next(); | |
| if (typeToken.type !== TokenType.KEYWORD || !SCALARS.has(typeToken.value)) { | |
| throw new ParseError(`'${typeToken.value}' is not a supported parameter type`, typeToken.line); | |
| } | |
| const scalar = typeToken.value; | |
| let isBuffer = false; | |
| if (at('*')) { | |
| next(); | |
| isBuffer = true; | |
| } | |
| return { kind: isBuffer ? 'buffer' : 'scalar', scalar, isConst }; | |
| } | |
| function parseParams() { | |
| expect('('); | |
| const params = []; | |
| if (at(')')) { | |
| next(); | |
| return params; | |
| } | |
| for (;;) { | |
| const type = parseParamType(); | |
| const name = expectIdent(); | |
| params.push({ name, type }); | |
| if (at(',')) { | |
| next(); | |
| if (at(')')) break; | |
| continue; | |
| } | |
| break; | |
| } | |
| expect(')'); | |
| return params; | |
| } | |
| // ---- primary expressions ---- | |
| function parsePrimary() { | |
| const token = peek(); | |
| if (at('(')) { | |
| next(); | |
| const inner = parseExpr(); | |
| expect(')'); | |
| return inner; | |
| } | |
| if (token.type === TokenType.INT) { | |
| next(); | |
| const value = token.value.startsWith('0x') || token.value.startsWith('0X') | |
| ? parseInt(token.value, 16) | |
| : parseInt(token.value, 10); | |
| return { kind: 'int', value }; | |
| } | |
| if (token.type === TokenType.FLOAT) { | |
| next(); | |
| // As in C, an unsuffixed literal is a double and `f` makes it a float. | |
| const single = /[fF]$/.test(token.value); | |
| return { kind: 'float', value: Number.parseFloat(token.value), single }; | |
| } | |
| // Functional cast: int(x), float(x). | |
| if (atScalar() && peek(1).type === TokenType.PUNCT && peek(1).value === '(') { | |
| const to = next().value; | |
| next(); | |
| const inner = parseExpr(); | |
| expect(')'); | |
| return { kind: 'cast', to, operand: inner }; | |
| } | |
| if (atKeyword('true')) { | |
| next(); | |
| return { kind: 'bool', value: true }; | |
| } | |
| if (atKeyword('false')) { | |
| next(); | |
| return { kind: 'bool', value: false }; | |
| } | |
| if (token.type === TokenType.IDENT) { | |
| next(); | |
| if (BUILTINS.has(token.value)) return { kind: 'builtin', name: token.value }; | |
| if (at('(')) throw new ParseError(`Unknown function '${token.value}'`, token.line); | |
| return { kind: 'var', name: token.value }; | |
| } | |
| throw new ParseError(`Unexpected '${token.value}' in expression`, token.line); | |
| } | |
| function parsePostfix() { | |
| let left = parsePrimary(); | |
| for (;;) { | |
| if (at('[')) { | |
| next(); | |
| const index = parseExpr(); | |
| expect(']'); | |
| left = { kind: 'index', base: left, index }; | |
| continue; | |
| } | |
| if (at('.') && left.kind === 'builtin') { | |
| next(); | |
| const member = expectIdent(); | |
| const mapped = BUILTIN_MEMBERS[left.name]; | |
| if (!mapped || !(member in mapped)) { | |
| throw new ParseError(`'${left.name}' has no member '${member}'`, peek().line); | |
| } | |
| left = { kind: 'call', callee: mapped[member], args: [] }; | |
| continue; | |
| } | |
| if (at('++') || at('--')) { | |
| const op = next().value; | |
| left = { kind: 'assign', op: op === '++' ? '+=' : '-=', target: left, value: { kind: 'int', value: 1 } }; | |
| continue; | |
| } | |
| break; | |
| } | |
| return left; | |
| } | |
| function parseUnary() { | |
| if (at('!') || at('-') || at('~')) { | |
| const op = next().value; | |
| return { kind: 'unary', op, operand: parseUnary() }; | |
| } | |
| // C-style cast: (float)x. | |
| if (at('(') && peek(1).type === TokenType.KEYWORD && SCALARS.has(peek(1).value) && peek(1).value !== 'void' | |
| && peek(2).type === TokenType.PUNCT && peek(2).value === ')') { | |
| next(); | |
| const to = next().value; | |
| next(); | |
| return { kind: 'cast', to, operand: parseUnary() }; | |
| } | |
| return parsePostfix(); | |
| } | |
| function parseBinary(minPrecedence) { | |
| let left = parseUnary(); | |
| for (;;) { | |
| const token = peek(); | |
| if (token.type !== TokenType.PUNCT) break; | |
| const precedence = PRECEDENCE[token.value]; | |
| if (precedence === undefined || precedence < minPrecedence) break; | |
| next(); | |
| const right = parseBinary(precedence + 1); | |
| left = { kind: 'binary', op: token.value, left, right }; | |
| } | |
| return left; | |
| } | |
| function parseExpr() { | |
| const left = parseBinary(1); | |
| if (ASSIGN_OPS.has(peek().value) && peek().type === TokenType.PUNCT) { | |
| const op = next().value; | |
| const value = parseExpr(); | |
| return { kind: 'assign', op, target: left, value }; | |
| } | |
| return left; | |
| } | |
| // ---- statements ---- | |
| function parseBlock() { | |
| expect('{'); | |
| const body = []; | |
| while (!at('}')) { | |
| if (peek().type === TokenType.EOF) throw new ParseError("Unterminated block: expected '}'", peek().line); | |
| body.push(parseStmt()); | |
| } | |
| expect('}'); | |
| return body; | |
| } | |
| function parseStmt() { | |
| if (at('{')) return { kind: 'block', body: parseBlock() }; | |
| if (atKeyword('if')) { | |
| next(); | |
| expect('('); | |
| const cond = parseExpr(); | |
| expect(')'); | |
| // A braceless body is valid CUDA and is used by the reference kernel, so a single | |
| // statement is accepted and wrapped rather than rejected. | |
| const then = at('{') ? parseBlock() : [parseStmt()]; | |
| let otherwise = []; | |
| if (atKeyword('else')) { | |
| next(); | |
| otherwise = at('{') ? parseBlock() : [parseStmt()]; | |
| } | |
| return { kind: 'if', cond, then, otherwise }; | |
| } | |
| if (atKeyword('while')) { | |
| next(); | |
| expect('('); | |
| const cond = parseExpr(); | |
| expect(')'); | |
| const body = at('{') ? parseBlock() : [parseStmt()]; | |
| return { kind: 'while', cond, body }; | |
| } | |
| if (atKeyword('for')) { | |
| next(); | |
| expect('('); | |
| let init = null; | |
| if (at(';')) next(); | |
| else init = parseSimpleStmt(); | |
| const cond = at(';') ? null : parseExpr(); | |
| expect(';'); | |
| const step = at(')') ? null : parseExpr(); | |
| expect(')'); | |
| const body = at('{') ? parseBlock() : [parseStmt()]; | |
| return { kind: 'for', init, cond, step, body }; | |
| } | |
| if (atKeyword('break') || atKeyword('continue')) { | |
| const kind = next().value; | |
| expect(';'); | |
| return { kind }; | |
| } | |
| if (atKeyword('return')) { | |
| next(); | |
| if (at(';')) { | |
| next(); | |
| return { kind: 'return', value: null }; | |
| } | |
| const value = parseExpr(); | |
| expect(';'); | |
| return { kind: 'return', value }; | |
| } | |
| return parseSimpleStmt(); | |
| } | |
| /** A declaration or expression statement, including its ';'. Also the init clause of `for`. */ | |
| function parseSimpleStmt() { | |
| // Local scalar declaration: int i = 0; (a functional cast like int(x) is an expression) | |
| if (atScalar() && !(peek(1).type === TokenType.PUNCT && peek(1).value === '(')) { | |
| const type = next().value; | |
| const name = expectIdent(); | |
| let init = null; | |
| if (at('=')) { | |
| next(); | |
| init = parseExpr(); | |
| } | |
| expect(';'); | |
| return { kind: 'decl', name, type, init }; | |
| } | |
| const expr = parseExpr(); | |
| expect(';'); | |
| // An assignment statement keeps the assign node itself, matching the Stmt kinds in ast.js. | |
| if (expr.kind === 'assign') return expr; | |
| return { kind: 'expr', value: expr }; | |
| } | |
| // ---- top level: a single __global__ kernel ---- | |
| if (!atKeyword('__global__')) { | |
| throw new ParseError(`A kernel must be declared __global__ (found '${peek().value}')`, peek().line); | |
| } | |
| next(); | |
| const returnToken = next(); | |
| if (returnToken.type !== TokenType.KEYWORD || !SCALARS.has(returnToken.value)) { | |
| throw new ParseError(`'${returnToken.value}' is not a supported return type`, returnToken.line); | |
| } | |
| const returnType = returnToken.value; | |
| const name = expectIdent(); | |
| const params = parseParams(); | |
| const body = parseBlock(); | |
| if (peek().type !== TokenType.EOF) { | |
| throw new ParseError(`Unexpected trailing input '${peek().value}'`, peek().line); | |
| } | |
| return { name, returnType, params, body, isGlobal: true }; | |
| } | |