Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 16 additions & 0 deletions lang/en.txt
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,22 @@ E109=not an lvalue (prefix)
E110=unknown field in lvalue
E111=not an lvalue

# ---- Triton frontend (E112-E125) ----
E112=indentation too deeply nested
E113=inconsistent dedent
E114=unterminated string literal
E115=unexpected character '!'
E116=unexpected character
E117=lexer made no progress (internal bug)
E118=unterminated triple-quoted string
E119=too many AST children
E120=expected identifier
E121=expected identifier after '.'
E122=expected parameter name
E123=expected identifier after 'import'
E124=expected attribute name after '.'
E125=unrecognised expression

# ---- ABEND Messages (hex IDs) ----
A0C1=illegal GPU instruction
A0C4=memory access violation (page not mapped)
Expand Down
17 changes: 16 additions & 1 deletion src/fe/bc_err.c
Original file line number Diff line number Diff line change
Expand Up @@ -78,7 +78,22 @@ static const char *bc_dflt[BC_EID_MAX] = {
/* E109 */ "not an lvalue (prefix)",
/* E110 */ "unknown field in lvalue",
/* E111 */ "not an lvalue",
/* E112-E129 */
/* ---- Triton frontend (E112-E125): lexer + parser, fixed messages ---- */
/* E112 */ "indentation too deeply nested",
/* E113 */ "inconsistent dedent",
/* E114 */ "unterminated string literal",
/* E115 */ "unexpected character '!'",
/* E116 */ "unexpected character",
/* E117 */ "lexer made no progress (internal bug)",
/* E118 */ "unterminated triple-quoted string",
/* E119 */ "too many AST children",
/* E120 */ "expected identifier",
/* E121 */ "expected identifier after '.'",
/* E122 */ "expected parameter name",
/* E123 */ "expected identifier after 'import'",
/* E124 */ "expected attribute name after '.'",
/* E125 */ "unrecognised expression",
/* E126-E149 */
};

/* ---- ABEND compiled-in defaults ----
Expand Down
22 changes: 20 additions & 2 deletions src/fe/bc_err.h
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
* E100-E129 Lowering
*/

#define BC_EID_MAX 130
#define BC_EID_MAX 150

typedef enum {
BC_E000 = 0, /* internal compiler error — the "this shouldn't happen" */
Expand Down Expand Up @@ -81,7 +81,25 @@ typedef enum {
BC_E108 = 108, /* parameter not addressable */
BC_E109 = 109, /* not an lvalue (prefix) */
BC_E110 = 110, /* unknown field in lvalue */
BC_E111 = 111 /* not an lvalue */
BC_E111 = 111, /* not an lvalue */

/* ---- Triton frontend (E112-E140), block reserved to avoid the FE
collisions the Python path used to have. E112-E125 are catalogued
below; E126-E140 (sema/lower) carry runtime-built messages. ---- */
BC_E112 = 112, /* indentation too deeply nested */
BC_E113 = 113, /* inconsistent dedent */
BC_E114 = 114, /* unterminated string literal (triton) */
BC_E115 = 115, /* unexpected character '!' */
BC_E116 = 116, /* unexpected character (triton) */
BC_E117 = 117, /* lexer made no progress */
BC_E118 = 118, /* unterminated triple-quoted string */
BC_E119 = 119, /* too many AST children */
BC_E120 = 120, /* expected identifier */
BC_E121 = 121, /* expected identifier after '.' */
BC_E122 = 122, /* expected parameter name */
BC_E123 = 123, /* expected identifier after 'import' */
BC_E124 = 124, /* expected attribute name after '.' */
BC_E125 = 125 /* unrecognised expression */
} bc_eid_t;

/* Returns format string for eid -- loaded translation or compiled-in English */
Expand Down
21 changes: 13 additions & 8 deletions src/triton/lex.c
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
* positions and emits INDENT/DEDENT at block boundaries. */

#include "triton.h"
#include "bc_err.h"

#include <string.h>
#include <stdio.h>
Expand Down Expand Up @@ -122,7 +123,11 @@ static void tn_err(tn_lex_t *L, uint16_t eid, const char *msg)
e->eid = eid;
e->loc.line = L->line;
e->loc.col = (uint16_t)(L->pos - L->line_start);
snprintf(e->msg, sizeof(e->msg), "%s", msg);
/* Prefer the shared catalogue (so --lang can translate); fall back to the
* inline text for any code not yet registered. */
const char *cat = bc_efmt((bc_eid_t)eid);
snprintf(e->msg, sizeof(e->msg), "%s",
(cat && strcmp(cat, "unknown error") != 0) ? cat : msg);
}

