d4r/kernels/native/ctile.py

"""Color-tap LDS tile pass for the hiluma output kernels (used by ptx2hip.py --ctile).

Path A of hiluma_engine_output_* (taken when a block's render footprint is wider than path B's 24-texel LDS tile,
e.g. at Quality) fetches 16 color texels per thread and converts each one (YCoCg, exposure, tonemap weight) right
after the fetch. The conversion of a texel depends only on the texel and the exposure %f1, and neighbouring
threads convert the same texels (~9x at ratio 1.5). This pass:

1. finds every tex.base on the color texture inside path A,
2. slices the instructions that compute the tap's packed result (cvt.rn.f16x2 x2 or cvt.rn.f16.f32) from the
   fetched texel (forward taint + backward need, including the tonemap branch diamond), and checks that no other
   value escapes the slice,
3. replaces the slice with `d4r.ctap.G` (tile lookup, falling back to fetch + the same conversion), and
4. emits one conversion function per slice shape (generated from the slice's own instructions, so the tile holds
   bit-identical values) and a cooperative tile fill at the start of path A.

The tile is keyed by the final (clamped) texel coordinates, so any tap outside it just takes the fallback.
"""
import re

REG = re.compile(r'%[A-Za-z_$][\w$]*')
CAP = 1024          # tile texels (TW*TH); a larger footprint falls back to direct fetches


class CtileErr(Exception):
    pass


def parse(body, tokenize):
    """-> list of dicts: kind ('stmt'|'label'|'open'|'close'), text, plus op/dst/src/pred for stmts"""
    items = []
    for kind, text in tokenize(body):
        it = {'kind': kind, 'text': text}
        if kind == 'stmt':
            s = text
            m = re.match(r'@(!?)(%p\d+)\s+(.*)$', s)
            it['pred'] = m.group(2) if m else None
            if m:
                s = m.group(3)
            parts = s.split(None, 1)
            it['op'] = parts[0]
            rest = parts[1] if len(parts) > 1 else ''
            it['rest'] = rest
            dst, src = [], []
            op = it['op']
            if op.startswith('.') or op.startswith('bra') or op.startswith('bar') or op.startswith('ret'):
                src = REG.findall(rest)
            elif op.startswith('st.') or op.startswith('sust'):
                src = REG.findall(rest)
            else:
                # first operand = destination(s)
                depth, i = 0, 0
                while i < len(rest):
                    c = rest[i]
                    if c in '{[':
                        depth += 1
                    elif c in '}]':
                        depth -= 1
                    elif c == ',' and depth == 0:
                        break
                    i += 1
                dst = REG.findall(rest[:i])
                src = REG.findall(rest[i:])
            if it['pred']:
                src = src + [it['pred']]
            it['dst'], it['src'] = dst, src
        items.append(it)
    return items


def branch_target(it):
    return it['rest'].strip().rstrip(';').strip().lstrip('$') if it['kind'] == 'stmt' and it['op'].startswith('bra') else None


def find_taps(items, color_aliases):
    return [i for i, it in enumerate(items) if it['kind'] == 'stmt' and it['op'] == 'tex.base.2d.v4.f32.s32'
            and re.search(r'\[\s*(%rd\d+)', it['rest']).group(1) in color_aliases]


