// IR optimisation (constant folding), IR printing, and WASM code generation. #include "kernel.h" #include "../wasm/wasm.h" #include #include #include namespace unify::kernel { namespace { float as_f(uint32_t b) { float f; std::memcpy(&f, &b, 4); return f; } uint32_t fb(float f) { uint32_t b; std::memcpy(&b, &f, 4); return b; } // ---------------------------------------------------------------------------- folding bool fold_expr(IRExprP& e, uint32_t& count) { for (auto& a : e->args) fold_expr(a, count); auto is_const = [](const IRExprP& x) { return x->kind == IRExpr::Const; }; if (e->kind == IRExpr::Convert && is_const(e->args[0])) { IRExpr& a = *e->args[0]; uint32_t bits; if (e->type == IRType::F32) bits = fb(float(int32_t(a.bits))); else { float f = as_f(a.bits); if (!(f > -2147483648.0f && f < 2147483648.0f)) return false; // would trap at runtime; keep it bits = uint32_t(int32_t(f)); } auto c = std::make_unique(); c->kind = IRExpr::Const; c->type = e->type; c->bits = bits; e = std::move(c); count++; return true; } if (e->kind == IRExpr::Binary && is_const(e->args[0]) && is_const(e->args[1])) { const std::string& op = e->op; IRType operand = e->args[0]->type; uint32_t r = 0; if (operand == IRType::I32) { int32_t a = int32_t(e->args[0]->bits), b = int32_t(e->args[1]->bits); if ((op == "/" || op == "%") && (b == 0 || (a == INT32_MIN && b == -1))) return false; // runtime trap semantics if (op == "+") r = uint32_t(a) + uint32_t(b); else if (op == "-") r = uint32_t(a) - uint32_t(b); else if (op == "*") r = uint32_t(a) * uint32_t(b); else if (op == "/") r = uint32_t(a / b); else if (op == "%") r = uint32_t(a % b); else if (op == "<") r = a < b; else if (op == "<=") r = a <= b; else if (op == ">") r = a > b; else if (op == ">=") r = a >= b; else if (op == "==") r = a == b; else if (op == "!=") r = a != b; else if (op == "&&") r = (a != 0) & (b != 0); else if (op == "||") r = (a != 0) | (b != 0); else return false; } else { float a = as_f(e->args[0]->bits), b = as_f(e->args[1]->bits); if (op == "+") r = fb(a + b); else if (op == "-") r = fb(a - b); else if (op == "*") r = fb(a * b); else if (op == "/") r = fb(a / b); else if (op == "<") r = a < b; else if (op == "<=") r = a <= b; else if (op == ">") r = a > b; else if (op == ">=") r = a >= b; else if (op == "==") r = a == b; else if (op == "!=") r = a != b; else return false; } auto c = std::make_unique(); c->kind = IRExpr::Const; c->type = e->type; c->bits = r; e = std::move(c); count++; return true; } if (e->kind == IRExpr::Unary && is_const(e->args[0])) { uint32_t v = e->args[0]->bits, r; if (e->op == "!") r = v == 0; else if (e->type == IRType::I32) r = uint32_t(-int32_t(v)); else r = fb(-as_f(v)); auto c = std::make_unique(); c->kind = IRExpr::Const; c->type = e->type; c->bits = r; e = std::move(c); count++; return true; } return false; } void fold_block(std::vector& body, uint32_t& count) { std::vector out; for (auto& s : body) { if (s->value) fold_expr(s->value, count); if (s->index) fold_expr(s->index, count); if (s->cond) fold_expr(s->cond, count); fold_block(s->then_body, count); fold_block(s->else_body, count); fold_block(s->body, count); fold_block(s->step, count); // Dead-branch elimination for constant conditions. if (s->kind == IRStmt::If && s->cond->kind == IRExpr::Const) { auto& taken = s->cond->bits ? s->then_body : s->else_body; for (auto& t : taken) out.push_back(std::move(t)); count++; continue; } if (s->kind == IRStmt::Loop && s->cond->kind == IRExpr::Const && s->cond->bits == 0) { count++; continue; } out.push_back(std::move(s)); } body = std::move(out); } // ---------------------------------------------------------------------------- printing const char* tn(IRType t) { return t == IRType::F32 ? "f32" : "i32"; } void print_expr(std::ostream& o, const IRExpr& e, const IRFunction& fn) { auto local_name = [&](uint32_t i) -> std::string { static const char* b[] = {"threadIdx.x", "blockIdx.x", "blockDim.x", "gridDim.x"}; if (i < 4) return b[i]; if (i < fn.param_types.size()) return "%p" + std::to_string(i - 4); return "%" + fn.local_names[i - fn.param_types.size()]; }; switch (e.kind) { case IRExpr::Const: if (e.type == IRType::F32) o << as_f(e.bits) << "f"; else o << int32_t(e.bits); break; case IRExpr::Local: o << local_name(e.local); break; case IRExpr::Load: o << "load." << tn(e.type) << " " << local_name(e.local) << "["; print_expr(o, *e.args[0], fn); o << "]"; break; case IRExpr::Binary: o << "("; print_expr(o, *e.args[0], fn); o << " " << e.op << "." << tn(e.args[0]->type) << " "; print_expr(o, *e.args[1], fn); o << ")"; break; case IRExpr::Unary: o << e.op << "("; print_expr(o, *e.args[0], fn); o << ")"; break; case IRExpr::Convert: o << "convert." << tn(e.type) << "("; print_expr(o, *e.args[0], fn); o << ")"; break; case IRExpr::Intrinsic: o << e.op << "("; for (size_t i = 0; i < e.args.size(); i++) { if (i) o << ", "; print_expr(o, *e.args[i], fn); } o << ")"; break; } } void print_block(std::ostream& o, const std::vector& body, const IRFunction& fn, int indent) { std::string pad(size_t(indent) * 2, ' '); for (auto& s : body) { switch (s->kind) { case IRStmt::SetLocal: { IRExpr l; l.kind = IRExpr::Local; l.local = s->local; o << pad; print_expr(o, l, fn); o << " = "; print_expr(o, *s->value, fn); o << "\n"; break; } case IRStmt::Store: { IRExpr l; l.kind = IRExpr::Local; l.local = s->local; o << pad << "store." << tn(s->type) << " "; print_expr(o, l, fn); o << "["; print_expr(o, *s->index, fn); o << "] = "; print_expr(o, *s->value, fn); o << "\n"; break; } case IRStmt::If: o << pad << "if "; print_expr(o, *s->cond, fn); o << " {\n"; print_block(o, s->then_body, fn, indent + 1); if (!s->else_body.empty()) { o << pad << "} else {\n"; print_block(o, s->else_body, fn, indent + 1); } o << pad << "}\n"; break; case IRStmt::Loop: o << pad << "loop while "; print_expr(o, *s->cond, fn); o << " {\n"; print_block(o, s->body, fn, indent + 1); if (!s->step.empty()) { o << pad << "} step {\n"; print_block(o, s->step, fn, indent + 1); } o << pad << "}\n"; break; case IRStmt::Return: o << pad << "return\n"; break; case IRStmt::Break: o << pad << "break\n"; break; case IRStmt::Continue: o << pad << "continue\n"; break; } } } // ---------------------------------------------------------------------------- WASM emission using wasm::write_sleb; using wasm::write_uleb; struct Emitter { std::vector code; // Branch depths: how many structured blocks deep we are relative to the innermost loop's // break (outer block) and continue (inner block) targets. std::vector break_depth, continue_depth; uint32_t depth = 0; void op(uint8_t b) { code.push_back(b); } void local_get(uint32_t i) { op(0x20); write_uleb(code, i); } void local_set(uint32_t i) { op(0x21); write_uleb(code, i); } void i32_const(int32_t v) { op(0x41); write_sleb(code, v); } void f32_const(uint32_t bits) { op(0x43); for (int i = 0; i < 4; i++) code.push_back(uint8_t(bits >> (8 * i))); } void begin(uint8_t kind) { op(kind); op(0x40); depth++; } void end() { op(0x0B); depth--; } void address(uint32_t base_local, const IRExpr& index) { local_get(base_local); expr(index); i32_const(4); op(0x6C); // i32.mul op(0x6A); // i32.add } void expr(const IRExpr& e) { bool f = e.kind == IRExpr::Binary ? e.args[0]->type == IRType::F32 : e.type == IRType::F32; switch (e.kind) { case IRExpr::Const: if (e.type == IRType::F32) f32_const(e.bits); else i32_const(int32_t(e.bits)); break; case IRExpr::Local: local_get(e.local); break; case IRExpr::Load: address(e.local, *e.args[0]); op(e.type == IRType::F32 ? 0x2A : 0x28); write_uleb(code, 2); write_uleb(code, 0); break; case IRExpr::Binary: { expr(*e.args[0]); expr(*e.args[1]); const std::string& o = e.op; if (o == "&&") { op(0x71); break; } if (o == "||") { op(0x72); break; } if (f) { if (o == "+") op(0x92); else if (o == "-") op(0x93); else if (o == "*") op(0x94); else if (o == "/") op(0x95); else if (o == "==") op(0x5B); else if (o == "!=") op(0x5C); else if (o == "<") op(0x5D); else if (o == ">") op(0x5E); else if (o == "<=") op(0x5F); else if (o == ">=") op(0x60); } else { if (o == "+") op(0x6A); else if (o == "-") op(0x6B); else if (o == "*") op(0x6C); else if (o == "/") op(0x6D); else if (o == "%") op(0x6F); else if (o == "==") op(0x46); else if (o == "!=") op(0x47); else if (o == "<") op(0x48); else if (o == ">") op(0x4A); else if (o == "<=") op(0x4C); else if (o == ">=") op(0x4E); } break; } case IRExpr::Unary: if (e.op == "!") { expr(*e.args[0]); op(0x45); } else if (e.type == IRType::F32) { expr(*e.args[0]); op(0x8C); } else { i32_const(0); expr(*e.args[0]); op(0x6B); } break; case IRExpr::Convert: expr(*e.args[0]); op(e.type == IRType::F32 ? 0xB2 : 0xA8); // f32.convert_i32_s / i32.trunc_f32_s break; case IRExpr::Intrinsic: { const std::string& n = e.op; if (e.type == IRType::I32 && (n == "min" || n == "max")) { // select(a, b, a < b) for min; a > b for max. Operands are side-effect free. expr(*e.args[0]); expr(*e.args[1]); expr(*e.args[0]); expr(*e.args[1]); op(n == "min" ? 0x48 : 0x4A); op(0x1B); break; } for (auto& a : e.args) expr(*a); if (n == "sqrtf") op(0x91); else if (n == "fabsf") op(0x8B); else if (n == "floorf") op(0x8E); else if (n == "ceilf") op(0x8D); else if (n == "fminf" || n == "min") op(0x96); else if (n == "fmaxf" || n == "max") op(0x97); break; } } } void stmts(const std::vector& body) { for (auto& s : body) stmt(*s); } void stmt(const IRStmt& s) { switch (s.kind) { case IRStmt::SetLocal: expr(*s.value); local_set(s.local); break; case IRStmt::Store: address(s.local, *s.index); expr(*s.value); op(s.type == IRType::F32 ? 0x38 : 0x36); write_uleb(code, 2); write_uleb(code, 0); break; case IRStmt::If: expr(*s.cond); begin(0x04); stmts(s.then_body); if (!s.else_body.empty()) { op(0x05); stmts(s.else_body); } end(); break; case IRStmt::Loop: { // block $break { loop $top { br_if $break !cond; block $continue { body } step; br $top } } begin(0x02); break_depth.push_back(depth); begin(0x03); uint32_t top = depth; expr(*s.cond); op(0x45); op(0x0D); write_uleb(code, depth - break_depth.back()); begin(0x02); continue_depth.push_back(depth); stmts(s.body); continue_depth.pop_back(); end(); stmts(s.step); op(0x0C); write_uleb(code, depth - top); end(); break_depth.pop_back(); end(); break; } case IRStmt::Return: op(0x0F); break; case IRStmt::Break: op(0x0C); write_uleb(code, depth - break_depth.back()); break; case IRStmt::Continue: op(0x0C); write_uleb(code, depth - continue_depth.back()); break; } } }; void section(std::vector& out, uint8_t id, const std::vector& body) { out.push_back(id); write_uleb(out, body.size()); out.insert(out.end(), body.begin(), body.end()); } void name(std::vector& out, const std::string& s) { write_uleb(out, s.size()); out.insert(out.end(), s.begin(), s.end()); } uint8_t vt(IRType t) { return t == IRType::F32 ? 0x7D : 0x7F; } } // namespace uint32_t fold_constants(IRFunction& fn) { uint32_t count = 0; fold_block(fn.body, count); return count; } std::string print_ir(const IRFunction& fn) { std::ostringstream o; o << "kernel " << fn.name << "("; for (size_t i = 4; i < fn.param_types.size(); i++) o << (i > 4 ? ", " : "") << "%p" << (i - 4) << ": " << tn(fn.param_types[i]); o << ") {\n"; for (size_t i = 0; i < fn.local_types.size(); i++) o << " local %" << fn.local_names[i] << ": " << tn(fn.local_types[i]) << "\n"; print_block(o, fn.body, fn, 1); o << "}\n"; return o.str(); } std::vector emit_wasm(const IRFunction& fn) { std::vector m = {0x00, 'a', 's', 'm', 0x01, 0x00, 0x00, 0x00}; const size_t user = fn.param_types.size() - 4; // Types: 0 = kernel(threadIdx, blockIdx, blockDim, gridDim, params...), 1 = dispatch(grid, block, params...) std::vector types; write_uleb(types, 2); types.push_back(0x60); write_uleb(types, fn.param_types.size()); for (IRType t : fn.param_types) types.push_back(vt(t)); write_uleb(types, 0); types.push_back(0x60); write_uleb(types, 2 + user); types.push_back(0x7F); types.push_back(0x7F); for (size_t i = 4; i < fn.param_types.size(); i++) types.push_back(vt(fn.param_types[i])); write_uleb(types, 0); section(m, 1, types); // Import the shared linear memory so every kernel sees the same buffers. std::vector imports; write_uleb(imports, 1); name(imports, "env"); name(imports, "memory"); imports.push_back(0x02); imports.push_back(0x00); write_uleb(imports, 1); section(m, 2, imports); std::vector funcs; write_uleb(funcs, 2); write_uleb(funcs, 0); write_uleb(funcs, 1); section(m, 3, funcs); std::vector exports; write_uleb(exports, 2); name(exports, fn.name); exports.push_back(0x00); write_uleb(exports, 0); name(exports, fn.name + "__dispatch"); exports.push_back(0x00); write_uleb(exports, 1); section(m, 7, exports); std::vector code; write_uleb(code, 2); { // kernel body Emitter e; std::vector body; write_uleb(body, fn.local_types.size()); for (IRType t : fn.local_types) { write_uleb(body, 1); body.push_back(vt(t)); } e.stmts(fn.body); e.op(0x0B); body.insert(body.end(), e.code.begin(), e.code.end()); write_uleb(code, body.size()); code.insert(code.end(), body.begin(), body.end()); } { // dispatch: for b in [0, grid) for t in [0, block) kernel(t, b, block, grid, params...) Emitter e; const uint32_t grid = 0, block = 1, b = uint32_t(2 + user), t = b + 1; std::vector body; write_uleb(body, 1); write_uleb(body, 2); body.push_back(0x7F); e.i32_const(0); e.local_set(b); e.begin(0x02); e.begin(0x03); e.local_get(b); e.local_get(grid); e.op(0x4F); e.op(0x0D); write_uleb(e.code, 1); // br_if b >= grid (unsigned) e.i32_const(0); e.local_set(t); e.begin(0x02); e.begin(0x03); e.local_get(t); e.local_get(block); e.op(0x4F); e.op(0x0D); write_uleb(e.code, 1); e.local_get(t); e.local_get(b); e.local_get(block); e.local_get(grid); for (uint32_t i = 0; i < user; i++) e.local_get(2 + i); e.op(0x10); write_uleb(e.code, 0); e.local_get(t); e.i32_const(1); e.op(0x6A); e.local_set(t); e.op(0x0C); write_uleb(e.code, 0); e.end(); e.end(); e.local_get(b); e.i32_const(1); e.op(0x6A); e.local_set(b); e.op(0x0C); write_uleb(e.code, 0); e.end(); e.end(); e.op(0x0B); body.insert(body.end(), e.code.begin(), e.code.end()); write_uleb(code, body.size()); code.insert(code.end(), body.begin(), body.end()); } section(m, 10, code); return m; } CompiledKernel compile(const std::string& source) { CompiledKernel out; KernelAST ast; if (!parse(source, ast, out.error)) return out; IRFunction fn; if (!lower(ast, fn, out.error)) return out; out.folded_constants = fold_constants(fn); out.name = fn.name; for (auto& p : ast.params) out.params.push_back({p.type, p.name}); out.ir_text = print_ir(fn); out.wasm = emit_wasm(fn); out.ok = true; return out; } } // namespace unify::kernel