static void tn_emit(tn_lex_t *L, int kind, uint32_t off, uint32_t len)
Expand Down Expand Up @@ -215,7 +220,7 @@ static int tn_indent(tn_lex_t *L)
: 0;
if (col > top) {
if (L->indent_depth >= TN_MAX_INDENTS) {
tn_err(L, 50, "indentation too deeply nested");
tn_err(L, 112, "indentation too deeply nested");
return 0;
}
L->indents[L->indent_depth++] = col;
Expand All @@ -230,7 +235,7 @@ static int tn_indent(tn_lex_t *L)
? L->indents[L->indent_depth - 1]
: 0;
if (new_top != col) {
tn_err(L, 51, "inconsistent dedent");
tn_err(L, 113, "inconsistent dedent");
return 0;
}
}
Expand Down Expand Up @@ -373,7 +378,7 @@ static void tn_scan_str_at(tn_lex_t *L, uint32_t start)
L->pos++;
guard++;
}
tn_err(L, 56, "unterminated triple-quoted string");
tn_err(L, 118, "unterminated triple-quoted string");
tn_emit_at(L, TN_TOK_STRING, start, L->pos - start,
start_line, start_line_start);
return;
Expand All @@ -393,7 +398,7 @@ static void tn_scan_str_at(tn_lex_t *L, uint32_t start)
continue;
}
if (L->src[L->pos] == '\n') {
tn_err(L, 52, "unterminated string literal");
tn_err(L, 114, "unterminated string literal");
tn_emit_at(L, TN_TOK_STRING, start, L->pos - start,
start_line, start_line_start);
return;
Expand Down Expand Up @@ -563,11 +568,11 @@ static void tn_scan_op(tn_lex_t *L)
if (tn_match2(L, '=')) { tn_emit_op(L, TN_TOK_NE, 2); return; }
/* A bare ! is not valid Python. Report it and advance so the
* parser can keep going. */
tn_err(L, 53, "unexpected character '!'");
tn_err(L, 115, "unexpected character '!'");
L->pos++;
return;
default:
tn_err(L, 54, "unexpected character");
tn_err(L, 116, "unexpected character");
L->pos++;
return;
}
Expand Down Expand Up @@ -600,7 +605,7 @@ int tn_tokenize(tn_lex_t *L)
uint32_t guard = 0;
while (L->pos < L->src_len) {
if (++guard > L->src_len * 4) {
tn_err(L, 55, "lexer made no progress (internal bug)");
tn_err(L, 117, "lexer made no progress (internal bug)");
return BC_ERR_TRITON;
}

Expand Down
40 changes: 20 additions & 20 deletions src/triton/lower.c
Original file line number Diff line number Diff line change
Expand Up @@ -211,7 +211,7 @@ static int l_shape_supported(tn_lower_t *L, uint32_t node_idx)
snprintf(msg, sizeof(msg),
"rank-2 tile (%s) needs matrix codegen "
"(MFMA / mma.sync), not yet lowered", sbuf);
l_err(L, 99, l_tok(L, node_idx), msg);
l_err(L, 138, l_tok(L, node_idx), msg);
return 0;
}