def slice_tap(items, t):
    """-> (removed index set, results list, kind 'in'|'out', externals set)"""
    tex = items[t]
    taint = set(tex['dst'])
    tainted_idx = [t]
    results, kind = [], None
    i = t + 1
    while i < len(items):
        it = items[i]
        if it['kind'] == 'stmt' and any(s in taint for s in it['src']):
            op = it['op']
            if op.startswith('cvt.rn.f16x2.f32') and kind in (None, 'in'):
                kind = 'in'
                results.append((i, it['dst'][0]))
                tainted_idx.append(i)
                if len(results) == 2:
                    break
            elif op.startswith('cvt.rn.f16.f32') and kind is None:
                kind = 'out'
                results.append((i, it['dst'][0]))
                tainted_idx.append(i)
                break
            else:
                taint.update(it['dst'])
                tainted_idx.append(i)
        i += 1
    else:
        raise CtileErr(f'tap at {t}: no packed result found')
    last = results[-1][0]
    # backward need from the results (within [t, last])
    need_regs = set()
    for ri, _ in results:
        need_regs.update(items[ri]['src'])
    needed = {ri for ri, _ in results}
    # conditional branches on tainted predicates (the tonemap branch) are part of the slice
    for j in tainted_idx:
        it = items[j]
        if it['op'].startswith('bra') and it['pred'] is not None:
            needed.add(j)
            need_regs.add(it['pred'])
    for j in range(last, t - 1, -1):
        it = items[j]
        if j in needed:
            continue
        if it['kind'] == 'stmt' and j in tainted_idx and any(d in need_regs for d in it['dst']):
            needed.add(j)
            need_regs.update(it['src'])
    extra = [j for j in tainted_idx if j not in needed]
    if extra:
        raise CtileErr(f'tap at {t}: tainted but not needed: {[items[j]["text"] for j in extra]}')
    removed = set(needed)
    # branch diamonds: region from each conditional branch to its merge label
    for j in sorted(needed):
        it = items[j]
        if not it['op'].startswith('bra'):
            continue
        targets, seen = {branch_target(it)}, set()
        k = j + 1
        while True:
            x = items[k]
            if x['kind'] == 'label':
                lab = x['text'].lstrip('$')
                if lab in targets and targets - seen == {lab} and lab not in {branch_target(it)}:
                    removed.add(k)
                    break
                seen.add(lab)
                removed.add(k)
            elif x['kind'] == 'stmt':
                if x['op'].startswith('bra'):
                    if x['pred'] is not None and k not in needed:
                        raise CtileErr(f'tap at {t}: foreign conditional branch in diamond')
                    targets.add(branch_target(x))
                    removed.add(k)
                elif k not in needed:
                    # constant moves used only inside the slice are part of it
                    if x['op'].startswith('mov') and not any(REG.match(s) for s in x['src']):
                        removed.add(k)
                    else:
                        raise CtileErr(f'tap at {t}: foreign instruction in diamond: {x["text"]}')
            k += 1
            if k > last:
                raise CtileErr(f'tap at {t}: diamond does not close before the result')
    # escape check: registers defined in the slice (except results) used outside it
    defs = set()
    for j in removed:
        if items[j]['kind'] == 'stmt':
            defs.update(items[j]['dst'])
    res_regs = {r for _, r in results}
    inner_defs = defs - res_regs
    for j, it in enumerate(items):
        if j in removed or it['kind'] != 'stmt':
            continue
        bad = inner_defs.intersection(it['src'])
        if bad:
            raise CtileErr(f'tap at {t}: {bad} escapes to {it["text"]}')
        if defs.intersection(it['dst']) - res_regs or res_regs.intersection(it['dst']):
            raise CtileErr(f'tap at {t}: slice register redefined by {it["text"]}')
    # labels removed must not be targeted from outside
    labels = {items[j]['text'].lstrip('$') for j in removed if items[j]['kind'] == 'label'}
    for j, it in enumerate(items):
        if j not in removed and branch_target(it) in labels:
            raise CtileErr(f'tap at {t}: label targeted from outside the slice')
    used = set()
    for j in removed:
        if items[j]['kind'] == 'stmt' and j != t:
            used.update(items[j]['src'])
    externals = used - defs - set(tex['dst'])
    externals.discard(None)
    return removed, results, kind, externals


