From 2043f3bd40f3843b8691fd7a93ceb4940157ebc3 Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Mon, 27 Jul 2026 00:51:41 +1200 Subject: [PATCH 1/2] fix: parse function pointer declarators --- src/fe/parser.c | 59 ++++++++++++++++++++++++++++++++++++++++++++++++- tests/fnptr.cu | 18 +++++++++++++++ tests/tphase.c | 14 ++++++++++++ 3 files changed, 90 insertions(+), 1 deletion(-) create mode 100644 tests/fnptr.cu diff --git a/src/fe/parser.c b/src/fe/parser.c index fadb40a..762b7e4 100644 --- a/src/fe/parser.c +++ b/src/fe/parser.c @@ -318,6 +318,8 @@ static uint32_t parse_expr(parser_t *P, int min_prec); static uint32_t parse_stmt(parser_t *P); static uint32_t parse_type_spec(parser_t *P, uint16_t *quals, uint16_t *cuda); static uint32_t parse_decl_or_stmt(parser_t *P); +static uint32_t fnptr(parser_t *P, int *depth); +static int is_fnptr(parser_t *P); static int64_t parse_int_text(const char *s, int len) { @@ -894,7 +896,12 @@ static uint32_t parse_param_list(parser_t *P) P->nodes[param].d.oper.flags = ptr_depth; } - if (cur_type(P) == TOK_IDENT) { + if (is_fnptr(P)) { + int d = P->nodes[param].d.oper.flags; + uint32_t fname = fnptr(P, &d); + P->nodes[param].d.oper.flags = d; + if (fname) add_child(P, param, fname); + } else if (cur_type(P) == TOK_IDENT) { uint32_t name = alloc_node(P, AST_IDENT); P->nodes[name].d.text.offset = cur(P)->offset; P->nodes[name].d.text.len = cur(P)->len; @@ -922,6 +929,37 @@ static uint32_t parse_param_list(parser_t *P) return first; } +/* T (*name)(sig) — the inner stars add to *depth. Booth has no indirect + calls on device, so the signature is parsed and dropped and what is left + is an ordinary pointer. Returns the name node, 0 for an abstract one. */ +static uint32_t fnptr(parser_t *P, int *depth) +{ + uint32_t name = 0; + advance(P); /* ( */ + while (cur_type(P) == TOK_STAR) { advance(P); (*depth)++; } + while (cur_type(P) == TOK_CONST || cur_type(P) == TOK_CU_RESTRICT) + advance(P); + if (cur_type(P) == TOK_IDENT) { + name = alloc_node(P, AST_IDENT); + P->nodes[name].d.text.offset = cur(P)->offset; + P->nodes[name].d.text.len = cur(P)->len; + advance(P); + } + expect(P, TOK_RPAREN); + if (cur_type(P) == TOK_LPAREN) { + advance(P); + parse_param_list(P); + expect(P, TOK_RPAREN); + } + return name; +} + +/* Distinguishes T (*f)(...) from a parenthesised expression. */ +static int is_fnptr(parser_t *P) +{ + return cur_type(P) == TOK_LPAREN && peek_type(P, 1) == TOK_STAR; +} + static int starts_declaration(parser_t *P) { int t = cur_type(P); @@ -1163,6 +1201,25 @@ static uint32_t parse_declaration(parser_t *P) } } + if (is_fnptr(P)) { + uint32_t fname = fnptr(P, &ptr_depth); + if (fname) { + decl_node = alloc_node(P, AST_VAR_DECL); + P->nodes[decl_node].qualifiers = quals; + P->nodes[decl_node].cuda_flags = cuda; + P->nodes[decl_node].d.oper.flags = ptr_depth; + add_child(P, decl_node, type_node); + add_child(P, decl_node, fname); + if (quals & QUAL_TYPEDEF) + reg_tname(P, P->nodes[fname].d.text.offset, + (uint16_t)P->nodes[fname].d.text.len); + if (!expect(P, TOK_SEMI)) sync_past_semi(P); + return decl_node; + } + match(P, TOK_SEMI); + return type_node; + } + if (cur_type(P) == TOK_IDENT || cur_type(P) == TOK_TILDE || cur_type(P) == TOK_OPERATOR) { int is_dtor = (cur_type(P) == TOK_TILDE); diff --git a/tests/fnptr.cu b/tests/fnptr.cu new file mode 100644 index 0000000..5a7ba7a --- /dev/null +++ b/tests/fnptr.cu @@ -0,0 +1,18 @@ +/* Function pointer declarators in each position they can turn up. */ + +typedef void (*cb_t)(int); +typedef int (*cmp_t)(int, int); +typedef void (*void_t)(); + +struct disp { + void (*hook)(float); + int n; +}; + +cb_t g_cb; +cmp_t g_cmp; + +__global__ void k(int *x, void (*unused)(int)) +{ + x[0] = 1; +} diff --git a/tests/tphase.c b/tests/tphase.c index a4234f7..3be8815 100644 --- a/tests/tphase.c +++ b/tests/tphase.c @@ -45,6 +45,20 @@ static void pha_ast(void) } TH_REG("phase", pha_ast) +/* A function pointer declarator used to parse as a call expression, which + cost nothing at parse time and everything later. */ +static void pha_fnptr(void) +{ + int rc = th_run(BC_BIN " --parse tests/fnptr.cu", obuf, TH_BUFSZ); + CHEQ(rc, 0); + CHECK(strstr(obuf, "0 parse error(s)") != NULL); + CHECK(strstr(obuf, "(ident cb_t)") != NULL); + CHECK(strstr(obuf, "(ident hook)") != NULL); + CHECK(strstr(obuf, "call") == NULL); + PASS(); +} +TH_REG("phase", pha_fnptr) + /* ---- phase: IR ---- */ static void pha_ir(void) From 470665e1845481cfbc30242c6533f50802af4e8a Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Mon, 27 Jul 2026 00:58:30 +1200 Subject: [PATCH 2/2] fix: parse constructors and destructors --- src/fe/parser.c | 57 +++++++++++++++++++++++++++++++++++++++++-------- src/fe/parser.h | 5 +++++ tests/ctor.cu | 20 +++++++++++++++++ tests/tphase.c | 14 ++++++++++++ 4 files changed, 87 insertions(+), 9 deletions(-) create mode 100644 tests/ctor.cu diff --git a/src/fe/parser.c b/src/fe/parser.c index 762b7e4..346a732 100644 --- a/src/fe/parser.c +++ b/src/fe/parser.c @@ -929,9 +929,20 @@ static uint32_t parse_param_list(parser_t *P) return first; } -/* T (*name)(sig) — the inner stars add to *depth. Booth has no indirect - calls on device, so the signature is parsed and dropped and what is left - is an ordinary pointer. Returns the name node, 0 for an abstract one. */ +/* Steps over a balanced (...) without building nodes for it. */ +static void skippar(parser_t *P) +{ + int depth = 0; + while (cur_type(P) != TOK_EOF) { + int t = cur_type(P); + if (t == TOK_LPAREN) depth++; + else if (t == TOK_RPAREN && --depth == 0) { advance(P); return; } + advance(P); + } +} + +/* T (*name)(sig) — inner stars add to *depth, and the signature is skipped + because Booth has no indirect calls to lower it to. 0 if abstract. */ static uint32_t fnptr(parser_t *P, int *depth) { uint32_t name = 0; @@ -946,11 +957,7 @@ static uint32_t fnptr(parser_t *P, int *depth) advance(P); } expect(P, TOK_RPAREN); - if (cur_type(P) == TOK_LPAREN) { - advance(P); - parse_param_list(P); - expect(P, TOK_RPAREN); - } + if (cur_type(P) == TOK_LPAREN) skippar(P); return name; } @@ -960,6 +967,19 @@ static int is_fnptr(parser_t *P) return cur_type(P) == TOK_LPAREN && peek_type(P, 1) == TOK_STAR; } +/* S() or ~S() directly inside struct S. Both lack a return type, so the + type spec would otherwise eat the name and leave a stray (. */ +static int is_ctor(parser_t *P) +{ + if (!P->cs_len) return 0; + if (cur_type(P) == TOK_TILDE) + return peek_type(P, 1) == TOK_IDENT && peek_type(P, 2) == TOK_LPAREN; + if (cur_type(P) != TOK_IDENT || peek_type(P, 1) != TOK_LPAREN) return 0; + return cur(P)->len == P->cs_len + && memcmp(P->src + cur(P)->offset, + P->src + P->cs_off, P->cs_len) == 0; +} + static int starts_declaration(parser_t *P) { int t = cur_type(P); @@ -1071,7 +1091,13 @@ static uint32_t parse_declaration(parser_t *P) return u; } - uint32_t type_node = parse_type_spec(P, &quals, &cuda); + uint32_t type_node; + if (is_ctor(P)) { + type_node = alloc_node(P, AST_TYPE_SPEC); + P->nodes[type_node].d.btype.kind = TYPE_VOID; + } else { + type_node = parse_type_spec(P, &quals, &cuda); + } if (!type_node) { parse_error(P, BC_E024); advance(P); @@ -1112,6 +1138,17 @@ static uint32_t parse_declaration(parser_t *P) if (!match(P, TOK_COMMA)) break; } } else { + /* Saved and restored so nested structs each get their own. + * Synthetic names live outside src, so memcmp must skip those. */ + uint32_t sv_off = P->cs_off; + uint16_t sv_len = P->cs_len; + uint32_t nm = P->nodes[type_node].first_child; + P->cs_len = 0; + if (nm && P->nodes[nm].type == AST_IDENT && + P->nodes[nm].d.text.offset < BC_ANON_BASE) { + P->cs_off = P->nodes[nm].d.text.offset; + P->cs_len = (uint16_t)P->nodes[nm].d.text.len; + } while (cur_type(P) != TOK_RBRACE && cur_type(P) != TOK_EOF) { if (cur_type(P) == TOK_PUBLIC || cur_type(P) == TOK_PRIVATE || cur_type(P) == TOK_PROTECTED) { @@ -1121,6 +1158,8 @@ static uint32_t parse_declaration(parser_t *P) uint32_t member = parse_declaration(P); if (member) add_child(P, def, member); } + P->cs_off = sv_off; + P->cs_len = sv_len; } expect(P, TOK_RBRACE); diff --git a/src/fe/parser.h b/src/fe/parser.h index 7c6ee3d..ef4038f 100644 --- a/src/fe/parser.h +++ b/src/fe/parser.h @@ -32,6 +32,11 @@ typedef struct { uint32_t anon_len; uint32_t anon_cnt; + /* Enclosing struct name, so a constructor can be told apart from a + * declaration that happens to start with a type name. len 0 = not in one. */ + uint32_t cs_off; + uint16_t cs_len; + bc_error_t errors[BC_MAX_ERRORS]; int num_errors; } parser_t; diff --git a/tests/ctor.cu b/tests/ctor.cu new file mode 100644 index 0000000..5d6aa8c --- /dev/null +++ b/tests/ctor.cu @@ -0,0 +1,20 @@ +/* Constructors and destructors, which carry no return type. */ + +struct vec3 { + float x, y, z; + vec3() { x = 0.0f; y = 0.0f; z = 0.0f; } + vec3(float v) { x = v; y = v; z = v; } + ~vec3() { } + float sum() { return x + y + z; } +}; + +struct outer { + struct inner { inner() { } }; + int n; + outer() { n = 0; } +}; + +__global__ void k(float *out) +{ + out[0] = 1.0f; +} diff --git a/tests/tphase.c b/tests/tphase.c index 3be8815..9b8f2d8 100644 --- a/tests/tphase.c +++ b/tests/tphase.c @@ -59,6 +59,20 @@ static void pha_fnptr(void) } TH_REG("phase", pha_fnptr) +/* Nested structs matter here: the enclosing name is saved and restored, + so inner's constructor must not be mistaken for outer's. */ +static void pha_ctor(void) +{ + int rc = th_run(BC_BIN " --parse tests/ctor.cu", obuf, TH_BUFSZ); + CHEQ(rc, 0); + CHECK(strstr(obuf, "0 parse error(s)") != NULL); + CHECK(strstr(obuf, "(ident vec3)") != NULL); + CHECK(strstr(obuf, "(ident inner)") != NULL); + CHECK(strstr(obuf, "(ident sum)") != NULL); + PASS(); +} +TH_REG("phase", pha_ctor) + /* ---- phase: IR ---- */ static void pha_ir(void)