Expand All @@ -237,7 +237,7 @@ static uint32_t l_lit(tn_lower_t *L, uint32_t node_idx)
case TN_LIT_NONE:
return BIR_MAKE_CONST(bir_const_int(L->bir, L->t_i32, 0));
case TN_LIT_STRING:
l_err(L, 90, t, "string literals are not lowerable in sitting one");
l_err(L, 129, t, "string literals are not lowerable in sitting one");
return BIR_VAL_NONE;
default:
return BIR_VAL_NONE;
Expand Down Expand Up @@ -272,7 +272,7 @@ static uint32_t l_name(tn_lower_t *L, uint32_t node_idx)
uint32_t v = L->node_val[aux];
if (v != BIR_VAL_NONE) return v;
}
l_err(L, 91, l_tok(L, node_idx),
l_err(L, 130, l_tok(L, node_idx),
"local referenced before it produced a BIR value");
return BIR_VAL_NONE;
}
Expand All @@ -282,7 +282,7 @@ static uint32_t l_name(tn_lower_t *L, uint32_t node_idx)
/* Bare module / intrinsic / type names should not appear as
* value expressions on their own; usually means the kernel
* uses something not yet lowered. Report and continue. */
l_err(L, 92, l_tok(L, node_idx),
l_err(L, 131, l_tok(L, node_idx),
"bare module or intrinsic name in expression position");
return BIR_VAL_NONE;
case TN_SYM_UNBOUND:
Expand Down Expand Up @@ -409,7 +409,7 @@ static uint32_t l_binop(tn_lower_t *L, uint32_t node_idx)
uint32_t rhs = l_expr(L, l_kid(L, node_idx, 1));
uint32_t r = l_emit_binop(L, (int)n->flags, lhs, rhs);
if (r == BIR_VAL_NONE && lhs != BIR_VAL_NONE && rhs != BIR_VAL_NONE)
l_err(L, 93, l_tok(L, node_idx), "binop not yet lowered");
l_err(L, 132, l_tok(L, node_idx), "binop not yet lowered");
return r;
}