def canon(items, removed, t, results):
    """canonical text of the slice (registers by first appearance, texel by position)"""
    names, labels = {}, {}
    for k, r in enumerate(items[t]['dst']):
        names[r] = f'T{k}'
    for k, (_, r) in enumerate(results):
        names[r] = f'R{k}'
    out = []
    for j in sorted(removed):
        if j == t:
            continue
        it = items[j]
        if it['kind'] == 'label':
            lab = it['text'].lstrip('$')
            labels.setdefault(lab, f'L{len(labels)}')
            out.append(labels[lab] + ':')
            continue
        if it['kind'] != 'stmt':
            continue
        s = it['text']

        def rn(m):
            r = m.group(0)
            if r not in names:
                names[r] = f'V{len(names)}'
            return names[r]
        s = REG.sub(rn, s)
        s = re.sub(r'\$(L__\w+)', lambda m: labels.setdefault(m.group(1), f'L{len(labels)}'), s)
        out.append(s)
    return '\n'.join(out), names


def apply(tr, tokenize):
    """Rewrite tr.body in place for the color tile. Returns the HIP prelude (conversion functions) and info."""
    body = tr.body
    items = parse(body, tokenize)
    # color texture handle: ld.param.u64 [%rd1+-56] (param+200, %rd10 / %rd69); taps use it or its mov aliases
    color = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'ld.param.u64'
             and re.search(r'\[%rd1\+-56\]', it['rest'])}
    aliases = set(color)
    for it in items:
        if it['kind'] == 'stmt' and it['op'] == 'mov.u64' and len(it['src']) == 1 and it['src'][0] in color:
            aliases.add(it['dst'][0])
    # path A's color taps lie between $L__BB0_1 and $L__BB0_77 (the per-tap blocks); path B's are elsewhere
    labels = {it['text'].lstrip('$'): i for i, it in enumerate(items) if it['kind'] == 'label'}
    a0, a1 = labels['L__BB0_1'], labels['L__BB0_77']
    taps = [t for t in find_taps(items, aliases) if a0 < t < a1]
    if len(taps) != 16:
        raise CtileErr(f'expected 16 path A color taps, found {len(taps)}')
    groups, per_tap = {}, []
    all_removed = set()
    for t in taps:
        removed, results, kind, ext = slice_tap(items, t)
        if removed & all_removed:
            raise CtileErr('overlapping slices')
        all_removed |= removed
        text, names = canon(items, removed, t, results)
        ext_names = sorted(ext, key=lambda r: names.get(r, r))
        key = (kind, text, tuple(names.get(r) for r in ext_names))
        g = groups.setdefault(key, len(groups))
        per_tap.append((t, removed, results, kind, ext_names, g))
    if len(groups) > 3:
        raise CtileErr(f'{len(groups)} slice shapes')
    # externals: %f1 (exposure) or constant registers (single mov-immediate definition)
    const_def = {}
    for it in items:
        if it['kind'] == 'stmt' and it['op'].startswith('mov') and len(it['dst']) == 1 and not it['src']:
            const_def.setdefault(it['dst'][0], []).append(it)
    group_info = {}
    for (t, removed, results, kind, ext_names, g) in per_tap:
        if g in group_info:
            continue
        params, consts = [], []
        for r in ext_names:
            if r == '%f1':
                params.append(r)
            elif r in const_def and len(const_def[r]) == 1:
                consts.append(const_def[r][0]['text'])
            else:
                raise CtileErr(f'external {r} of tap {t}')
        group_info[g] = dict(t=t, removed=removed, results=results, kind=kind, params=params, consts=consts)
    return items, per_tap, group_info, color


