ayo
a627439025
Shim - Adaptive output-event wait: sleep to ~120 us before the predicted completion (smoothed per wait site), spin near it, back off when late. D4R_SHIM_OUTPUT_POLL_US restores a fixed interval. - Defer the input-marker wait into the worker thread with VRAM interop (D4R_SHIM_DEFER_MARKER, default on), taking the host setup off the critical path. MicroCUDA - Experimental ZLUDA-free backend in wine_nvcuda_bridge, plus tools/microcuda and scripts/build_microcuda.sh. Serves accuracy, fast and native kernels; unregistered modules fall back to the reference path. Unclassified writes invalidate prepared weights. M kernels - SWIN_PREP_IDENTITY_AT: key weight prep by allocation identity (MicroCUDA); enc3 tube uses it with 8 slots. ZLUDA ignores it. - dec2 (gfx1101): SWIN_PACKED_RESIDUAL, byte-identical, ~1.6% faster. - Opt-in gfx12 hardware fp8 conversion (SWIN_FP8_HWCVT), unverified on hardware. Native K translators - ctile: handle variants that re-materialize the parameter base (mvlo). - rewrite_output: check live-in registers, not just outputs. - build.sh: native input covers all 8 flag variants, output the 4 HDR "max" variants; depth-tile skip stays mvhi-only. Tools - dump_runner and bench_dlss_kernel harness additions.
73 lines
3.5 KiB
Python
Executable File
73 lines
3.5 KiB
Python
Executable File
#!/usr/bin/env python3
|
|
"""Build the staged output rewrite from the exact translated kernel.
|
|
|
|
Replaces path B's cooperative tile fill with output_tile.h. All other phases and
|
|
the DLAA path remain intact. Refuses to generate if the region's register contract
|
|
changes, so new translator output cannot silently reuse an incompatible mapping.
|
|
"""
|
|
import argparse
|
|
import re
|
|
from pathlib import Path
|
|
|
|
REGISTER = re.compile(r"\b(?:rd|rs|r|f|fd|p)\d+\b")
|
|
ASSIGN = re.compile(r"^\s*((?:rd|rs|r|f|fd|p)\d+)\s*=")
|
|
LABEL = re.compile(r"^(L__BB0_\d+):;$")
|
|
|
|
|
|
def rewrite(source):
|
|
head, main = source.split('extern "C" __attribute__((global))', 1)
|
|
lines = main.splitlines(keepends=True)
|
|
labels = [(i, LABEL.match(line.strip()).group(1)) for i, line in enumerate(lines) if LABEL.match(line.strip())]
|
|
blocks = {name: (start, labels[k + 1][0] if k + 1 < len(labels) else len(lines))
|
|
for k, (start, name) in enumerate(labels)}
|
|
pending, region = ['L__BB0_124'], set()
|
|
while pending:
|
|
name = pending.pop()
|
|
if name in region or name == 'L__BB0_138':
|
|
continue
|
|
region.add(name)
|
|
start, end = blocks[name]
|
|
body = ''.join(lines[start + 1:end])
|
|
pending += re.findall(r'goto (L__BB0_\d+);', body)
|
|
if not re.search(r'^\s*goto L__BB0_\d+;\s*$', lines[end - 1]):
|
|
# This region's fallthroughs are only its internal branch joins.
|
|
pending += [n for i, n in labels if i == end]
|
|
expected = {'L__BB0_124'} | {f'L__BB0_{k}' for k in range(127, 138)}
|
|
if region != expected:
|
|
raise ValueError(f'tile CFG changed: {sorted(region)}')
|
|
removed = {i for name in region for i in range(*blocks[name])}
|
|
definitions = {m.group(1) for i in removed if (m := ASSIGN.match(lines[i]))}
|
|
# Definitions declared at function scope count as uses. Ignore declaration-only lines.
|
|
uses_outside = set(REGISTER.findall(''.join(line for i, line in enumerate(lines)
|
|
if i not in removed and not re.match(r'^\s*(?:uint\d+_t|float|double|bool)\s+(?:rd|rs|r|f|fd|p)\d+(?:,|;)', line))))
|
|
live_out = definitions & uses_outside
|
|
if live_out != {'r215', 'r217'}:
|
|
raise ValueError(f'tile register contract changed: {sorted(live_out)}')
|
|
# The replacement passes r3, r4, r7, r8 and f1 by name: the region's inputs must be exactly the ones it was
|
|
# written for (other flag variants of the kernel reuse this region; the checks above alone only cover outputs).
|
|
region_text = ''.join(lines[i] for i in sorted(removed))
|
|
live_in = set(REGISTER.findall(region_text)) - definitions
|
|
if live_in != {'r3', 'r4', 'r5', 'r6', 'r7', 'r8', 'rd1', 'f1', 'f4'}:
|
|
raise ValueError(f'tile input contract changed: {sorted(live_in)}')
|
|
replacement = '''L__BB0_124:;
|
|
output_tile_fill(KARGP, (int32_t)r3, (int32_t)r4, (int32_t)r7, (int32_t)(r8 + 1u),
|
|
f1, (AS3 uint8_t*)sharedColor, (AS3 uint8_t*)sharedDepth);
|
|
r215 = (uint32_t)__builtin_amdgcn_workitem_id_y();
|
|
r217 = (uint32_t)__builtin_amdgcn_workitem_id_x();
|
|
goto L__BB0_138;
|
|
'''
|
|
result = []
|
|
for i, line in enumerate(lines):
|
|
if i == blocks['L__BB0_124'][0]:
|
|
result.append(replacement)
|
|
if i not in removed:
|
|
result.append(line)
|
|
return head + '#include "output_tile.h"\n' + 'extern "C" __attribute__((global))' + ''.join(result)
|
|
|
|
|
|
if __name__ == '__main__':
|
|
parser = argparse.ArgumentParser(description=__doc__)
|
|
parser.add_argument('source', type=Path)
|
|
args = parser.parse_args()
|
|
print(rewrite(args.source.read_text()), end='')
|