Expand Down Expand Up @@ -438,7 +438,7 @@ static uint32_t l_compare(tn_lower_t *L, uint32_t node_idx)
case TN_CMP_EQ: pred = BIR_FCMP_OEQ; break;
case TN_CMP_NE: pred = BIR_FCMP_ONE; break;
default:
l_err(L, 96, l_tok(L, node_idx),
l_err(L, 135, l_tok(L, node_idx),
"comparison predicate not yet lowered for floats");
return BIR_VAL_NONE;
}
Expand All @@ -456,7 +456,7 @@ static uint32_t l_compare(tn_lower_t *L, uint32_t node_idx)
case TN_CMP_EQ: pred = BIR_ICMP_EQ; break;
case TN_CMP_NE: pred = BIR_ICMP_NE; break;
default:
l_err(L, 96, l_tok(L, node_idx),
l_err(L, 135, l_tok(L, node_idx),
"comparison predicate not yet lowered");
return BIR_VAL_NONE;
}
Expand Down Expand Up @@ -492,7 +492,7 @@ static uint32_t l_unop(tn_lower_t *L, uint32_t node_idx)
return inst;
}
default:
l_err(L, 94, l_tok(L, node_idx),
l_err(L, 133, l_tok(L, node_idx),
"unary operator not yet lowered");
return BIR_VAL_NONE;
}
Expand Down Expand Up @@ -782,7 +782,7 @@ static uint32_t l_intrinsic_call(tn_lower_t *L, uint32_t call_idx,
uint32_t tn = l_pos_arg(L, call_idx, 1);
uint32_t fn = l_pos_arg(L, call_idx, 2);
if (cn == 0 || tn == 0 || fn == 0) {
l_err(L, 95, l_tok(L, call_idx),
l_err(L, 134, l_tok(L, call_idx),
"tl.where expects condition, true value, false value");
return BIR_VAL_NONE;
}
Expand All @@ -791,7 +791,7 @@ static uint32_t l_intrinsic_call(tn_lower_t *L, uint32_t call_idx,
if (cond == BIR_VAL_NONE) return BIR_VAL_NONE;
int ck = l_type_kind(L, l_val_type(L, cond));
if (ck != BIR_TYPE_INT) {
l_err(L, 95, l_tok(L, call_idx),
l_err(L, 134, l_tok(L, call_idx),
"tl.where condition must lower to an integer/bool mask");
return BIR_VAL_NONE;
}
Expand All @@ -804,7 +804,7 @@ static uint32_t l_intrinsic_call(tn_lower_t *L, uint32_t call_idx,
uint32_t fty = l_val_type(L, fv);
uint32_t sty = l_tjoin(L, tty, fty);
if (sty == BIR_VAL_NONE) {
l_err(L, 95, l_tok(L, call_idx),
l_err(L, 134, l_tok(L, call_idx),
"tl.where value types not yet lowerable");
return BIR_VAL_NONE;
}
Expand Down Expand Up @@ -843,7 +843,7 @@ static uint32_t l_intrinsic_call(tn_lower_t *L, uint32_t call_idx,
snprintf(msg, sizeof(msg),
"intrinsic tl.%s not yet lowered in sitting one",
tn_intrinsic_name(intrinsic_id));
l_err(L, 95, l_tok(L, call_idx), msg);
l_err(L, 134, l_tok(L, call_idx), msg);
return BIR_VAL_NONE;
}
}
Expand Down Expand Up @@ -902,7 +902,7 @@ static int l_tile(tn_lower_t *L, uint32_t node, int *out){
int rows = (sh.rank>=1) ? sh.dims[0] : 1;
int cols = (sh.rank>=2) ? sh.dims[1] : 1;
if (rows<=0 || cols<=0 || rows*cols>TN_TILE_MAX){
l_err(L, 99, l_tok(L,node), "tile shape unsupported or too large");
l_err(L, 138, l_tok(L,node), "tile shape unsupported or too large");
return -1;
}
switch (n->kind){
Expand All @@ -912,7 +912,7 @@ static int l_tile(tn_lower_t *L, uint32_t node, int *out){
uint32_t aux = L->sema->node_sym_aux[node];
if (aux<TN_MAX_NODES && L->node_tile[aux]>=0){ *out=L->node_tile[aux]; return 0; }
}
l_err(L, 91, l_tok(L,node), "tile local referenced before defined");
l_err(L, 130, l_tok(L,node), "tile local referenced before defined");
return -1;
}
case TN_NK_SUBSCRIPT: {
Expand Down Expand Up @@ -1065,7 +1065,7 @@ static uint32_t l_call(tn_lower_t *L, uint32_t node_idx)
return l_intrinsic_call(L, node_idx, id);
}

l_err(L, 96, l_tok(L, node_idx),
l_err(L, 135, l_tok(L, node_idx),
"non-intrinsic call not yet lowered");
return BIR_VAL_NONE;
}
Expand All @@ -1090,11 +1090,11 @@ static uint32_t l_expr(tn_lower_t *L, uint32_t node_idx)
case TN_NK_ATTR:
case TN_NK_IFEXPR:
case TN_NK_BOOLOP:
l_err(L, 97, l_tok(L, node_idx),
l_err(L, 136, l_tok(L, node_idx),
"expression kind not yet lowered in sitting two");
return BIR_VAL_NONE;
case TN_NK_EXPR_SPAN:
l_err(L, 98, l_tok(L, node_idx),
l_err(L, 137, l_tok(L, node_idx),
"opaque expression span left over from parser");
return BIR_VAL_NONE;
default:
Expand Down Expand Up @@ -1236,7 +1236,7 @@ static void l_stmt(tn_lower_t *L, uint32_t node_idx)
uint32_t rhs_idx = l_kid(L, node_idx, 1);
const tn_node_t *tn = &L->parser->nodes[target_idx];
if (tn->kind != TN_NK_NAME) {
l_err(L, 102, l_tok(L, target_idx),
l_err(L, 139, l_tok(L, target_idx),
"AugAssign target kind not yet lowered");
break;
}
Expand Down Expand Up @@ -1286,7 +1286,7 @@ static void l_stmt(tn_lower_t *L, uint32_t node_idx)
case TN_AUG_SHL: bop = TN_BOP_SHL; break;
case TN_AUG_SHR: bop = TN_BOP_SHR; break;
default:
l_err(L, 103, l_tok(L, node_idx),
l_err(L, 140, l_tok(L, node_idx),
"AugAssign operator not yet lowered");
break;
}
Expand Down Expand Up @@ -1323,7 +1323,7 @@ static void l_stmt(tn_lower_t *L, uint32_t node_idx)
l_for(L, node_idx);
break;
case TN_NK_WHILE:
l_err(L, 99, l_tok(L, node_idx),
l_err(L, 138, l_tok(L, node_idx),
"while loops not yet lowered");
break;
default:
Expand Down
Loading
Loading