def emit(tr, tokenize, Translator):
    """Apply the pass to translator tr (before translate()). Installs the pseudo-op handler; returns prelude text."""
    items, per_tap, group_info, color = apply(tr, tokenize)
    # ---- conversion functions, one per group, translated from the group's first slice
    prelude = []
    for g, info in sorted(group_info.items()):
        t, removed, results = info['t'], info['removed'], info['results']
        sub = Translator.__new__(Translator)
        sub.kernel = tr.kernel
        sub.param_name = tr.param_name
        sub.regs, sub.shared, sub.locals, sub.lines = {}, [], [], []
        sub.scope_names, sub.scope_id = [{}], 0
        tex = items[t]
        # texel inputs
        for k, r in enumerate(tex['dst']):
            sub.reg(r)
            sub.lines.append(f'    {sub.reg(r)} = t_[{k}];')
        for c in info['consts']:
            sub.stmt(c)
        for j in sorted(removed):
            if j == t:
                continue
            it = items[j]
            if it['kind'] == 'label':
                sub.lines.append(f'{it["text"].replace("$", "")}:;')
            elif it['kind'] == 'stmt':
                sub.stmt(it['text'])
        outs = [sub.reg(r) for _, r in results]
        decl = []
        for name, ty in sorted(sub.regs.items()):
            if name in [sub.reg(p) for p in info['params']]:
                continue
            decl.append(f'    {ty} {name};')
        args = ', '.join(['f4 t_'] + [f'float {sub.reg(p)}' for p in info['params']] + [f'uint32_t* o{k}' for k in range(len(outs))])
        prelude.append(f'RT void d4r_ctap_conv{g}({args})\n{{')
        prelude.extend(decl)
        prelude.extend(sub.lines)
        for k, o in enumerate(outs):
            prelude.append(f'    *o{k} = (uint32_t){o};')
        prelude.append('}')
    # ---- tile storage (aliases the kernel's LDS through d4r_lds)
    offs, off = {}, 0
    for g, info in sorted(group_info.items()):
        for k in range(len(info['results'])):
            offs[(g, k)] = off
            off += CAP * 4
    # ---- rewrite the body: removed statements dropped, pseudo-op at each tap, fill at the start of path A
    tap_at = {t: (results, ext, g) for (t, removed, results, kind, ext, g) in per_tap}
    removed_all = set()
    for (t, removed, *_r) in per_tap:
        removed_all |= removed
    out = []
    for i, it in enumerate(items):
        if it['kind'] == 'open':
            out.append('{'); continue
        if it['kind'] == 'close':
            out.append('}'); continue
        if it['kind'] == 'label':
            if i in removed_all:
                continue
            out.append(it['text'] + ':')
            if it['text'].lstrip('$') == 'L__BB0_1':
                out.append('d4r.ctile_fill;')
            continue
        if i in tap_at:
            results, ext, g = tap_at[i]
            m = re.search(r'\[\s*(%rd\d+)\s*,\s*\{([^}]*)\}\s*\]', it['rest'])
            regs = ', '.join(r for _, r in results)
            out.append(f'd4r.ctap.{g} {regs}, {m.group(1)}, {m.group(2)};')
            continue
        if i in removed_all:
            continue
        out.append(it['text'] + ';')
    tr.body = '\n'.join(out)
    summary = dict(groups=group_info, offs=offs, lds_bytes=off, taps=len(per_tap))

    orig_op = tr.op

    def op(opname, a):
        if opname == 'd4r.ctile_fill':
            tr.emit('d4r_ctile_fill(KARGP, (int32_t)r3, (int32_t)r4, (int32_t)r5, (int32_t)r6, f1, &ct_x0, &ct_y0, &ct_w, &ct_h);')
            return
        if opname.startswith('d4r.ctap.'):
            g = int(opname.split('.')[2])
            nres = len(group_info[g]['results'])
            res = [tr.reg(x) for x in a[:nres]]
            h = tr.val(a[nres], 'u64')
            x, y = tr.val(a[nres + 1], 'u32'), tr.val(a[nres + 2], 'u32')
            prm = ''.join(f', {tr.val(p, "f32")}' for p in group_info[g]['params'])
            tr.emit(f'{{ int32_t dx_ = (int32_t){x} - ct_x0, dy_ = (int32_t){y} - ct_y0;')
            tr.emit(f'  if ((uint32_t)dx_ < (uint32_t)ct_w && (uint32_t)dy_ < (uint32_t)ct_h) {{ uint32_t i_ = (uint32_t)(dy_ * ct_w + dx_);')
            for k, r in enumerate(res):
                tr.emit(f'    {r} = (uint32_t)((const AS3 uint32_t*)(d4r_lds + {offs[(g, k)]}))[i_];')
            outs = ', '.join(f'o_ + {k}' for k in range(nres))
            tr.emit(f'  }} else {{ uint32_t o_[2]; d4r_ctap_conv{g}(tex_fetch({h}, (int32_t){x}, (int32_t){y}){prm}, {outs});')
            for k, r in enumerate(res):
                tr.emit(f'    {r} = o_[{k}];')
            tr.emit('  } }')
            return
        orig_op(opname, a)
    tr.op = op
    # result register types: f16x2 packs are u32 regs, f16 packs are u16 regs: both assigned via (uint32_t) casts
    # fill function (needs the conversion functions and offsets)
    fill = ['RT void d4r_ctile_fill(const AS4 uint8_t* KARGP, int32_t x0l, int32_t y0l, int32_t x1l, int32_t y1l, float f1,',
            '                       int32_t* ox0, int32_t* oy0, int32_t* ow, int32_t* oh)',
            '{',
            '    // final texel coordinate of a tile-relative logical coordinate u: u > 0 ? min(lo + u, hi) : lo',
            '    // (the clamp every color tap applies; params +208/+212 = lo, +216/+220 = hi)',
            '    int32_t lox = *(const AS4 int32_t*)(KARGP + 208), loy = *(const AS4 int32_t*)(KARGP + 212);',
            '    int32_t hix = *(const AS4 int32_t*)(KARGP + 216), hiy = *(const AS4 int32_t*)(KARGP + 220);',
            '    int32_t fx0 = x0l > 0 ? __builtin_elementwise_min(lox + x0l, hix) : lox;',
            '    int32_t fy0 = y0l > 0 ? __builtin_elementwise_min(loy + y0l, hiy) : loy;',
            '    int32_t fx1 = x1l > 0 ? __builtin_elementwise_min(lox + x1l, hix) : lox;',
            '    int32_t fy1 = y1l > 0 ? __builtin_elementwise_min(loy + y1l, hiy) : loy;',
            '    int32_t w = fx1 - fx0 + 1, h = fy1 - fy0 + 1;',
            f'    if (w <= 0 || h <= 0 || w * h > {CAP}) {{ w = 0; h = 0; }}',
            '    uint64_t tex = *(const AS4 uint64_t*)(KARGP + 200);',
            '    int32_t n = w * h;',
            '    for (int32_t i = (int32_t)(__builtin_amdgcn_workitem_id_y() * 16u + __builtin_amdgcn_workitem_id_x()); i < n; i += 256)',
            '    {',
            '        int32_t ty = i / w, tx = i - ty * w;',
            '        f4 t_ = tex_fetch(tex, fx0 + tx, fy0 + ty);']
    for g, info in sorted(group_info.items()):
        nres = len(info['results'])
        prm = ''.join(', f1' for _ in info['params'])
        outs = ', '.join(f'o_ + {k}' for k in range(nres))
        fill.append(f'        {{ uint32_t o_[2]; d4r_ctap_conv{g}(t_{prm}, {outs});')
        for k in range(nres):
            fill.append(f'          ((AS3 uint32_t*)(d4r_lds + {offs[(g, k)]}))[i] = o_[{k}];')
        fill.append('        }')
    fill += ['    }',
             '    bar_sync();',
             '    *ox0 = fx0; *oy0 = fy0; *ow = w; *oh = h;',
             '}']
    for p in sum((info['params'] for info in group_info.values()), []):
        if p != '%f1':
            raise CtileErr('only %f1 is supported as a conversion parameter')
    return prelude, fill, summary