diff options
Diffstat (limited to 'src/parser')
| -rw-r--r-- | src/parser/parser_core.c | 24 | ||||
| -rw-r--r-- | src/parser/parser_stmt.c | 111 |
2 files changed, 134 insertions, 1 deletions
diff --git a/src/parser/parser_core.c b/src/parser/parser_core.c index 3e683fb..c3c91fe 100644 --- a/src/parser/parser_core.c +++ b/src/parser/parser_core.c @@ -63,6 +63,9 @@ ASTNode *parse_program_nodes(ParserContext *ctx, Lexer *l) int attr_weak = 0; int attr_export = 0; int attr_comptime = 0; + int attr_cuda_global = 0; // @global -> __global__ + int attr_cuda_device = 0; // @device -> __device__ + int attr_cuda_host = 0; // @host -> __host__ char *deprecated_msg = NULL; char *attr_section = NULL; @@ -232,7 +235,23 @@ ASTNode *parse_program_nodes(ParserContext *ctx, Lexer *l) } else { - zwarn_at(attr, "Unknown attribute: %.*s", attr.len, attr.start); + // Checking for CUDA attributes... + if (0 == strncmp(attr.start, "global", 6) && 6 == attr.len) + { + attr_cuda_global = 1; + } + else if (0 == strncmp(attr.start, "device", 6) && 6 == attr.len) + { + attr_cuda_device = 1; + } + else if (0 == strncmp(attr.start, "host", 4) && 4 == attr.len) + { + attr_cuda_host = 1; + } + else + { + zwarn_at(attr, "Unknown attribute: %.*s", attr.len, attr.start); + } } t = lexer_peek(l); @@ -469,6 +488,9 @@ ASTNode *parse_program_nodes(ParserContext *ctx, Lexer *l) s->func.pure = attr_pure; s->func.section = attr_section; s->func.is_comptime = attr_comptime; + s->func.cuda_global = attr_cuda_global; + s->func.cuda_device = attr_cuda_device; + s->func.cuda_host = attr_cuda_host; if (attr_deprecated && s->func.name) { diff --git a/src/parser/parser_stmt.c b/src/parser/parser_stmt.c index 5307768..daf3f72 100644 --- a/src/parser/parser_stmt.c +++ b/src/parser/parser_stmt.c @@ -2503,6 +2503,117 @@ ASTNode *parse_statement(ParserContext *ctx, Lexer *l) return parse_guard(ctx, l); } + // CUDA launch: launch kernel(args) with { grid: X, block: Y }; + if (strncmp(tk.start, "launch", 6) == 0 && tk.len == 6) + { + Token launch_tok = lexer_next(l); // eat 'launch' + + // Parse the kernel call expression + ASTNode *call = parse_expression(ctx, l); + if (!call || call->type != NODE_EXPR_CALL) + { + zpanic_at(launch_tok, "Expected kernel call after 'launch'"); + } + + // Expect 'with' + Token with_tok = lexer_peek(l); + if (with_tok.type != TOK_IDENT || strncmp(with_tok.start, "with", 4) != 0 || + with_tok.len != 4) + { + zpanic_at(with_tok, "Expected 'with' after kernel call in launch statement"); + } + lexer_next(l); // eat 'with' + + // Expect '{' for configuration block + if (lexer_peek(l).type != TOK_LBRACE) + { + zpanic_at(lexer_peek(l), "Expected '{' after 'with' in launch statement"); + } + lexer_next(l); // eat '{' + + ASTNode *grid = NULL; + ASTNode *block = NULL; + ASTNode *shared_mem = NULL; + ASTNode *stream = NULL; + + // Parse configuration fields + while (lexer_peek(l).type != TOK_RBRACE && lexer_peek(l).type != TOK_EOF) + { + Token field_name = lexer_next(l); + if (field_name.type != TOK_IDENT) + { + zpanic_at(field_name, "Expected field name in launch configuration"); + } + + // Expect ':' + if (lexer_peek(l).type != TOK_COLON) + { + zpanic_at(lexer_peek(l), "Expected ':' after field name"); + } + lexer_next(l); // eat ':' + + // Parse value expression + ASTNode *value = parse_expression(ctx, l); + + // Assign to appropriate field + if (strncmp(field_name.start, "grid", 4) == 0 && field_name.len == 4) + { + grid = value; + } + else if (strncmp(field_name.start, "block", 5) == 0 && field_name.len == 5) + { + block = value; + } + else if (strncmp(field_name.start, "shared_mem", 10) == 0 && field_name.len == 10) + { + shared_mem = value; + } + else if (strncmp(field_name.start, "stream", 6) == 0 && field_name.len == 6) + { + stream = value; + } + else + { + zpanic_at(field_name, "Unknown launch configuration field (expected: grid, " + "block, shared_mem, stream)"); + } + + // Optional comma + if (lexer_peek(l).type == TOK_COMMA) + { + lexer_next(l); + } + } + + // Expect '}' + if (lexer_peek(l).type != TOK_RBRACE) + { + zpanic_at(lexer_peek(l), "Expected '}' to close launch configuration"); + } + lexer_next(l); // eat '}' + + // Expect ';' + if (lexer_peek(l).type == TOK_SEMICOLON) + { + lexer_next(l); + } + + // Require at least grid and block + if (!grid || !block) + { + zpanic_at(launch_tok, "Launch configuration requires at least 'grid' and 'block'"); + } + + ASTNode *n = ast_create(NODE_CUDA_LAUNCH); + n->cuda_launch.call = call; + n->cuda_launch.grid = grid; + n->cuda_launch.block = block; + n->cuda_launch.shared_mem = shared_mem; + n->cuda_launch.stream = stream; + n->token = launch_tok; + return n; + } + // Do-while loop: do { body } while condition; if (strncmp(tk.start, "do", 2) == 0 && tk.len == 2) { |