From d7f447c36e66ff4e98b2c42bdff4e30f8eed7460 Mon Sep 17 00:00:00 2001 From: Chen-Yu Yang Date: Mon, 17 Aug 2026 07:39:42 -0400 Subject: [PATCH] consts are weak [pr] --- test/amd/test_emu2_pcode.py | 36 +++---- test/backend/test_const_folding.py | 6 +- test/backend/test_dtype_alu.py | 24 +++++ test/backend/test_isel.py | 5 +- test/backend/test_linearizer.py | 12 +-- test/mockgpu/amd/emu.py | 2 +- test/mockgpu/amd/pcode.py | 46 ++++----- test/null/test_const_folding.py | 10 +- test/null/test_graph_rewrite.py | 55 +++++------ test/null/test_simplify_valid_idx.py | 18 ++-- test/null/test_uop_graph.py | 52 ++++++----- test/null/test_uop_symbolic.py | 6 +- test/null/test_uops.py | 40 +++++--- test/null/test_viz.py | 6 +- test/unit/test_dtype_weak.py | 16 ++-- tinygrad/codegen/__init__.py | 17 ++-- tinygrad/codegen/decomp/dtype.py | 12 ++- tinygrad/codegen/decomp/op.py | 7 +- tinygrad/codegen/late/coalesce.py | 8 ++ tinygrad/codegen/late/linearizer.py | 3 +- tinygrad/codegen/late/regalloc.py | 3 +- tinygrad/renderer/cstyle.py | 60 ++++++------ tinygrad/renderer/isa/x86.py | 93 ++++++++++-------- tinygrad/renderer/llvmir.py | 6 +- tinygrad/renderer/nir.py | 15 ++- tinygrad/renderer/ptx.py | 13 +-- tinygrad/renderer/wgsl.py | 10 +- tinygrad/runtime/ops_python.py | 5 +- tinygrad/runtime/support/hcq2.py | 12 +-- tinygrad/schedule/rangeify.py | 4 +- tinygrad/tensor.py | 3 +- tinygrad/uop/ops.py | 29 +++--- tinygrad/uop/render.py | 9 +- tinygrad/uop/spec.py | 21 +++-- tinygrad/uop/symbolic.py | 66 +++++++------ tinygrad/uop/validate.py | 4 +- tinygrad/uop/weak.py | 135 ++++++++++++++++----------- 37 files changed, 508 insertions(+), 361 deletions(-) diff --git a/test/amd/test_emu2_pcode.py b/test/amd/test_emu2_pcode.py index 29dcfefdff..a83164028a 100644 --- a/test/amd/test_emu2_pcode.py +++ b/test/amd/test_emu2_pcode.py @@ -67,32 +67,32 @@ class TestParseExpr(unittest.TestCase): def test_integer_literals(self): """Test parsing integer literals.""" - self.assertEqual(parse_expr('0', {}).val, 0) - self.assertEqual(parse_expr('42', {}).val, 42) - self.assertEqual(parse_expr('42U', {}).val, 42) + self.assertEqual(parse_expr('0', {}).src[0].val, 0) + self.assertEqual(parse_expr('42', {}).src[0].val, 42) + self.assertEqual(parse_expr('42U', {}).src[0].val, 42) def test_negative_integers(self): """Test parsing negative integer literals.""" result = parse_expr('-1', {}) - self.assertEqual(result.val, -1) + self.assertEqual(result.src[0].val, -1) self.assertEqual(result.dtype, dtypes.int) def test_float_literals(self): """Test parsing float literals.""" result = parse_expr('1.0F', {}) - self.assertEqual(result.val, 1.0) + self.assertEqual(result.src[0].val, 1.0) self.assertEqual(result.dtype, dtypes.float32) def test_hex_literals(self): """Test parsing hex literals.""" result = parse_expr('0xFF', {}) - self.assertEqual(result.val, 255) + self.assertEqual(result.src[0].val, 255) def test_variable_lookup(self): """Test variable lookup in parse_expr.""" vrs = {'x': UOp.const(42, dtypes.uint32)} result = parse_expr('x', vrs) - self.assertEqual(result.val, 42) + self.assertEqual(result.src[0].val, 42) def test_binary_ops(self): """Test parsing binary operations.""" @@ -104,8 +104,8 @@ class TestParseExpr(unittest.TestCase): # Subtraction with constant folding result = parse_expr('10 - 5', {}) - self.assertEqual(result.op, Ops.CONST) - self.assertEqual(result.val, 5) + self.assertTrue(result.op is Ops.CAST and result.src[0].op is Ops.CONST) + self.assertEqual(result.src[0].val, 5) def test_ternary(self): """Test parsing ternary expressions.""" @@ -262,8 +262,8 @@ class TestDSPcodePatterns(unittest.TestCase): _, assigns = parse_pcode(pcode, srcs) # Check addresses: 100 + 2*4 = 108, 100 + 5*4 = 120 # assigns[i][1] is (addr, val) tuple for MEM writes; mypy sees UOp - self.assertEqual(assigns[0][1][0].simplify().val, 108) # type: ignore[index] - self.assertEqual(assigns[1][1][0].simplify().val, 120) # type: ignore[index] + self.assertEqual(assigns[0][1][0].simplify().src[0].val, 108) # type: ignore[index] + self.assertEqual(assigns[1][1][0].simplify().src[0].val, 120) # type: ignore[index] def test_ds_store_data_values(self): """Test DS_STORE_2ADDR_B32 uses correct data values.""" @@ -280,8 +280,8 @@ class TestDSPcodePatterns(unittest.TestCase): _, assigns = parse_pcode(pcode, srcs) # assigns[i][1] is (addr, val) tuple for MEM writes; mypy sees UOp # DATA[31:0] should preserve the value - self.assertEqual(assigns[0][1][1].simplify().val, 0xAAAAAAAA) # type: ignore[index] - self.assertEqual(assigns[1][1][1].simplify().val, 0xBBBBBBBB) # type: ignore[index] + self.assertEqual(assigns[0][1][1].simplify().src[0].val, 0xAAAAAAAA) # type: ignore[index] + self.assertEqual(assigns[1][1][1].simplify().src[0].val, 0xBBBBBBBB) # type: ignore[index] class TestConditionalParsing(unittest.TestCase): """Test conditional (if/elsif/else) pcode parsing.""" @@ -306,12 +306,12 @@ class TestConcatWidthParsing(unittest.TestCase): def test_permlanex16_altrow_concat(self): for row, expected in [(0, 1), (1, 0), (2, 3), (3, 2)]: parsed = parse_expr('{ row[1], ~row[0] }', {'row': UOp.const(row, dtypes.uint32)}) - self.assertEqual(parsed.simplify().val, expected) + self.assertEqual(parsed.simplify().src[0].val, expected) def test_permlane64_altlane_concat(self): for lane, expected in [(0, 32), (1, 33), (31, 63), (32, 0), (63, 31)]: parsed = parse_expr('{ ~lane[5], lane[4:0] }', {'lane': UOp.const(lane, dtypes.uint32)}) - self.assertEqual(parsed.simplify().val, expected) + self.assertEqual(parsed.simplify().src[0].val, expected) def test_permlane64_wave64_pcode_indices(self): vgpr = UOp.param(0, dtypes.uint32, (256,)) @@ -332,13 +332,13 @@ class TestConcatWidthParsing(unittest.TestCase): self.assertEqual(simp.op, Ops.LOAD) self.assertEqual(simp.src[0].op, Ops.INDEX) idx = simp.src[0].src[1].simplify() - self.assertEqual(idx.op, Ops.CONST) - return idx.val + self.assertTrue(idx.op is Ops.CAST and idx.src[0].op is Ops.CONST) + return idx.src[0].val _, assigns = parse_pcode(PCODE[VOP1Op.V_PERMLANE64_B32_E32], srcs) self.assertEqual(len(assigns), 64) for lane, (dst_idx, src_idx) in {0: (64, 32), 31: (95, 63), 32: (96, 0), 63: (127, 31)}.items(): - self.assertEqual(assigns[lane][1][0].simplify().val, dst_idx) # type: ignore[index] + self.assertEqual(assigns[lane][1][0].simplify().src[0].val, dst_idx) # type: ignore[index] self.assertEqual(load_idx(assigns[lane][1][1]), src_idx) # type: ignore[index] class TestAllPcode(unittest.TestCase): diff --git a/test/backend/test_const_folding.py b/test/backend/test_const_folding.py index 0c2c801560..15f9d01928 100644 --- a/test/backend/test_const_folding.py +++ b/test/backend/test_const_folding.py @@ -17,7 +17,7 @@ def _check_ast_count(desired_count:int, t:Tensor): class TestMovedConstFolding(unittest.TestCase): def test_contiguous_deviceless_const(self): t = Tensor(UOp.const(2.0, dtypes.float)).contiguous() - self.assertIs(t.uop.op, Ops.CONST) + self.assertTrue(t.uop.op is Ops.CAST and t.uop.src[0].op is Ops.CONST) self.assertIsNone(t.uop.device) def test_add_shrunk_zero(self): @@ -169,8 +169,8 @@ class TestMultiConstFolding(unittest.TestCase): class TestThreefryConstFolding(unittest.TestCase): def test_threefry(self): # THREEFRY(const,const) folds to a const once decomposed - x = threefry2x32(UOp.const(5, dtypes.uint64), UOp.const(10, dtypes.uint64)) - self.assertIs(x.simplify().op, Ops.CONST) + x = threefry2x32(UOp.const(5, dtypes.uint64), UOp.const(10, dtypes.uint64)).simplify() + self.assertTrue(x.op is Ops.CAST and x.src[0].op is Ops.CONST) class TestTautologicalCompare(unittest.TestCase): # without const folding, these would have triggered -Wtautological-compare in clang diff --git a/test/backend/test_dtype_alu.py b/test/backend/test_dtype_alu.py index 6a182616c8..34f8c66b0b 100644 --- a/test/backend/test_dtype_alu.py +++ b/test/backend/test_dtype_alu.py @@ -7,6 +7,8 @@ from tinygrad.runtime.ops_python import from_storage_scalar from tinygrad.renderer.ptx import PTXRenderer from tinygrad.renderer.nir import NIRRenderer from tinygrad.uop import Ops +from tinygrad.codegen import to_program +from tinygrad.renderer.isa import ISARenderer import numpy as np import pytest from hypothesis import assume, given, strategies as strat, settings @@ -399,6 +401,28 @@ class TestDTypeALU(unittest.TestCase): if float_dtype not in supported_dtypes: float_dtype = dtypes.float32 universal_test_cast(a, float_dtype, unsigned_dtype) + # an ISA renderer is excluded because its program is machine instructions, so no Ops.FDIV survives to be counted. + # only the op count is asserted: the value is not bit-exact anywhere, LLVM emits fdiv with arcp/afn and AMD then + # lowers it to an unrefined v_rcp, which is 1 ulp off the correctly rounded quotient + @unittest.skipIf(isinstance(Device[Device.DEFAULT].renderer, ISARenderer) or + Ops.FDIV not in Device[Device.DEFAULT].renderer.code_for_op, "renderer divides by reciprocal") + def test_float_div_is_not_a_reciprocal_multiply(self): + # a*(1/b) is 1-2 ulp worse than a/b, so the contraction must fire even with the bare 1 the decomp now mints + def ops(t): return [u.op for u in to_program(t.contiguous().schedule_linear().src[-1].src[0], Device[Device.DEFAULT].renderer).src[1].src] + strong = ops(Tensor([1e10], dtype=dtypes.float32)/Tensor([1.5], dtype=dtypes.float32)) + self.assertEqual((strong.count(Ops.FDIV), strong.count(Ops.MUL)), (1, 0)) + # with a weak numerator the 2.0 rides in bare, so the contraction sees a different shape on the way in + weak = ops(2.0/Tensor([1.5], dtype=dtypes.float32)) + self.assertEqual((weak.count(Ops.FDIV), weak.count(Ops.MUL)), (1, 0)) + + def test_out_of_range_unsigned_literal_keeps_its_width(self): + # -1000000 stays MATHEMATICAL even though it does not fit uint64 (wrapping to 2**64-1000000 would miscompile + # in-range programs), and the width comes from the strong sibling's promotion -- a uint64 multiply, not a bounds default + b = Tensor([True], dtype=dtypes.bool).contiguous() + out_of_range = (b*Tensor(-1000000, dtype=dtypes.uint64)+1e10).item() + self.assertEqual(out_of_range, (b*Tensor(2**64-1000000, dtype=dtypes.uint64)+1e10).item()) + self.assertEqual(out_of_range, 1.8446744073709552e+19) + def test_unsafe_cast_float_to_int(self): # the value is off the float32 grid but rounds in-range: the buffer and const-fold paths must agree # (out-of-range float->int cast stays undefined: hardware may saturate where the fold wraps) diff --git a/test/backend/test_isel.py b/test/backend/test_isel.py index ed38be7c3b..6129a71e5e 100644 --- a/test/backend/test_isel.py +++ b/test/backend/test_isel.py @@ -46,10 +46,11 @@ class TestIselX86(unittest.TestCase): # complex address is [base + index*scale + displacement] def test_complex_address(self): a = UOp.variable("a", 0, 0, dtypes.int32) - load = UOp.param(0, dtypes.int32, (16,)).index(a + 1).load() + # a program states every width, so the offset that folds into the displacement is a literal, not a bare const + load = UOp.param(0, dtypes.int32, (16,)).index(a + UOp.const(1, dtypes.int32)).load() n = self.isel_rewrite(load) # displacement is the constant in "a" scaled to the buffer element size, dtype is int8 when the value fits otherwise int32 - self.assertTrue(n.src[2].op is Ops.CONST and n.src[2].dtype is dtypes.int8 and n.src[2].val == 4) + self.assertTrue(n.src[2].op is Ops.CAST and n.src[2].dtype is dtypes.int8 and n.src[2].src[0].val == 4) if __name__ == "__main__": unittest.main() diff --git a/test/backend/test_linearizer.py b/test/backend/test_linearizer.py index 0f36811baa..bc76afb5b9 100644 --- a/test/backend/test_linearizer.py +++ b/test/backend/test_linearizer.py @@ -16,7 +16,7 @@ from test.helpers import replace_opts, check_schedule from test.backend.test_softmax_fusion import single_kernel_softmax MOCKGPU = DEV.interface.startswith("MOCK") -from tinygrad.uop.render import print_uops # noqa: F401 # pylint: disable=unused-import +from tinygrad.uop.render import print_uops @unittest.skipIf(isinstance(Device[Device.DEFAULT].renderer, ISARenderer), "isa backends don't preserve the op spec when lowering") class TestLinearizer(unittest.TestCase): @@ -248,11 +248,11 @@ class TestLinearizer(unittest.TestCase): uops = tuple(to_program(replace_opts(ast, opt), renderer=Device[Device.DEFAULT].renderer).src[1].src) begin_range = [i for i, x in enumerate(uops) if x.op is Ops.RANGE][-1] end_range = [i for i, x in enumerate(uops) if x.op is Ops.END][0] - for i,u in enumerate(uops): print(i, u.op, [uops.index(s) for s in u.src], u.arg, u.dtype) + print_uops(list(uops)) for u in uops: if u.op is Ops.STORE and u.src[0].addrspace is AddrSpace.REG: if uops.index(u) < begin_range: - assert u.src[1].op is Ops.CONST + assert u.src[1].op is Ops.CAST and u.src[1].src[0].op is Ops.CONST else: assert u.src[1].op in GroupOp.ALU assert begin_range < uops.index(u) < end_range @@ -268,9 +268,9 @@ class TestLinearizer(unittest.TestCase): uops = tuple(to_program(replace_opts(ast, []), renderer=Device[Device.DEFAULT].renderer).src[1].src) idxs = dedup([uop for uop in uops if uop.op is Ops.SPECIAL]) idxs = sorted(idxs, key=lambda uop: uop.arg) - assert (idxs[0].arg, idxs[0].src[0].val) == ('gidx0', 6), idxs[0] - assert (idxs[1].arg, idxs[1].src[0].val) == ('gidx1', 5), idxs[1].arg - assert (idxs[2].arg, idxs[2].src[0].val) == ('gidx2', 4), idxs[2].arg + assert (idxs[0].arg, idxs[0].src[0].src[0].val) == ('gidx0', 6), idxs[0] + assert (idxs[1].arg, idxs[1].src[0].src[0].val) == ('gidx1', 5), idxs[1].arg + assert (idxs[2].arg, idxs[2].src[0].src[0].val) == ('gidx2', 4), idxs[2].arg def test_sum_collapse(self): t = Tensor([2]).reshape(1, 1).expand(256, 256).sum() diff --git a/test/mockgpu/amd/emu.py b/test/mockgpu/amd/emu.py index 7c6fe9038d..6b66778006 100644 --- a/test/mockgpu/amd/emu.py +++ b/test/mockgpu/amd/emu.py @@ -709,7 +709,7 @@ class _Ctx: # VGPR bit-slice assignment: VGPR[lane][reg][hi:lo] = (vgpr_idx, rhs_val, hi, lo[, cond]) -> read-modify-write if dest.startswith('VGPR[') and re.search(r'\[\d+:\d+\]', dest): # VGPR bit-slice: (vgpr_idx, rhs_val, hi_bit, lo_bit) - hi/lo are UOp constants - hi_bit, lo_bit = int(val[2].val), int(val[3].val) + hi_bit, lo_bit = int(val[2].src[0].val), int(val[3].src[0].val) width = hi_bit - lo_bit + 1 old = self.vgpr.index(val[0]).load() new_val = _set_bits(old, _val_to_bits(val[1]), width, lo_bit).cast(dtypes.uint32) diff --git a/test/mockgpu/amd/pcode.py b/test/mockgpu/amd/pcode.py index bb66ffe03f..b6af2d9e52 100644 --- a/test/mockgpu/amd/pcode.py +++ b/test/mockgpu/amd/pcode.py @@ -8,6 +8,10 @@ from tinygrad.codegen.decomp.dtype import f2f VarVal = UOp | tuple[str, list[str], str] def _const(dt, v): return UOp.const(v, dt) +# every literal this parser mints is the pair CAST(dt, CONST(v)): _lit is its value at the CAST's width, or None +def _lit(u:UOp): return u.src[0].val if u.op is Ops.CAST and u.src[0].op is Ops.CONST else None +# restating a literal's width RE-MINTS it: casting the pair would stack two width statements on one literal +def _restate(u:UOp, dt): return _const(dt, v) if (v:=_lit(u)) is not None else u.cast(dt) def _u32(v): return _const(dtypes.uint32, v) def _u64(v): return _const(dtypes.uint64, v) def _to_u32(v): return v if v.dtype == dtypes.uint32 else v.bitcast(dtypes.uint32) if v.dtype.itemsize == 4 else v.cast(dtypes.uint32) @@ -55,8 +59,7 @@ def _expr_bits(v: UOp) -> int: if v.op in (Ops.AND, Ops.XOR): widths: list[int] = [] for src in v.src: - if src.op == Ops.CONST and isinstance(src.val, int) and src.val > 0 and (src.val & (src.val + 1)) == 0: - widths.append(src.val.bit_length()) + if isinstance(sv:=_lit(src), int) and sv > 0 and (sv & (sv + 1)) == 0: widths.append(sv.bit_length()) if widths: return max(widths) return v.dtype.bitsize @@ -144,9 +147,9 @@ def _minmax_reduce(is_max: bool, dt, *args: UOp) -> UOp: def _find_two_pi_mul(x): if x.op != Ops.MUL or len(x.src) != 2: return None for i, s in enumerate(x.src): - if s.op == Ops.CONST and abs(s.val - 6.283185307179586) < 1e-5: return (x.src[1-i], 6.283185307179586) + if (v:=_lit(s)) is not None and abs(v - 6.283185307179586) < 1e-5: return (x.src[1-i], 6.283185307179586) if s.op == Ops.MUL and len(s.src) == 2: - vals = [ss.val for ss in s.src if ss.op == Ops.CONST] + [ss.src[0].val for ss in s.src if ss.op == Ops.CAST and ss.src[0].op == Ops.CONST] + vals = [v for ss in s.src if (v:=_lit(ss)) is not None] if len(vals) == 2 and abs(vals[0] * vals[1] - 6.283185307179586) < 1e-5: return (x.src[1-i], vals[0] * vals[1]) return None @@ -163,7 +166,7 @@ def _trig_reduce(x, phase=0.0): def _signext(val: UOp) -> UOp: for bits, mask, ext in [(4, 0xF, 0xFFFFFFF0), (8, 0xFF, 0xFFFFFF00), (16, 0xFFFF, 0xFFFF0000)]: - if (val.op == Ops.AND and len(val.src) == 2 and val.src[1].op == Ops.CONST and val.src[1].val == mask) or val.dtype.itemsize == bits // 8: + if (val.op == Ops.AND and len(val.src) == 2 and _lit(val.src[1]) == mask) or val.dtype.itemsize == bits // 8: v32 = val.cast(dtypes.uint32) if val.dtype != dtypes.uint32 else val sb = (v32 >> _u32(bits - 1)) & _u32(1) return sb.ne(_u32(0)).where(v32 | _u32(ext), v32).cast(dtypes.int) @@ -482,7 +485,7 @@ class Parser: def _apply_binop(self, left, right, op): if op in ('||', '&&', '|', '^', '&'): left, right = self._coerce_bitwise(left, right) elif op in ('>=', '<=', '>', '<', '==', '!=', '<>', '>>', '<<'): left, right = self._coerce_cmp(left, right) - elif left.dtype != right.dtype: right = right.cast(left.dtype) + elif left.dtype != right.dtype: right = _restate(right, left.dtype) match op: case '||' | '|': return left | right case '&&' | '&': return left & right @@ -497,7 +500,7 @@ class Parser: if not dtypes.is_int(right.dtype): right = right.cast(dtypes.uint32) return (left >> right) if op == '>>' else (left << right) case '+' | '-': - if op == '-' and left.op == Ops.CONST and right.op == Ops.CONST: return _const(left.dtype, left.val - right.val) + if op == '-' and (lv:=_lit(left)) is not None and (rv:=_lit(right)) is not None: return _const(left.dtype, lv - rv) return (left + right) if op == '+' else (left - right) case '*' | '/': # Integer promotion: promote 16-bit integers to 32-bit before multiply to avoid overflow @@ -507,7 +510,7 @@ class Parser: left, right = left.cast(pdt), right.cast(pdt) if op == '*': return left * right return (left // right) if dtypes.is_int(left.dtype) else (left / right) - case '**': return UOp(Ops.EXP2, src=(right.cast(left.dtype),)) if left.op == Ops.CONST and left.val == 2.0 else left + case '**': return UOp(Ops.EXP2, src=(right.cast(left.dtype),)) if _lit(left) == 2.0 else left _PREC = [('||',), ('&&',), ('|',), ('^',), ('&',), ('==', '!=', '<>'), ('>=', '<=', '>', '<'), ('>>', '<<'), ('+', '-'), ('*', '/'), ('**',)] @@ -529,8 +532,8 @@ class Parser: return inner.eq(_const(inner.dtype, 0)) if self.try_eat_val('-', 'OP'): inner = self.unary() - if inner.op == Ops.CONST: - return _const(dtypes.int if inner.dtype == dtypes.uint32 else inner.dtype, -inner.val) + if (v:=_lit(inner)) is not None: + return _const(dtypes.int if inner.dtype == dtypes.uint32 else inner.dtype, -v) return inner.neg() if self.try_eat_val('+', 'OP'): return self.unary() return self.postfix() @@ -669,15 +672,15 @@ class Parser: self.eat('OP') width = self.parse() self.eat('RBRACKET') - if width.op == Ops.CONST: - w = int(width.val) + if (wv:=_lit(width)) is not None: + w = int(wv) return (base >> _to_u32(first)) & _const(base.dtype, (1 << w) - 1) return base if self.try_eat('COLON'): second = self.parse() self.eat('RBRACKET') - if first.op == Ops.CONST and second.op == Ops.CONST: - a, b = int(first.val), int(second.val) + if (fv:=_lit(first)) is not None and (sv:=_lit(second)) is not None: + a, b = int(fv), int(sv) if a < b: return _bitreverse(base, b - a + 1) hi, lo = a, b if lo >= base.dtype.itemsize * 8: @@ -698,8 +701,8 @@ class Parser: dt_suffix = DTYPES.get(self.eat('IDENT').val, dtypes.uint32) if var_name is None: var_name = self._find_var_name(base) - if first.op == Ops.CONST: - idx = int(first.val) + if (fv:=_lit(first)) is not None: + idx = int(fv) # Check for array element (var@idx) if var_name and f'{var_name}@{idx}' in self.vars: v = self.vars[f'{var_name}@{idx}'] @@ -758,7 +761,7 @@ class Parser: if type_char == 'F' and inner.dtype in (dtypes.uint32, dtypes.uint64, dtypes.ulong, dtypes.int, dtypes.int64): if inner.dtype.itemsize != dt.itemsize: inner = inner.cast(dtypes.uint32 if dt.itemsize == 4 else dtypes.uint64) return inner.bitcast(dt) - return inner.cast(dt) + return _restate(inner, dt) if self.at('IDENT'): ident = self.peek().val fmt = ident[0].lower() @@ -872,8 +875,8 @@ class Parser: def _coerce_cmp(self, l: UOp, r: UOp) -> tuple[UOp, UOp]: if l.dtype != r.dtype: - if r.dtype == dtypes.int and r.op == Ops.CONST and r.val < 0: l = l.cast(dtypes.int) - else: r = r.cast(l.dtype) + if r.dtype == dtypes.int and (rv:=_lit(r)) is not None and rv < 0: l = l.cast(dtypes.int) + else: r = _restate(r, l.dtype) return l, r def _coerce_bitwise(self, l: UOp, r: UOp) -> tuple[UOp, UOp]: @@ -969,8 +972,9 @@ def parse_block(lines: list[str], start: int, env: dict[str, VarVal], funcs: dic p.eat('QUOTE') if p.at('NUM'): return int(p.eat('NUM').val.rstrip('UuLl')) expr = p.parse().simplify() - assert expr.op == Ops.CONST, f"loop bound must be constant, got {expr}" - return int(expr.val) + v = _lit(expr) + assert v is not None, f"loop bound must be constant, got {expr}" + return int(v) start_val = parse_bound() p.eat('COLON') end_val = parse_bound() diff --git a/test/null/test_const_folding.py b/test/null/test_const_folding.py index a4ac45e350..1c673f6534 100644 --- a/test/null/test_const_folding.py +++ b/test/null/test_const_folding.py @@ -55,7 +55,7 @@ class TestWeakConstFolding(unittest.TestCase): # log10 backward folds log10(2)/log(2) = 1/log(10) in one rounding, not the double-rounded 1/float32(log(10)) x = Tensor([1.0, 2.0, 3.0]) ast = next(s.src[0] for s in x.log10().sum().gradient(x)[0].schedule_linear().src if s.src[0].op is Ops.SINK) - const = next(u.arg for u in full_rewrite(ast).toposort() if u.op is Ops.CONST and u.dtype is dtypes.float32) + const = next(u.src[0].val for u in full_rewrite(ast).toposort() if u.dtype is dtypes.float32 and u.op is Ops.CAST and u.src[0].op is Ops.CONST) # correctly rounded: within half a float32 ulp of the exact value (folding at float32 lands 0.66 ulp off) self.assertLess(abs(const - 1/math.log(10)), 2**-26) @@ -129,9 +129,9 @@ class TestBitcastConstFolding(unittest.TestCase): for (from_dt, from_v), (to_dt, to_v) in itertools.product(cases.items(), cases.items()): if not math.isnan(from_v): r = full_rewrite(UOp.const(from_v, from_dt).bitcast(to_dt).sink()).src[0] - self.assertEqual(r.op, Ops.CONST, msg:=f"{from_dt} -> {to_dt} ({from_v} -> {to_v})") + self.assertTrue(r.op is Ops.CAST and r.src[0].op is Ops.CONST, msg:=f"{from_dt} -> {to_dt} ({from_v} -> {to_v})") self.assertEqual(r.dtype, to_dt, msg) - np.testing.assert_equal(r.val, to_v, msg) + np.testing.assert_equal(r.src[0].val, to_v, msg) t({dtypes.int8: 0, dtypes.uint8: 0, dtypes.bool: False}) t({dtypes.int8: 1, dtypes.uint8: 1, dtypes.bool: True}) @@ -153,8 +153,8 @@ class TestBitcastConstFolding(unittest.TestCase): def test_vec_bitcast(self): with Context(SPEC=0): srcs = full_rewrite(UOp.const((-1, -2**31, 75), dtypes.int32).bitcast(dtypes.uint32).sink()).src - self.assertTrue(all(r.op is Ops.CONST and r.dtype == dtypes.uint32 for r in srcs)) - self.assertEqual(tuple(x.val for x in srcs), (2**32-1, 2**31, 75)) + self.assertTrue(all(r.op is Ops.CAST and r.src[0].op is Ops.CONST and r.dtype == dtypes.uint32 for r in srcs)) + self.assertEqual(tuple(x.src[0].val for x in srcs), (2**32-1, 2**31, 75)) # folds advance indexing into basic indexing class TestIndexingConstFolding(unittest.TestCase): diff --git a/test/null/test_graph_rewrite.py b/test/null/test_graph_rewrite.py index 06aa090359..f536574a77 100644 --- a/test/null/test_graph_rewrite.py +++ b/test/null/test_graph_rewrite.py @@ -14,14 +14,15 @@ def apply_rewrite(expr): @Context(SPEC=0) def apply_rewrite_values(expr): srcs = full_rewrite(expr.sink()).src - if len(srcs) == 1: - if srcs[0].op is Ops.CONST: return (srcs[0].val,) - if srcs[0].op is Ops.STACK: return tuple(s.val for s in srcs[0].src) - return tuple(s.val for s in srcs) + if len(srcs) == 1 and srcs[0].op is Ops.STACK: srcs = srcs[0].src + return tuple(s.src[0].val for s in srcs) +# reads both literal spellings: a pre-rewrite graph has bare CONSTs, a post-rewrite one has the pair def evaluate_uop(uop, variables): - if uop.op == Ops.CONST: + if uop.op is Ops.CONST: return uop.val + elif uop.op is Ops.CAST and uop.src[0].op is Ops.CONST: + return uop.src[0].val elif uop.op == Ops.PARAM and uop.arg.addrspace is AddrSpace.ALU: return variables[uop.expr] elif uop.op in GroupOp.ALU: @@ -33,31 +34,31 @@ def evaluate_uop(uop, variables): class TestArithmeticSimplifications(unittest.TestCase): def test_full_graph_rewrite_division_by_zero(self): optimized_div_uop = apply_rewrite(UOp.const(10.0) / UOp.const(0.0)) - self.assertEqual(optimized_div_uop.op, Ops.CONST) - self.assertTrue(math.isinf(optimized_div_uop.val) or math.isnan(optimized_div_uop.val)) + self.assertTrue(optimized_div_uop.op is Ops.CAST and optimized_div_uop.src[0].op is Ops.CONST) + self.assertTrue(math.isinf(optimized_div_uop.src[0].val) or math.isnan(optimized_div_uop.src[0].val)) def test_full_graph_rewrite_redundant_operations(self): optimized_uop = apply_rewrite((UOp.const(10.0) + UOp.const(0.0)) * UOp.const(1.0)) - self.assertEqual(optimized_uop.op, Ops.CONST) - self.assertEqual(optimized_uop.val, 10.0) + self.assertTrue(optimized_uop.op is Ops.CAST and optimized_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_uop.src[0].val, 10.0) def test_full_graph_rewrite_large_graph(self): prev_uop = UOp.const(0) for i in range(1, 101): prev_uop += UOp.const(i) optimized_uop = apply_rewrite(prev_uop) - self.assertEqual(optimized_uop.op, Ops.CONST) - self.assertEqual(optimized_uop.val, sum(range(1, 101))) + self.assertTrue(optimized_uop.op is Ops.CAST and optimized_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_uop.src[0].val, sum(range(1, 101))) def test_full_graph_rewrite_division_by_one(self): optimized_uop = apply_rewrite(UOp.const(42.0) / UOp.const(1.0)) - self.assertEqual(optimized_uop.op, Ops.CONST) - self.assertEqual(optimized_uop.val, 42.0) + self.assertTrue(optimized_uop.op is Ops.CAST and optimized_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_uop.src[0].val, 42.0) def test_full_graph_rewrite_modulo_by_one(self): optimized_uop = apply_rewrite(UOp.const(42) % UOp.const(1)) - self.assertEqual(optimized_uop.op, Ops.CONST) - self.assertEqual(optimized_uop.val, 0) + self.assertTrue(optimized_uop.op is Ops.CAST and optimized_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_uop.src[0].val, 0) class TestFoldingAndReduction(unittest.TestCase): @@ -100,8 +101,8 @@ class TestFoldingAndReduction(unittest.TestCase): inner_range = UOp.range(4, 1) expr = (outer_range * 10) + inner_range optimized_reduce_uop = apply_rewrite(expr.reduce(Ops.ADD, outer_range, inner_range)) - self.assertEqual(optimized_reduce_uop.op, Ops.CONST) - self.assertEqual(optimized_reduce_uop.val, sum((i * 10) + j for i in range(8) for j in range(4))) + self.assertTrue(optimized_reduce_uop.op is Ops.CAST and optimized_reduce_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_reduce_uop.src[0].val, sum((i * 10) + j for i in range(8) for j in range(4))) class TestModuloAndDivisionFolding(unittest.TestCase): @@ -109,22 +110,22 @@ class TestModuloAndDivisionFolding(unittest.TestCase): # index dtype because div-mod rules only work on index x_var_uop = UOp.variable('x', 0, 100).cast(dtypes.weakint) optimized_mod_uop = apply_rewrite(((x_var_uop * 4) + 2) % 4) - self.assertEqual(optimized_mod_uop.op, Ops.CONST) - self.assertEqual(optimized_mod_uop.val, 2) + self.assertTrue(optimized_mod_uop.op is Ops.CAST and optimized_mod_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_mod_uop.src[0].val, 2) def test_full_graph_rewrite_division_folding_with_define_var(self): # index dtype because div-mod rules only work on index n_var_uop = UOp.variable('n', 1, 1000).cast(dtypes.weakint) optimized_div_uop = apply_rewrite((n_var_uop * 6) // 3) self.assertEqual(optimized_div_uop.op, Ops.MUL) - self.assertEqual(optimized_div_uop.src[1].val, 2) + self.assertEqual(optimized_div_uop.src[1].src[0].val, 2) def test_full_graph_rewrite_complex_mod_div_folding(self): # index dtype because div-mod rules only work on index k_var_uop = UOp.variable('k', 0, 50).cast(dtypes.weakint) optimized_div_uop = apply_rewrite(((k_var_uop * 12 + 8) % 6) // 2) - self.assertEqual(optimized_div_uop.op, Ops.CONST) - self.assertEqual(optimized_div_uop.val, 1) + self.assertTrue(optimized_div_uop.op is Ops.CAST and optimized_div_uop.src[0].op is Ops.CONST) + self.assertEqual(optimized_div_uop.src[0].val, 1) def test_graph_rewrite_div_folding_bug(self): lhs = UOp(Ops.ADD, src=( @@ -161,9 +162,9 @@ class TestEdgeCasesAndSpecialOperations(unittest.TestCase): def test_full_graph_rewrite_transcendental_edge_cases(self): optimized_sink = full_rewrite(UOp.const(-1.0).log2().sink(UOp.const(0.0).reciprocal())) optimized_log2_neg, optimized_recip_zero = optimized_sink.src - self.assertTrue(math.isnan(optimized_log2_neg.val), f"Expected NaN for log2(-1.0), got {optimized_log2_neg.val}") - self.assertTrue(math.isinf(optimized_recip_zero.val) and optimized_recip_zero.val > 0, - f"Expected +inf for reciprocal(0.0), got {optimized_recip_zero.val}") + self.assertTrue(math.isnan(optimized_log2_neg.src[0].val), f"Expected NaN for log2(-1.0), got {optimized_log2_neg.src[0].val}") + self.assertTrue(math.isinf(optimized_recip_zero.src[0].val) and optimized_recip_zero.src[0].val > 0, + f"Expected +inf for reciprocal(0.0), got {optimized_recip_zero.src[0].val}") @unittest.skip("broken") def test_full_graph_rewrite_modulo_negative_dividend(self): @@ -183,7 +184,7 @@ class TestGEPAndVectorizeRewrite(unittest.TestCase): def test_gep_single_element_extraction(self): # GEP on a vector dtype to extract a single element base_vector = UOp.const((1.0, 2.0, 3.0, 4.0)) - self.assertEqual(apply_rewrite(base_vector.index(2)).val, 3.0) + self.assertEqual(apply_rewrite(base_vector.index(2)).src[0].val, 3.0) def test_gep_tuple_extraction(self): # GEP on a vector dtype to extract multiple elements as a vector @@ -193,7 +194,7 @@ class TestGEPAndVectorizeRewrite(unittest.TestCase): def test_gep_on_const_stack(self): # GEP on a const STACK to extract a single element const_stack = UOp.const((1.0, 2.0, 3.0, 4.0)) - self.assertEqual(apply_rewrite(const_stack.index(2)).val, 3.0) + self.assertEqual(apply_rewrite(const_stack.index(2)).src[0].val, 3.0) def test_gep_tuple_on_const_stack(self): # GEP on a const STACK using a tuple to extract multiple elements diff --git a/test/null/test_simplify_valid_idx.py b/test/null/test_simplify_valid_idx.py index a5dc1d83d2..92d1da10ce 100644 --- a/test/null/test_simplify_valid_idx.py +++ b/test/null/test_simplify_valid_idx.py @@ -3,7 +3,7 @@ import unittest, itertools from tinygrad.codegen.late.coalesce import indexing_simplify from tinygrad.dtype import dtypes from tinygrad.uop.ops import UOp, Ops, graph_rewrite -from tinygrad.uop.weak import pm_lower_index_dtype +from tinygrad.uop.weak import pm_commit_weak from tinygrad.uop.symbolic import simplify_valid, sym, pm_move_where_on_load from tinygrad.helpers import Context from test.helpers import full_rewrite @@ -500,7 +500,7 @@ class TestImageSimplification(unittest.TestCase): idx_y = (f + UOp.const(1.0)).cast(dtypes.int) load = get_load_image_uop((10, 10, 4), (UOp.const(-1) < idx_y) & (idx_y < UOp.const(10)), (Special("gidx0", 10), idx_y)) - off = graph_rewrite(load.sink(), pm_lower_index_dtype+indexing_simplify, ctx={}).src[0].src[0] + off = graph_rewrite(load.sink(), pm_commit_weak+indexing_simplify).src[0].src[0] self.assertEqual(off.src[1].get_valid(), UOp.const(True)) class TestDropTrueGate(unittest.TestCase): @@ -530,7 +530,7 @@ class TestRangeShrink(unittest.TestCase): load = get_gated_load_uop(r < UOp.const(4), r) ranges = self.get_ranges(load.sink()) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 4) + self.assertEqual(ranges[0].src[0].src[0].val, 4) def test_range_shrink_picks_max_guard(self): # two loads guard the same range with r < 4 and r < 8 -> shrink to max(4, 8) = 8 @@ -539,7 +539,7 @@ class TestRangeShrink(unittest.TestCase): load2 = get_gated_load_uop(r < UOp.const(8), r) ranges = self.get_ranges(UOp.sink(load1, load2)) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 8) + self.assertEqual(ranges[0].src[0].src[0].val, 8) def test_range_no_shrink_guard_ge_max(self): # guard r < 300 with range max 204 -> no shrink (guard doesn't constrain) @@ -547,7 +547,7 @@ class TestRangeShrink(unittest.TestCase): load = get_gated_load_uop(r < UOp.const(300), r) ranges = self.get_ranges(load.sink()) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 204) + self.assertEqual(ranges[0].src[0].src[0].val, 204) def test_range_no_shrink_when_unguarded_elsewhere(self): # one load guards r < 4, but another load uses r without a gate -> no shrink @@ -556,7 +556,7 @@ class TestRangeShrink(unittest.TestCase): load2 = UOp(Ops.LOAD, src=(UOp.param(1, dtypes.float, (204,)).index(r),)) ranges = self.get_ranges(UOp.sink(load1, load2)) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 204) + self.assertEqual(ranges[0].src[0].src[0].val, 204) def test_range_no_shrink_when_used_in_reduce(self): # range used in both a gated load AND directly in the reduce expression -> no shrink @@ -565,7 +565,7 @@ class TestRangeShrink(unittest.TestCase): red = (r.cast(dtypes.float) + gated_load).reduce(r, arg=Ops.ADD) ranges = self.get_ranges(red.sink()) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 204) + self.assertEqual(ranges[0].src[0].src[0].val, 204) def test_range_shrink_to_single_iteration(self): # guard r < 1 shrinks range to 1 -> single iteration, range eliminated entirely @@ -581,7 +581,7 @@ class TestRangeShrink(unittest.TestCase): x = (r < 4).where(UOp.const(1.0), Invalid) ranges = self.get_ranges(UOp.param(0, dtypes.float, (204,)).index(r).store((r < 4).where(x, Invalid)).sink()) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 4) + self.assertEqual(ranges[0].src[0].src[0].val, 4) def test_range_shrink_store_where_invalid_flipped(self): # above, but flipped @@ -590,7 +590,7 @@ class TestRangeShrink(unittest.TestCase): x = (r < 4).where(UOp.const(1.0), Invalid) ranges = self.get_ranges(UOp.param(0, dtypes.float, (204,)).index(r).store((r >= 4).where(Invalid, x)).sink()) self.assertEqual(len(ranges), 1) - self.assertEqual(ranges[0].src[0].val, 4) + self.assertEqual(ranges[0].src[0].src[0].val, 4) if __name__ == '__main__': unittest.main() diff --git a/test/null/test_uop_graph.py b/test/null/test_uop_graph.py index 6e77bcec12..ec755f1bae 100644 --- a/test/null/test_uop_graph.py +++ b/test/null/test_uop_graph.py @@ -14,9 +14,14 @@ simple_pm = PatternMatcher([ ((UPat.var('x') + UPat.cvar('c1')) + UPat.cvar('c2'), lambda x,c1,c2: x + (c1.val+c2.val)), ]) +def lit_value(u:UOp): + if u.op is Ops.CONST: return u.val + return u.src[0].val + def const_values(u:UOp): if u.op is Ops.CONST: return (u.val,) - if u.op is Ops.STACK: return tuple(x.val for x in u.src) + if u.op is Ops.CAST: return (u.src[0].val,) + if u.op is Ops.STACK: return tuple(lit_value(x) for x in u.src) raise AssertionError(f"expected const-like UOp, got {u.op}") class TestGraphRewriteConst(unittest.TestCase): @@ -25,7 +30,7 @@ class TestGraphRewriteConst(unittest.TestCase): v2 = v1.index(1) ret = graph_rewrite(v2, sym) self.assertEqual(ret.dtype, dtypes.int) - self.assertEqual(ret.val, 1) + self.assertEqual(ret.src[0].val, 1) def test_add_const(self): v1 = UOp.const((0,1,2)) @@ -196,10 +201,10 @@ class TestUOpGraph(unittest.TestCase): c2 = UOp.const(2.0, dtypes.float) out = c1+c2 uops = to_uops_list([out]) - self.assertEqual(len(uops), 2) # +1 for SINK + self.assertEqual(len(uops), 2) # +1 for SINK (a literal's inner CONST is not linearized, only its CAST) out = uops[-2] - self.assertEqual(out.op, Ops.CONST) - self.assertEqual(out.val, 3.0) + self.assertTrue(out.op is Ops.CAST and out.src[0].op is Ops.CONST) + self.assertEqual(out.src[0].val, 3.0) def test_where_same_fold(self): v = UOp.variable('tmp', 0, 1) @@ -210,8 +215,8 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out]) self.assertEqual(len(uops), 2) # +1 for SINK out = uops[-2] - self.assertEqual(out.op, Ops.CONST) - self.assertEqual(out.val, 1.0) + self.assertTrue(out.op is Ops.CAST and out.src[0].op is Ops.CONST) + self.assertEqual(out.src[0].val, 1.0) def test_where_const_fold(self): bf = UOp.const(False) @@ -221,8 +226,8 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out]) self.assertEqual(len(uops), 2) # +1 for SINK out = uops[-2] - self.assertEqual(out.op, Ops.CONST) - self.assertEqual(out.val, 2.0) + self.assertTrue(out.op is Ops.CAST and out.src[0].op is Ops.CONST) + self.assertEqual(out.src[0].val, 2.0) def test_const_cast(self): bf = UOp.const(False) @@ -230,8 +235,8 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out]) self.assertEqual(len(uops), 2) # +1 for SINK out = uops[-2] - self.assertEqual(out.op, Ops.CONST) - self.assertEqual(out.val, 0) + self.assertTrue(out.op is Ops.CAST and out.src[0].op is Ops.CONST) + self.assertEqual(out.src[0].val, 0) def test_const_bitcast(self): bf = UOp.const(1.0, dtypes.float) @@ -239,8 +244,8 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out]) self.assertEqual(len(uops), 2) # +1 for SINK out = uops[-2] - self.assertEqual(out.op, Ops.CONST) - self.assertEqual(out.val, 0x3F800000) + self.assertTrue(out.op is Ops.CAST and out.src[0].op is Ops.CONST) + self.assertEqual(out.src[0].val, 0x3F800000) @unittest.expectedFailure def test_const_shape_change_bitcast(self): @@ -315,7 +320,8 @@ class TestUOpGraph(unittest.TestCase): vec = UOp(Ops.STACK, src=tuple(consts)) with Context(SPEC=0): uops = to_uops_list([vec.index(i) for i in range(vec_size)]) - for uop, const in zip(uops, consts): + # each element folds to a two-uop literal, so read the outputs off the SINK + for uop, const in zip(uops[-1].src, consts): self.assertEqual(uop, const) def test_cast_alu_fold(self): @@ -326,7 +332,7 @@ class TestUOpGraph(unittest.TestCase): alu = (ld<1).cast(dtypes.bool) out = d0.index(idx).store(alu) uops = to_uops_list([out]) - self.assertEqual(len([x for x in uops if x.op is Ops.CAST]), 0) + self.assertEqual(len([x for x in uops if x.op is Ops.CAST and x.src[0].op is not Ops.CONST]), 0) def test_double_cast_fold(self): d0 = UOp.param(0, dtypes.float, (1,)) @@ -336,7 +342,7 @@ class TestUOpGraph(unittest.TestCase): alu = ld.cast(dtypes.float).cast(dtypes.float) out = d0.index(idx).store(alu) uops = to_uops_list([out]) - self.assertEqual(len([x for x in uops if x.op is Ops.CAST]), 1) + self.assertEqual(len([x for x in uops if x.op is Ops.CAST and x.src[0].op is not Ops.CONST]), 1) def test_depth_2_const_fold(self): v = UOp.variable("tmp", 0, 1, dtypes.int, param=True) @@ -348,8 +354,8 @@ class TestUOpGraph(unittest.TestCase): self.assertEqual(len(uops), 5) # +1 for SINK, +1 for the PARAM shape STACK out = uops[-2] # -2 to skip SINK self.assertEqual(out.op, Ops.ADD) - self.assertEqual(out.src[1].op, Ops.CONST) - self.assertEqual(out.src[1].val, 6) + self.assertTrue(out.src[1].op is Ops.CAST and out.src[1].src[0].op is Ops.CONST) + self.assertEqual(out.src[1].src[0].val, 6) def test_bitcast_to_same_dtype_fold(self): for dt in dtypes.ints + dtypes.floats + (dtypes.bool,): @@ -361,7 +367,7 @@ class TestUOpGraph(unittest.TestCase): def test_sub_with_cast_folds(self): a = Variable("a", 0, 5) uops = to_uops_list([a.cast(dtypes.int)+(-a).cast(dtypes.int)]) - assert uops[0] == UOp.const(0, dtypes.int) + assert uops[-2] == UOp.const(0, dtypes.int) assert uops[-1].op == Ops.SINK def test_where_on_gated_load_fold(self): @@ -373,7 +379,7 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out.index(ridx0).store(w)]) for u in uops: assert u.op is not Ops.WHERE - if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].val==5 + if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].src[0].val==5 def test_where_on_gated_load_folds_swapped_branches(self): ridx0 = UOp.range(100, 0) @@ -383,7 +389,7 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([w]) for u in uops: assert u.op is not Ops.WHERE - if u.op is Ops.LOAD: assert u.src[1].val==5 + if u.op is Ops.LOAD: assert u.src[1].src[0].val==5 def test_where_on_gated_load_with_cast(self): ridx0 = UOp.range(100, 0) @@ -395,7 +401,7 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([out.index(ridx0).store(w)]) for u in uops: assert u.op is not Ops.WHERE - if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].val == 5 + if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].src[0].val == 5 def test_where_on_casted_gated_load_extra_cond(self): ridx0 = UOp.range(100, 0) @@ -427,7 +433,7 @@ class TestUOpGraph(unittest.TestCase): uops = to_uops_list([st]) for u in uops: assert u.op is not Ops.WHERE - if u.op is Ops.STORE: assert u.src[1].val==5 + if u.op is Ops.STORE: assert u.src[1].src[0].val==5 def test_load_idx_becomes_int(self): # mnist indexing with split reduceop diff --git a/test/null/test_uop_symbolic.py b/test/null/test_uop_symbolic.py index 369ea785fe..bc2ec42423 100644 --- a/test/null/test_uop_symbolic.py +++ b/test/null/test_uop_symbolic.py @@ -6,7 +6,7 @@ from tinygrad.dtype import dtypes, ConstType, DType, Invalid from tinygrad.uop.ops import UOp, Ops, graph_rewrite, sym_infer from tinygrad.uop.spec import spec_shared, type_verify from tinygrad.uop.symbolic import sym, commutative, pm_simplify_valid, pm_move_where_on_load -from tinygrad.uop.weak import pm_cast_weak +from tinygrad.uop.weak import pm_commit_weak from tinygrad.uop.validate import uops_to_z3 def check_uop_against_string(self, v:UOp, s:str): @@ -36,7 +36,7 @@ class TestSymbolic(unittest.TestCase): self.assertEqual(solver.check(expr1 != expr2), z3.unsat, "simplified expression not equal to original") def helper_test_variable(self, v, n, m, s, test_z3:bool=True): - v_simplified = graph_rewrite(v, sym+pm_cast_weak, name="simplify symbolic uop") + v_simplified = graph_rewrite(v, sym+pm_commit_weak, name="simplify symbolic uop") if test_z3: self.check_equal_z3(v, v_simplified) nmin, nmax = v_simplified.vmin, v_simplified.vmax check_uop_against_string(self, v_simplified, s) @@ -1017,7 +1017,7 @@ class TestSymbolic(unittest.TestCase): cond = Variable("s", 0, 3, dtypes.int) < 2 a = Variable("a", 0, 3, dtypes.int) self.assertIs(graph_rewrite(cond.where(a, a+1).cast(dtypes.half), sym), cond.where(a.cast(dtypes.half), (a+1).cast(dtypes.half))) - self.assertIs(graph_rewrite(cond.where(a, uconst(2)).cast(dtypes.half), sym), cond.where(a.cast(dtypes.half), UOp.const(2, dtypes.half))) + self.assertIs(graph_rewrite(cond.where(a, uconst(2)).cast(dtypes.half), sym), cond.where(a.cast(dtypes.half), uconst(2.0))) self.assertIs(graph_rewrite(cond.where(a, UOp.invalid()).cast(dtypes.half), sym), cond.where(a.cast(dtypes.half), UOp.invalid())) def test_where_merge_branches(self): diff --git a/test/null/test_uops.py b/test/null/test_uops.py index c1d6931fd6..a50025d7f7 100644 --- a/test/null/test_uops.py +++ b/test/null/test_uops.py @@ -6,7 +6,7 @@ from tinygrad.helpers import Timing, Context, cdiv from tinygrad.dtype import dtypes, AddrSpace, ConstFloat, Invalid # noqa: F401 from tinygrad.device import Device from tinygrad.uop.ops import Ops, ParamArg, PatternMatcher, UOp, UPat, dtype_from_uop, exec_alu, graph_rewrite # noqa: F401 # ParamArg used by eval(str(uop)) roundtrip tests -from tinygrad.uop.weak import pm_lower_index_dtype +from tinygrad.uop.weak import pm_commit_weak, pm_lower_weak from tinygrad.uop.spec import spec_program, spec_shared, type_verify from tinygrad.uop.symbolic import sym, pm_remove_invalid from test.helpers import eval_uop, to_uops_list @@ -68,6 +68,11 @@ class TestDTypeFromUOp(unittest.TestCase): self.assertIs((out:=graph_rewrite(gate.where(value, UOp.invalid()), pm_remove_invalid)).src[2], UOp.const(0, dtypes.float)) type_verify(out.sink(), spec_program) + def test_bitcast_rejects_a_weak_operand(self): + # a bitcast reinterprets its operand's bits, so the operand must state a width: a raw build raises + with self.assertRaises(RuntimeError): UOp(Ops.BITCAST, src=(UOp.const(1.0),), arg=dtypes.uint32) + self.assertEqual(UOp(Ops.BITCAST, src=(UOp.const(1.0, dtypes.float32),), arg=dtypes.uint32).dtype, dtypes.uint32) + def test_remove_invalid_stack_lanes(self): stack = UOp(Ops.STACK, dtypes.half, (UOp.const(1, dtypes.half), UOp.invalid())) out = graph_rewrite(stack, pm_remove_invalid) @@ -75,22 +80,26 @@ class TestDTypeFromUOp(unittest.TestCase): type_verify(out.sink(), spec_program) class TestLowerIndexDtype(unittest.TestCase): + def lower(self, u:UOp) -> UOp: return graph_rewrite(graph_rewrite(u, pm_commit_weak), pm_lower_weak) + def test_gated_shrink_lowers_to_selected_width(self): # coalesce builds gated SHRINKs for masked vectorized loads; lowering must resolve them at the # width the offset bounds select (this one needs long) buf = UOp.param(0, dtypes.float, (2**31+64,)) i = UOp.variable("i", 0, 2**28) shrink = UOp(Ops.SHRINK, src=(buf, (i*24).valid(i < 2**28), UOp.const(4))) - lowered = graph_rewrite(shrink.sink(), pm_lower_index_dtype) - self.assertTrue(all(u.dtype != dtypes.weakint for u in lowered.backward_slice_with_self), "lowering must resolve all weakint") + lowered = self.lower(shrink.sink()) + self.assertTrue(all(u.op is Ops.CONST for u in lowered.backward_slice_with_self if u.dtype in dtypes.weaks), + "lowering must resolve every weak width, except a typed literal's value half") sh = next(u for u in lowered.backward_slice_with_self if u.op is Ops.SHRINK) self.assertEqual(sh.src[1].dtype, dtypes.long) def test_reg_buffer_size_lowers(self): reg = UOp.placeholder((4,), dtypes.float, 0, addrspace=AddrSpace.REG) self.assertEqual(reg.src[0].dtype, dtypes.weakint) - lowered = graph_rewrite(reg.sink(), pm_lower_index_dtype) - self.assertTrue(all(u.dtype != dtypes.weakint for u in lowered.backward_slice_with_self), "lowering must resolve all weakint") + lowered = self.lower(reg.sink()) + self.assertTrue(all(u.op is Ops.CONST for u in lowered.backward_slice_with_self if u.dtype in dtypes.weaks), + "lowering must resolve every weak width, except a typed literal's value half") self.assertEqual(next(u for u in lowered.backward_slice_with_self if u.op is Ops.BUFFER).src[0].dtype, dtypes.int) class TestSafeCast(unittest.TestCase): @@ -134,7 +143,7 @@ class TestConstFloatEq(unittest.TestCase): self.assertFalse(Invalid != HoldsInvalid()) def test_matchers_agree_on_nan(self): - n = UOp.const(math.nan, dtypes.float32) + n = UOp.const(math.nan) for compiled in (False, True): pm = PatternMatcher([(UPat(Ops.CONST, arg=math.nan), lambda: True)], compiled=compiled) self.assertTrue(pm.rewrite(n), f"{compiled=}") @@ -279,6 +288,9 @@ class TestGatedStoreRewrite(unittest.TestCase): for x in gated_uops: self.assertIs(x.op, Ops.STORE) for x in gated_uops: self.assertEqual(len(x.src), 2) +# assertions below count arithmetic ops: a literal's CAST is filtered out (one test asserts a CAST must NOT appear) +def alu_ops(uops): return [x.op for x in uops if x.op is not Ops.CAST or x.src[0].op is not Ops.CONST] + @unittest.skipIf(Device.DEFAULT == "METAL", "compiler bug") @unittest.skipUnless(Ops.SHR in Device[Device.DEFAULT].renderer.code_for_op, "fast_idiv requires SHR") class TestFastIdiv(unittest.TestCase): @@ -290,7 +302,7 @@ class TestFastIdiv(unittest.TestCase): a = UOp(Ops.CDIV, dt, (l, c)) uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) Device[Device.DEFAULT].renderer.render(uops) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.SHR, ops, f"For dtype={dt} divison by power of two did not simplify to shift") self.assertNotIn(Ops.CDIV, ops, f"For dtype={dt} divison by power of two did not simplify to shift") @@ -301,7 +313,7 @@ class TestFastIdiv(unittest.TestCase): c = UOp.const(8).cast(dt) a = UOp(Ops.FLOORMOD, dt, (g.index(c), c)) uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.AND, ops, f"For dtype={dt} FLOORMOD by pow2 did not simplify to AND") self.assertNotIn(Ops.CMOD, ops, f"For dtype={dt} FLOORMOD by pow2 left a MOD") self.assertNotIn(Ops.FLOORMOD, ops, f"For dtype={dt} FLOORMOD survived past late rewrite") @@ -313,7 +325,7 @@ class TestFastIdiv(unittest.TestCase): c = UOp.const(2).cast(dt) a = UOp(Ops.FLOORDIV, dt, (g.index(c), c)) uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.SHR, ops, f"For dtype={dt} FLOORDIV by power of two did not simplify to shift") self.assertNotIn(Ops.CDIV, ops, f"For dtype={dt} FLOORDIV by power of two did not simplify to shift") self.assertNotIn(Ops.FLOORDIV, ops, f"For dtype={dt} FLOORDIV survived past late rewrite") @@ -327,14 +339,14 @@ class TestFastIdiv(unittest.TestCase): a = UOp(Ops.CDIV, src=(l, c)) uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) Device[Device.DEFAULT].renderer.render(uops) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.SHR, ops) self.assertNotIn(Ops.CDIV, ops) b = UOp(Ops.CMOD, src=(l, c)) uops = to_uops_list([b], ren=Device[Device.DEFAULT].renderer) Device[Device.DEFAULT].renderer.render(uops) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.SHR, ops) self.assertNotIn(Ops.CMOD, ops) @@ -348,7 +360,7 @@ class TestFastIdiv(unittest.TestCase): def test_fast_idiv_remove_powers_of_two(self): ridx = UOp.range(2**20, 0) uops = to_uops_list([ridx//(7*64)], ren=Device[Device.DEFAULT].renderer) - ops = [x.op for x in uops] + ops = alu_ops(uops) # this requires shifting out the powers of two before doing fast_idiv # (((ridx0>>6)*18725)>>17) instead of (int)((((long)(ridx0)*1198373)>>29)) self.assertNotIn(Ops.CAST, ops) @@ -362,7 +374,7 @@ class TestFastIdiv(unittest.TestCase): a = UOp(Ops.CDIV, src=(l, c)) uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) Device[Device.DEFAULT].renderer.render(uops) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertIn(Ops.SHR, ops) self.assertNotIn(Ops.CDIV, ops) @@ -373,7 +385,7 @@ class TestFastIdiv(unittest.TestCase): a = UOp(Ops.CDIV, src=(l, c)) with Context(DISABLE_FAST_IDIV=1): uops = to_uops_list([a], ren=Device[Device.DEFAULT].renderer) - ops = [x.op for x in uops] + ops = alu_ops(uops) self.assertNotIn(Ops.SHR, ops) self.assertIn(Ops.CDIV, ops) diff --git a/test/null/test_viz.py b/test/null/test_viz.py index 49fbff6140..4947af57c7 100644 --- a/test/null/test_viz.py +++ b/test/null/test_viz.py @@ -236,8 +236,8 @@ class TestViz(unittest.TestCase): def test_const_node_visibility(self): with save_viz() as viz: a = UOp.variable("a", 0, 10, dtype=dtypes.int) - z = UOp.const(0, a.dtype) - y = UOp.const(math.pi, dtypes.float) + z = UOp.const(0) + y = UOp.const(math.pi) alu = a*z ret = exec_rewrite(sink:=UOp.sink(alu, y), [sym]) lst = viz.list_items() @@ -249,7 +249,7 @@ class TestViz(unittest.TestCase): self.assertTrue(graphs[0][id(y)]["exclude"]) self.assertFalse(graphs[0][id(alu)]["exclude"]) self.assertEqual(graphs[0][id(y)]["label"].split("\n")[:2], ["CONST", "3.14159"]) - self.assertEqual(list(graphs[1]), [id(z), id(y), id(ret)]) + self.assertEqual(list(graphs[1]), [id(z), id(ret.src[0]), id(y), id(ret)]) # a*0 folds to CAST(int, z): both nodes are in the graph def test_const_reshape_expand_folded(self): # CONST->EXPAND should be folded into the ALU node, not shown as separate EXPAND nodes diff --git a/test/unit/test_dtype_weak.py b/test/unit/test_dtype_weak.py index 90b8d16b90..a2f23afc9d 100644 --- a/test/unit/test_dtype_weak.py +++ b/test/unit/test_dtype_weak.py @@ -4,7 +4,7 @@ from tinygrad import Tensor, dtypes, TinyJit from tinygrad.helpers import Context from tinygrad.dtype import least_upper_float from tinygrad.uop.ops import UOp, Ops, dtype_from_uop, graph_rewrite -from tinygrad.uop.weak import pm_lower_index_dtype, pm_commit_weak +from tinygrad.uop.weak import pm_commit_weak from tinygrad.uop.symbolic import symbolic_simple from tinygrad.uop.spec import spec_shared, type_verify from tinygrad.engine.jit import JitError @@ -73,7 +73,7 @@ class TestWeakPromotion(unittest.TestCase): recips = [u for u in (x / y)._uop.toposort() if u.op is Ops.RECIPROCAL] self.assertEqual([(u.dtype, u.src[0].dtype) for u in recips], [(dtypes.float32, dtypes.float32)]) with Context(DEFAULT_FLOAT=dtypes.float16): - committed = graph_rewrite((UOp.const(1).cast(dtypes.int32) + UOp.const(1.0)).cast(dtypes.float32), pm_lower_index_dtype, ctx={}) + committed = graph_rewrite((UOp.const(1).cast(dtypes.int32) + UOp.const(1.0)).cast(dtypes.float32), pm_commit_weak) self.assertEqual([u.dtype for u in committed.toposort() if u.op is Ops.ADD], [dtypes.float32]) def test_div_sub_operand_kept_weak(self): @@ -84,7 +84,7 @@ class TestWeakPromotion(unittest.TestCase): def test_cast_weak_expression_commits_at_cast_floor(self): # the floor never narrows: a cast BELOW the default does not pull the compute width down with it with Context(DEFAULT_FLOAT=dtypes.float32): - narrowed = graph_rewrite((UOp.const(1.0) + UOp.const(2.0)).cast(dtypes.float16), pm_lower_index_dtype, ctx={}) + narrowed = graph_rewrite((UOp.const(1.0) + UOp.const(2.0)).cast(dtypes.float16), pm_commit_weak) self.assertEqual((narrowed.dtype, narrowed.src[0].dtype), (dtypes.float16, dtypes.float32)) def test_cast_weak_expression_value_uses_cast_floor(self): @@ -124,16 +124,17 @@ class TestWeakPromotion(unittest.TestCase): with Context(DEFAULT_FLOAT=dtypes.float16): dst = UOp.param(0, dtypes.bfloat16, (1,)).index(UOp.const(0).cast(dtypes.int32)) gate = UOp.const(True) - out = graph_rewrite(dst.store(UOp.const(5.0), gate), pm_lower_index_dtype, ctx={}) + out = graph_rewrite(dst.store(UOp.const(5.0), gate), pm_commit_weak) # a bare weak CONST commits directly: the pass runs without symbolic, so a CAST here would survive it self.assertEqual((out.src[1], out.src[2]), (UOp.const(5.0, dtypes.bfloat16), gate)) def test_weak_srcs_commit_only_at_a_concrete_lub(self): weak_lub = UOp(Ops.ADD, src=(UOp.const(1), UOp.const(1.0))) - self.assertIs(graph_rewrite(weak_lub, pm_lower_index_dtype, ctx={}), weak_lub) + self.assertIs(graph_rewrite(weak_lub, pm_commit_weak), weak_lub) concrete = UOp.const(2.0).cast(dtypes.float16) - where = graph_rewrite(UOp(Ops.WHERE, src=(UOp.const(True), concrete, UOp.const(1.0))), pm_lower_index_dtype, ctx={}) - self.assertEqual(tuple(x.dtype for x in where.src), (dtypes.bool, dtypes.float16, dtypes.float16)) + # the weak arm stays bare: its sibling states the width, so the WHERE already derives float16 for it + where = graph_rewrite(UOp(Ops.WHERE, src=(UOp.const(True), concrete, UOp.const(1.0))), pm_commit_weak) + self.assertEqual((where.dtype, tuple(x.dtype for x in where.src)), (dtypes.float16, (dtypes.bool, dtypes.float16, dtypes.weakfloat))) def test_weak_shift_lhs_commits_the_node(self): # a shift derives its lhs's dtype, so committing the lhs restates the root (WGSL's packed store writes `mask << shift_am`) @@ -287,6 +288,5 @@ class TestSignedUint64Weakfloat(unittest.TestCase): self.assertEqual((i64 < u64).item(), True) # comparison meets at float self.assertAlmostEqual((i64 + u64).sin().item(), math.sin(2), places=5) # Unary lowers before transcendental - if __name__ == "__main__": unittest.main() diff --git a/tinygrad/codegen/__init__.py b/tinygrad/codegen/__init__.py index 05d666a724..44ac2608f8 100644 --- a/tinygrad/codegen/__init__.py +++ b/tinygrad/codegen/__init__.py @@ -3,7 +3,7 @@ import itertools, functools from tinygrad.helpers import DISABLE_FAST_IDIV, TRANSCENDENTAL, SPEC, DEBUG, VIZ, IMAGE, NOOPT, EMULATED_DTYPES, NOLOCALS, USE_TC from tinygrad.helpers import ALLOW_TF32, DEFAULT_FLOAT, DEFAULT_INT, TracingKey, Context, panic from tinygrad.uop.ops import PatternMatcher, graph_rewrite, UOp, Ops, UPat, rewrite_group, KernelInfo, ProgramInfo, GroupOp, AxisType -from tinygrad.uop.weak import pm_lower_index_dtype, pm_commit_weak, pm_cast_weak +from tinygrad.uop.weak import pm_lower_weak, pm_commit_weak, pm_cast_const from tinygrad.uop.render import pyrender from tinygrad.uop.spec import type_verify, spec_tensor, spec_program from tinygrad.renderer import Renderer, Estimates @@ -17,7 +17,7 @@ from tinygrad.uop.movement import mop_cleanup from tinygrad.codegen.decomp.dtype import pm_dtype_decomps from tinygrad.codegen.decomp.op import get_late_rewrite_patterns, get_simplifying_rewrite_patterns from tinygrad.codegen.decomp.transcendental import get_transcendental_patterns -from tinygrad.codegen.late.coalesce import indexing_simplify +from tinygrad.codegen.late.coalesce import indexing_simplify, pm_narrow_index from tinygrad.codegen.opt.postrange import apply_opts from tinygrad.codegen.late.gater import pm_move_gates_from_index from tinygrad.codegen.simplify import pm_simplify_ranges, pm_flatten_range, pm_split_ranges, pm_load_collapse, pm_reduce_unparented @@ -342,11 +342,13 @@ def full_rewrite_to_sink(ast:UOp, ren:Renderer, optimize:bool=True) -> UOp: # extra symbolic before decomp. crashes without this? # NOTE: also run indexing_simplify here, while the index is still weakint and (x+y)*c -> x*c+y*c applies - sink = graph_rewrite(sink, sym+indexing_simplify, name="extra symbolic") + # pm_commit_weak rides along: a width minted in a round commits in that same fixpoint + sink = graph_rewrite(sink, sym+indexing_simplify+pm_commit_weak, name="extra symbolic") - # lower index dtype + # THE BOUNDARY: below here every width is stated. pm_narrow_index rides along because index widths settle here # NOTE: we need indexing_simplify to remove the cast to long using the Invalid - sink = graph_rewrite(sink, symbolic_simple+pm_lower_index_dtype+indexing_simplify, ctx={}, name="lower all index dtypes") + # NOTE: symbolic must NOT be composed here -- its cast collapse eats the weak CAST a lowered node wears, and it cycles + sink = graph_rewrite(sink, pm_lower_weak+pm_narrow_index+indexing_simplify, name="lower all index dtypes") # final symbolic before decomp sink = graph_rewrite(sink, symbolic, name="final symbolic") @@ -370,7 +372,7 @@ def full_rewrite_to_sink(ast:UOp, ren:Renderer, optimize:bool=True) -> UOp: # final rules for the renderer (without sym) extra_matcher = ren.extra_matcher if ren.extra_matcher is not None else PatternMatcher([]) - pm_final_rewrite = pm_commit_weak+pm_cast_weak+pm_decomp+extra_matcher+pm_split_ends + pm_final_rewrite = pm_commit_weak+pm_decomp+extra_matcher+pm_split_ends sink = graph_rewrite(sink, pm_final_rewrite+pm_remove_invalid, ctx=ren, name="final rewrite") # add implicit barriers (stores/loads through LOCAL memory ordered by AFTER or across loop iterations need workgroup barriers) @@ -383,6 +385,9 @@ def full_rewrite_to_sink(ast:UOp, ren:Renderer, optimize:bool=True) -> UOp: num_params = len([x for x in sink.toposort() if x.op is Ops.PARAM and x.arg.slot != -1]) sink = graph_rewrite(sink, pm_number_params, ctx=[num_params], name="number params with -1", walk=True) + # wrap every const still riding bare, so no renderer reads one + sink = graph_rewrite(sink, pm_cast_const, name="state literal widths") + if VIZ: graph_rewrite(sink, PatternMatcher([]), name="View Output AST") if SPEC: type_verify(sink, spec_program) diff --git a/tinygrad/codegen/decomp/dtype.py b/tinygrad/codegen/decomp/dtype.py index db85526c83..5ec8ad9f0a 100644 --- a/tinygrad/codegen/decomp/dtype.py +++ b/tinygrad/codegen/decomp/dtype.py @@ -3,12 +3,14 @@ from tinygrad.dtype import dtypes, DType, truncate from tinygrad.helpers import flatten, DEBUG, EMULATED_DTYPES, Context, SPEC from tinygrad.uop import GroupOp from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher, graph_rewrite, ParamArg +from tinygrad.uop.weak import commit_weak from tinygrad.renderer import Renderer from tinygrad.codegen.decomp.transcendental import exponent_bias, shl, shr # ***** long as 2 ints ***** l2i_dt = {dtypes.long: dtypes.int, dtypes.ulong: dtypes.uint} +long_word_tags = {(w, dt) for w in (0, 1) for dt in l2i_dt.values()} def unpack32(v:UOp) -> tuple[UOp, UOp]: return v.bitcast(dtypes.uint) & 0xFFFF, shr(v.bitcast(dtypes.uint), 16) def reindex(idx:UOp, off:int, mul=2) -> UOp: if idx.op is Ops.SHRINK: @@ -134,8 +136,14 @@ def f2f_store(st, idx, val, fr:DType, to:DType): if (n:=val.max_numel()) == 1: return st.replace(src=(idx, f2f(val.bitcast(f2f_dt[to]), to, fr))) return UOp.group(*(st.replace(src=(reindex(idx, i, 1), f2f(val.index(i).bitcast(f2f_dt[to]), to, fr))) for i in range(n))) +def split_word(x:UOp, v:int) -> UOp: return UOp.const(truncate[x.tag[1]]((v >> 32) if x.tag[0] == 1 else (v & 0xFFFFFFFF)), x.tag[1]) + # tag is the 32-bit word this node becomes - (0 for the low word, 1 for the high, the dtype the consumer wants) pm_long_decomp = PatternMatcher([ + # a weak CONST states no width and cannot be split into words: commit it at the long dtype a sibling src states + (UPat(GroupOp.All, name='x'), lambda x: x.replace(src=tuple(commit_weak(s, dt) if s.op is Ops.CONST and s.dtype in dtypes.weaks else s + for s in x.src)) if any(s.op is Ops.CONST and s.dtype in dtypes.weaks for s in x.src) + and (dt:=next((s.dtype for s in x.src if s.dtype in l2i_dt), None)) is not None else None), (UPat(GroupOp.Defines, src=(UPat.var("sz"),), name="x"), lambda x,sz: x.replace(dtype=l2i_dt[x.dtype], arg=replace(x.arg, dtype=l2i_dt[x.dtype]), src=(sz*2,)) if x.dtype in l2i_dt else None), (UPat(Ops.INDEX, tuple(l2i_dt.keys()), name='x'), lambda x: @@ -147,6 +155,8 @@ pm_long_decomp = PatternMatcher([ split_l2i(ctx, x.op, dt:=l2i_dt[a.dtype], *flatten((s.rtag((0, dt)), s.rtag((1, dt))) for s in x.src))), (UPat(Ops.CAST, tuple(l2i_dt.keys()), src=(UPat.var('a', tuple(l2i_dt.keys())),), name="x"), lambda ctx,a,x: split_l2i(ctx, Ops.BITCAST, l2i_dt[x.dtype], a.rtag((0, dt:=l2i_dt[a.dtype])), a.rtag((1, dt)))[x.tag[0]]), + # a literal splits by value; the general CAST arm below would drop its high word + (UPat(Ops.CAST, src=(UPat(Ops.CONST, name='c'),), tag=long_word_tags, name='x'), lambda x,c: split_word(x, c.val)), (UPat(Ops.CAST, tuple(l2i_dt.keys()), src=(UPat.var('a'),), name="x"), lambda ctx,a,x: split_l2i(ctx, x.op, x.dtype, a)[x.tag[0]] if x.tag is not None else None), (UPat(Ops.CAST, src=(UPat.var('a', tuple(l2i_dt.keys())),), name="x"), lambda ctx,a,x: @@ -161,8 +171,6 @@ pm_long_decomp = PatternMatcher([ if x.tag is not None else None), (UPat(Ops.LOAD, tuple(l2i_dt.keys()), src=(UPat.var('idx'),), name='x'), lambda x,idx: x.replace(dtype=l2i_dt[x.dtype], src=(reindex(idx, x.tag[0]).replace(dtype=l2i_dt[x.dtype], tag=None),), tag=None) if x.tag is not None else None), - (UPat(Ops.CONST, tag={(w, dt) for w in (0, 1) for dt in l2i_dt.values()}, name='x'), lambda x: - UOp.const(truncate[x.tag[1]]((x.val >> 32) if x.tag[0] == 1 else (x.val & 0xFFFFFFFF)), x.tag[1])) ]) # float decomposition patterns - ctx is (fr, to) tuple diff --git a/tinygrad/codegen/decomp/op.py b/tinygrad/codegen/decomp/op.py index 6a48cdca53..98f4e12c95 100644 --- a/tinygrad/codegen/decomp/op.py +++ b/tinygrad/codegen/decomp/op.py @@ -1,6 +1,6 @@ from typing import Callable import functools -from tinygrad.dtype import dtypes +from tinygrad.dtype import dtypes, weak_dtype from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher from tinygrad.renderer import Renderer @@ -128,6 +128,7 @@ def get_late_rewrite_patterns(ops:tuple[Ops, ...], disable_fast_idiv:bool) -> Pa if Ops.SHL in ops: pat += [(UPat.var('x').alu(Ops.SHL, UPat.cvar('n'))+UPat.var('c'), lambda x,n,c: x.alu(Ops.MULACC, x.const_like(1< a/b if Ops.FDIV in ops: - pat += [(UPat.var("x").reciprocal(), lambda x: x.const_like(1).alu(Ops.FDIV, x))] - pat += [(UPat.var("a", dtypes.floats) * UPat(Ops.FDIV, dtypes.floats, src=(UPat.const(1), UPat.var("b"))), lambda a,b: a.alu(Ops.FDIV, b))] + # mint the 1 bare: the contraction rule below matches it as a bare CONST + pat += [(UPat.var("x").reciprocal(), lambda x: x.const_like(1, weak_dtype(x.dtype)).alu(Ops.FDIV, x))] + pat += [(UPat.var("a") * UPat(Ops.FDIV, dtypes.floats, src=(UPat.const(1), UPat.var("b"))), lambda a,b: a.alu(Ops.FDIV, b))] return PatternMatcher(pat) diff --git a/tinygrad/codegen/late/coalesce.py b/tinygrad/codegen/late/coalesce.py index 64e65f2320..19f19d8481 100644 --- a/tinygrad/codegen/late/coalesce.py +++ b/tinygrad/codegen/late/coalesce.py @@ -61,6 +61,14 @@ indexing_simplify = PatternMatcher([ UPat.var("valid").where(UPat.var("idx_x"), UPat(arg=Invalid)))), simplify_valid_image_load), ]) +# a valid index into an n-element buffer lives in [0,n): a gated long index narrows when n-1 fits int32 (out-of-gate wraps, discarded) +# TODO: more generic +pm_narrow_index = PatternMatcher([ + (UPat((Ops.INDEX, Ops.SHRINK), src=(UPat.var("buf"), UPat.var("gate").where(UPat.var("idx", dtypes.long), UPat(Ops.CONST, arg=Invalid))), + allow_any_len=True, name="u"), + lambda u,buf,gate,idx: u.replace(src=(buf, idx.cast(dtypes.int).valid(gate))+u.src[2:]) if buf.max_numel()-1 <= dtypes.int32.max else None), +]) + # get list of (height, width) that do not require pitch padding def image_valid_dims(base:DType, size:int, arch:str) -> list[tuple[int,int]]: if (ALIGN:=next((int(p.split('=')[1]) for p in arch.split(',') if p.startswith("IMAGE_PITCH_ALIGNMENT=")), 0)) == 0: return [] diff --git a/tinygrad/codegen/late/linearizer.py b/tinygrad/codegen/late/linearizer.py index e16778a367..e9f7254db8 100644 --- a/tinygrad/codegen/late/linearizer.py +++ b/tinygrad/codegen/late/linearizer.py @@ -43,7 +43,8 @@ def linearize(sink:UOp) -> list[UOp]: for v in u.src: out_degree[v] -= 1 if out_degree[v] == 0: heapq.heappush(heap, (-nkey[v],v)) - newlst = newlst[::-1] + # a renderer never emits a CONST, only the CAST over it: drop the value halves + newlst = [u for u in newlst[::-1] if u.op is not Ops.CONST] if getenv("DEBUG_LINEARIZE"): for i,u in enumerate(newlst): diff --git a/tinygrad/codegen/late/regalloc.py b/tinygrad/codegen/late/regalloc.py index 0d675e62ee..33aab5f5de 100644 --- a/tinygrad/codegen/late/regalloc.py +++ b/tinygrad/codegen/late/regalloc.py @@ -4,7 +4,8 @@ from tinygrad.uop.ops import UOp, Ops, PatternMatcher, UPat from tinygrad.renderer.isa import ISARenderer, Register, greg from tinygrad.dtype import dtypes -PSEUDO_OPS = {Ops.CONST, Ops.NOOP, Ops.AFTER, Ops.BARRIER, Ops.GROUP, Ops.STACK} +# CAST replaces CONST here: the linearizer drops every CONST, so a literal reaches regalloc as the CAST over it +PSEUDO_OPS = {Ops.CAST, Ops.NOOP, Ops.AFTER, Ops.BARRIER, Ops.GROUP, Ops.STACK} class LinearScanRegallocContext: # returns the uop that defines the virtual register diff --git a/tinygrad/renderer/cstyle.py b/tinygrad/renderer/cstyle.py index 0014952bdf..4a97a8c99b 100644 --- a/tinygrad/renderer/cstyle.py +++ b/tinygrad/renderer/cstyle.py @@ -3,6 +3,7 @@ import math, sys, struct from collections import defaultdict, Counter from tinygrad.codegen.opt import tc from tinygrad.uop.ops import GroupOp, Ops, UOp, PatternMatcher, UPat, range_str, axis_letters +from tinygrad.uop.weak import commit_weak from tinygrad.helpers import strip_parens, getenv, prod, dedup, Target, NUM_CPU_THREADS, IMAGE, FLOAT16, is_image_shape from tinygrad.dtype import dtypes, DType, AddrSpace, truncate, float_to_bf16 from tinygrad.renderer import Renderer @@ -19,20 +20,20 @@ base_rewrite = PatternMatcher([ (UPat(Ops.IF, name="x"), lambda ctx,x: f"if ({ctx[x.src[0]]}) {{"), (UPat((Ops.ENDIF, Ops.END)), lambda ctx: "}"), - # const - (UPat(Ops.CONST, dtype=dtypes.floats, name="x"), lambda ctx,x: None if math.isfinite(v:=x.val) else \ + # const, above the casting rules: a literal is the CAST that states its width over the CONST that carries its value + (UPat.cvar("c").cast(dtypes.floats, name="x"), lambda ctx,x,c: None if math.isfinite(v:=c.val) else \ f"({ctx.render_cast(x, ctx.nan if math.isnan(v) else ctx.infinity if v > 0 else f'-{ctx.infinity}')})"), - (UPat(Ops.CONST, dtype=dtypes.float, name="x"), lambda ctx,x: f"{x.val}f"), - (UPat(Ops.CONST, dtype=dtypes.int64, name="x"), lambda ctx,x: f"{x.val}l"), - (UPat(Ops.CONST, dtype=dtypes.uint64, name="x"), lambda ctx,x: f"{truncate[x.dtype](x.val)}ul"), - (UPat(Ops.CONST, dtype=dtypes.uint32, name="x"), lambda ctx,x: f"{truncate[x.dtype](x.val)}u"), - (UPat(Ops.CONST, dtype=dtypes.bool, name="x"), lambda ctx,x: "1" if x.val else "0"), + (UPat.cvar("c").cast(dtypes.float), lambda ctx,c: f"{c.val}f"), + (UPat.cvar("c").cast(dtypes.int64), lambda ctx,c: f"{c.val}l"), + (UPat.cvar("c").cast(dtypes.uint64, name="x"), lambda ctx,x,c: f"{truncate[x.dtype](c.val)}ul"), + (UPat.cvar("c").cast(dtypes.uint32, name="x"), lambda ctx,x,c: f"{truncate[x.dtype](c.val)}u"), + (UPat.cvar("c").cast(dtypes.bool), lambda ctx,c: "1" if c.val else "0"), # consts are rendered to larger type and casted - (UPat(Ops.CONST, (*dtypes.fp8s, dtypes.bfloat16, dtypes.half), name="x"), lambda ctx,x: f"({ctx.render_cast(x, f'{x.val}f')})"), - (UPat(Ops.CONST, (dtypes.uint8, dtypes.uint16), name="x"), lambda ctx,x: f"({ctx.render_cast(x, f'{x.val}u')})"), - (UPat(Ops.CONST, (dtypes.int8, dtypes.int16), name="x"), lambda ctx,x: f"({ctx.render_cast(x, str(x.val))})"), + (UPat.cvar("c").cast((*dtypes.fp8s, dtypes.bfloat16, dtypes.half), name="x"), lambda ctx,x,c: f"({ctx.render_cast(x, f'{c.val}f')})"), + (UPat.cvar("c").cast((dtypes.uint8, dtypes.uint16), name="x"), lambda ctx,x,c: f"({ctx.render_cast(x, f'{c.val}u')})"), + (UPat.cvar("c").cast((dtypes.int8, dtypes.int16), name="x"), lambda ctx,x,c: f"({ctx.render_cast(x, str(c.val))})"), # default const render - (UPat(Ops.CONST, name="x"), lambda ctx,x: str(x.val)), + (UPat.cvar("c").cast(), lambda ctx,c: str(c.val)), # casting (UPat(Ops.CAST, name="x"), lambda ctx,x: f"__builtin_convertvector({ctx[x.src[0]]}, {ctx.render_type(x)})" \ @@ -47,7 +48,7 @@ base_rewrite = PatternMatcher([ # SHRINK/INDEX (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var('idx')), name="x"), lambda ctx,**kwargs: ctx.render_index(**kwargs)), - (UPat(Ops.SHRINK, src=(UPat.var("buf"), UPat.var('idx'), UPat.cvar()), name="x"), lambda ctx,**kwargs: ctx.render_index(**kwargs)), + (UPat(Ops.SHRINK, src=(UPat.var("buf"), UPat.var('idx'), UPat.cvar().cast()), name="x"), lambda ctx,**kwargs: ctx.render_index(**kwargs)), (UPat(Ops.STACK, name="x"), lambda ctx,x: f"{ctx.float4.replace('float4', ctx.render_type(x))}" + \ f"{ctx.float4_style[0]}{','.join([ctx[y] for y in x.src])}{ctx.float4_style[1]}"), @@ -74,16 +75,22 @@ base_rewrite = PatternMatcher([ def create_non_native_float_pats(dts:tuple[DType, ...], casting:bool=True): patterns = PatternMatcher([ + # a weak operand commits at the emulated width first: the rules below move the computation to float + (UPat(GroupOp.ALU, dtype=dts, name="x"), lambda x: x.replace(src=tuple(commit_weak(s, x.dtype) if s.dtype in dtypes.weaks else s for s in x.src)) + if any(s.dtype in dtypes.weaks for s in x.src) else None), (UPat(Ops.WHERE, dtype=dts, src=(UPat.var("b"), UPat.var("x"), UPat.var("y")), name="w"), lambda w,b,x,y: b.where(x.cast(dtypes.float), y.cast(dtypes.float)).cast(w.dtype)), (UPat(GroupOp.ALU-{Ops.WHERE}, dtype=dts, name="x"), lambda x: UOp(x.op, src=tuple(vv.cast(dtypes.float) for vv in x.src), arg=x.arg).cast(x.dtype)), - (UPat(GroupOp.ALU, dtypes.bool, name="alu", src=(UPat.var("x", dtype=dts), UPat.var("y", dtype=dts))), - lambda alu,x,y: UOp(alu.op, src=(x.cast(dtypes.float), y.cast(dtypes.float)), arg=alu.arg))]) + # a comparison's dtype is bool, so the rule above cannot reach it. src=[...] only GATES; the rebuild takes alu.src in its OWN order + (UPat(GroupOp.ALU, dtypes.bool, name="alu", src=[UPat(dtype=dts), UPat(dtype=(*dts, *dtypes.weaks))]), + lambda alu: UOp(alu.op, src=tuple(s.cast(dtypes.float) for s in alu.src), arg=alu.arg))]) if casting: # add float intermediate casting patterns += PatternMatcher([ - (UPat(Ops.CAST, dts, (UPat.var("x"),), name="y"), lambda x,y: x.cast(dtypes.float).cast(y.dtype) if x.dtype!=dtypes.float else None), + # a literal declines: routing it through float nests a second cast + (UPat(Ops.CAST, dts, (UPat.var("x"),), name="y"), + lambda x,y: x.cast(dtypes.float).cast(y.dtype) if x.dtype!=dtypes.float and x.op is not Ops.CONST else None), (UPat(Ops.CAST, name="x", src=(UPat.var("y", dts),)), lambda x,y: y.cast(dtypes.float).cast(x.dtype) if x.dtype!=dtypes.float else None)]) return patterns @@ -161,8 +168,8 @@ class CStyleLanguage(Renderer): def render_index(self, x:UOp, buf:UOp, idx:UOp): if buf.addrspace == AddrSpace.ALU: # this is lane access in C - if idx.op is not Ops.CONST: return f"({self[buf]})[{self[idx]}]" - return self[buf]+(f"[{idx.val}]" if buf.max_numel() > self.gep_arr_threshold else f".{'xyzwabcd'[idx.val]}") + if idx.op is not Ops.CAST or idx.src[0].op is not Ops.CONST: return f"({self[buf]})[{self[idx]}]" + return self[buf]+(f"[{idx.src[0].val}]" if buf.max_numel() > self.gep_arr_threshold else f".{'xyzwabcd'[idx.src[0].val]}") return f"({self[buf]}+{strip_parens(self[idx]) if idx.arg == Ops.ADD else self[idx]})" def render_buffer(self, x:UOp): @@ -226,7 +233,7 @@ class CStyleLanguage(Renderer): if u.op is Ops.SPECIAL: r[u] = u.arg elif u.op is Ops.RANGE: r[u] = f"{axis_letters[u.arg[-1]]}idx"+range_str(u) else: - prefix = {Ops.WMMA: "wmma", Ops.CONST: "const", Ops.BUFFER: "buf", Ops.CAST: "cast", Ops.BITCAST: "cast", Ops.STACK: "cast", + prefix = {Ops.WMMA: "wmma", Ops.BUFFER: "buf", Ops.CAST: "cast", Ops.BITCAST: "cast", Ops.STACK: "cast", Ops.INDEX: "bidx", Ops.LOAD: "val"}.get(u.op, "alu") r[u] = f"{prefix}{c[prefix]}" @@ -234,7 +241,9 @@ class CStyleLanguage(Renderer): assert l is not None, f"failed to render {u.op} {u.dtype} {[(x.op,x.dtype) for x in u.src]} {u.arg}" if u.op in {Ops.ENDIF, Ops.END}: depth -= 1 - if (u.op is not Ops.CAST or u.max_numel() == 1) and (u.op in {Ops.CONST, Ops.INDEX, Ops.SHRINK, Ops.CUSTOMI} or \ + # a literal is data, not an instruction: it inlines at every use and never takes an SSA name + if (u.op is not Ops.CAST or u.max_numel() == 1) and ((u.op is Ops.CAST and u.src[0].op is Ops.CONST) or \ + u.op in {Ops.INDEX, Ops.SHRINK, Ops.CUSTOMI} or \ (u.op is Ops.LOAD and u.src[0].addrspace == AddrSpace.REG and child_count[u] == 1) or \ (u.op is Ops.CAST and u.addrspace in (AddrSpace.GLOBAL, AddrSpace.LOCAL)) or \ (u.op in {Ops.STACK, *(GroupOp.ALU-{Ops.WHERE}), Ops.CAST, Ops.BITCAST} and child_count[u] == 1 and not getenv("EXPAND_SSA"))): @@ -318,8 +327,7 @@ class OpenCLRenderer(CStyleLanguage): string_rewrite = PatternMatcher([ (UPat(Ops.BITCAST, name="x"), lambda ctx,x: f"as_{ctx.render_dtype(x.dtype)}(({ctx.render_dtype(x.src[0].dtype)})({ctx[x.src[0]]}))"), # bfloat16 constants need to be rendered as their bit pattern since bf16 is stored as ushort - (UPat(Ops.CONST, dtypes.bfloat16, name="x"), - lambda ctx,x: f"{(struct.unpack('I', struct.pack('f', float_to_bf16(x.val)))[0] >> 16)}u"), + (UPat.cvar("c").cast(dtypes.bfloat16), lambda ctx,c: f"{(struct.unpack('I', struct.pack('f', float_to_bf16(c.val)))[0] >> 16)}u"), # load/store image (OpenCL) (UPat.var('buf').index(UPat.var('idx_y'), UPat.var('idx_x')), lambda ctx,buf,idx_y,idx_x: f"IMAGE<{ctx[buf]}, {ctx[idx_y]}, {ctx[idx_x]}>"), (UPat(Ops.LOAD, dtype=dtypes.float, src=(UPat.var('buf').index(UPat.var('idx_y'), UPat.var('idx_x')), UPat.var("var"), UPat.var("gate"))), @@ -493,8 +501,8 @@ class HIPRenderer(CStyleLanguage): (UPat(Ops.WMMA, name="x"), lambda ctx,x: f"__{_wmma_name(x)}({ctx[x.src[0]]}, {ctx[x.src[1]]}, {ctx[x.src[2]]}," f" {fp8_index(x.src[0].dtype)}, {fp8_index(x.src[0].dtype)}, 0, 0, 0, 0)" if x.arg[0][2] == 128 else None), (UPat(Ops.WMMA, name="x"), lambda ctx,x: f"__{_wmma_name(x)}({ctx[x.src[0]]}, {ctx[x.src[1]]}, {ctx[x.src[2]]}, 0, 0, 0)"), - (UPat(Ops.CONST, dtypes.fp8s, name="x"), lambda ctx,x: - f"f32_to_fp8({ctx.nan if math.isnan(v:=x.val) else ctx.infinity if v == math.inf else f'-{ctx.infinity}' if v == -math.inf else f'{v}f'}," + (UPat.cvar("c").cast(dtypes.fp8s, name="x"), lambda ctx,x,c: + f"f32_to_fp8({ctx.nan if math.isnan(v:=c.val) else ctx.infinity if v == math.inf else f'-{ctx.infinity}' if v == -math.inf else f'{v}f'}," f" {fp8_index(x.dtype)})"), (UPat(Ops.CAST, dtypes.fp8s, (UPat(dtype=dtypes.float),), name="x",), lambda ctx,x: f"f32_to_fp8({ctx[x.src[0]]}, {fp8_index(x.dtype)})"), @@ -520,7 +528,7 @@ class HIPRenderer(CStyleLanguage): lambda x: x.replace(src=(x.src[0].bitcast(dtypes.uint64), x.src[1].bitcast(dtypes.uint64), x.src[2])) if x.src[0].max_numel() == 8 and x.src[0].dtype in dtypes.fp8_ocp else None), # bfloat16 constant casting - (UPat.cvar('x', dtypes.bfloat16), lambda x: cast_float_to_bf16(UOp.const(x.val, dtypes.float))), + (UPat.cvar('x').cast(dtypes.bfloat16), lambda x: cast_float_to_bf16(UOp.const(x.val, dtypes.float))), ]) def asm(self, prg:UOp, lin:UOp) -> bytes: @@ -536,7 +544,7 @@ class HIPRenderer(CStyleLanguage): prefix, ockl = [], [] type_map = { dtypes.bfloat16: "bf16", dtypes.float: "f32", dtypes.half: "f16", dtypes.fp8e4m3: "_fp8_fp8", dtypes.fp8e5m2: "_bf8_bf8" } used_dtypes = uops_to_dtypes(uops) - if any(u.op is Ops.CONST and not math.isfinite(u.val) for u in uops): + if any(u.op is Ops.CAST and u.src[0].op is Ops.CONST and not math.isfinite(u.src[0].val) for u in uops): prefix += ["#define INFINITY (__builtin_inff())", "#define NAN (__builtin_nanf(\"\"))"] if any(u.op is Ops.SPECIAL for u in uops): prefix.append("typedef long unsigned int size_t;") @@ -550,7 +558,7 @@ class HIPRenderer(CStyleLanguage): if any(dt in dtypes.fp8s for dt, _ in used_dtypes): prefix += ["typedef unsigned char hip_bf8;", "typedef unsigned char hip_fp8;"] if any((u.op is Ops.CAST and u.dtype in dtypes.fp8s and u.src[0].dtype == dtypes.float) or - (u.op is Ops.CONST and u.dtype in dtypes.fp8s) for u in uops): + (u.op is Ops.CAST and u.src[0].op is Ops.CONST and u.dtype in dtypes.fp8s) for u in uops): prefix.append("""static inline __attribute__((device)) unsigned char f32_to_fp8(float v, int is_bf8) { v = (((*(unsigned*)&v)&0x7F800000)!=0x7F800000)?__builtin_amdgcn_fmed3f(v,is_bf8?57344.0f:448.0f,is_bf8?-57344.0f:-448.0f) : v; return (unsigned char)(is_bf8?__builtin_amdgcn_cvt_pk_bf8_f32(v,v,0,false):__builtin_amdgcn_cvt_pk_fp8_f32(v,v,0,false));\n}""") diff --git a/tinygrad/renderer/isa/x86.py b/tinygrad/renderer/isa/x86.py index 85b4ece0d2..5291fc7902 100644 --- a/tinygrad/renderer/isa/x86.py +++ b/tinygrad/renderer/isa/x86.py @@ -221,15 +221,17 @@ reg_strs = {"rax": {4:"eax", 2:"ax", 1:"al"}, "rcx": {4:"ecx", 2:"cx", 1:"cl"}, # ***** X86 instruction selection ***** def base(x:UOp, i:int) -> UOp: return s.src[0] if (s:=x.src[i]).op is Ops.INDEX else s -def lane(x:UOp, i:int) -> int: return s.src[1].val if (s:=x.src[i]).op is Ops.INDEX else 0 +def lane(x:UOp, i:int) -> int: return s.src[1].src[0].val if (s:=x.src[i]).op is Ops.INDEX else 0 def to_int(dt:DType): return {dtypes.float16: dtypes.int16, dtypes.float32: dtypes.int32, dtypes.float64: dtypes.int64}[dt] def def_reg(dt:DType, reg:Register|None=None) -> UOp: return UOp(Ops.INS, dt, arg=X86Ops.DEFINE, tag=None if reg is None else (reg,)) -def imm(dt:DType, v:int) -> UOp: return UOp.const(truncate[dt](v), dt).rtag() +# an immediate is a literal, so it wears a literal's spelling: value on the CONST, width on the CAST. +# built raw because .cast(dt) folds away at the dtypes a CONST derives, and imm(bool) is one of them +def imm(dt:DType, v:int) -> UOp: return UOp(Ops.CAST, src=(UOp(Ops.CONST, arg=truncate[dt](v)),), arg=dt).rtag() def to_imm(c:UOp) -> UOp|None: - if c.op is not Ops.CONST: return None - if c.dtype is dtypes.int64: return imm(dtypes.int32, c.val) if not c.overflows(dtypes.int32) else None - if c.dtype is dtypes.uint64: return imm(dtypes.uint32, c.val) if not c.overflows(dtypes.uint32) else None - if c.dtype in dtypes.ints+(dtypes.bool,): return imm(c.dtype, c.val) + if c.op is not Ops.CAST or (v:=c.src[0]).op is not Ops.CONST: return None + if c.dtype is dtypes.int64: return imm(dtypes.int32, v.val) if not v.overflows(dtypes.int32) else None + if c.dtype is dtypes.uint64: return imm(dtypes.uint32, v.val) if not v.overflows(dtypes.uint32) else None + if c.dtype in dtypes.ints+(dtypes.bool,): return imm(c.dtype, v.val) return None def cmp(x:UOp) -> UOp: if x.src[0].dtype is dtypes.float32: return x.ins(X86Ops.VUCOMISS, dtype=dtypes.void) @@ -289,8 +291,10 @@ def fold_address(x:UOp) -> tuple[UOp, UOp, UOp, UOp]: # buffers are indexed by element, everything else (the stack pointer) by byte scale = base.dtype.itemsize if base.op in {Ops.PARAM, Ops.BUFFER, Ops.AFTER} else 1 sz = imm(dtypes.uint8, base.dtype.itemsize) - if idx.op is Ops.ADD and idx.src[1].op is Ops.CONST: return (base, _cast(idx.src[0]), _disp(idx.src[1].val * scale), sz) - if idx.op is Ops.CONST: return (base, UOp(Ops.NOOP), _disp(idx.val * scale), sz) + # a literal index folds into the displacement. this runs mid isel, where a CAST is a real cast unless it wraps a CONST + if idx.op is Ops.ADD and (c:=idx.src[1]).op is Ops.CAST and c.src[0].op is Ops.CONST: + return (base, _cast(idx.src[0]), _disp(c.src[0].val * scale), sz) + if idx.op is Ops.CAST and idx.src[0].op is Ops.CONST: return (base, UOp(Ops.NOOP), _disp(idx.src[0].val * scale), sz) return (base, _cast(idx), _disp(0), sz) def abi(ctx:IselContext, x:UOp) -> UOp|None: @@ -353,7 +357,8 @@ isel_matcher = PatternMatcher([ # cast of void is a noop (UPat.var("y").cast(name="x"), lambda y,x: y if y.dtype == dtypes.void else None), # range is lowered to acc, cmp, jmp after regalloc - (UPat(Ops.RANGE, src=(UPat.cvar("c"),), allow_any_len=True, name="x"), lambda c,x: x.replace(src=(imm(c.dtype, c.val),) + x.src[1:])), + (UPat(Ops.RANGE, src=(UPat.cvar("c").cast(name="b"),), allow_any_len=True, name="x"), + lambda b,c,x: x.replace(src=(imm(b.dtype, c.val),) + x.src[1:])), (UPat(Ops.RANGE, name="x"), lambda ctx,x: x.replace(tag=(ctx.vreg(WGPR),)) if not isinstance(x.tag, tuple) else None), # really all a backedge END is is an IF with a tag referencing the RANGE start label (UPat(Ops.END, src=(UPat(), UPat(), UPat(GroupOp.Comparison, name="cond")), name="x"), @@ -366,11 +371,11 @@ isel_matcher = PatternMatcher([ if not x.src or x.src[0].arg is not X86Ops.RET else None), # function abi constraints (UPat((Ops.PARAM, Ops.SPECIAL), name="x"), abi), - # constants that can't be immediates, move them to registers - (UPat.cvar("x", dtypes.int64s), lambda x: x.ins(X86Ops.MOVABS, src=(imm(x.dtype, x.val),)) if not x.tag else None), - (UPat.cvar("x", dtypes.ints+(dtypes.bool,)), lambda x: x.ins(X86Ops.MOVi, src=(imm(x.dtype, x.val),)) if not x.tag else None), - (UPat.cvar("x", dtypes.floats), lambda x: - UOp.const(struct.unpack((dt:=to_int(x.dtype)).fmt, struct.pack(x.dtype.fmt, x.val))[0], dt).bitcast(x.dtype) if not x.tag else None), + # constants that can't be immediates, move them to registers. the tag marks an already selected immediate, it rides on the width half + (UPat.cvar("c").cast(dtypes.int64s, name="x"), lambda c,x: x.ins(X86Ops.MOVABS, src=(imm(x.dtype, c.val),)) if not x.tag else None), + (UPat.cvar("c").cast(dtypes.ints+(dtypes.bool,), name="x"), lambda c,x: x.ins(X86Ops.MOVi, src=(imm(x.dtype, c.val),)) if not x.tag else None), + (UPat.cvar("c").cast(dtypes.floats, name="x"), lambda c,x: + UOp.const(struct.unpack((dt:=to_int(x.dtype)).fmt, struct.pack(x.dtype.fmt, c.val))[0], dt).bitcast(x.dtype) if not x.tag else None), # conditional moves that use masks NOTE: these currently assume a mask producing cmp exists (UPat.var("m").where(UPat.var("a", dtypes.int8s+dtypes.int16s+dtypes.int32s+(dtypes.int64,)), UPat.var("b")), lambda m,a,b: a.ins(X86Ops.VPBLENDVB, src=(b, a, m.replace(dtype=m.src[0].dtype))) if a.max_numel() > 1 else None), @@ -420,15 +425,15 @@ isel_matcher = PatternMatcher([ (UPat(Ops.STACK, dtypes.float32, name="x"), vinsertps), (UPat(Ops.STACK, dtypes.ints+(dtypes.bool,), name="x"), vpins), # INDEX on a vector register value extracts a single element - (UPat.var("y", dtypes.int8s+(dtypes.bool,)).index(UPat.cvar("c"), name="x"), + (UPat.var("y", dtypes.int8s+(dtypes.bool,)).index(UPat.cvar("c").cast(), name="x"), lambda y,c,x: x.ins(X86Ops.VPEXTRB, src=(y, imm(dtypes.uint8, c.val))) if _is_vec_xmm(y) else None), - (UPat.var("y", dtypes.int16s).index(UPat.cvar("c"), name="x"), + (UPat.var("y", dtypes.int16s).index(UPat.cvar("c").cast(), name="x"), lambda y,c,x: x.ins(X86Ops.VPEXTRW, src=(y, imm(dtypes.uint8, c.val))) if _is_vec_xmm(y) else None), - (UPat.var("y", dtypes.int32s).index(UPat.cvar("c"), name="x"), + (UPat.var("y", dtypes.int32s).index(UPat.cvar("c").cast(), name="x"), lambda y,c,x: x.ins(X86Ops.VPEXTRD, src=(y, imm(dtypes.uint8, c.val))) if _is_vec_xmm(y) else None), - (UPat.var("y", dtypes.int64s).index(UPat.cvar("c"), name="x"), + (UPat.var("y", dtypes.int64s).index(UPat.cvar("c").cast(), name="x"), lambda y,c,x: x.ins(X86Ops.VPEXTRQ, src=(y, imm(dtypes.uint8, c.val))) if _is_vec_xmm(y) else None), - (UPat.var("y", dtypes.floats).index(UPat.cvar("c"), name="x"), + (UPat.var("y", dtypes.floats).index(UPat.cvar("c").cast(), name="x"), lambda y,c,x: x.ins(X86Ops.VPSRLDQ, src=(y, imm(dtypes.uint8, c.val * x.dtype.itemsize))) if _is_vec_xmm(y) else None), # packed bitwise ((UPat() & UPat()).named("x"), lambda x: x.ins(X86Ops.VPAND) if x.max_numel() > 1 else None), @@ -453,15 +458,19 @@ isel_matcher = PatternMatcher([ # scalar int binary ((UPat(dtype=dtypes.ints).alu(Ops.CDIV, UPat())).named("x"), idiv), # scalar int binary with immediate - (UPat.var("a", dtypes.ints) << UPat.cvar("c"), lambda a,c: a.ins(X86Ops.SHLi, src=(a, imm(dtypes.uint8, c.val)))), - (UPat.var("a", dtypes.uints) >> UPat.cvar("c"), lambda a,c: a.ins(X86Ops.SHRi, src=(a, imm(dtypes.uint8, c.val)))), - (UPat.var("a", dtypes.sints) >> UPat.cvar("c"), lambda a,c: a.ins(X86Ops.SARi, src=(a, imm(dtypes.uint8, c.val)))), - (UPat.var("a", dtypes.ints) + UPat.cvar("c"), lambda a,c: a.ins(X86Ops.ADDi, src=(a, i)) if (i:=to_imm(c)) is not None else None), - (UPat.var("a", dtypes.ints) * UPat.cvar("c"), lambda a,c: a.ins(X86Ops.IMULi, src=(a, i)) if (i:=to_imm(c)) is not None else None), - (UPat.var("a", dtypes.ints+(dtypes.bool,)) & UPat.cvar("c"), lambda a,c: a.ins(X86Ops.ANDi, src=(a, i)) if (i:=to_imm(c)) is not None else None), - (UPat.var("a", dtypes.ints+(dtypes.bool,)) | UPat.cvar("c"), lambda a,c: a.ins(X86Ops.ORi, src=(a, i)) if (i:=to_imm(c)) is not None else None), - (UPat.var("a", dtypes.ints+(dtypes.bool,)) ^ UPat.cvar("c"), lambda a,c: a.ins(X86Ops.XORi, src=(a, i)) if (i:=to_imm(c)) is not None else None), - (UPat(Ops.SUB, dtypes.ints, (UPat.var("a"), UPat.cvar("c"))), lambda a,c: a.ins(X86Ops.SUBi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat.var("a", dtypes.ints) << UPat.cvar("c").cast(), lambda a,c: a.ins(X86Ops.SHLi, src=(a, imm(dtypes.uint8, c.val)))), + (UPat.var("a", dtypes.uints) >> UPat.cvar("c").cast(), lambda a,c: a.ins(X86Ops.SHRi, src=(a, imm(dtypes.uint8, c.val)))), + (UPat.var("a", dtypes.sints) >> UPat.cvar("c").cast(), lambda a,c: a.ins(X86Ops.SARi, src=(a, imm(dtypes.uint8, c.val)))), + (UPat.var("a", dtypes.ints) + UPat.cvar().cast(name="c"), lambda a,c: a.ins(X86Ops.ADDi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat.var("a", dtypes.ints) * UPat.cvar().cast(name="c"), lambda a,c: a.ins(X86Ops.IMULi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat.var("a", dtypes.ints+(dtypes.bool,)) & UPat.cvar().cast(name="c"), + lambda a,c: a.ins(X86Ops.ANDi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat.var("a", dtypes.ints+(dtypes.bool,)) | UPat.cvar().cast(name="c"), + lambda a,c: a.ins(X86Ops.ORi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat.var("a", dtypes.ints+(dtypes.bool,)) ^ UPat.cvar().cast(name="c"), + lambda a,c: a.ins(X86Ops.XORi, src=(a, i)) if (i:=to_imm(c)) is not None else None), + (UPat(Ops.SUB, dtypes.ints, (UPat.var("a"), UPat.cvar().cast(name="c"))), + lambda a,c: a.ins(X86Ops.SUBi, src=(a, i)) if (i:=to_imm(c)) is not None else None), # scalar int binary with register ((UPat(dtype=dtypes.ints) << UPat()).named("x"), lambda x: shift(x, X86Ops.SHL)), ((UPat(dtype=dtypes.uints) >> UPat()).named("x"), lambda x: shift(x, X86Ops.SHR)), @@ -572,7 +581,7 @@ def lower_range(ctx, x:UOp) -> tuple[UOp, list[UOp]]: if x.dtype is dtypes.void: return (label, [label]) else: acc = x.ins(X86Ops.MOVi, src=(imm(x.dtype, 0),) + x.src[1:]) - cmp = UOp(Ops.INS, arg=X86Ops.CMPi if x.src[0].op is Ops.CONST else X86Ops.CMP, src=(acc, x.src[0])) + cmp = UOp(Ops.INS, arg=X86Ops.CMPi if x.src[0].op is Ops.CAST else X86Ops.CMP, src=(acc, x.src[0])) jump_out = UOp(Ops.INS, arg=X86Ops.JGE, src=(cmp,), tag=f".LOOP_OUT_{loop_label}") ctx.loop_label[acc] = loop_label return (acc, [acc, label, cmp, jump_out]) @@ -605,6 +614,7 @@ post_regalloc_matcher = PatternMatcher([ # ***** X86 instruction encoding ***** +# isel turns every real CAST into an instruction, so a CAST reaching here is a literal: value on its CONST, width on itself def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> bytes|None: def _encode(reg_uop:UOp|None, rm_uop:UOp, idx_uop:UOp|None=None, disp_uop:UOp|None=None, sz_uop:UOp|None=None, vvvv_uop:UOp|None=None, imm_uop:UOp|None=None) -> bytes: @@ -614,7 +624,7 @@ def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> rm = cast(Register, greg(rm_uop)).index idx = cast(Register, greg(idx_uop)).index if idx_uop is not None and greg(idx_uop) is not None else 4 # for a memory operand the rm size is the element size from the address, otherwise it's the size of the value in the register - rm_sz = sz_uop.val if sz_uop is not None else rm_uop.dtype.itemsize + rm_sz = sz_uop.src[0].val if sz_uop is not None else rm_uop.dtype.itemsize reg_sz = reg_uop.dtype.itemsize if reg_uop is not None else 0 sz = reg_sz or rm_sz @@ -647,10 +657,10 @@ def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> # 0b10 -- signals memory access with 32bit displacement # 0b11 -- signals no memory access if disp_uop is not None: - assert disp_uop.op is Ops.CONST, "displacement must be a constant" + assert disp_uop.op is Ops.CAST, "displacement must be a literal" assert disp_uop.dtype in (dtypes.int8, dtypes.int32), "displacement can only be 1 or 4 byte signed int" # rbp/r13 always require a displacement - if disp_uop.val != 0 or rm == 0b101: mod = 0b01 if disp_uop.dtype.itemsize == 1 else 0b10 + if disp_uop.src[0].val != 0 or rm == 0b101: mod = 0b01 if disp_uop.dtype.itemsize == 1 else 0b10 else: mod = 0b00 else: mod = 0b11 # x 0b0 and idx 0b100 means rsp which means no index exists @@ -664,10 +674,10 @@ def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> # DISP byte if mod == 0b01 or mod == 0b10: assert disp_uop is not None - inst += struct.pack(unwrap(disp_uop.dtype.fmt), disp_uop.val) + inst += struct.pack(unwrap(disp_uop.dtype.fmt), disp_uop.src[0].val) # IMM byte if imm_uop is not None: - if imm_uop.op is Ops.CONST: inst += struct.pack(unwrap(imm_uop.dtype.fmt), imm_uop.val) + if imm_uop.op is Ops.CAST: inst += struct.pack(unwrap(imm_uop.dtype.fmt), imm_uop.src[0].val) elif isinstance(greg(imm_uop), Register): inst += bytes([(greg(imm_uop).index & 0b1111) << 4 | 0b0000]) return inst @@ -677,13 +687,13 @@ def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> if x.arg in X86GroupOp.WriteMem: if len(x.src) > 4: address, rest = x.src[:4], x.src[4:] else: address, rest = (x, None, None, None), x.src - imm_uop = rest[:1] if rest and rest[0].op is Ops.CONST else (None,) + imm_uop = rest[:1] if rest and rest[0].op is Ops.CAST else (None,) return _encode(rest[0], *address, *(None, *rest[1:])) if reg is None else _encode(None, *address, *(None, *imm_uop)) if x.arg in X86GroupOp.Rm1st: if len(x.src) > 3: address, rest = x.src[:4], x.src[4:] else: address, rest = (x.src[0], None, None, None), x.src[1:] - imm_uop = rest[:1] if rest and rest[0].op is Ops.CONST else (None,) + imm_uop = rest[:1] if rest and rest[0].op is Ops.CAST else (None,) return _encode(x, *address, *(None, *imm_uop)) if reg is None else _encode(None, *address, *(x if sel else None, *imm_uop)) if x.arg in X86GroupOp.Rm2nd: @@ -701,7 +711,7 @@ def encode(x:UOp, opc:int, reg:int|None=None, pp:int=0, sel:int=0, we:int=0) -> encodings = { # moves X86Ops.MOVABS: lambda x: - bytes([0b0100 << 4 | 0b1 << 3 | 0b00 << 2 | greg(x).index >> 3, 0xB8 + (greg(x).index & 0b111)]) + struct.pack(x.dtype.fmt, x.src[0].val), + bytes([0b0100 << 4 | 0b1 << 3 | 0b00 << 2 | greg(x).index >> 3, 0xB8 + (greg(x).index & 0b111)]) + struct.pack(x.dtype.fmt, x.src[0].src[0].val), X86Ops.MOV: lambda x: encode(x, 0x8B), X86Ops.MOVi: lambda x: encode(x, 0xC7, reg=0), X86Ops.MOVm: lambda x: encode(x, 0x89), X86Ops.LEA: lambda x: encode(x, 0x8D), X86Ops.VMOVSS: lambda x: encode(x, 0x10, pp=2, sel=1), X86Ops.VMOVSSm: lambda x: encode(x, 0x11, pp=2, sel=1), @@ -724,8 +734,8 @@ encodings = { X86Ops.VCVTPS2PD: lambda x: encode(x, 0x5A, pp=0, sel=1), X86Ops.VCVTPD2PS: lambda x: encode(x, 0x5A, pp=1, sel=1), X86Ops.VCVTTPS2DQ: lambda x: encode(x, 0x5B, pp=2, sel=1), X86Ops.VCVTTPD2DQ: lambda x: encode(x, 0xE6, pp=1, sel=1), # the int src is the 2nd src (the rm field), if it was folded into a memory operand its width is the element size of the address - X86Ops.VCVTSI2SS: lambda x: encode(x, 0x2A, pp=2, sel=1, we=(x.src[4].val if len(x.src) > 4 else x.src[1].dtype.itemsize) == 8), - X86Ops.VCVTSI2SD: lambda x: encode(x, 0x2A, pp=3, sel=1, we=(x.src[4].val if len(x.src) > 4 else x.src[1].dtype.itemsize) == 8), + X86Ops.VCVTSI2SS: lambda x: encode(x, 0x2A, pp=2, sel=1, we=(x.src[4].src[0].val if len(x.src) > 4 else x.src[1].dtype.itemsize) == 8), + X86Ops.VCVTSI2SD: lambda x: encode(x, 0x2A, pp=3, sel=1, we=(x.src[4].src[0].val if len(x.src) > 4 else x.src[1].dtype.itemsize) == 8), X86Ops.VCVTTSS2SI: lambda x: encode(x, 0x2C, pp=2, sel=1, we=x.dtype.itemsize == 8), X86Ops.VCVTTSD2SI: lambda x: encode(x, 0x2C, pp=3, sel=1, we=x.dtype.itemsize == 8), # int division @@ -840,10 +850,11 @@ class X86Renderer(ISARenderer): def _format_op(x:UOp) -> str: return f" {(o[7:-1] if (o:=str(x.arg))[-1] in ('i', 'm') else o[7:]).lower():7s}" def _format_operands(x:UOp) -> str: def _format(src:tuple[UOp, ...]) -> list[str]: - return [str(s.val) if s.op is Ops.CONST else reg_strs[o].get(s.dtype.itemsize, o) if \ + return [str(s.src[0].val) if s.op is Ops.CAST else reg_strs[o].get(s.dtype.itemsize, o) if \ (o:=str(greg(s))) in reg_strs else o for s in src if greg(s) is not None] def _mem_adress(base:UOp, idx:UOp, disp:UOp, sz:UOp) -> list[str]: - return [f"[{greg(base)}" + (f" + {greg(idx)}*{sz.val}" if greg(idx) else "") + (f" + {disp.val}" if disp.val else "") + "]"] + return [f"[{greg(base)}" + (f" + {greg(idx)}*{sz.src[0].val}" if greg(idx) else "") + \ + (f" + {disp.src[0].val}" if disp.src[0].val else "") + "]"] if len(x.src) > 4 and x.arg in X86GroupOp.WriteMem: ret = _mem_adress(*x.src[:4]) + _format(x.src[4:]) elif len(x.src) > 3 and x.arg in X86GroupOp.Rm1st: ret = _format((x,)) + _mem_adress(*x.src[:4]) + _format(x.src[4:]) diff --git a/tinygrad/renderer/llvmir.py b/tinygrad/renderer/llvmir.py index 3ce6fc8919..a1571a7b62 100644 --- a/tinygrad/renderer/llvmir.py +++ b/tinygrad/renderer/llvmir.py @@ -81,8 +81,8 @@ base_rewrite = PatternMatcher([ (UPat((Ops.INDEX, Ops.SHRINK), src=(UPat((Ops.BUFFER, Ops.PARAM, Ops.AFTER)),), allow_any_len=True, name="x"), lambda ctx,x: f" {ctx[x]} = getelementptr inbounds {ldt(x.dtype)}, {ldt(x.dtype, ptr=True)} {ctx[x.src[0]]}, {ldt(x.src[1].dtype)} {ctx[x.src[1]]}"), # register index - (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.cvar("idx")), name="x"), lambda ctx,buf,idx,x: - f" {ctx[x]} = extractelement {ldt(buf.dtype, buf.max_numel())} {ctx[buf]}, i32 {idx.val}" if buf.addrspace == AddrSpace.ALU else None), + (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.cvar("c").cast()), name="x"), lambda ctx,buf,c,x: + f" {ctx[x]} = extractelement {ldt(buf.dtype, buf.max_numel())} {ctx[buf]}, i32 {c.val}" if buf.addrspace == AddrSpace.ALU else None), # load/store (UPat(Ops.LOAD, src=(UPat.var("idx"), UPat.var("alt"), UPat.var("mask")), name="x"), @@ -185,7 +185,7 @@ class LLVMRenderer(Renderer): kernel.append(f" {r[u]} = addrspacecast [{size} x {ldt(u.dtype)}] addrspace(3)* @{r[u][1:]} to [{size} x {ldt(u.dtype)}]*") else: kernel.append(f" {r[u]} = alloca [{size} x {ldt(u.dtype)}], align 16") - elif u.op is Ops.CONST: r[u] = lconst(u.val, u.dtype) + elif u.op is Ops.CAST and u.src[0].op is Ops.CONST: r[u] = lconst(u.src[0].val, u.dtype) elif u.op is Ops.CAST and ldt(u.dtype) == ldt(u.src[0].dtype): r[u] = r[u.src[0]] # cast from signed to unsigned of the same size is a noop, or pointer cast else: diff --git a/tinygrad/renderer/nir.py b/tinygrad/renderer/nir.py index a2153394f6..da92c02afc 100644 --- a/tinygrad/renderer/nir.py +++ b/tinygrad/renderer/nir.py @@ -122,7 +122,8 @@ class NIRRenderer(Renderer): extra_matcher = PatternMatcher([ # handle negative unsigned CONST - (UPat.cvar("x", dtypes.uints), lambda x: UOp.const(x.dtype.max+x.val+1, x.dtype) if x.val < 0 else None), + (UPat(Ops.CAST, dtypes.uints, src=(UPat(Ops.CONST, name="x"),), name="c"), + lambda x,c: UOp.const(c.dtype.max+x.val+1, c.dtype) if x.val < 0 else None), # from ptx (UPat.var('x', dtype=dtypes.bool) uint8 @@ -137,7 +138,7 @@ class NIRRenderer(Renderer): (UPat(Ops.CAST, (dtypes.uchar, dtypes.ushort), src=(UPat.var("x", dtypes.floats),), name="c"), lambda x,c: x.cast(dtypes.int32).cast(c.dtype)), # load/store use pointer arithmetic, and the cast does nothing. NOTE: this doesn't apply to image indexing cause it's 1-D (UPat((Ops.INDEX, Ops.SHRINK), src=(UPat.var("buf"), UPat.var("off")), allow_any_len=True, name="x"), lambda x,buf,off: x.replace( - src=(buf,UOp.const(off.val, dtypes.long) if off.op is Ops.CONST else off.cast(dtypes.long))+x.src[2:]) + src=(buf,UOp.const(off.src[0].val, dtypes.long) if off.op is Ops.CAST and off.src[0].op is Ops.CONST else off.cast(dtypes.long))+x.src[2:]) if buf.addrspace != AddrSpace.REG and not is_image_shape(buf._shape) else None), # images need index to be int for nir (coordinates only: the INDEX keeps its access dtype) (UPat.var("buf").index(UPat.var("idx_y"), UPat.var("idx_x"), name="x"), @@ -145,7 +146,7 @@ class NIRRenderer(Renderer): ]) def_rewrite = PatternMatcher([ - (UPat(Ops.CONST, name="x"), lambda ctx,x: nimm(ctx.b, x.val, x.dtype)), + (UPat(Ops.CAST, src=(UPat(Ops.CONST, name="c"),), name="x"), lambda ctx,x,c: nimm(ctx.b, c.val, x.dtype)), (UPat(Ops.PARAM, name="x"), lambda ctx,x: ctx.param(ctx.b, x, x.dtype.itemsize if x.addrspace is AddrSpace.ALU else 8)), (UPat(Ops.SPECIAL, name="x"), lambda ctx,x: nchannel(ctx.b, {'g':ngid, 'l':nlid, 'i': nid}[x.arg[0]](ctx.b), int(x.arg[-1]))), (UPat(Ops.STORE, src=(UPat((Ops.INDEX, Ops.SHRINK), src=(UPat.var("buf"),UPat.var("off")), allow_any_len=True), UPat.var("val"))), @@ -186,7 +187,8 @@ class NIRRenderer(Renderer): def render(self, uops:list[UOp]): self.prerender(uops) - for u in [u for u in uops if u.op is Ops.SPECIAL and u.arg[0] == "l"]: self.b.shader.contents.info.workgroup_size[int(u.arg[-1])] = u.src[0].val + for u in [u for u in uops if u.op is Ops.SPECIAL and u.arg[0] == "l"]: + self.b.shader.contents.info.workgroup_size[int(u.arg[-1])] = u.src[0].src[0].val self.r: dict[UOp, Any] = {} self.param_idx = 0 ranges: list[mesa.nir_def|None] = [] @@ -195,7 +197,10 @@ class NIRRenderer(Renderer): if u.op in {Ops.NOOP, Ops.GROUP} or (u.op is Ops.STACK and len(u.src) == 0): pass elif u.op in {Ops.INDEX, Ops.SHRINK}: # INDEX on a register value picks the element, memory INDEX is handled in the LOAD/STORE patterns - if u.src[0].op not in {Ops.PARAM, Ops.BUFFER, Ops.AFTER}: self.r[u] = nchannel(self.b, self.r[u.src[0]], u.src[1].val) + if u.src[0].op not in {Ops.PARAM, Ops.BUFFER, Ops.AFTER}: + off = u.src[1] # a literal under width casts: the value lives on the bare CONST at the end of the CAST chain + while off.op is Ops.CAST: off = off.src[0] + self.r[u] = nchannel(self.b, self.r[u.src[0]], off.val) elif u.op is Ops.AFTER: self.r[u] = self.r[u.src[0]] elif u.op == Ops.SINK: diff --git a/tinygrad/renderer/ptx.py b/tinygrad/renderer/ptx.py index 30c73c2f93..f62f887581 100644 --- a/tinygrad/renderer/ptx.py +++ b/tinygrad/renderer/ptx.py @@ -79,8 +79,9 @@ def modifier(a: DType, b: DType): return '.rzi' if dtypes.is_int(a) and dtypes.i (a.itemsize < b.itemsize or dtypes.is_int(b) or b == dtypes.bool) else '' string_rewrite = PatternMatcher([ - (UPat.cvar("x", dtypes.bool), lambda ctx, x: f"setp.ne.s16 {ctx.r[x]}, {render_val(x.val, x.dtype)}, 0;"), - (UPat.cvar("x"), lambda ctx, x: f"mov.b{ctx.types[x.dtype][1:]} {ctx.r[x]}, {render_val(x.val, x.dtype)};"), + # a literal is CAST(dt, CONST(v)): above the generic CAST arms, which would read the dropped CONST's register + (UPat.cvar("c").cast(dtypes.bool, name="x"), lambda ctx, x, c: f"setp.ne.s16 {ctx.r[x]}, {render_val(c.val, x.dtype)}, 0;"), + (UPat.cvar("c").cast(name="x"), lambda ctx, x, c: f"mov.b{ctx.types[x.dtype][1:]} {ctx.r[x]}, {render_val(c.val, x.dtype)};"), (UPat(Ops.SPECIAL, name="x"), lambda ctx,x: f"mov.u32 %{x.arg}, %{'ctaid' if x.arg[0] == 'g' else 'tid'}.{chr(120+int(x.arg[-1]))};"), (UPat(Ops.PARAM, name="x"), lambda ctx, x: f"ld.param.{ctx.types[dtypes.ulong] if x.addrspace is AddrSpace.GLOBAL else ctx.mem_types[x.dtype]} {ctx.r[x]}, [data{x.arg.slot}+0];"), @@ -200,10 +201,10 @@ class PTXRenderer(Renderer): r[u] = [ssa("reg", u, self.types[u.dtype]) for _ in range(u.max_numel())] continue if u.op in {Ops.INDEX, Ops.SHRINK, Ops.LOAD} and u.src[0].addrspace in (AddrSpace.REG, AddrSpace.ALU): - # on REG, INDEX/SHRINK pick the register (must be CONST) and LOAD is a noop - if u.op is not Ops.LOAD and u.src[1].op is not Ops.CONST: + # on REG, INDEX/SHRINK pick the register (must be a literal) and LOAD is a noop + if u.op is not Ops.LOAD and (u.src[1].op is not Ops.CAST or u.src[1].src[0].op is not Ops.CONST): raise RuntimeError(f"PTX does not support dynamic register indexing: {u}") - r[u] = r[u.src[0]] if u.op is Ops.LOAD else r[u.src[0]][u.src[1].val] + r[u] = r[u.src[0]] if u.op is Ops.LOAD else r[u.src[0]][u.src[1].src[0].val] continue if u.op is Ops.SPECIAL: r[u] = "%" + u.arg elif u.op is Ops.LOAD: @@ -216,7 +217,7 @@ class PTXRenderer(Renderer): [ssa("wmma_acc", dtype="b32") for _ in range(0, len(r[u.src[2]]), 4 // u.dtype.itemsize)]] r[u] = [ssa("wmma", dtype=self.types[u.dtype]) for _ in range(u.max_numel())] prefix, dtype = {Ops.CAST: ("cast", None), Ops.BITCAST: ("cast", None), Ops.END: ("pred", "pred"), Ops.RANGE: ("ridx", None), - Ops.CONST: ("const", None), Ops.BUFFER: ("local", "u64"), Ops.INDEX: ("bidx", "u64"), Ops.SHRINK: ("bidx", "u64"), + Ops.BUFFER: ("local", "u64"), Ops.INDEX: ("bidx", "u64"), Ops.SHRINK: ("bidx", "u64"), Ops.PARAM: ("dat", "u64" if u.addrspace is AddrSpace.GLOBAL else None), **{op: ("alu", None) for op in GroupOp.ALU}}.get(u.op, (None, None)) if u.op is Ops.RANGE and u.dtype == dtypes.void: prefix = None # loop headers don't have a register if prefix: r[u] = ssa(prefix, u, dtype) diff --git a/tinygrad/renderer/wgsl.py b/tinygrad/renderer/wgsl.py index 83a9355e6b..66c265d792 100644 --- a/tinygrad/renderer/wgsl.py +++ b/tinygrad/renderer/wgsl.py @@ -69,10 +69,12 @@ class WGSLRenderer(CStyleLanguage): string_rewrite = PatternMatcher([ (UPat(Ops.NEG, dtypes.uints, src=(UPat.var('x'))), lambda ctx,x: f"(0-{ctx[x]})"), - (UPat.cvar("x", dtype=dtypes.bool), lambda x: "true" if x.val else "false"), - (UPat(Ops.CONST, dtype=(dtypes.uchar, dtypes.ushort, dtypes.uint32), name="x"), - lambda x: f"bitcast({x.val})" if x.val < 0 else f"{x.val&0xFFFFFFFF}u"), - (UPat(Ops.CONST, dtype=dtypes.int32, name="x"), lambda ctx,x: f"{truncate[x.dtype](x.val)}"), + (UPat.cvar("c").cast(dtypes.bool), lambda c: "true" if c.val else "false"), + (UPat.cvar("c").cast((dtypes.uchar, dtypes.ushort, dtypes.uint32)), + lambda c: f"bitcast({c.val})" if c.val < 0 else f"{c.val&0xFFFFFFFF}u"), + # type the i32 literal: a u32(...) conversion of a bare abstract int rejects any negative value. the conversion + # (not an `i` suffix) because the suffix binds tighter than unary minus, so INT_MIN would overflow while parsing + (UPat.cvar("c").cast(dtypes.int32, name="x"), lambda ctx,x,c: f"i32({truncate[x.dtype](c.val)})"), (UPat(Ops.BUFFER, name="x"), lambda ctx,x: f"var{'' if x.addrspace == AddrSpace.LOCAL else ''} {ctx[x]}: array<{ctx.buf_map(x)},{_packed_size(x)}>;"), (UPat(Ops.BITCAST, dtype=dtypes.half, name="x", src=(UPat(dtype=(dtypes.short, dtypes.ushort, dtypes.uint32),),)), diff --git a/tinygrad/runtime/ops_python.py b/tinygrad/runtime/ops_python.py index 0b2126deaa..f9be6f2d2b 100644 --- a/tinygrad/runtime/ops_python.py +++ b/tinygrad/runtime/ops_python.py @@ -56,7 +56,8 @@ class PythonProgram(Program['PythonDevice']): i = 0 while i < len(self.uops): u = self.uops[i] - src_values = [values[v] for v in u.src if v.dtype is not dtypes.void] + # a literal's CONST half is never executed, so it has no entry in values + src_values = [values[v] for v in u.src if v.dtype is not dtypes.void and v.op is not Ops.CONST] src_dtypes = [v.dtype for v in u.src if v.dtype is not dtypes.void] if getenv("TRACE"): print(i, u.op, u.dtype, u.arg, src_values, src_dtypes) if u.op is Ops.END: @@ -102,7 +103,7 @@ class PythonProgram(Program['PythonDevice']): elif u.op is Ops.SPECIAL: if u.arg[0] == 'g': values[u] = [idxs[2-int(u.arg[-1])]] * warp_size elif u.arg[0] == 'l': values[u] = [x[2-int(u.arg[-1])] for x in warp] - elif u.op is Ops.CONST: values[u] = [u.val] * warp_size + elif u.op is Ops.CAST and u.src[0].op is Ops.CONST: values[u] = [u.src[0].val] * warp_size elif u.op in {Ops.INDEX, Ops.SHRINK}: ret:list = [] if u.src[0].addrspace == AddrSpace.ALU: diff --git a/tinygrad/runtime/support/hcq2.py b/tinygrad/runtime/support/hcq2.py index 3655834410..877ab28ada 100644 --- a/tinygrad/runtime/support/hcq2.py +++ b/tinygrad/runtime/support/hcq2.py @@ -60,8 +60,8 @@ def make_binary_patch(buf:UOp, blob:bytes) -> UOp: def make_cmdbuf(lin, devs, buf:UOp|None=None): blob, patches = bytearray(), [] for s in (s for ins in lin.src for s in ins.src): - if s.op is not Ops.CONST: patches.append((len(blob), s)) - blob.extend(struct.pack(f'<{s.dtype.fmt}', s.val if s.op is Ops.CONST else 0x0)) + if not (lit:=(s.op is Ops.CAST and s.src[0].op is Ops.CONST)): patches.append((len(blob), s)) + blob.extend(struct.pack(f'<{s.dtype.fmt}', s.src[0].val if lit else 0x0)) cmdbuf = buf if buf is not None else UOp.placeholder((len(blob) // 4,), dtypes.uint32, next(UOp.unique_num), device=devs).rtag("cmdbuf") return cmdbuf.after(make_binary_patch(cmdbuf, bytes(blob)), *make_patches(cmdbuf, patches)) @@ -317,8 +317,8 @@ def make_scatter_loops(patches:list[UOp], inputs_table:tuple, lt_patches:list[UO data = [] for p in patches: words = [(off, val, get_getaddrs(val)) for off,val in zip(p.src[0].src[1].src, p.src[1].src)] - data += [(off.val, slots[gaddrs[0]]) for off,_,gaddrs in words if gaddrs][::2] - scalars = [(off.val*dst.dtype.itemsize, val) for off,val,gaddrs in words if not gaddrs] + data += [(off.src[0].val, slots[gaddrs[0]]) for off,_,gaddrs in words if gaddrs][::2] + scalars = [(off.src[0].val*dst.dtype.itemsize, val) for off,val,gaddrs in words if not gaddrs] subs[p] = UOp.group(*make_patches(dst, scalars)) if scalars else UOp(Ops.NOOP) word_table, slot_table = (UOp.placeholder((len(data),), dtypes.uint32, next(UOp.unique_num), device=dst.device).rtag("systems") for _ in range(2)) @@ -502,8 +502,8 @@ def fold_binary(buf:UOp, blob:UOp) -> UOp: def fold_const_store(buf:UOp, off:UOp, val:UOp) -> UOp: for off,val in zip(off.src, val.src): for b,v in zip((bs:=mb.bufs if isinstance((mb:=buf.buffer), MultiBuffer) else (mb,)), val.src if val.op is Ops.STACK else (val,)*len(bs)): - data = struct.pack(f'<{v.dtype.fmt}', truncate[v.dtype]((v.src[0] if v.op is Ops.CAST else v).val)) - b.ensure_allocated().as_memoryview(force_zero_copy=True, no_sync=True).cast('B')[(bo:=off.val*buf.dtype.itemsize):bo+len(data)] = data + data = struct.pack(f'<{v.dtype.fmt}', truncate[v.dtype](v.src[0].val)) + b.ensure_allocated().as_memoryview(force_zero_copy=True, no_sync=True).cast('B')[(bo:=off.src[0].val*buf.dtype.itemsize):bo+len(data)] = data return UOp(Ops.NOOP) def resolve_getaddr(buf:UOp, g:UOp) -> UOp: diff --git a/tinygrad/schedule/rangeify.py b/tinygrad/schedule/rangeify.py index 430fb52a6a..e2e09d4869 100644 --- a/tinygrad/schedule/rangeify.py +++ b/tinygrad/schedule/rangeify.py @@ -304,10 +304,12 @@ pm_const_buffer_folding = pm_mops+PatternMatcher([ (UPat(Ops.INDEX, name="idx").f(Ops.STAGE, allow_any_len=True, name="b2"), remove_noop_bufferize), (UPat(Ops.INDEX, src=(UPat(Ops.STAGE),), allow_any_len=True, name="idx").f(Ops.NOOP).f(Ops.STAGE, allow_any_len=True, name="b2"), remove_noop_bufferize), - # no buffers for const (ranges don't matter for const - it's the same value everywhere) + # no buffers for const (ranges don't matter for const - it's the same value everywhere). one rule per literal spelling (UPat(Ops.CONST, name='c').f(Ops.STAGE, allow_any_len=True, name="b"), lambda c,b: b.const_like(c.val)), + (UPat.cvar('c').cast().f(Ops.STAGE, allow_any_len=True, name="b"), lambda c,b: b.const_like(c.val)), # indexing a const is a const (UPat(Ops.INDEX, src=(UPat(Ops.CONST, name="c"),),), lambda c: c), + (UPat(Ops.INDEX, src=(UPat.cvar().cast(name="c"),),), lambda c: c), # indexing an after with all fully invalid stores is invalid (UPat(Ops.INDEX, src=(UPat(Ops.AFTER, name="after"),), allow_any_len=True, name="idx"), lambda idx,after: idx.const_like(Invalid) if after_all_invalid(after) else None), diff --git a/tinygrad/tensor.py b/tinygrad/tensor.py index d23c6ce33a..e11504465f 100644 --- a/tinygrad/tensor.py +++ b/tinygrad/tensor.py @@ -744,7 +744,8 @@ class Tensor(RandMixin): ref_frames = [x.contiguous() for x in ref_frames or []] assert frame_pos.is_bound_var, "frame_pos must be a bound Variable" srcs = (out:=Tensor.empty(*shape, device=self.device, dtype=self.dtype), self.contiguous(), state.contiguous(), *ref_frames) - fn = UOp(Ops.CUSTOM_FUNCTION, src=(frame_pos.src[0], *[UOp.const(s, dtypes.int) for s in shape]), arg="encdec") + # the dims are metadata, never emitted code. they must stay BARE: exec_encdec reads them back with `s.op is Ops.CONST` + fn = UOp(Ops.CUSTOM_FUNCTION, src=(frame_pos.src[0], *[UOp.const(s) for s in shape]), arg="encdec") return Tensor(out.uop.after(fn.call(*[s.uop for s in srcs], frame_pos))) P = ParamSpec("P") diff --git a/tinygrad/uop/ops.py b/tinygrad/uop/ops.py index f1f641fe54..a6847d561f 100644 --- a/tinygrad/uop/ops.py +++ b/tinygrad/uop/ops.py @@ -96,9 +96,9 @@ def multirange_str(rngs:Iterable[UOp], color=False, pad=None) -> str: def shape_to_shape_arg(arg:tuple[sint, ...]) -> UOp: for x in arg: if isinstance(x, UOp) and not dtypes.is_int(x.dtype): raise RuntimeError(f"shape must be int, got {x.dtype} in {arg}") - if len(arg) == 0: return UOp(Ops.STACK) - elif len(arg) == 1: return UOp.const(arg[0], dtypes.weakint) - else: return UOp(Ops.STACK, src=tuple(UOp.const(x) if isinstance(x, int) else x for x in arg)) + # a dim is metadata, never emitted code, so an int dim states no width + src = tuple(UOp.const(x) if isinstance(x, int) else x for x in arg) + return src[0] if len(src) == 1 else UOp(Ops.STACK, src=src) def consumer_map_from_toposort(lst:Iterable[UOp]): ret: dict[UOp, dict[UOp, None]] = {} @@ -169,6 +169,8 @@ def dtype_from_uop(op:Ops, src:tuple[UOp,...], arg:Any) -> DType|None: return dtypes.uint8 case Ops.CAST | Ops.BITCAST: assert isinstance(arg, DType), f"CAST/BITCAST arg must be DType, got {arg}" + # a BITCAST reinterprets its operand's BITS, so that operand must state a width and a weak value states none + if op is Ops.BITCAST and src[0].dtype in dtypes.weaks: raise RuntimeError(f"cannot bitcast a weak value, got {src[0].dtype}") return arg case Ops.CONST: # derived from the value. order matters: bool is an int subclass, ConstFloat is a float subclass @@ -187,14 +189,16 @@ class UOpMetaClass(type): ucache:dict[tuple, weakref.ReferenceType[UOp]] = {} def __call__(cls, op:Ops, dtype:DType|None=None, src:tuple[UOp,...]=tuple(), arg:Any=None, tag:Any=None, metadata:tuple[Metadata,...]|None=None, _buffer:Buffer|None=None): - if dtype is None: dtype = dtype_from_uop(op, src, arg) or dtypes.void - # CONST derives its dtype by value only when the constructor omits one + # ONE derivation, two duties: it FILLS an omitted dtype and is the SPEC=2 answer (asking with an Invalid src can RAISE) + derived = dtype_from_uop(op, src, arg) if dtype is None or (SPEC == 2 and not any(s.base.is_invalid for s in src)) else None + if dtype is None: dtype = derived or dtypes.void + # a literal's value belongs on its CAST's grid: re-mint it there, unless the re-mint has the cast's own dtype + if op is Ops.CAST and src[0].op is Ops.CONST and (c:=UOp(Ops.CONST, arg=dtype.const(src[0].arg))).dtype is not dtype: src = (c,) # TODO: delete this once the dtype field is removed, for now it just re-implements spec.py # an INDEX presents its access dtype, which a still-weak source matches up to weakness - if SPEC == 2 and op is not Ops.CONST and \ - not any(s.base.is_invalid for s in src) and (expected_dtype:=dtype_from_uop(op, src, arg)) is not None and expected_dtype != dtype and \ - not (op is Ops.INDEX and weak_dtype(expected_dtype) == weak_dtype(dtype)): - raise RuntimeError(f"bad dtype {dtype}, expected {expected_dtype} on {op}") + if SPEC == 2 and derived is not None and derived != dtype and \ + not (op is Ops.INDEX and weak_dtype(derived) == weak_dtype(dtype)): + raise RuntimeError(f"bad dtype {dtype}, expected {derived} on {op}") if (wret:=UOpMetaClass.ucache.get(key:=(op, dtype, src, arg, tag), None)) is not None and (ret:=wret()) is not None: return ret UOpMetaClass.ucache[key] = weakref.ref(created:=super().__call__(*key)) if metadata is not None: all_metadata[created] = metadata @@ -257,6 +261,7 @@ class UOp(RandMixin, metaclass=UOpMetaClass): if (self.op, self.dtype, self.src, self.arg, self.tag) == new_args: return self return UOp(*new_args) def rtag(self, tag=True): return self.replace(tag=tag) + # the value is minted on the cast's grid, so .val on the CONST reads the literal at its stated width in either spelling @property def val(self): assert self.op is Ops.CONST, f"val is only for CONST, got {self.op}" @@ -614,7 +619,8 @@ class UOp(RandMixin, metaclass=UOpMetaClass): if isinstance(b, UOp): return b.cast(dtype) # NOTE: it always has to be STACK now, even if they are all the same if isinstance(b, tuple): return UOp.stack(*[UOp.const(c, dtype) for c in b]) - return UOp(Ops.CONST, dtype, arg=dtype.const(b), src=()) + # .cast folds away at exactly the dtypes a CONST derives (bool/weakint/weakfloat): bare there, the pair everywhere else + return UOp(Ops.CONST, arg=dtype.const(b), src=()).cast(dtype) @staticmethod def range(end:sint, axis_id, axis_type=AxisType.WEAK, *arg, dtype=dtypes.weakint, src=(), **kwargs): return UOp(Ops.RANGE, src=(sint_to_uop(end, dtype),)+src, arg=(axis_id, axis_type)+arg, **kwargs) @@ -987,7 +993,8 @@ class UOp(RandMixin, metaclass=UOpMetaClass): return unwrap(self.arg.name) def bind(self, val:int|UOp): assert self.is_variable, f"op is {self.op}, need Variable" - uval = self.const_like(val) if isinstance(val, int) else val + # the Variable states the width, so the bound value stays BARE: is_bound_var tests for a CONST there, unbind reads .val + uval = UOp.const(val) if isinstance(val, int) else val assert self.vmin <= uval.vmin and uval.vmax <= self.vmax, f"bind {val} not in range [{self.vmin}, {self.vmax}]" assert uval.divides(self.arg.multiple_of) is not None, f"bind {val} not divisible by {self.arg.multiple_of}" return self.after(self.store(uval)) diff --git a/tinygrad/uop/render.py b/tinygrad/uop/render.py index 6718d4636e..d474f72235 100644 --- a/tinygrad/uop/render.py +++ b/tinygrad/uop/render.py @@ -18,7 +18,8 @@ def pretty_print(x:UOp, cache=None, d=0)->str: def print_uops(uops:list[UOp]): uops_index = {u:i for i,u in enumerate(uops)} for i,u in enumerate(uops): - formatted_srcs = [(uops_index[x] if x.op is not Ops.CONST else f"{x.val}") if x in uops else "--" for x in u.src] + # a literal is data, not a program point: print its value + formatted_srcs = [(f"{x.src[0].val}" if x.op is Ops.CAST and x.src[0].op is Ops.CONST else uops_index[x]) if x in uops else "--" for x in u.src] print(f"{i:4d} {str(u.op):20s}: {multirange_str(u.ranges, color=True, pad=10)} {str(u.dtype):40s} " f"{str(formatted_srcs):32s} {u.arg}") # for debug @@ -38,6 +39,7 @@ renderer = PatternMatcher([ (UPat((Ops.SPECIAL), name="x"), lambda x: x.arg), (UPat(Ops.RANGE, dtypes.void, name="x"), lambda x: f"loop{x.arg[0]}"), (UPat(Ops.RANGE, name="x"), lambda x: f"r{range_str(x)}"), + (UPat.cvar("c").cast(), lambda c: str(c.val)), (UPat(Ops.CONST, name="x"), lambda x: str(x.val)), (UPat(Ops.CAST, name="x"), lambda ctx,x: f"({str(x.dtype)[7:]})({ctx[x.src[0]]})"), (UPat(Ops.NEG, name="x"), lambda ctx,x: f"(-{ctx[x.src[0]]})"), @@ -80,8 +82,9 @@ def render_marg(ctx,x:UOp): sugar = {Ops.SINK, Ops.END, Ops.STORE, Ops.LOAD, Ops.SQRT, Ops.INDEX, Ops.REDUCE, Ops.AFTER, Ops.THREEFRY, Ops.RECIPROCAL, Ops.EXP2, Ops.LOG2, Ops.SIN, Ops.CONTIGUOUS, Ops.BARRIER, Ops.DETACH} pm_pyrender_extra = PatternMatcher([ - (UPat(Ops.CONST, src=(), name="x"), lambda x: f"UOp.const({x.val}, {x.dtype})"), - (UPat((Ops.CAST, Ops.BITCAST), name="x"), lambda ctx,x: f"{ctx[x.src[0]]}.{x.op.name.lower()}({x.dtype})"), + (UPat(Ops.CONST, src=(), name="x"), lambda x: f"UOp.const({x.val})"), # a bare CONST derives its dtype from the value + # a same-dtype cast does not round-trip (.cast folds it away), so fall through to the raw UOp(...) form + (UPat((Ops.CAST, Ops.BITCAST), name="x"), lambda ctx,x: f"{ctx[x.src[0]]}.{x.op.name.lower()}({x.dtype})" if x.dtype != x.src[0].dtype else None), (UPat(Ops.SPECIAL, src=(UPat(Ops.CONST),), name="x"), lambda x: f"UOp.special({x.src[0].val}, {repr(x.arg)}, dtype={x.dtype})"), (UPat(Ops.BUFFER, src=(UPat(),), name="x"), lambda x: f"UOp.new_buffer({repr(x.arg.device)}, {x.max_numel()}, {x.dtype}, {x.arg.slot})" diff --git a/tinygrad/uop/spec.py b/tinygrad/uop/spec.py index 35fa4bb94f..791a9fefb0 100644 --- a/tinygrad/uop/spec.py +++ b/tinygrad/uop/spec.py @@ -1,6 +1,6 @@ import math from typing import Any -from tinygrad.uop.ops import PatternMatcher, UPat, GroupOp, Ops, UOp, AxisType, KernelInfo, ParamArg +from tinygrad.uop.ops import PatternMatcher, UPat, GroupOp, Ops, UOp, AxisType, KernelInfo, ParamArg, dtype_from_uop from tinygrad.uop.render import print_uops, pyrender from tinygrad.dtype import DType, dtypes, AddrSpace, Invalid, ConstFloat from tinygrad.helpers import DEBUG, Context, SPEC, Metadata, panic, CHECK_OOB, all_same, is_image_shape @@ -53,8 +53,8 @@ spec_shared = PatternMatcher([ # NOOP. TODO: remove this (UPat(Ops.NOOP), lambda: True), - # CONST is everywhere; Invalid is a bool const - (UPat(Ops.CONST, src=(), name="x"), lambda x: x.dtype is dtypes.bool if x.is_invalid else type(x.val) is type(x.dtype.const(x.val))), + # a CONST states no width: its dtype must be the one DERIVED from its arg. any other dtype is spelled CAST(dt, CONST(v)) + (UPat(Ops.CONST, src=(), name="x"), lambda x: x.dtype is dtype_from_uop(Ops.CONST, x.src, x.arg)), # STACK is everywhere too (UPat(Ops.STACK, dtype=dtypes.void, src=()), lambda: True), @@ -73,7 +73,7 @@ spec_shared = PatternMatcher([ (UPat(GroupOp.ALU, name="x"), lambda x: all(matches_dtype(y, x.dtype) or y.dtype in dtypes.weaks for y in x.src)), # CAST - (UPat((Ops.BITCAST, Ops.CAST), src=(UPat(),), name="x"), lambda x: isinstance(x.arg, DType)), + (UPat((Ops.BITCAST, Ops.CAST), src=(UPat(),), name="x"), lambda x: isinstance(x.arg, DType) and x.dtype is x.arg), # RANGE can be in the big graph now. a void RANGE is a bound-less loop header, the arg is an axis id like RANGE (UPat(Ops.RANGE, src=(UPat.var("x"),), allow_any_len=True, name="rng"), lambda rng,x: @@ -200,11 +200,12 @@ spec_tensor = PatternMatcher([ # these ops can exist in programs but not the tensor spec. example: LOAD spec_program = PatternMatcher([ - # index and weak dtypes are not allowed in programs - (UPat(GroupOp.All, (dtypes.weakint, dtypes.weakfloat)), lambda: False), + # every width in a program is stated: a CONST appears only under the CAST stating its width, and is the only weak node + (UPat(GroupOp.All, name="x"), lambda x: False if x.op is not Ops.CAST and any(s.op is Ops.CONST for s in x.src) else None), + (UPat(GroupOp.All, dtypes.weaks, name="x"), lambda x: None if x.op is Ops.CONST else False), # allow special SHRINK - (UPat(Ops.SHRINK, src=(UPat((Ops.PARAM, Ops.BUFFER, Ops.AFTER)), UPat(), UPat(Ops.CONST))), lambda: True), + (UPat(Ops.SHRINK, src=(UPat((Ops.PARAM, Ops.BUFFER, Ops.AFTER)), UPat(), UPat(Ops.CAST, src=(UPat(Ops.CONST),)))), lambda: True), # movement ops are not allowed in programs (UPat(GroupOp.Movement), lambda: False), @@ -249,9 +250,11 @@ spec_kernel_graph = PatternMatcher([ (UPat(Ops.SINK, dtypes.void), lambda: True), # the store of a bound Variable binds it: AFTER(BUFFER, STORE(BUFFER, CONST)) in call args (UPat(Ops.STORE, dtypes.void, (UPat(Ops.BUFFER, name="b"), UPat(Ops.CONST))), lambda b: b.is_variable), - # const + stack to make vconsts and shape args + # const + stack to make vconsts and shape args, in both spellings -- each at the dtypes its mint fixes (UPat(Ops.CONST, src=()), lambda: True), - (UPat(Ops.STACK, name="s"), lambda s: all(x.op in (Ops.CONST, Ops.PARAM) or x.is_variable or x.is_bound_var for x in s.src) or None), + (UPat(Ops.CAST, src=(UPat(Ops.CONST, src=()),)), lambda: True), + (UPat(Ops.STACK, name="s"), lambda s: all(x.op in (Ops.CONST, Ops.PARAM) or (x.op is Ops.CAST and x.src[0].op is Ops.CONST) + or x.is_variable or x.is_bound_var for x in s.src) or None), # linear for more kernels (TODO: we should enter non sink calls) #(UPat(Ops.LINEAR), lambda: True), # param is outside buffer, buffer is local buffer diff --git a/tinygrad/uop/symbolic.py b/tinygrad/uop/symbolic.py index 647460639d..4897cb7387 100644 --- a/tinygrad/uop/symbolic.py +++ b/tinygrad/uop/symbolic.py @@ -1,12 +1,12 @@ # all of symbolic lives here now import math from collections import defaultdict -from tinygrad.uop.ops import Ops, PatternMatcher, UPat, UOp, GroupOp, exec_alu -from tinygrad.dtype import PyConst, ConstType, dtypes, can_lossless_cast, Invalid, bitcast +from tinygrad.uop.ops import Ops, PatternMatcher, UPat, UOp, GroupOp, exec_alu, promo_dtype +from tinygrad.dtype import PyConst, ConstType, dtypes, can_lossless_cast, Invalid, bitcast, weak_dtype, truncate from tinygrad.helpers import partition, all_same, prod, flatten, unwrap, IMAGE, dedup from tinygrad.uop.divandmod import div_and_mod_symbolic from tinygrad.uop.movement import mop_cleanup -from tinygrad.uop.weak import commit_weak +from tinygrad.uop.weak import pm_uncast_const, commit_weak # TODO: symbolic shouldn't be importing from codegen from tinygrad.codegen.decomp.transcendental import xpow @@ -15,27 +15,18 @@ from tinygrad.codegen.decomp.transcendental import xpow def simplify_pow(x:UOp, c:UOp) -> UOp|None: if c.val < 0: return x.reciprocal().pow(-c.val) - if c.val == 0: return x.const_like(1) + # mint the 1 bare so x*1 -> x fires in the recursion below; the consumer restores the width + if c.val == 0: return x.const_like(1, weak_dtype(x.dtype)) if int(c.val-0.5)+0.5 == c.val: return x.pow(c.val-0.5) * x.sqrt() if int(c.val) == c.val: return (y := x.pow(c.val//2)) * y * (x if c.val%2 == 1 else 1) return None -def fold_bitcast(root:UOp, c:UOp) -> UOp|None: +def fold_bitcast(root:UOp, c:UOp, v:ConstType) -> UOp|None: if c.dtype.fmt is None or root.dtype.fmt is None or c.dtype.itemsize != root.dtype.itemsize: return None - return root.const_like(bitcast(c.val, c.dtype, root.dtype)) + # the value is MATHEMATICAL and may not fit: reading it as bits is the emission that pins it to the stated width + return root.const_like(bitcast(truncate[c.dtype](v), c.dtype, root.dtype)) -# const folding works for CONST, STACK, and casted CONST -const_folding_pat = UPat.any(UPat((Ops.CONST, Ops.STACK)), UPat(Ops.CAST, src=(UPat(Ops.CONST),))) - -def const_arg(u:UOp) -> ConstType|tuple[ConstType, ...]|None: - if u.op is Ops.CONST: return u.val - if u.op is Ops.CAST and u.src[0].op is Ops.CONST: return u.dtype.const(u.src[0].val) - if u.op is Ops.STACK and all(s.op is Ops.CONST for s in u.src): return tuple(s.val for s in u.src) - return None - -def fold_const_alu(a:UOp) -> UOp|None: - vals = [const_arg(s) for s in a.src] - return None if any(v is None for v in vals) else a.const_like(exec_alu(a.op, a.dtype, vals, False)) +def fold_const_alu(a:UOp, vals) -> UOp: return a.const_like(exec_alu(a.op, a.dtype, vals, False)) def _quotient_base(q:UOp, base:UOp, div:int) -> UOp|None: # the B with q == B//div and B%div == base%div, or None. only such congruence is needed to recombine, and canonicalization @@ -71,6 +62,12 @@ def fold_add_divmod_recombine(x:UOp) -> UOp|None: # this needs to be before symbolic so that 0*something_that_might_be_invalid doesnt become 0 invalid_pat = UPat(Ops.CONST, arg=Invalid, name="i") invalid_gate = UPat.var("cond").where(UPat.var("x"), invalid_pat) + +# the two literal spellings and the value reader for each: Invalid carries no width, so it rides bare inside either +bare_lit = UPat.any(UPat(Ops.CONST), UPat(Ops.STACK, src=UPat(Ops.CONST))) +pair_lit = UPat.any(p:=UPat(Ops.CAST, src=(UPat(Ops.CONST),)), UPat(Ops.STACK, src=UPat.any(p, UPat(Ops.CONST, arg=Invalid)))) +def bare_arg(u:UOp): return tuple(s.val for s in u.src) if u.op is Ops.STACK else u.val +def pair_arg(u:UOp): return tuple(s.val if s.is_invalid else s.src[0].val for s in u.src) if u.op is Ops.STACK else u.src[0].val pm_data_invalid = PatternMatcher([ (invalid_pat.broadcast(), lambda i: i), (UPat(GroupOp.Unary|{Ops.CAST, Ops.BITCAST}, src=(invalid_pat,)), lambda i: i), @@ -136,11 +133,15 @@ symbolic_simple = pm_data_invalid + PatternMatcher([ lambda x,mask,c: x // c.val if c.val > 0 and c.val & (c.val-1) == 0 and mask.val | (c.val-1) == -1 else None), (UPat.var("x", dtype=dtypes.ints+(dtypes.bool, dtypes.weakint)) != UPat.var("x"), lambda x: x.const_like(False, dtypes.bool)), # x != x -> False (only ints) - # ** constant folding ** - (UPat(GroupOp.Unary, src=(const_folding_pat,), name="a"), fold_const_alu), + # ** constant folding ** one rule per spelling combination: all-bare folds with no width involved, all-pair evaluates + # each operand at its stated width, mixed normalizes the bare operand to the promotion (then all-pair takes it). + # these arms mint pairs; stripping a redundant one is pm_uncast_const's duty in `symbolic` below # NOTE: THREEFRY(const,const) folds via its decomposition - (UPat(GroupOp.Binary-{Ops.THREEFRY}, src=(const_folding_pat,)*2, name="a"), fold_const_alu), - (UPat(GroupOp.Ternary, src=(const_folding_pat,)*3, name="a"), fold_const_alu), + (UPat(GroupOp.ALU-{Ops.THREEFRY}, src=bare_lit, name="a"), lambda a: fold_const_alu(a, [bare_arg(s) for s in a.src])), + (UPat(GroupOp.ALU-{Ops.THREEFRY}, src=pair_lit, name="a"), lambda a: fold_const_alu(a, [pair_arg(s) for s in a.src])), + (UPat(GroupOp.Binary-{Ops.THREEFRY}, src=[pair_lit, bare_lit], name="a"), lambda a: + a.replace(dtype=None, src=tuple(UOp.const(bare_arg(s), dt) if s.dtype in dtypes.weaks else s for s in a.src)) + if (dt:=promo_dtype(a.src)) not in dtypes.weaks else None), # bool MUL is AND, ADD/MAX is OR. prevents other rules to rewrite bool ADD/MUL incorrectly (UPat.var('x', dtype=dtypes.bool) * UPat.var('y', dtype=dtypes.bool), lambda x,y: x&y), (UPat.var('x', dtype=dtypes.bool) + UPat.var('y', dtype=dtypes.bool), lambda x,y: x|y), @@ -158,7 +159,10 @@ symbolic_simple = pm_data_invalid + PatternMatcher([ and isinstance(x.val, float) and (math.isnan(x.val) or math.isinf(x.val)) else 0)), # *** cast/bitcast *** (UPat((Ops.CAST, Ops.BITCAST), name="root"), lambda root: root.src[0] if root.dtype == root.src[0].dtype else None), - (UPat(Ops.BITCAST, name="root", src=(UPat.cvar("c"),)), fold_bitcast), + # a BITCAST reads its operand at the width it STATES, so a weak const is nonsense here: the bare arm is bool only + (UPat(Ops.BITCAST, name="root", src=(UPat(Ops.CONST, dtypes.bool, name="c"),)), lambda root,c: fold_bitcast(root, c, c.val)), + (UPat(Ops.BITCAST, name="root", src=(UPat(Ops.CAST, src=(UPat(Ops.CONST, name="c"),), name="lit"),)), + lambda root,lit,c: fold_bitcast(root, lit, c.val)), # b.cast(a).cast(b) -> b if a preserves all values in b (UPat.var('x').cast(name="a").cast(name="b"), lambda x,a,b: x if x.dtype == b.dtype and can_lossless_cast(b.dtype, a.dtype) else None), # bitcast twice @@ -284,13 +288,12 @@ symbolic = symbolic_simple+commutative+PatternMatcher([ # cast/long folding # if the intermediate cast doesnt narrow we can do it in one cast (UPat.var('x').cast(name="a").cast(name="b"), lambda x,a,b: x.cast(b.dtype) if can_lossless_cast(x.dtype, a.dtype) else None), + # commit_weak, not .cast: a weak b.dtype is not a literal spelling, and a CAST(weakfloat, CONST) reaches no commit round (UPat.var('x', dtypes.ints+(dtypes.weakint,)).cast(dtypes.ints+(dtypes.weakint,), name="a").cast(name="b"), - lambda x,a,b: x.cast(b.dtype) if a.dtype.min<=x.vmin and x.vmax<=a.dtype.max else None), - # try to do math in int instead of long, keep weak const weak + lambda x,a,b: commit_weak(x, b.dtype) if a.dtype.min<=x.vmin and x.vmax<=a.dtype.max else None), + # try to do math in int instead of long (UPat(GroupOp.Binary, src=(UPat.var("x", dtypes.long), UPat.var("y", dtypes.long)), name="u"), lambda u,x,y: - (UOp.const(x.val) if x.op is Ops.CONST else x.cast(dtypes.int)).alu(u.op, - UOp.const(y.val) if y.op is Ops.CONST else y.cast(dtypes.int)).cast(u.dtype) - if not any(v.overflows(dtypes.int) for v in (u,x,y)) else None), + x.cast(dtypes.int).alu(u.op, y.cast(dtypes.int)).cast(u.dtype) if not any(v.overflows(dtypes.int) for v in (u,x,y)) else None), ((UPat.var("x", dtypes.weakint) + UPat.cvar("c")).cast(dtypes.sints, name="cast"), lambda x,c,cast:x.cast(cast.dtype)+cast.const_like(c.val)), # only RANGE/IF/STORE/KERNEL have side effects (UPat(Ops.AFTER, name="x"), lambda x: x.replace(src=(x.src[0],)+ @@ -298,7 +301,12 @@ symbolic = symbolic_simple+commutative+PatternMatcher([ else y.src for y in x.src[1:]]))))), # after/end with 1 src is just src[0] (UPat((Ops.AFTER, Ops.END), src=(UPat.var("s"),)), lambda s: s), -])+div_and_mod_symbolic + # a CAST over a typed literal is a value conversion: evaluate at the stated width, keep the outer cast. NOT in + # symbolic_simple -- the emulated-float rounds compose that with a producer that re-expands the pair (bf16 cycle) + (UPat(Ops.CAST, dtypes.all, name="root", src=(UPat(Ops.CAST, dtypes.all, src=(UPat(Ops.CONST, name="c"),)),)), + lambda root,c: root.const_like(c.val)), + # the rules above key on bare CONSTs, so a const wearing a redundant width statement has to be unwrapped in the same fixpoint +])+div_and_mod_symbolic+pm_uncast_const # ******** we take a small aside to "simplify_valid" to rewrite valids ******** diff --git a/tinygrad/uop/validate.py b/tinygrad/uop/validate.py index 2dfb9025ea..a408c28e4c 100644 --- a/tinygrad/uop/validate.py +++ b/tinygrad/uop/validate.py @@ -45,13 +45,15 @@ z3_renderer = PatternMatcher([ (UPat((Ops.LOAD, Ops.INDEX), dtypes.bool), lambda ctx: (z3.Bool(f"load{len(ctx[1])}", ctx=ctx[0]), None)), # constants (UPat(Ops.CONST, arg=Invalid), lambda ctx: (z3.Int("Invalid", ctx=ctx[0]), None)), - (UPat(Ops.CONST, dtypes.ints+(dtypes.weakint,), name="x"), lambda x,ctx: (z3.IntVal(x.val, ctx=ctx[0]), None)), + (UPat(Ops.CONST, dtypes.weakint, name="x"), lambda x,ctx: (z3.IntVal(x.val, ctx=ctx[0]), None)), (UPat(Ops.CONST, dtypes.bool, name="x"), lambda x,ctx: (z3.BoolVal(x.val, ctx=ctx[0]), None)), # casts from floats create new variables (UPat(Ops.CAST, dtypes.ints+(dtypes.weakint,), src=(UPat(dtype=dtypes.floats),), name="x"), lambda x,ctx: create_bounded(f"cast{len(ctx[1])}", x.dtype.min, x.dtype.max, ctx[0])), # A comparison between floats introduces a new bool variable (UPat(GroupOp.Comparison, src=UPat(dtype=dtypes.floats)), lambda ctx: (z3.Bool(f"float_cmp{len(ctx[1])}", ctx=ctx[0]), None)), + # a same-dtype cast states a width, which z3 does not model: identity. must precede the rules below (bool->bool) + (UPat(Ops.CAST, name="x"), lambda x,ctx: (ctx[1][x.src[0]], None) if x.dtype == x.src[0].dtype else None), # casts from bool/int to int/bool (UPat(Ops.CAST, dtypes.ints+(dtypes.weakint,),src=(UPat.var("x", dtypes.bool),)), lambda x,ctx: (z3.If(ctx[1][x], 1, 0), None)), (UPat(Ops.CAST, dtypes.ints+(dtypes.weakint,), src=(UPat.var("x", dtypes.ints+(dtypes.weakint,)),)), lambda x,ctx: (ctx[1][x], None)), diff --git a/tinygrad/uop/weak.py b/tinygrad/uop/weak.py index 23001e80e8..b4899e6f1c 100644 --- a/tinygrad/uop/weak.py +++ b/tinygrad/uop/weak.py @@ -1,82 +1,103 @@ from dataclasses import replace -from tinygrad.dtype import dtypes, DType, AddrSpace, Invalid, least_upper_dtype, strong_dtype, weak_dtype +from tinygrad.dtype import dtypes, DType, AddrSpace, least_upper_dtype, strong_dtype, weak_dtype from tinygrad.helpers import unwrap -from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher, GroupOp, graph_rewrite, dtype_from_uop +from tinygrad.uop.ops import UOp, UPat, Ops, PatternMatcher, GroupOp, dtype_from_uop, promo_dtype def default_dtype(u:UOp): if u.dtype is dtypes.weakfloat: return dtypes.default_float return dtypes.long if u.overflows(dtypes.int32) else dtypes.int +def commit_weak(s:UOp, dt:DType) -> UOp: + # a CONST re-mints, never takes a cast: at bool/weakint/weakfloat a CAST would be a second spelling of one literal + return UOp.const(s.val, dt) if s.op is Ops.CONST else s.cast(dt) + +# s may stay bare only where u derives a width from src anyway. both checks bind: a comparison's dtype is bool, a shift's is its lhs's +def bare_ok(u:UOp, s:UOp, src:tuple[UOp, ...]) -> bool: + return s.op is Ops.CONST and u.op in GroupOp.Broadcastable and promo_dtype(src) not in dtypes.weaks \ + and dtype_from_uop(u.op, src, u.arg) not in dtypes.weaks + +# absorb the weak CAST off each src (the consumer states that width instead), THEN give the still-bare weak CONSTs their own +# default. bare_ok is asked at the ABSORBED srcs -- the widths this node ends up meeting at, not the ones it came in with +def take_widths(u:UOp) -> tuple[UOp, ...]: + src = tuple(s.src[0] if s.op is Ops.CAST and s.dtype in dtypes.weaks else s for s in u.src) + return tuple(commit_weak(s, default_dtype(s)) if s.op is Ops.CONST and s.dtype in dtypes.weaks and not bare_ok(u, s, src) else s + for s in src) + +def commit_srcs_at(u:UOp, dt:DType) -> UOp|None: + # the root re-derives: a shift's dtype is its lhs's, so committing the lhs commits the node too + ret = u.replace(dtype=None, src=tuple(commit_weak(s, dt) if s.dtype in dtypes.weaks and not bare_ok(u, s, u.src) else s for s in u.src)) + return None if ret is u else ret + +def commit_weak_srcs(u:UOp) -> UOp|None: + if not any(s.dtype in dtypes.weaks for s in u.src) or (dt:=least_upper_dtype(*(s.dtype for s in u.src))) in dtypes.weaks: return None + return commit_srcs_at(u, dt) + +# a concrete CAST over a weak node states the width the value will live at. that width is a floor, never a narrowing +def cast_weak_srcs(c:UOp, u:UOp) -> UOp|None: + # only within the kind: an int cast of a weakfloat node is a value conversion, not a statement about the node's width + if c.dtype in dtypes.weaks or weak_dtype(c.dtype) is not u.dtype: return None + return None if (ret:=commit_srcs_at(u, least_upper_dtype(c.dtype, default_dtype(u)))) is None else ret.cast(c.dtype) + +# rides every round that can mint a weak const, and must reach fixpoint before pm_lower_weak below hands one its own default +pm_commit_weak = PatternMatcher([ + (UPat(GroupOp.Broadcastable, name="u"), commit_weak_srcs), + (UPat(Ops.STORE, src=(UPat(), UPat(dtype=dtypes.weaks)), allow_any_len=True, name="u"), + lambda u: u.replace(src=(u.src[0], commit_weak(u.src[1], u.src[0].dtype), *u.src[2:]))), + # NOTE: no CONST arm. a concrete CAST over a weak CONST is already the committed pair, minted that way by UOp.const + (UPat(Ops.CAST, name="c", src=(UPat(GroupOp.ALU, dtype=dtypes.weaks, name="u"),)), cast_weak_srcs), +]) + +# a lowered node keeps its old weak dtype as a CAST, which only take_widths absorbs +# NOTE: the round running this must not compose symbolic -- its cast collapse eats that CAST and the round cycles def lower_weak_node(u:UOp) -> UOp|None: - start, src = (1 if u.op is Ops.WHERE else 0), tuple(s.src[0] if s.op is Ops.CAST and s.dtype in dtypes.weaks else s for s in u.src) - if src == u.src or any(s.dtype in dtypes.weaks for s in src[start:]): return None + start = 1 if u.op is Ops.WHERE else 0 # WHERE's cond is bool, never part of the width unification + src = take_widths(u) + # a weak CONST take_widths kept bare is settled -- this node derives its width; any OTHER weak src is not ours to lower yet + if src == u.src or any(s.dtype in dtypes.weaks and s.op is not Ops.CONST for s in src[start:]): return None + # a Binary widens from its own bounds as well as the lowered srcs, every other op derives from the lowered srcs alone dt = strong_dtype(least_upper_dtype(default_dtype(u), *(s.dtype for s in src)) if u.op in GroupOp.Binary else unwrap(dtype_from_uop(u.op, src, u.arg))) - return u.replace(dtype=None, src=src[:start]+tuple(s if s.base.is_invalid else commit_weak(s, dt) for s in src[start:])).cast(u.dtype) + return u.replace(dtype=None, src=src[:start]+tuple(s if s.base.is_invalid or s.dtype in dtypes.weaks else commit_weak(s, dt) + for s in src[start:])).cast(u.dtype) + +# every other consumer edge. a CONST gets no rule of its own: CONST(v) -> CAST(dt, CONST(v)) contains its own input, which deadlocks +def absorb_weak_srcs(u:UOp) -> UOp|None: + if u.op is Ops.CAST and u.src[0].op is Ops.CONST: return None # a literal, not a consumer: its CONST is the value half + return None if (src:=take_widths(u)) == u.src else u.replace(dtype=None, src=src) pm_lower_weak = PatternMatcher([ - (UPat(Ops.CONST, dtype=dtypes.weaks, name="u"), lambda u: UOp.const(u.val, default_dtype(u)).cast(u.dtype)), # two stacked weak casts are two kind conversions: each resolves at its own kind's default - # a SINGLE weak cast is never rewritten here, each consumer absorbs it on its own edge (see lower_weak_srcs) (UPat(Ops.CAST, dtype=dtypes.weaks, src=(UPat(Ops.CAST, dtype=dtypes.weaks, src=(UPat.var("x"),)),), name="u"), lambda u,x: x.cast(default_dtype(u.src[0])).cast(default_dtype(u)).cast(u.dtype) if x.dtype not in dtypes.weaks else None), - # Binary can widen from the bounds, all other nodes derive from the lowered sources. # a weakfloat Unary (sin/exp2/...) must resolve here, before the transcendental decomposition (UPat(GroupOp.Binary|GroupOp.Unary|{Ops.WHERE, Ops.RANGE, Ops.STACK, Ops.SPECIAL}, name="u"), lower_weak_node), (UPat((Ops.PARAM, Ops.BUFFER), dtype=dtypes.weakint, name="u"), lambda u: u.replace(dtype=None, arg=replace(u.arg, dtype=default_dtype(u))).cast(dtypes.weakint) if u.addrspace == AddrSpace.ALU else None), + (UPat(GroupOp.All, name="u"), absorb_weak_srcs), ]) -def lower_weak_srcs(ctx:dict[UOp, UOp]|None, u:UOp) -> UOp|None: - if ctx is None: ctx = {} - def lower(s:UOp) -> UOp: - if (r:=ctx.get(s)) is None: - r = graph_rewrite(s, pm_lower_weak) - # the consumer absorbs the cast on its own edge - ctx[s] = r = r.src[0] if r.op is Ops.CAST and r.dtype in dtypes.weaks else r - return r - # a comparison demands a common operand width: lower it whole so the Binary rule unifies its operands - ret = lower(u) if u.op in GroupOp.Comparison else u.replace(src=tuple(lower(s) if s.dtype in dtypes.weaks else s for s in u.src)) - return None if ret is u else ret +# drop a CAST over a weak const where the consumer restores the width anyway, so bare-CONST rules keep matching. two +# statements must survive the drop: the width the operands MEET at, and the node's own DERIVED dtype +def uncast_const(u:UOp) -> UOp|None: + src = tuple(s.src[0] if s.op is Ops.CAST and s.dtype not in dtypes.weaks and s.src[0].op is Ops.CONST + and s.src[0].dtype in dtypes.weaks else s for s in u.src) + if src == u.src or promo_dtype(src) != promo_dtype(u.src) or dtype_from_uop(u.op, src, u.arg) is not u.dtype: return None + return u.replace(src=src) -def commit_weak(s:UOp, dt:DType) -> UOp: - # a CONST commits directly at dt (the value stays mathematical, emission truncates), a non-const src takes the cast - return UOp.const(s.val, dt) if s.op is Ops.CONST else s.cast(dt) +# the inverse of pm_cast_const: drop a width statement the consumer re-derives, so bare-keyed rules keep matching. +# composes into symbolic and never symbolic_simple: the widths it drops are first stated by the round above it +pm_uncast_const = PatternMatcher([(UPat(GroupOp.Broadcastable, name="u"), uncast_const)]) -def commit_weak_srcs(u:UOp) -> UOp|None: - if not any(s.dtype in dtypes.weaks for s in u.src): return None - if (dt:=least_upper_dtype(*(s.dtype for s in u.src))) in dtypes.weaks: return None - # the root re-derives: a shift's dtype is its lhs's, so committing the lhs commits the node too - return u.replace(dtype=None, src=tuple(commit_weak(s, dt) if s.dtype in dtypes.weaks else s for s in u.src)) +def cast_const(u:UOp, s:UOp) -> UOp: + if s.op is not Ops.CONST or s.is_invalid: return s + # a bool const already has a strong dtype, so its CAST is built raw: .cast(bool) would fold at construction + if s.dtype not in dtypes.weaks: return UOp(Ops.CAST, src=(s,), arg=s.dtype) + # the width its consumer derives; where nothing does, commit_weak is the identity and spec_program rejects it LOUDLY + return commit_weak(s, promo_dtype(u.src)) if u.op in GroupOp.Broadcastable else s -# runs in index lowering and in the decomps: a rule that mints a weak const commits it in the same rewrite, so none reaches the renderer -pm_commit_weak = PatternMatcher([ - (UPat(GroupOp.Broadcastable, name="u"), commit_weak_srcs), - # demand from the destination: a STORE's weak value commits at the destination's dtype - (UPat(Ops.STORE, src=(UPat(), UPat(dtype=dtypes.weaks)), allow_any_len=True, name="u"), - lambda u: u.replace(src=(u.src[0], commit_weak(u.src[1], u.src[0].dtype), *u.src[2:]))), -]) +# THE DOOR: wrap every remaining bare const in its CAST. keyed on the CONSUMER -- "bare" is a property of the edge +def cast_consts(u:UOp) -> UOp|None: + if u.op is Ops.CAST and u.src[0].op is Ops.CONST: return None + return None if (src:=tuple(cast_const(u, s) for s in u.src)) == u.src else u.replace(src=src) -# a concrete CAST over a weak node states the width the value will live at. that width is a floor, never a narrowing -def cast_weak_srcs(c:UOp, u:UOp) -> UOp|None: - if c.dtype in dtypes.weaks or weak_dtype(c.dtype) is not u.dtype: return None - dt = least_upper_dtype(c.dtype, default_dtype(u)) - return u.replace(dtype=None, src=tuple(commit_weak(s, dt) if s.dtype in dtypes.weaks else s for s in u.src)).cast(c.dtype) - -pm_cast_weak = PatternMatcher([ - (UPat(Ops.CAST, name="c", src=(UPat(GroupOp.ALU, dtype=dtypes.weaks, name="u"),)), cast_weak_srcs), - (UPat(Ops.CAST, name="c", src=(UPat(Ops.CONST, dtype=dtypes.weaks, name="u"),)), lambda c,u: commit_weak(u, c.dtype)), -]) - -pm_lower_index_dtype = pm_commit_weak+pm_cast_weak+PatternMatcher([ - # a CAST between two concrete dtypes over a CONST is a value conversion: evaluate it once, at the width the CAST states - # TODO: delete this once CONST has no dtype - (UPat(Ops.CAST, dtypes.all, name="root", src=(UPat.cvar("c", dtypes.all),)), lambda root, c: root.const_like(c.val)), - (UPat(GroupOp.All, name="u"), - lambda ctx,u: lower_weak_srcs(ctx, u) if u.dtype not in dtypes.weaks and any(s.dtype in dtypes.weaks for s in u.src) else None), - # a valid index into an n-element buffer lives in [0,n): a gated long index narrows when n-1 fits int32 (out-of-gate wraps, discarded) - # TODO: more generic - (UPat((Ops.INDEX, Ops.SHRINK), src=(UPat.var("buf"), UPat.var("gate").where(UPat.var("idx", dtypes.long), UPat(Ops.CONST, arg=Invalid))), - allow_any_len=True, name="u"), - lambda u,buf,gate,idx: u.replace(src=(buf, idx.cast(dtypes.int).valid(gate))+u.src[2:]) if buf.max_numel()-1 <= dtypes.int32.max else None), -]) +pm_cast_const = PatternMatcher([(UPat(GroupOp.All, name="u"), cast_consts)])