From a6274390250951b7c7db4213640b143407b333f0 Mon Sep 17 00:00:00 2001 From: ayo <93222141+countervolts@users.noreply.github.com> Date: Wed, 7 Oct 2026 15:54:38 -0400 Subject: [PATCH] Add MicroCUDA backend, adaptive output wait, wider native K coverage (improve perf) 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. --- kernels/build.sh | 52 +- kernels/m/rrlite_dec2_4x4.hip | 6 +- kernels/m/rrlite_enc3_tube_4x4.hip | 8 +- kernels/m/swin_block.h | 9 +- kernels/m/swin_common.h | 14 + kernels/native/ctile.py | 9 +- kernels/native/in_lds2reg.py | 192 +++++++ kernels/native/rewrite_output.py | 6 + kernels/tools/dump_runner.cpp | 92 +++- scripts/build_microcuda.sh | 15 + tools/bench_dlss_kernel.cpp | 66 ++- tools/d4r_nvngx_shim.cpp | 101 +++- tools/microcuda/arch_spoof.c | 35 ++ tools/microcuda/check_runtime.py | 202 +++++++ tools/microcuda/make_pack.py | 134 +++++ tools/microcuda/make_registry.py | 99 ++++ tools/microcuda/merge_pack.py | 74 +++ tools/microcuda/runtime.cpp | 786 ++++++++++++++++++++++++++++ tools/microcuda/zluda_precompile.py | 95 ++++ tools/wine_nvcuda_bridge.c | 235 ++++++++- 20 files changed, 2186 insertions(+), 44 deletions(-) create mode 100644 kernels/native/in_lds2reg.py create mode 100755 scripts/build_microcuda.sh create mode 100644 tools/microcuda/arch_spoof.c create mode 100644 tools/microcuda/check_runtime.py create mode 100644 tools/microcuda/make_pack.py create mode 100644 tools/microcuda/make_registry.py create mode 100644 tools/microcuda/merge_pack.py create mode 100644 tools/microcuda/runtime.cpp create mode 100644 tools/microcuda/zluda_precompile.py diff --git a/kernels/build.sh b/kernels/build.sh index 2fae61a..0229fbb 100755 --- a/kernels/build.sh +++ b/kernels/build.sh @@ -125,13 +125,23 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then for mode in static dynamic; do for range in hdr ldr; do specs+=("rrlite_downsample_kernel_${mode}_${range}:sust_only${downsample_mode}"); done done - # The fast set serves two K kernels from native code translated from the PTX (kernels/native): these builds - # replace the sust_only build of the same name. The accuracy set keeps ZLUDA's compile (denormals preserved). - native_out=hiluma_engine_output_depthinv_mvhi_hdr_max_v2_rel - native_in=hiluma_engine_input_depthinv_mvhi_hdr_v2_rel + # The fast set serves K's hiluma input/output kernels from native code translated from the PTX + # (kernels/native): these builds replace the sust_only build of the same name. Every flag variant whose + # code the translators prove compatible is covered (they refuse anything else): input = all 8 + # depth x motion x range variants; output = the 4 HDR "max" variants (the LDR max kernels are a different + # kernel structure and keep the sust_only build). D4R_K_NO_DEPTH_TILE (skip the depth tile fill) is proven + # for mvhi only: mvlo kernels read the depth tile, so they keep the fill (measured: byte-exact only then). + # The accuracy set keeps ZLUDA's compile (denormals preserved). + native_outs=() native_ins=() + for depth in depthinv depthreg; do + for mv in mvhi mvlo; do + native_outs+=("hiluma_engine_output_${depth}_${mv}_hdr_max_v2_rel") + for range in hdr ldr; do native_ins+=("hiluma_engine_input_${depth}_${mv}_${range}_v2_rel"); done + done + done for spec in "${specs[@]}"; do IFS=: read -r kernel src mode <<< "$spec" - [[ "$ACCURACY" == 0 && "$kernel" == "$native_out" ]] && continue + [[ "$ACCURACY" == 0 && " ${native_outs[*]} " == *" $kernel "* ]] && continue # accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the # denormal handling and gave the same image on an RX 9070 XT (post 0.79 -> 0.67 ms, downsample 0.43 -> 0.39 ms) w64=0 @@ -145,24 +155,28 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then done if [[ "$WHAT" != l && "$ACCURACY" == 0 ]]; then # NVIDIA's PTX translated to HIP (native/ptx2hip.py) and rebuilt natively; the generated sources stay in a - # temporary directory. Both kernels run through the same texture/surface helpers as the other texture - # kernels and are byte-identical to ZLUDA's compile of the PTX. + # temporary directory. They run through the same texture/surface helpers as the other texture kernels + # and are byte-identical to ZLUDA's compile of the PTX (validated per variant in the harness). ptx_of() { grep -l -E "\.entry[[:space:]]+$1[[:space:]]*\(" "$PTX_DIR"/*.ptx | head -1; } NAT="$HERE/native" NT="$(mktemp -d)" trap 'rm -rf "$NT"' EXIT - ptx="$(ptx_of "$native_out")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_out" >&2; exit 2; } - # hiluma output (path B for Quality..UltraPerf), LDS tile as a function, global-memory redirect stores - python3 "$NAT/ptx2hip.py" --ctile --as1 "$ptx" "$native_out" > "$NT/output_translated.hip" - python3 "$NAT/rewrite_output.py" "$NT/output_translated.hip" > "$NT/$native_out.hip" - D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" "$OUT/$native_out.hsaco" \ - -DD4R_NATIVE_F2I -DD4R_K_NO_DEPTH_TILE -DD4R_AS1_STORES - # hiluma input: scratch arrays in registers, wave32 - ptx="$(ptx_of "$native_in")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_in" >&2; exit 2; } - python3 "$NAT/ptx2hip.py" --ifconv --sink "$ptx" "$native_in" > "$NT/in_ifs.hip" - python3 "$NAT/in_fuse.py" "$NT/in_ifs.hip" > "$NT/in_unit.h" - cp "$NAT/in_regs.hip" "$NT/in_regs.hip" - D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" WAVE=32 "$NAT/build_native.sh" "$NT/in_regs.hip" "$OUT/$native_in.hsaco" + for native_out in "${native_outs[@]}"; do + ptx="$(ptx_of "$native_out")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_out" >&2; exit 2; } + # hiluma output (path B for Quality..UltraPerf), LDS tile as a function, global-memory redirect stores + python3 "$NAT/ptx2hip.py" --ctile --as1 "$ptx" "$native_out" > "$NT/output_translated.hip" + python3 "$NAT/rewrite_output.py" "$NT/output_translated.hip" > "$NT/$native_out.hip" + depth_tile=(); [[ "$native_out" == *_mvhi_* ]] && depth_tile=(-DD4R_K_NO_DEPTH_TILE) + D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" "$OUT/$native_out.hsaco" \ + -DD4R_NATIVE_F2I "${depth_tile[@]}" -DD4R_AS1_STORES + done + for native_in in "${native_ins[@]}"; do + ptx="$(ptx_of "$native_in")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_in" >&2; exit 2; } + # hiluma input: per-thread LDS scratch and local array in registers (in_lds2reg.py), wave32 + python3 "$NAT/ptx2hip.py" --ifconv --sink "$ptx" "$native_in" > "$NT/in_ifs.hip" + python3 "$NAT/in_lds2reg.py" "$NT/in_ifs.hip" > "$NT/$native_in.hip" + D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" WAVE=32 "$NAT/build_native.sh" "$NT/$native_in.hip" "$OUT/$native_in.hsaco" + done fi fi [[ "$ACCURACY" == 1 ]] && printf '1\n' > "$OUT/d4r-accuracy.txt" diff --git a/kernels/m/rrlite_dec2_4x4.hip b/kernels/m/rrlite_dec2_4x4.hip index 22458f8..a91a92c 100644 --- a/kernels/m/rrlite_dec2_4x4.hip +++ b/kernels/m/rrlite_dec2_4x4.hip @@ -1,6 +1,8 @@ // fast numerics: intermediate activations f16 (no FP8 re-quantisation), f32 accumulation through GEMM chains -// RX 7700 XT: memory-clause scheduling at the 192 VGPR cap; byte-identical to the default build, gfx1101 only -// d4r-build-flags-gfx1101: -mllvm -misched=gcn-max-memory-clause +// RX 7700 XT: memory-clause scheduling at the 192 VGPR cap; byte-identical to the default build, gfx1101 only. +// SWIN_PACKED_RESIDUAL (residual kept as already-rounded halves until the output projection): byte-identical, +// -1.6% on the captured dec2 launch (0.768 -> 0.756 ms, 3 alternating reps of 100); enc1/enc2/dec1/tube: no gain +// d4r-build-flags-gfx1101: -mllvm -misched=gcn-max-memory-clause -DSWIN_PACKED_RESIDUAL #define SWIN_NO_Q8 #define SWIN_F32ACC #ifndef SWIN_MAX_VGPR diff --git a/kernels/m/rrlite_enc3_tube_4x4.hip b/kernels/m/rrlite_enc3_tube_4x4.hip index 6a3c222..29c5c38 100644 --- a/kernels/m/rrlite_enc3_tube_4x4.hip +++ b/kernels/m/rrlite_enc3_tube_4x4.hip @@ -8,9 +8,13 @@ #define SWIN_MLP_R(NW) 2 #endif // NGX can recycle a weights address for another tube block after feature -// recreation. Address-only caching then serves stale packed weights. Prepare -// each launch; the default single image is sufficient on the launch stream. +// recreation. Address-only caching then serves stale packed weights, so +// d4r_prep_key_at stays 0 (ZLUDA prepares every launch). MicroCUDA instead keys +// preparation by allocation identity (d4r_prep_identity_at): one image per tube +// weight set (6 per evaluation), re-prepared on any free/reallocation or write. #define SWIN_PREP_KEY_AT 0 +#define SWIN_PREP_IDENTITY_AT 1 +#define SWIN_PREP_SLOTS 8 #include "swin_block.h" #ifndef SWIN_NW #define SWIN_NW 8 diff --git a/kernels/m/swin_block.h b/kernels/m/swin_block.h index a3b2505..adde34d 100644 --- a/kernels/m/swin_block.h +++ b/kernels/m/swin_block.h @@ -1050,6 +1050,12 @@ __device__ __forceinline__ void swin_block(const CommonParams& p, const TubePara #ifndef SWIN_PREP_KEY_AT #define SWIN_PREP_KEY_AT 1 // weights pointer at offset 0; 0 forces preparation on every launch #endif +#ifndef SWIN_PREP_IDENTITY_AT +// weights pointer at offset (value - 1) for runtimes that key preparation by allocation identity (pointer + +// allocation generation + write epoch: MicroCUDA), never by address alone; 0 = no such key. Runtimes that only +// know d4r_prep_key_at (ZLUDA) ignore it. +#define SWIN_PREP_IDENTITY_AT 0 +#endif // Optional weight-image slots for several weight sets with stable address identity. The ZLUDA // hook can run prep only for a weights pointer it has not seen (d4r_prep_key_at / d4r_prep_key_slots); @@ -1090,13 +1096,14 @@ template __device__ __forceinline__ int prep_slot_claim(uint64_t* keys, using NAME##_L = SwinLayout; \ constexpr int NAME##_S = SWIN_PREP_SLOTS; \ __device__ wslot g_w16[(NAME##_S + 1) * NAME##_L::TOTAL]; \ - __device__ uint64_t g_prep_keys[NAME##_S + 1]; \ + extern "C" __device__ __attribute__((used)) uint64_t g_prep_keys[NAME##_S + 1] = {}; \ __constant__ SwinDescs g_descs = make_descs(); \ extern "C" __device__ __attribute__((used)) uint32_t d4r_block_z = NWAVES; \ extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_blocks = (NAME##_L::TOTAL + 127) / 128; \ /* enc3 disables address-only reuse: a recreated feature may recycle a weights address */ \ extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_at = SWIN_PREP_KEY_AT; \ extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_key_slots = NAME##_S; \ + extern "C" __device__ __attribute__((used)) uint32_t d4r_prep_identity_at = SWIN_PREP_IDENTITY_AT; \ extern "C" __global__ void __launch_bounds__(128) NAME##_prep(PARAMS p) \ { \ const int idx = blockIdx.x * 128 + threadIdx.x; \ diff --git a/kernels/m/swin_common.h b/kernels/m/swin_common.h index 34335d6..6106228 100644 --- a/kernels/m/swin_common.h +++ b/kernels/m/swin_common.h @@ -337,6 +337,20 @@ __device__ __forceinline__ sop lds16(const half_t* p) // e4m3 bytes of an f16 operand whose values are e4m3 values (K order kept: byte j = element j) __device__ __forceinline__ gop to_fp8(sop v) { +#if defined(SWIN_FP8_HWCVT) && defined(__GFX12__) + // Opt-in, UNVERIFIED ON HARDWARE: the values already are e4m3 values (finite, |x| <= 448), so the hardware + // conversion has nothing to round or saturate; byte j = element j as below. Open question before enabling + // it by default: whether e4m3 subnormals (k * 2^-9) survive the gfx12 denormal mode in effect. + uint32_t w[2]; +#pragma unroll + for (int j = 0; j < 2; ++j) + { + int packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * j], (float)v[4 * j + 1], 0, false); + packed = __builtin_amdgcn_cvt_pk_fp8_f32((float)v[4 * j + 2], (float)v[4 * j + 3], packed, true); + w[j] = (uint32_t)packed; + } + return (u2v){w[0], w[1]}; +#endif uint32_t c[4]; #pragma unroll for (int j = 0; j < 4; ++j) diff --git a/kernels/native/ctile.py b/kernels/native/ctile.py index e72f935..c0642a3 100755 --- a/kernels/native/ctile.py +++ b/kernels/native/ctile.py @@ -224,8 +224,15 @@ def apply(tr, tokenize): 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 + # (some variants, e.g. mvlo, re-materialize the parameter base: %rdN = param_0 + 256, the same address as %rd1) + param0 = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'mov.b64' + and re.search(r',\s*\w+_param_0\s*$', it['rest'])} + bases = {it['dst'][0] for it in items if it['kind'] == 'stmt' and it['op'] == 'add.s64' + and len(it['src']) == 1 and it['src'][0] in param0 and re.search(r',\s*256\s*$', it['rest'])} + if '%rd1' not in bases: + raise CtileErr(f'parameter base %rd1 not found ({sorted(bases)})') 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'])} + and any(re.search(r'\[' + re.escape(b) + r'\+-56\]', it['rest']) for b in bases)} 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: diff --git a/kernels/native/in_lds2reg.py b/kernels/native/in_lds2reg.py new file mode 100644 index 0000000..05008e7 --- /dev/null +++ b/kernels/native/in_lds2reg.py @@ -0,0 +1,192 @@ +#!/usr/bin/env python3 +"""in_lds2reg.py IN.hip > OUT.hip: a hiluma input kernel (any variant, translated by ptx2hip.py --ifconv --sink) +with its per-thread scratch kept in registers. Generalizes in_fuse.py + in_regs.hip, which matched one variant's +register numbers and its output address formula (mvlo variants store to a different address). + +The kernel stays the translated kernel: same thread ids, same arithmetic, same final stores. Only two +per-thread scratch areas change, both value-preserving: +- LDS: 4 planes x 3 floats at S + 3072 p + 4 j, where S = T * 12 + base is the thread's own slot, written in + straight-line code before any branch, read back at fixed offsets and once as plane k (address + (k << 8) * 12 + S) -> 12 registers + a 4-way select. The shared array is removed (12 KB LDS per block). +- the 32-byte local depot: 4 float pairs at depot + 0..28 written before any branch, read once as the pair at + depot + 8 k -> 8 registers + a 4-way select (it compiled to scratch memory). +The plane index k must be the same register for both reads, and is proven to lie in 0..3 by evaluating its +defining expression (which may depend only on workgroup/workitem ids and constants) for every thread of +grids up to 65536 x 16384 threads. Anything else is refused. +""" +import re +import sys + +import numpy as np + +src = open(sys.argv[1]).read() +lines = src.split('\n') +k0 = next(i for i, l in enumerate(lines) if l.startswith('extern "C" __attribute__((global))')) +assert lines[k0 + 1] == '{' +body_start = k0 + 2 +first_label = next(i for i in range(body_start, len(lines)) if re.match(r'^L__BB0_\d+:;$', lines[i])) + +defs = {} +for i, l in enumerate(lines): + m = re.match(r'^ ((?:r|rd)\d+) = (.*);$', l) + if m: + defs.setdefault(m.group(1), []).append((i, m.group(2))) + + +def single(reg): + d = defs.get(reg, []) + assert len(d) == 1, f'{reg}: {len(d)} definitions' + return d[0] + + +# ---- LDS slot S, base, plane address P = (k << 8) * 12 + S +ST = re.compile(r'^ \*\(AS3 float\*\)\(uintptr_t\)\((r\d+) \+ (\d+)u\) = \(float\)\((f\d+)\);$') +LD = re.compile(r'^ (f\d+) = \*\(AS3 float\*\)\(uintptr_t\)\((r\d+) \+ (\d+)u\);$') +stores = [(i, ST.match(l)) for i, l in enumerate(lines) if ST.match(l)] +slot = {m.group(1) for _, m in stores} +assert len(slot) == 1, slot +S = slot.pop() +_, sdef = single(S) +m = re.fullmatch(r'\(uint32_t\)\((r\d+) \* 12u \+ (r\d+)\)', sdef) +assert m, sdef +T, BASE = m.groups() +_, bdef = single(BASE) +shared = re.fullmatch(r'\(uint32_t\)\(\(uint32_t\)\(uintptr_t\)\(AS3 uint8_t\*\)(\w+)\)', bdef) +assert shared, bdef +shared_name = shared.group(1) +offsets = {} +for i, m in stores: + assert i < first_label, f'LDS store inside control flow: {lines[i]}' + off = int(m.group(2)) + assert off % 4 == 0 and off // 3072 < 4 and (off % 3072) // 4 < 3 and off not in offsets, lines[i] + offsets[off] = m.group(3) +assert len(offsets) == 12, sorted(offsets) +loads = [(i, LD.match(l)) for i, l in enumerate(lines) if LD.match(l)] +plane_regs = {m.group(2) for _, m in loads if m.group(2) != S} +assert len(plane_regs) == 1, plane_regs +P = plane_regs.pop() +_, pdef = single(P) +m = re.fullmatch(r'\(uint32_t\)\((r\d+) \* 12u \+ ' + S + r'\)', pdef) +assert m, pdef +_, shdef = single(m.group(1)) +m = re.fullmatch(r'\(uint32_t\)\((r\d+) << \(8u & 31\)\)', shdef) +assert m, shdef +K = m.group(1) +for i, m in loads: + if m.group(2) == P: + assert m.group(3) == '0', lines[i] + +# ---- local depot +DBASE = {r for r, d in defs.items() if any(e == '(uint64_t)(l2_SPL + (uint64_t)0ull)' for _, e in d)} +DST = re.compile(r'^ \*\(AS5 float\*\)\(uintptr_t\)\(uint32_t\)\((rd\d+) \+ (\d+)\) = \(float\)\((f\d+)\);$') +DLD = re.compile(r'^ (f\d+) = \*\(AS5 float\*\)\(uintptr_t\)\(uint32_t\)\((rd\d+) \+ ([04])\);$') +dstores = {} +for i, l in enumerate(lines): + m = DST.match(l) + if m: + assert m.group(1) in DBASE and i < first_label, l + off = int(m.group(2)) + assert off % 4 == 0 and off < 32 and off not in dstores, l + dstores[off] = m.group(3) +assert len(dstores) == 8, sorted(dstores) +dloads = [(i, DLD.match(l)) for i, l in enumerate(lines) if DLD.match(l)] +assert len(dloads) == 2, len(dloads) +dl = {m.group(2) for _, m in dloads} +assert len(dl) == 1, dl +DL = dl.pop() +_, ddef = single(DL) +m = re.fullmatch(r'\(uint64_t\)\((rd\d+) \+ (rd\d+)\)', ddef) +assert m and m.group(1) in DBASE, ddef +_, kdef = single(m.group(2)) +assert kdef == f'(uint64_t)((int64_t)(int32_t){K} * (int64_t)(int32_t)8)', kdef +assert all(re.search(r'\bAS5\b', lines[i]) is None or i in [j for j, _ in dloads] or DST.match(lines[i]) + or 'l2_SPL' in lines[i] or '__local_depot0' in lines[i] for i in range(len(lines))), 'other depot use' +assert sum('AS3' in l for l in lines) == len(stores) + len(loads) + 1, 'other LDS use' # + the base address + +# ---- prove k in 0..3 for every thread: evaluate its expression over the thread-id range +IDS = {'workgroup_id_x': 'wgx', 'workitem_id_x': 'wix', 'workgroup_id_y': 'wgy', 'workitem_id_y': 'wiy'} + + +def to_np(reg, env, seen=()): + assert reg not in seen + if reg in env: + return env[reg] + _, e = single(reg) + m = re.fullmatch(r'\(uint32_t\)\(\(uint32_t\)__builtin_amdgcn_(\w+)\(\)\)', e) + if m: + env[reg] = env[IDS[m.group(1)]] + return env[reg] + expr = e + for r in sorted(set(re.findall(r'\br\d+\b', e)), key=len, reverse=True): + to_np(r, env, seen + (reg,)) + expr = re.sub(r'\(uint32_t\)', '', expr) + expr = re.sub(r'\(int32_t\)(r\d+)', r's32(\1)', expr) + expr = re.sub(r'(\d+)u\b', r'np.uint32(\1)', expr) + expr = re.sub(r'\(np\.uint32\((\d+)\) & np\.uint32\(31\)\)', r'np.uint32(\1)', expr) + assert re.fullmatch(r'[\w\s().,+\-*&|<>]*', expr), e + env[reg] = u32(eval(expr, {'np': np, 's32': s32, 'bfi32': bfi32}, env)) + return env[reg] + + +def u32(x): + return np.asarray(x).astype(np.int64).astype(np.uint32) if np.asarray(x).dtype.kind == 'i' else np.asarray(x, np.uint32) + + +def s32(x): + return x.view(np.int32) if isinstance(x, np.ndarray) else np.int32(x) + + +def bfi32(a, b, offset, width): + mask = np.uint32(((1 << int(width)) - 1) << int(offset)) + return (u32(a) << np.uint32(offset)) & mask | (u32(b) & ~mask) + + +gx = np.arange(65536, dtype=np.uint32) +for gy0 in range(0, 16384, 256): + gy = np.arange(gy0, gy0 + 256, dtype=np.uint32)[:, None] + env = {'wgx': np.broadcast_to(gx >> 4, (256, 65536)).copy(), 'wix': np.broadcast_to(gx & 15, (256, 65536)).copy(), + 'wgy': np.broadcast_to(gy >> 4, (256, 65536)).copy(), 'wiy': np.broadcast_to(gy & 15, (256, 65536)).copy()} + k = to_np(K, env) + assert k.max() <= 3, f'plane index {K} reaches {k.max()} at rows {gy0}..' + +# ---- emit +out = [] +sel = lambda name: (f'{K} == 0u ? {name(0)} : {K} == 1u ? {name(1)} : {K} == 2u ? {name(2)} : {name(3)}') +for i, l in enumerate(lines): + if re.match(rf'^__attribute__\(\(shared\)\).*\b{shared_name}\[', l): + continue # the LDS scratch array + if i == k0 + 1: + out.append(l) + out.append(' float ' + ', '.join(f'shv{p}_{j}' for p in range(4) for j in range(3)) + ';') + out.append(' float ' + ', '.join(f'dpv{p}_{j}' for p in range(4) for j in range(2)) + ';') + continue + if i == single(BASE)[0]: + out.append(f' {BASE} = 0u; // (LDS scratch kept in registers)') + continue + m = ST.match(l) + if m: + off = int(m.group(2)) + out.append(f' shv{off // 3072}_{(off % 3072) // 4} = {m.group(3)};') + continue + m = LD.match(l) + if m: + if m.group(2) == P: + out.append(f' {m.group(1)} = {sel(lambda p: f"shv{p}_0")}; // {K} in 0..3 (proven)') + else: + off = int(m.group(3)) + assert off in offsets, f'load before store: {l}' + out.append(f' {m.group(1)} = shv{off // 3072}_{(off % 3072) // 4};') + continue + m = DST.match(l) + if m: + off = int(m.group(2)) + out.append(f' dpv{off // 8}_{(off % 8) // 4} = {m.group(3)};') + continue + m = DLD.match(l) + if m: + j = int(m.group(3)) // 4 + out.append(f' {m.group(1)} = {sel(lambda p: f"dpv{p}_{j}")};') + continue + out.append(l) +print(f'// Generated by kernels/native/in_lds2reg.py from {sys.argv[1].split("/")[-1]} (do not edit).') +print('\n'.join(out), end='') diff --git a/kernels/native/rewrite_output.py b/kernels/native/rewrite_output.py index ae31309..3315ac6 100755 --- a/kernels/native/rewrite_output.py +++ b/kernels/native/rewrite_output.py @@ -43,6 +43,12 @@ def rewrite(source): 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); diff --git a/kernels/tools/dump_runner.cpp b/kernels/tools/dump_runner.cpp index 5d9f7aa..42cab1f 100644 --- a/kernels/tools/dump_runner.cpp +++ b/kernels/tools/dump_runner.cpp @@ -133,17 +133,87 @@ int main(int argc, char** argv) size_t argSize = args.size(); void* cfg[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, args.data(), HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END}; + // phase kernels NAME_post1..NAME_post8, launched after the main kernel with the same parameters, exactly as + // the ZLUDA hook and MicroCUDA do: u32 globals d4r_postK_grid_x/_grid_y (0 = the launch's grid; if either is + // set, grid z is 1) and d4r_postK_block_x/_y/_z (0 = 32/1/1) + struct Post { hipFunction_t fn; unsigned gx, gy, gz, bx, by, bz; }; + std::vector posts; + auto global_u32 = [&](const std::string& name) -> uint32_t { + hipDeviceptr_t p; + size_t n = 0; + uint32_t v = 0; + if (hipModuleGetGlobal(&p, &n, mod, name.c_str()) == hipSuccess && n == 4 && hipMemcpyDtoH(&v, p, 4) != hipSuccess) + v = 0; + return v; + }; + for (int k = 1; k <= 8; ++k) + { + Post q{}; + const std::string name = std::string(argv[2]) + "_post" + std::to_string(k); + if (hipModuleGetFunction(&q.fn, mod, name.c_str()) != hipSuccess) + break; + const std::string g = "d4r_post" + std::to_string(k) + "_"; + const uint32_t px = global_u32(g + "grid_x"), py = global_u32(g + "grid_y"); + q.gx = px ? px : grid[0]; + q.gy = py ? py : grid[1]; + q.gz = (px || py) ? 1 : grid[2]; + q.bx = global_u32(g + "block_x"); + q.by = global_u32(g + "block_y"); + q.bz = global_u32(g + "block_z"); + if (!q.bx) q.bx = 32; + if (!q.by) q.by = 1; + if (!q.bz) q.bz = 1; + posts.push_back(q); + } + if (!posts.empty()) + printf("%s: %zu post phase(s)\n", argv[2], posts.size()); const bool prepOnce = getenv("PREP_ONCE") != nullptr; int launches = 0; - auto launch = [&]() { + // PHASE_TIMES=1: an event before and after every phase (prep, main, post1..) for per-phase GPU times + const bool phaseTimes = getenv("PHASE_TIMES") != nullptr; + std::vector marks(posts.size() + 3); + for (auto& m : marks) + CHECK(hipEventCreate(&m)); + std::vector> phaseMs(posts.size() + 2); + bool recordPhases = false; + auto launch = [&]() -> hipError_t { + bool prepped = false; + if (recordPhases) (void)hipEventRecord(marks[0]); if (prep != nullptr && !(prepOnce && launches++ > 0)) { hipError_t e = hipModuleLaunchKernel(prep, prepGrid, 1, 1, 128, 1, 1, 0, nullptr, nullptr, cfg); if (e != hipSuccess) return e; + prepped = true; } - return hipModuleLaunchKernel(fn, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared, nullptr, - nullptr, cfg); + if (recordPhases) (void)hipEventRecord(marks[1]); + hipError_t e = hipModuleLaunchKernel(fn, grid[0], grid[1], grid[2], block[0], block[1], block[2], shared, + nullptr, nullptr, cfg); + if (e != hipSuccess) + return e; + if (recordPhases) (void)hipEventRecord(marks[2]); + for (size_t k = 0; k < posts.size(); ++k) + { + const Post& q = posts[k]; + e = hipModuleLaunchKernel(q.fn, q.gx, q.gy, q.gz, q.bx, q.by, q.bz, 0, nullptr, nullptr, cfg); + if (e != hipSuccess) + return e; + if (recordPhases) (void)hipEventRecord(marks[3 + k]); + } + if (recordPhases) + { + if (hipError_t w = hipEventSynchronize(marks.back()); w != hipSuccess) + return w; + for (size_t k = 0; k < phaseMs.size(); ++k) + { + if (k == 0 && !prepped) + continue; + float t; + if (hipEventElapsedTime(&t, marks[k], marks[k + 1]) == hipSuccess) + phaseMs[k].push_back(t); + } + } + return hipSuccess; }; CHECK(launch()); CHECK(hipDeviceSynchronize()); @@ -178,6 +248,22 @@ int main(int argc, char** argv) CHECK(hipEventElapsedTime(&t, a, b)); ms.push_back(t); } + if (phaseTimes) + { + recordPhases = true; + for (int i = 0; i < iters; ++i) + CHECK(launch()); + CHECK(hipDeviceSynchronize()); + for (size_t k = 0; k < phaseMs.size(); ++k) + { + auto& v = phaseMs[k]; + if (v.empty()) + continue; + std::sort(v.begin(), v.end()); + const std::string phase = k == 0 ? "prep" : k == 1 ? "main" : "post" + std::to_string(k - 1); + printf("%s %s: median %.4f ms, min %.4f ms\n", argv[2], phase.c_str(), v[v.size() / 2], v[0]); + } + } std::sort(ms.begin(), ms.end()); if (!ms.empty()) printf("%s: median %.4f ms, min %.4f ms over %d launches\n", argv[2], ms[ms.size() / 2], ms[0], iters); diff --git a/scripts/build_microcuda.sh b/scripts/build_microcuda.sh new file mode 100755 index 0000000..71422c9 --- /dev/null +++ b/scripts/build_microcuda.sh @@ -0,0 +1,15 @@ +#!/usr/bin/env bash +# Experimental backend; no installation or modification of the reference path. +# usage: build_microcuda.sh [OUTPUT_DIR] +# Modules come from a pack made separately with tools/microcuda/make_pack.py (D4R_MICROCUDA_PACK at run time). +set -euo pipefail +root="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +[[ $# -le 1 ]] || { echo "usage: $0 [OUTPUT_DIR]" >&2; exit 2; } +rocm="${D4R_ROCM_DIR:-/opt/rocm}" +out="$(realpath -m "${1:-$root/build/microcuda}")" +mkdir -p "$out" +"${CXX:-g++}" -std=c++20 -O2 -fPIC -shared -Wall -Wextra -Werror -Wno-deprecated-declarations \ + -I"$rocm/include" "$root/tools/microcuda/runtime.cpp" \ + -L"$rocm/lib" -Wl,-rpath,"$rocm/lib" -lamdhip64 -lcrypto -lpthread \ + -Wl,-z,defs -Wl,-soname,libd4r_microcuda.so -o "$out/libd4r_microcuda.so" +echo "Built $out/libd4r_microcuda.so (opt-in; modules from a make_pack.py pack)" diff --git a/tools/bench_dlss_kernel.cpp b/tools/bench_dlss_kernel.cpp index 4788580..5dbf389 100644 --- a/tools/bench_dlss_kernel.cpp +++ b/tools/bench_dlss_kernel.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -120,10 +121,15 @@ int main(int argc, char** argv) std::vector image = read_file(argv[1]); image.resize(image.size() + 16, 0); - void* library = dlopen("libcuda.so", RTLD_NOW | RTLD_LOCAL); + const char* library_path = std::getenv("D4R_BENCH_LIBCUDA"); + if (!library_path) library_path = "libcuda.so"; + const auto startup = std::chrono::steady_clock::now(); + auto us_since = [](auto start) { return std::chrono::duration( + std::chrono::steady_clock::now() - start).count(); }; + void* library = dlopen(library_path, RTLD_NOW | RTLD_LOCAL); if (library == nullptr) { - std::fprintf(stderr, "dlopen(libcuda.so): %s\n", dlerror()); + std::fprintf(stderr, "dlopen(%s): %s\n", library_path, dlerror()); return 1; } auto init = load(library, "cuInit"); @@ -149,8 +155,20 @@ int main(int argc, char** argv) CHECK(init(0)); CHECK(device_get(&device, 0)); CHECK(context_create(&context, 0, device)); + std::printf("startup_init_context_us=%.3f\n", us_since(startup)); + auto stage = std::chrono::steady_clock::now(); CHECK(module_load(&module, image.data())); + std::printf("module_load_us=%.3f\n", us_since(stage)); + stage = std::chrono::steady_clock::now(); CHECK(get_function(&function, module, kernel.c_str())); + std::printf("first_lookup_us=%.3f\n", us_since(stage)); + stage = std::chrono::steady_clock::now(); + for (int i = 0; i < 10000; ++i) { + void* again = nullptr; + CHECK(get_function(&again, module, kernel.c_str())); + if (again != function) return 3; + } + std::printf("cached_lookup_us=%.3f\n", us_since(stage) / 10000); for (Allocation& allocation : allocations) CHECK(mem_alloc(&allocation.device, allocation.bytes)); for (const Pointer& pointer : pointers) @@ -161,20 +179,39 @@ int main(int argc, char** argv) size_t argument_bytes = arguments.size(); void* extra[] = {reinterpret_cast(1), arguments.data(), reinterpret_cast(2), &argument_bytes, nullptr}; - auto restore = [&]() -> CUresult { - for (Allocation& allocation : allocations) + const bool immutable_weights = std::getenv("D4R_BENCH_IMMUTABLE_WEIGHTS") != nullptr; + size_t weight_allocation = allocations.size(); + if (immutable_weights) { + if (kernel != "dltss_pwin_enc0_layer") return 2; + for (const auto& pointer : pointers) if (pointer.offset == 64) weight_allocation = pointer.allocation; + if (weight_allocation == allocations.size()) return 2; + } + auto restore = [&](bool initial = true) -> CUresult { + for (size_t i = 0; i < allocations.size(); ++i) { + if (!initial && immutable_weights && i == weight_allocation) continue; + auto& allocation = allocations[i]; if (CUresult result = copy_h2d(allocation.device, allocation.data.data(), allocation.bytes)) return result; + } return 0; }; + void* parameter_array[] = {arguments.data()}; + const bool use_parameter_array = std::getenv("D4R_BENCH_PARAM_ARRAY") != nullptr; auto run = [&]() { return launch_kernel(function, launch[0], launch[1], launch[2], launch[3], launch[4], launch[5], - launch[6], nullptr, nullptr, extra); + launch[6], nullptr, use_parameter_array ? parameter_array : nullptr, + use_parameter_array ? nullptr : extra); }; + stage = std::chrono::steady_clock::now(); CHECK(restore()); + std::printf("restore_h2d_us=%.3f\n", us_since(stage)); + stage = std::chrono::steady_clock::now(); CHECK(run()); + std::printf("first_launch_cpu_us=%.3f\n", us_since(stage)); + stage = std::chrono::steady_clock::now(); CHECK(synchronize()); + std::printf("first_sync_wall_us=%.3f\n", us_since(stage)); for (size_t index = 0; index < allocations.size(); ++index) { std::vector result(allocations[index].bytes); @@ -201,18 +238,35 @@ int main(int argc, char** argv) CHECK(event_create(&start, 0)); CHECK(event_create(&end, 0)); std::vector times; + std::vector cpu_times, sync_times; for (int iteration = 0; iteration < iterations; ++iteration) { - CHECK(restore()); + CHECK(restore(false)); CHECK(synchronize()); CHECK(event_record(start, nullptr)); + stage = std::chrono::steady_clock::now(); CHECK(run()); + cpu_times.push_back(us_since(stage)); CHECK(event_record(end, nullptr)); + stage = std::chrono::steady_clock::now(); CHECK(synchronize()); + sync_times.push_back(us_since(stage)); float milliseconds = 0; CHECK(event_elapsed(&milliseconds, start, end)); times.push_back(milliseconds); } + using Stats = CUresult (*)(void*, uint64_t*, uint64_t*); + auto stats = reinterpret_cast(dlsym(library, "d4rMicrocudaGetDispatchStats")); + if (stats) { + uint64_t prep = 0, skips = 0; + CHECK(stats(function, &prep, &skips)); + std::printf("microcuda_prep_runs=%llu skips=%llu\n", (unsigned long long)prep, (unsigned long long)skips); + } + std::sort(cpu_times.begin(), cpu_times.end()); + std::sort(sync_times.begin(), sync_times.end()); + if (!cpu_times.empty()) + std::printf("launch_cpu_median_us=%.3f sync_wall_median_us=%.3f\n", + cpu_times[cpu_times.size()/2], sync_times[sync_times.size()/2]); std::sort(times.begin(), times.end()); if (!times.empty()) std::printf("kernel %s grid=%u,%u,%u block=%u,%u,%u iterations=%d min=%.4f median=%.4f max=%.4f ms\n", diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp index 8868f0b..0f46037 100644 --- a/tools/d4r_nvngx_shim.cpp +++ b/tools/d4r_nvngx_shim.cpp @@ -1615,6 +1615,7 @@ struct FrameTiming double prep = 0, worker = 0, workerWait = 0, finishWait = 0, markerWait = 0; double ngxHost = 0, ctxSync = 0, gpuEval = -1; double motionDilationWall = 0; + bool markerDeferred = false; // the worker waits for the input marker (D4R_SHIM_DEFER_MARKER) bool gpuEventsRecorded = false; double outputSyncWall = 0, outputSyncCpu = -1; bool outputSyncBlocking = false; @@ -1680,6 +1681,9 @@ struct Feature unsigned int preset = 0; CudaEvent profileStart = nullptr, profileEnd = nullptr; CudaEvent outputReadyEvent = nullptr; + // output wait pacing (D4R_SHIM_OUTPUT_POLL_US): smoothed time from the event record to its completion, kept + // separately for each wait site (they wait for very different amounts of GPU work) + double publishWaitEmaUs = 0.0, evalWaitEmaUs = 0.0; bool profileEventsReady = false; CudaImage color, depth, motion, exposure, output; @@ -3757,16 +3761,59 @@ static bool upload_inputs_vram(Feature& feature, InputSlot& slot, bool hasExposu // Only called on the CUDA worker. A completion event covers queued work on // the default stream; do not rely on HIP's blocking flag to release the CPU. -static int synchronize_default_stream(Feature& feature, bool* querySleep = nullptr) +static int synchronize_default_stream(Feature& feature, bool* querySleep = nullptr, double* waitEmaUs = nullptr) { if (feature.outputReadyEvent != nullptr) { const int recorded = g.cu.eventRecord(feature.outputReadyEvent, nullptr); if (recorded == 0) { + // D4R_SHIM_OUTPUT_POLL_US: N > 0 polls every N us (200 was the fixed interval before), 0 yields between + // queries (spins). Unset: adaptive - sleep until shortly before the predicted completion (smoothed + // wait of earlier frames), then poll every 20 us, backing off to 200 us when the GPU runs late. + static const int fixedPoll = [] { + const char* v = std::getenv("D4R_SHIM_OUTPUT_POLL_US"); + return v != nullptr && *v != '\0' ? static_cast(env_uint("D4R_SHIM_OUTPUT_POLL_US", 200)) : -1; + }(); + const auto waitStart = std::chrono::steady_clock::now(); + unsigned polls = 0; + double lastNotReadyUs = -1.0; // elapsed time of the last query that found the work unfinished const int result = d4r_wait_event( [&] { return g.cu.eventQuery(feature.outputReadyEvent); }, - [] { d4r_sleep_us(200); }); + [&] { + if (fixedPoll >= 0) + { + d4r_sleep_us(static_cast(fixedPoll)); + return; + } + constexpr double kMarginUs = 120.0, kStepUs = 20.0, kChunkUs = 200.0, kSpinAfterUs = 150.0; + const double elapsed = std::chrono::duration( + std::chrono::steady_clock::now() - waitStart).count(); + lastNotReadyUs = elapsed; // this sleep follows a query that returned NOT_READY + ++polls; + const double predicted = waitEmaUs != nullptr ? *waitEmaUs : 0.0; + const double remaining = predicted - kMarginUs - elapsed; + // far from the predicted completion: sleep in chunks of at most 200 us (a single long sleep + // can overshoot by far more under Wine); near it: 20 us polls; late: back off to 200 us + if (remaining > kStepUs) + d4r_sleep_us(static_cast(std::min(remaining, kChunkUs))); + else if (elapsed < predicted + kSpinAfterUs) + d4r_sleep_us(0); // final stretch around the predicted completion: spin (SwitchToThread); + // a Wine sleep rounds up to ~50 us, which would be the detection latency + else + d4r_sleep_us(static_cast(std::min( + kStepUs * (1.0 + (elapsed - predicted - kSpinAfterUs) / 160.0), kChunkUs))); + }); + if (result == 0 && fixedPoll < 0 && waitEmaUs != nullptr) + { + const double waited = std::chrono::duration( + std::chrono::steady_clock::now() - waitStart).count(); + // Learn the completion time, not our own wait: the work finished between the last NOT_READY query + // and the successful one (at most one 200 us chunk apart), so the midpoint cannot ratchet upwards + double sample = lastNotReadyUs >= 0.0 ? 0.5 * (lastNotReadyUs + waited) : waited; + sample = std::min(sample, 20000.0); + *waitEmaUs = *waitEmaUs == 0.0 ? sample : 0.8 * *waitEmaUs + 0.2 * sample; + } if (result == 0) { if (querySleep != nullptr) *querySleep = true; @@ -3811,7 +3858,7 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa { const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; const double cpuStart = timing.enabled ? profile_thread_cpu_ms() : -1.0; - result = synchronize_default_stream(*feature, &timing.outputSyncBlocking); + result = synchronize_default_stream(*feature, &timing.outputSyncBlocking, &feature->publishWaitEmaUs); if (timing.enabled) { timing.outputSyncWall = profile_ms(syncStart, ProfileClock::now()); @@ -3943,6 +3990,8 @@ static bool recreate_for_render_size(Feature& feature, FrameParams& params, uint return feature.cudaHandle != nullptr && width == params.renderWidth && height == params.renderHeight; } +static bool wait_for_input_marker(Feature* feature, int slotIndex, uint32_t frame); + static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, FrameParams params, FrameTiming timing) { InputSlot& slot = feature->inputs[slotIndex]; @@ -4039,6 +4088,14 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram } feature->outputRedirected = redirect; } + if (timing.markerDeferred) + { + const auto markerStart = ProfileClock::now(); + if (!wait_for_input_marker(feature, slotIndex, frame)) + return; + if (timing.enabled) + timing.markerWait = profile_ms(markerStart, ProfileClock::now()); + } const int eventStartResult = timing.enabled && feature->profileEventsReady ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; const auto evalStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -4057,7 +4114,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram // D4R_SHIM_EVAL_SYNC=0 (VRAM interop): no wait between NGX's kernels and the output copy queued // behind them on the same stream; publish_vram synchronises before the result is released. static const bool evalSync = env_uint("D4R_SHIM_EVAL_SYNC", 1) != 0; - const int syncResult = params.vram && !evalSync ? 0 : synchronize_default_stream(*feature); + const int syncResult = params.vram && !evalSync ? 0 : synchronize_default_stream(*feature, nullptr, &feature->evalWaitEmaUs); const auto evaluated = ProfileClock::now(); const bool gpuEventsRecorded = eventStartResult == 0 && eventEndResult == 0; float gpuEvalMs = -1.0f; @@ -4120,10 +4177,12 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram post_publish(feature, hostIndex, frame, params, timing); } -static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, FrameParams params, FrameTiming timing) +// Waits for this frame's own input marker (the GPU reached the game's input copies). false: the frame was +// dropped (retired feature, skipped submission, timeout) and has been retired here. +static bool wait_for_input_marker(Feature* feature, int slotIndex, uint32_t frame) { InputSlot& slot = feature->inputs[slotIndex]; - const auto start = ProfileClock::now(); + const auto start = std::chrono::steady_clock::now(); // A later input marker does not prove this slot's copies were submitted: // the game may have discarded this frame's command list. Require its own // tag, otherwise stale input data can enter DLSS's temporal history. @@ -4143,7 +4202,7 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram *inputMarker); slot.busy = false; frame_retired(feature, frame); - return; + return false; } // D4R_SHIM_MARKER_POLL_US: poll interval (0 = yield); detection latency is on the critical path static const unsigned markerPollUs = env_uint("D4R_SHIM_MARKER_POLL_US", 200); @@ -4152,6 +4211,34 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram else d4r_sleep_us(markerPollUs); } + return true; +} + +static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, FrameParams params, FrameTiming timing) +{ + InputSlot& slot = feature->inputs[slotIndex]; + const auto start = ProfileClock::now(); + // D4R_SHIM_DEFER_MARKER (default on): with VRAM interop nothing else happens here before the marker lands, so + // the worker thread waits for it itself, after the host-side setup of the evaluation (which no longer queues + // behind the marker: ~0.1 ms off the critical path) and without a thread wake-up between marker and NGX. + static const bool deferMarker = env_uint("D4R_SHIM_DEFER_MARKER", 1) != 0 && + env_uint("D4R_MOTION_DILATION", 0) == 0; + if (params.vram && deferMarker && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) + { + timing.markerDeferred = true; + if (timing.enabled) + { + timing.prepStart = start; + timing.prep = profile_ms(start, ProfileClock::now()); + timing.workerQueued = ProfileClock::now(); + } + g.worker.post([feature, slotIndex, frame, params, timing] { + run_evaluation(feature, slotIndex, frame, params, timing); + }); + return; + } + if (!wait_for_input_marker(feature, slotIndex, frame)) + return; const auto ready = ProfileClock::now(); // With VRAM interop the inputs are already on the GPU; host staging runs // only to cross-check them (D4R_SHIM_VRAM_VERIFY). diff --git a/tools/microcuda/arch_spoof.c b/tools/microcuda/arch_spoof.c new file mode 100644 index 0000000..7620845 --- /dev/null +++ b/tools/microcuda/arch_spoof.c @@ -0,0 +1,35 @@ +/* LD_PRELOAD for zluda_precompile.py: report D4R_SPOOF_GFX (e.g. gfx1201) as the device's gcnArchName, so + * ZLUDA compiles (and caches) code objects for a GPU that is not installed. Loading them then fails, which + * the precompile treats as expected; the cache entry is written before the load. Never use it for running. */ +#define _GNU_SOURCE +#define __HIP_PLATFORM_AMD__ +#include +#include +#include +#include + +hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* props, int device) +{ + static hipError_t (*real)(hipDeviceProp_tR0600*, int); + if (!real) + real = (hipError_t (*)(hipDeviceProp_tR0600*, int))dlsym(RTLD_NEXT, "hipGetDevicePropertiesR0600"); + if (!real) + { + /* ZLUDA loaded RTLD_LOCAL (ctypes): its HIP is not in the global scope RTLD_NEXT searches */ + void* hip = dlopen("libamdhip64.so.7", RTLD_NOW | RTLD_NOLOAD); + if (!hip) + hip = dlopen("libamdhip64.so", RTLD_NOW | RTLD_NOLOAD); + if (hip) + real = (hipError_t (*)(hipDeviceProp_tR0600*, int))dlsym(hip, "hipGetDevicePropertiesR0600"); + } + if (!real) + return hipErrorNotInitialized; + hipError_t result = real(props, device); + const char* arch = getenv("D4R_SPOOF_GFX"); + if (result == hipSuccess && props && arch && *arch) + { + memset(props->gcnArchName, 0, sizeof(props->gcnArchName)); + strncpy(props->gcnArchName, arch, sizeof(props->gcnArchName) - 1); + } + return result; +} diff --git a/tools/microcuda/check_runtime.py b/tools/microcuda/check_runtime.py new file mode 100644 index 0000000..b691317 --- /dev/null +++ b/tools/microcuda/check_runtime.py @@ -0,0 +1,202 @@ +#!/usr/bin/env python3 +"""GPU integration checks for the MicroCUDA runtime; no games or NGX evaluation. +usage: check_runtime.py LIBRARY PACK_DIR KNOWN_PTX NATIVE_DIR + +KNOWN_PTX is the K enc0 module text (a module in the pack with a native dltss_pwin_enc0_layer in NATIVE_DIR). +Pack-state failures (switch mismatch, tampered object) run in child processes: the pack is read once. +""" +import ctypes as c +import json +import os +from pathlib import Path +import shutil +import struct +import subprocess +import sys +import tempfile + +ptr, u64, size = c.c_void_p, c.c_uint64, c.c_size_t + + +def api(lib, name, args): + f = getattr(lib, name) + f.argtypes = args + f.restype = c.c_int + return f + + +def child(library, ptx): + """Prints the cuModuleLoadData result for ptx under the environment it was started with.""" + lib = c.CDLL(library) + assert api(lib, 'cuInit', [c.c_uint])(0) == 0 + context = ptr() + assert api(lib, 'cuCtxCreate_v2', [c.POINTER(ptr), c.c_uint, c.c_int])(c.byref(context), 0, 0) == 0 + module = ptr() + print(api(lib, 'cuModuleLoadData', [c.POINTER(ptr), ptr])(c.byref(module), c.create_string_buffer(Path(ptx).read_bytes()))) + + +def run_child(library, ptx, **env): + out = subprocess.run([sys.executable, __file__, '--child', library, ptx], env={**os.environ, **env}, + capture_output=True, text=True, check=True) + return int(out.stdout.split()[-1]) + + +def main(): + if sys.argv[1] == '--child': + return child(*sys.argv[2:]) + library, pack, ptx, native = sys.argv[1:] + os.environ['D4R_MICROCUDA_PACK'] = pack + os.environ['D4R_MICROCUDA_NATIVE_DIR'] = native + switches = next(l[9:].strip() for l in open(Path(pack) / 'microcuda-pack.txt') if l.startswith('switches')) + for key in [k for k in os.environ if k.startswith('D4R_ZLUDA_')]: + del os.environ[key] + for item in filter(None, switches.split(',')): + k, v = item.split('=', 1) + os.environ[k] = v + checks = [] + + def expect(label, result, code=0): + assert result == code, (label, result, code) + checks.append(label) + + expect('switch-set mismatch refuses the pack', run_child(library, ptx, D4R_ZLUDA_WMMA='0'), 200) + with tempfile.TemporaryDirectory(prefix='d4r-microcuda-check-') as directory: + shutil.copytree(pack, directory, dirs_exist_ok=True) + for co in (Path(directory) / 'modules').glob('*.co'): + b = bytearray(co.read_bytes()) + b[-1] ^= 1 + co.write_bytes(b) + expect('tampered pack object rejected', run_child(library, ptx, D4R_MICROCUDA_PACK=directory), 200) + expect('missing pack', run_child(library, ptx, D4R_MICROCUDA_PACK='/nonexistent-microcuda'), 500) + + lib = c.CDLL(library) + A = lambda name, args: api(lib, name, args) + init = A('cuInit', [c.c_uint]) + load = A('cuModuleLoadData', [c.POINTER(ptr), ptr]) + unload = A('cuModuleUnload', [ptr]) + get = A('cuModuleGetFunction', [c.POINTER(ptr), ptr, c.c_char_p]) + launch = A('cuLaunchKernel', [ptr] + [c.c_uint] * 7 + [ptr, c.POINTER(ptr), c.POINTER(ptr)]) + text = c.create_string_buffer(Path(ptx).read_bytes()) + module, function = ptr(), ptr() + expect('module before init', load(c.byref(module), text), 3) + expect('invalid init flags', init(1), 1) + expect('init', init(0)) + proc = A('cuGetProcAddress_v2', [c.c_char_p, c.POINTER(ptr), c.c_int, u64, c.POINTER(c.c_int)]) + address, status = ptr(), c.c_int() + expect('getProc modern alias', proc(b'cuMemAlloc', c.byref(address), 12080, 0, c.byref(status))) + assert address.value == c.cast(lib.cuMemAlloc_v2, ptr).value and status.value == 0 + expect('getProc array API', proc(b'cuArrayCreate', c.byref(address), 12080, 0, c.byref(status))) + expect('getProc unknown', proc(b'cuGraphLaunch', c.byref(address), 12080, 0, c.byref(status)), 500) + assert address.value is None and status.value == 1 + expect('getProc perthread flag rejected', proc(b'cuLaunchKernel', c.byref(address), 12080, 2, c.byref(status)), 801) + expect('getProc future version rejected', proc(b'cuLaunchKernel', c.byref(address), 13000, 0, c.byref(status)), 801) + expect('getProc null output', proc(b'cuInit', None, 12080, 0, None), 1) + device = c.c_int() + expect('device', A('cuDeviceGet', [c.POINTER(c.c_int), c.c_int])(c.byref(device), 0)) + context = ptr() + expect('context', A('cuCtxCreate_v2', [c.POINTER(ptr), c.c_uint, c.c_int])(c.byref(context), 0, device.value)) + current = ptr() + expect('get current', A('cuCtxGetCurrent', [c.POINTER(ptr)])(c.byref(current))) + assert current.value == context.value + + expect('unknown PTX', load(c.byref(module), c.create_string_buffer(b'.version 8.7\n.target sm_89\n.address_size 64\n')), 200) + bad = bytearray(Path(ptx).read_bytes()) + bad[0] ^= 1 + expect('changed PTX', load(c.byref(module), c.create_string_buffer(bytes(bad))), 200) + expect('known PTX', load(c.byref(module), text)) + expect('unknown function', get(c.byref(function), module, b'unknown'), 500) + assert function.value is None + expect('native function', get(c.byref(function), module, b'dltss_pwin_enc0_layer')) + another = ptr() + expect('cached lookup', get(c.byref(another), module, b'dltss_pwin_enc0_layer')) + assert another.value == function.value + os.environ['D4R_MICROCUDA_NATIVE_DIR'] = '/nonexistent-microcuda' + plain_module, plain = ptr(), ptr() + expect('second module instance', load(c.byref(plain_module), text)) + expect('pack function without native override', get(c.byref(plain), plain_module, b'dltss_pwin_enc0_layer')) + assert plain.value != function.value + expect('second module unload', unload(plain_module)) + os.environ['D4R_MICROCUDA_NATIVE_DIR'] = native + args = c.create_string_buffer(176) + n = size(176) + extra = (ptr * 5)(3, c.addressof(args), 2, c.addressof(n), None) + expect('HIP token rejected as CUDA token', launch(function, 1, 1, 1, 32, 1, 1, 0, None, None, extra), 801) + unterminated = (ptr * 8)(1, c.addressof(args), 2, c.addressof(n), 1, c.addressof(args), 2, c.addressof(n)) + expect('unterminated packed args', launch(function, 1, 1, 1, 32, 1, 1, 0, None, None, unterminated), 1) + expect('null function', launch(None, 1, 1, 1, 32, 1, 1, 0, None, None, extra), 400) + + allocate = A('cuMemAlloc_v2', [c.POINTER(u64), size]) + d2h = A('cuMemcpyDtoH_v2', [ptr, u64, size]) + source = c.create_string_buffer(bytes(range(256))) + result = c.create_string_buffer(256) + device_mem = u64() + expect('allocation', allocate(c.byref(device_mem), 256)) + expect('zero-filled allocation readback', d2h(result, device_mem, 256)) + assert result.raw == bytes(256) + stream = ptr() + expect('stream create', A('cuStreamCreate', [c.POINTER(ptr), c.c_uint])(c.byref(stream), 1)) + expect('async H2D', A('cuMemcpyHtoDAsync_v2', [u64, ptr, size, ptr])(device_mem, source, 256, stream)) + event = ptr() + expect('event create', A('cuEventCreate', [c.POINTER(ptr), c.c_uint])(c.byref(event), 0)) + expect('event record', A('cuEventRecord', [ptr, ptr])(event, stream)) + expect('event wait', A('cuEventSynchronize', [ptr])(event)) + expect('event query ready', A('cuEventQuery', [ptr])(event)) + expect('D2H', d2h(result, device_mem, 256)) + assert result.raw == source.raw[:256] + expect('event destroy', A('cuEventDestroy_v2', [ptr])(event)) + expect('stream destroy', A('cuStreamDestroy_v2', [ptr])(stream)) + expect('free', A('cuMemFree_v2', [u64])(device_mem)) + + # arrays: CUDA_ARRAY_DESCRIPTOR {size_t Width, Height; int Format; unsigned NumChannels} + array_create = A('cuArrayCreate_v2', [c.POINTER(ptr), ptr]) + array_desc = A('cuArrayGetDescriptor_v2', [ptr, ptr]) + array_destroy = A('cuArrayDestroy', [ptr]) + memcpy2d = A('cuMemcpy2D_v2', [ptr]) + unorm = ptr() + expect('UNORM8x4 array', array_create(c.byref(unorm), c.create_string_buffer(struct.pack(' array -> host + packed = ptr() + expect('packed 1010102 array', array_create(c.byref(packed), c.create_string_buffer(struct.pack('.co and OUT_DIR/microcuda-pack.txt, read by the runtime. Fatbins without +PTX become 'empty' entries: ZLUDA loads those as modules without functions. +Needs the blake3 Python package. No PTX or NVIDIA code is written; the objects are local build +products like ZLUDA's cache itself. +""" +import argparse +import collections +import hashlib +import json +import os +from pathlib import Path +import sqlite3 +import struct +import sys + +import blake3 + +sys.path.insert(0, str(Path(__file__).resolve().parents[2] / 'kernels/tools')) +from extract_dlss_ptx import fatbins, ptx_entries +from kernel_manifest import fnv1a64 + +# EF_AMDGPU_MACH (low byte of the ELF e_flags) of each supported target +MACH = {'gfx1100': 0x41, 'gfx1101': 0x46, 'gfx1102': 0x47, 'gfx1103': 0x44, 'gfx1200': 0x48, 'gfx1201': 0x4e} + +# switches that do not change generated code (see zluda_build_version in ZLUDA's module.rs) +IGNORED_SWITCHES = ('D4R_ZLUDA_CACHE_HOME=',) + + +def fingerprint(libcuda): + s = os.stat(libcuda) + return f'{s.st_size:x}.{int(s.st_mtime):x}.{s.st_mtime_ns % 1_000_000_000:x}' + + +def switches_of(version): + # "+[+]" + parts = version.split('+', 2) + raw = parts[2] if len(parts) == 3 else '' + return ','.join(s for s in raw.split(',') if s and not s.startswith(IGNORED_SWITCHES)) + + +def main(): + p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + p.add_argument('dll', type=Path) + p.add_argument('db', type=Path) + p.add_argument('libcuda', type=Path) + p.add_argument('out', type=Path) + p.add_argument('--switches', help='code-generation switch set to take when the cache holds several') + p.add_argument('--arch', default='gfx1101') + a = p.parse_args() + + fp = fingerprint(a.libcuda) + db = sqlite3.connect(f'file:{a.db}?mode=ro', uri=True) + rows = collections.defaultdict(list) + for h, version, attrs, rowid in db.execute( + 'select hash, zluda_version, backend_key, id from modules where device = ? and zluda_version like ?', + (a.arch, f'%+{fp}%')): + rows[h].append((switches_of(version), version, attrs, rowid)) + sets = collections.Counter(s for v in rows.values() for s, *_ in v) + if not sets: + sys.exit(f'no {a.arch} modules for ZLUDA build {fp} in {a.db}') + if a.switches is None: + if len(sets) > 1: + sys.exit('several switch sets in the cache; pick one with --switches:\n ' + + '\n '.join(f'{n:4} {s!r}' for s, n in sets.most_common())) + a.switches = next(iter(sets)) + + dll = a.dll.read_bytes() + (a.out / 'modules').mkdir(parents=True, exist_ok=True) + lines, missing, attrs_seen, version_seen = [], [], set(), set() + for index, (_, fatbin) in enumerate(fatbins(dll)): + ffnv = fnv1a64(fatbin) + # ZLUDA compiles the last PTX of a fatbin that it can use; take the last one it compiled + ptxs = list(ptx_entries(fatbin)) + if not ptxs: + # no PTX (only cubins for NVIDIA GPUs): ZLUDA's release build loads an empty module, so must we + lines.append(f'empty {ffnv:016x} {len(fatbin)}') + continue + chosen = None + for ptx in reversed(ptxs): + text = ptx.rstrip(b'\0') + for s, version, attrs, rowid in rows.get(blake3.blake3(text).hexdigest(), []): + if s == a.switches: + chosen = (text, version, attrs, rowid) + break + if chosen: + break + if chosen is None: + missing.append(f'{index}:{ffnv:016x}') + continue + text, version, attrs, rowid = chosen + attrs_seen.add(attrs) + version_seen.add(version) + binary = db.execute('select binary from modules where id = ?', (rowid,)).fetchone()[0] + if binary[:4] != b'\x7fELF' or struct.unpack_from(' 1 or len(version_seen) > 1: + sys.exit(f'inconsistent compile attributes in the chosen set: {attrs_seen} {version_seen}') + fp8 = json.loads(next(iter(attrs_seen))).get('wmma_fp8_native', False) if attrs_seen else False + header = [ + '# d4r MicroCUDA module pack (generated by tools/microcuda/make_pack.py; do not edit)', + f'# dll sha256 {hashlib.sha256(dll).hexdigest()}', + f'# zluda {next(iter(version_seen))}', + f'# attributes {next(iter(attrs_seen))}', + f'arch {a.arch}', + f'switches {a.switches}', + # RDNA4: compiled with native FP8 WMMA (d4r.ini NativeFp8); the runtime refuses a mismatched process + f'fp8native {int(fp8)}', + ] + (a.out / 'microcuda-pack.txt').write_text('\n'.join(header + sorted(lines)) + '\n') + print(f'{sum(l.startswith("module") for l in lines)} modules and {sum(l.startswith("empty") for l in lines)} empty modules packed for {a.arch}, switches {a.switches!r}; ' + f'{len(missing)} fatbins without a compiled object') + if missing: + (a.out / 'missing.txt').write_text('\n'.join(missing) + '\n') + + +if __name__ == '__main__': + main() diff --git a/tools/microcuda/make_registry.py b/tools/microcuda/make_registry.py new file mode 100644 index 0000000..fd45b07 --- /dev/null +++ b/tools/microcuda/make_registry.py @@ -0,0 +1,99 @@ +#!/usr/bin/env python3 +"""Offline inventory/registry. No proprietary PTX is emitted into the header. + +make_registry.py DLL NATIVE_DIR OUT_DIR [NAME ...] +Only audited K enc0 export is enabled for now. Inventory covers every PTX. +Hash identity: existing PTX FNV plus full fatbin FNV+size to avoid runtime LZ4. +""" +import argparse +import hashlib +import json +from pathlib import Path +import re +import subprocess +import sys + +sys.path.insert(0, str(Path(__file__).resolve().parents[2] / 'kernels/tools')) +from extract_dlss_ptx import fatbins, ptx_entries +from kernel_manifest import fnv1a64 + +ENTRY = re.compile(rb'\.entry\s+([\w$]+)\s*\((.*?)\)', re.S) +AUDITED = {'dltss_pwin_enc0_layer'} + + +def generate(dll, native, out, names): + if not set(names) <= AUDITED: + raise ValueError('Only K enc0 is ABI/phase audited; extend AUDITED after validation') + data = dll.read_bytes() + manifest = set() + for line in (native / 'd4r-kernels.txt').read_text().splitlines(): + fields = line.split() + if len(fields) == 2 and not line.startswith('#'): + manifest.add((fields[0], int(fields[1], 16))) + inventory, enabled = [], [] + for i, (_, fatbin) in enumerate(fatbins(data)): + for j, ptx in enumerate(ptx_entries(fatbin)): + ptx = ptx.rstrip(b'\0') + entries = ENTRY.findall(ptx) + phash, fhash = fnv1a64(ptx), fnv1a64(fatbin) + row = {'fatbin_index': i, 'ptx_index': j, 'ptx_hash': f'{phash:016x}', + 'fatbin_hash': f'{fhash:016x}', 'fatbin_bytes': len(fatbin), + 'kernels': []} + for raw_name, params in entries: + name = raw_name.decode() + m = re.fullmatch(rb'\s*\.param\s+\.align\s+(\d+)\s+\.b8\s+\w+\[(\d+)\]\s*', params) + abi = {'align': int(m[1]), 'bytes': int(m[2])} if m else None + row['kernels'].append({'name': name, 'aggregate_abi': abi, + 'native_manifest_match': (name, phash) in manifest, + 'native_file': (native / (name + '.hsaco')).is_file()}) + if name not in names: + continue + if abi != {'align': 8, 'bytes': 176}: + raise ValueError(f'{name}: unsupported module/ABI') + if (name, phash) not in manifest: + raise ValueError(f'{name}: PTX hash not verified by native manifest') + hsaco = native / (name + '.hsaco') + notes = subprocess.check_output(['llvm-readobj', '--notes', str(hsaco)], text=True) + # Verify both prep and main kernargs before generating a runnable registry. + for symbol in [name + '_prep', name]: + blocks = notes.split(' - .args:')[1:] + block = next((b for b in blocks if re.search(r'\.name:\s+' + re.escape(symbol) + r'\s', b)), '') + if not (re.search(r'\.size:\s+176\s', block) and + re.search(r'\.offset:\s+0\s', block) and + re.search(r'\.kernarg_segment_size:\s+176\s', block) and + block.count('.value_kind:') == 1 and + re.search(r'\.value_kind:\s+by_value\s', block) and + re.search(r'\.kernarg_segment_align:\s+8\s', block)): + raise ValueError(f'{symbol}: HSACO argument ABI mismatch') + arch = re.search(r'amdhsa.target:\s+amdgcn-amd-amdhsa--(gfx\d+)', notes) + if not arch: + raise ValueError('Missing exact GPU target in HSACO') + enabled.append((phash, len(ptx), fhash, len(fatbin), name, abi['bytes'], + arch[1], hashlib.sha256(hsaco.read_bytes()).hexdigest())) + inventory.append(row) + if {r[4] for r in enabled} != set(names): + raise ValueError('Requested kernel absent from DLL') + out.mkdir(parents=True, exist_ok=True) + (out / 'inventory.json').write_text(json.dumps({'dll_sha256': hashlib.sha256(data).hexdigest(), + 'modules': inventory}, indent=2) + '\n') + lines = ['// Generated offline; do not edit. Native bytes/ABI were verified at generation.', + 'static const RegistryEntry registry[] = {'] + for ph, ps, fh, fs, name, size, arch, digest in sorted(set(enabled)): + lines.append(f' {{0x{ph:016x}ULL, {ps}, 0x{fh:016x}ULL, {fs}, "{name}", {size}, "{arch}", "{digest}"}},') + lines += ['};', ''] + (out / 'registry.h').write_text('\n'.join(lines)) + print(f'{len(inventory)} PTX records inventoried; {len(enabled)} registry entries enabled') + + +def main(): + p = argparse.ArgumentParser(description=__doc__) + p.add_argument('dll', type=Path) + p.add_argument('native', type=Path) + p.add_argument('out', type=Path) + p.add_argument('names', nargs='*', default=['dltss_pwin_enc0_layer']) + a = p.parse_args() + generate(a.dll, a.native, a.out, a.names) + + +if __name__ == '__main__': + main() diff --git a/tools/microcuda/merge_pack.py b/tools/microcuda/merge_pack.py new file mode 100644 index 0000000..ea2a1e8 --- /dev/null +++ b/tools/microcuda/merge_pack.py @@ -0,0 +1,74 @@ +#!/usr/bin/env python3 +"""Per-module compile tuning: build a pack from several variant packs of the same DLL, ZLUDA build and target. + +merge_pack.py BASE_PACK OUT_PACK [--all VARIANT_PACK] [--choose CHOICES.json] + +BASE_PACK fixes the header (switch set the process must run with) and supplies every module not chosen +otherwise. --all takes every module from one variant pack (to measure that variant); --choose takes a JSON +map {fatbin_fnv_hex: variant_pack_dir}. Variants must only differ in code-generation choices that keep +results bit-exact (scheduler strategy, wave size); the caller validates the merged pack's output against the +reference before using it. Each module line keeps its own SHA256, and a comment records its variant. +""" +import argparse +import json +from pathlib import Path +import shutil +import sys + + +def read(pack): + header, modules = [], {} + for line in (pack / 'microcuda-pack.txt').read_text().splitlines(): + if line.startswith(('module ', 'empty ')): + modules[line.split()[1]] = line + else: + header.append(line) + return header, modules + + +def field(header, key): + return next((l.split(' ', 1)[1] if ' ' in l else '' for l in header if l.split(' ', 1)[0] == key), None) + + +def main(): + p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + p.add_argument('base', type=Path) + p.add_argument('out', type=Path) + p.add_argument('--all', type=Path) + p.add_argument('--choose', type=Path) + a = p.parse_args() + header, modules = read(a.base) + choice = {} + if a.all: + choice = {f: a.all for f in modules} + if a.choose: + choice.update({f: Path(v) for f, v in json.loads(a.choose.read_text()).items()}) + variants = {} + if a.out.exists(): + shutil.rmtree(a.out) + (a.out / 'modules').mkdir(parents=True) + lines, notes = [], [] + for f, line in sorted(modules.items()): + source = a.base + if f in choice and line.startswith('module '): + v = choice[f] + if v not in variants: + vh, vm = read(v) + if field(vh, 'arch') != field(header, 'arch') or \ + next(l for l in vh if l.startswith('# dll')) != next(l for l in header if l.startswith('# dll')): + sys.exit(f'{v}: different DLL or target') + variants[v] = (vh, vm) + vh, vm = variants[v] + if f in vm: + line, source = vm[f], v + notes.append(f'# variant {f} {field(vh, "switches")}') + if line.startswith('module '): + name = line.split()[6] + shutil.copyfile(source / 'modules' / name, a.out / 'modules' / name) + lines.append(line) + (a.out / 'microcuda-pack.txt').write_text('\n'.join(header + notes + lines) + '\n') + print(f'{len(lines)} entries, {len(notes)} from variant packs -> {a.out}') + + +if __name__ == '__main__': + main() diff --git a/tools/microcuda/runtime.cpp b/tools/microcuda/runtime.cpp new file mode 100644 index 0000000..dadffd0 --- /dev/null +++ b/tools/microcuda/runtime.cpp @@ -0,0 +1,786 @@ +// Experimental D4R-only Linux CUDA ABI over HIP, without ZLUDA's runtime or PTX translator. +// Modules come from an offline pack (tools/microcuda/make_pack.py: the code objects ZLUDA compiled +// for one exact DLSS DLL, ZLUDA build and switch set); anything outside it fails closed. Native +// kernel overrides, launch handling, arrays and texture/surface objects follow the patched ZLUDA +// runtime exactly (its module.rs, function.rs, memory.rs, array.rs, surf.rs) so that its code +// objects and d4r's native kernels see the same ABI. See ZLUDA_REPLACEMENT.md. +#define __HIP_PLATFORM_AMD__ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +extern char** environ; + +using CUresult = int; +using CUdeviceptr = uint64_t; +constexpr int SUCCESS=0, INVALID=1, OOM=2, NOT_INIT=3, INVALID_IMAGE=200, + INVALID_HANDLE=400, NOT_FOUND=500, NOT_READY=600, UNSUPPORTED=801, UNKNOWN=999; + +static int error(hipError_t e) { + // HIP errors are not globally interchangeable with CUDA driver errors. + switch (e) { + case hipSuccess: return SUCCESS; + case hipErrorInvalidValue: case hipErrorInvalidConfiguration: + case hipErrorInvalidDevicePointer: return INVALID; + case hipErrorOutOfMemory: return OOM; + case hipErrorNotInitialized: return NOT_INIT; + case hipErrorInvalidDevice: return 101; + case hipErrorNoDevice: return 100; + case hipErrorInvalidImage: return INVALID_IMAGE; + case hipErrorInvalidContext: return 201; + case hipErrorNoBinaryForGpu: return 209; + case hipErrorInvalidHandle: return INVALID_HANDLE; + case hipErrorNotFound: case hipErrorInvalidDeviceFunction: return NOT_FOUND; + case hipErrorNotReady: return NOT_READY; + case hipErrorNotSupported: return UNSUPPORTED; + case hipErrorIllegalAddress: return 700; + case hipErrorLaunchOutOfResources: return 701; + default: return UNKNOWN; + } +} +static std::atomic initialized{false}; +static uint64_t fnv(const unsigned char* p, size_t size) { + uint64_t h=0xcbf29ce484222325ULL; + for(size_t i=0;i allocations; +static uint64_t next_generation=1; +static std::atomic free_generation{0}; +static void track_alloc(uint64_t base,size_t bytes) { + std::lock_guard lock(allocations_lock); allocations[base]={bytes,next_generation++,0}; +} +static void track_free(uint64_t base) { + std::lock_guard lock(allocations_lock); allocations.erase(base); + free_generation.fetch_add(1,std::memory_order_release); +} +static void track_write(uint64_t dst,size_t bytes) { + if(!bytes) return; + std::lock_guard lock(allocations_lock); + auto it=allocations.upper_bound(dst); + if(it!=allocations.begin()) --it; + for(;it!=allocations.end()&&it->firstfirst+it->second.bytes>dst) ++it->second.epoch; +} +// -> false if p is not inside a tracked allocation +static bool identity_of(uint64_t p,uint64_t& generation,uint64_t& epoch) { + std::lock_guard lock(allocations_lock); + auto it=allocations.upper_bound(p); + if(it==allocations.begin()) return false; + --it; + if(p-it->first>=it->second.bytes) return false; + generation=it->second.generation; epoch=it->second.epoch; return true; +} + +extern "C" { +int cuInit(unsigned flags) { + if(flags) return INVALID; + int r=error(hipInit(0)); if(!r) initialized.store(true,std::memory_order_release); return r; +} +int cuDeviceGetCount(int* n) { return n?error(hipGetDeviceCount(n)):INVALID; } +int cuDeviceGet(int* d,int ordinal) { return d?error(hipDeviceGet(d,ordinal)):INVALID; } +// Contexts are bookkeeping over HIP's per-device state, as in ZLUDA (context.rs): no hipCtx API +// (deprecated in HIP), just hipSetDevice and a per-thread current-context stack. +struct Context { int device; }; +static thread_local std::vector context_stack; +static int make_current(Context* c) { return c?error(hipSetDevice(c->device)):SUCCESS; } +int cuCtxCreate_v2(Context** c,unsigned,int d) { + if(!c) return INVALID; + int count=0; if(int r=error(hipGetDeviceCount(&count))) return r; + if(d<0||d>=count) return 101; + *c=new Context{d}; + if(int r=make_current(*c)) { delete *c; *c=nullptr; return r; } + context_stack.push_back(*c); return SUCCESS; +} +int cuCtxDestroy_v2(Context* c) { + if(!c) return INVALID; + context_stack.erase(std::remove(context_stack.begin(),context_stack.end(),c),context_stack.end()); + delete c; return SUCCESS; +} +int cuCtxPushCurrent_v2(Context* c) { if(int r=make_current(c)) return r; context_stack.push_back(c); return SUCCESS; } +int cuCtxPopCurrent_v2(Context** c) { + if(context_stack.empty()) return 201; + if(c) *c=context_stack.back(); + context_stack.pop_back(); + return context_stack.empty()?SUCCESS:make_current(context_stack.back()); +} +int cuCtxSetCurrent(Context* c) { + if(int r=make_current(c)) return r; + if(context_stack.empty()) context_stack.push_back(c); else context_stack.back()=c; + if(!c) context_stack.pop_back(); + return SUCCESS; +} +int cuCtxGetCurrent(Context** c) { if(!c) return INVALID; *c=context_stack.empty()?nullptr:context_stack.back(); return SUCCESS; } +int cuCtxGetDevice(int* d) { + if(!d) return INVALID; + if(context_stack.empty()||!context_stack.back()) return 201; + *d=context_stack.back()->device; return SUCCESS; +} +int cuCtxSynchronize() { return error(hipDeviceSynchronize()); } +int cuDeviceGetName(char* p,int n,int d) { return p&&n>0?error(hipDeviceGetName(p,n,d)):INVALID; } +int cuDeviceGetUuid(hipUUID* uuid,int d) { return uuid?error(hipDeviceGetUuid(uuid,d)):INVALID; } +int cuDeviceGetLuid(char* luid,unsigned* mask,int d) { + if(!luid||!mask) return INVALID; + hipDeviceProp_t props{}; + int r=error(hipGetDeviceProperties(&props,d)); + if(!r) { std::memcpy(luid,props.luid,8); *mask=props.luidDeviceNodeMask; } + return r; +} +int cuDeviceTotalMem_v2(size_t* bytes,int d) { return bytes?error(hipDeviceTotalMem(bytes,d)):INVALID; } +int cuDeviceComputeCapability(int* major,int* minor,int d) { + int ignored; if(!major||!minor) return INVALID; + int r=error(hipDeviceGetAttribute(&ignored,hipDeviceAttributeMultiprocessorCount,d)); + if(!r) { *major=8; *minor=9; } return r; // same NGX sm_89 selection as D4R reference +} +int cuDriverGetVersion(int* v) { if(!v) return INVALID; *v=12080; return SUCCESS; } +int cuDeviceGetAttribute(int* out,int attribute,int d) { + if(!out) return INVALID; + if(attribute==75||attribute==76) { int ma,mi; int r=cuDeviceComputeCapability(&ma,&mi,d); + if(!r) *out=attribute==75?ma:mi; + return r; } + // Explicit CUDA->HIP enum map. Only the currently audited subset. + hipDeviceAttribute_t h; + switch(attribute) { + case 1: h=hipDeviceAttributeMaxThreadsPerBlock; break; + case 2: h=hipDeviceAttributeMaxBlockDimX; break; + case 3: h=hipDeviceAttributeMaxBlockDimY; break; + case 4: h=hipDeviceAttributeMaxBlockDimZ; break; + case 5: h=hipDeviceAttributeMaxGridDimX; break; + case 6: h=hipDeviceAttributeMaxGridDimY; break; + case 7: h=hipDeviceAttributeMaxGridDimZ; break; + case 8: h=hipDeviceAttributeMaxSharedMemoryPerBlock; break; + case 10: { int v; int r=error(hipDeviceGetAttribute(&v,hipDeviceAttributeWarpSize,d)); + if(!r) *out=32; + return r; } + case 13: h=hipDeviceAttributeClockRate; break; + case 16: h=hipDeviceAttributeMultiprocessorCount; break; + case 22: h=hipDeviceAttributeMaxTexture2DWidth; break; + case 23: h=hipDeviceAttributeMaxTexture2DHeight; break; + case 35: { int ignored; int r=error(hipDeviceGetAttribute(&ignored, + hipDeviceAttributeMultiprocessorCount,d)); if(!r) *out=0; return r; } // TCC_DRIVER + default: return UNSUPPORTED; + } + return error(hipDeviceGetAttribute(out,h,d)); +} +// cuMemAlloc zero-fills like ZLUDA (memory.rs alloc_v2): NGX output must not depend on stale VRAM. +int cuMemAlloc_v2(CUdeviceptr* p,size_t bytes) { + if(!p) return INVALID; + void* raw=nullptr; int r=error(hipMalloc(&raw,bytes)); if(r) return r; + if((r=error(hipMemset(raw,0,bytes)))) { auto ignored=hipFree(raw); (void)ignored; return r; } + *p=reinterpret_cast(raw); track_alloc(*p,bytes); return SUCCESS; +} +int cuMemFree_v2(CUdeviceptr p) { track_free(p); return error(hipFree(reinterpret_cast(p))); } +int cuMemHostAlloc(void** p,size_t bytes,unsigned flags) { return p?error(hipHostMalloc(p,bytes,flags)):INVALID; } +int cuMemAllocHost_v2(void** p,size_t bytes) { return cuMemHostAlloc(p,bytes,0); } +int cuMemFreeHost(void* p) { return error(hipHostFree(p)); } +int cuMemGetInfo_v2(size_t* free,size_t* total) { return free&&total?error(hipMemGetInfo(free,total)):INVALID; } +int cuMemcpyHtoD_v2(CUdeviceptr d,const void* s,size_t n) { track_write(d,n); return error(hipMemcpyHtoD(reinterpret_cast(d),const_cast(s),n)); } +int cuMemcpyDtoH_v2(void* d,CUdeviceptr s,size_t n) { return error(hipMemcpyDtoH(d,reinterpret_cast(s),n)); } +int cuMemcpyDtoD_v2(CUdeviceptr d,CUdeviceptr s,size_t n) { track_write(d,n); return error(hipMemcpyDtoD(reinterpret_cast(d),reinterpret_cast(s),n)); } +int cuMemcpyHtoDAsync_v2(CUdeviceptr d,const void* s,size_t n,hipStream_t st) { track_write(d,n); return error(hipMemcpyHtoDAsync(reinterpret_cast(d),const_cast(s),n,st)); } +int cuMemcpyDtoHAsync_v2(void* d,CUdeviceptr s,size_t n,hipStream_t st) { return error(hipMemcpyDtoHAsync(d,reinterpret_cast(s),n,st)); } +int cuMemcpyDtoDAsync_v2(CUdeviceptr d,CUdeviceptr s,size_t n,hipStream_t st) { track_write(d,n); return error(hipMemcpyDtoDAsync(reinterpret_cast(d),reinterpret_cast(s),n,st)); } +int cuMemsetD8_v2(CUdeviceptr d,unsigned char v,size_t n) { track_write(d,n); return error(hipMemset(reinterpret_cast(d),v,n)); } +int cuEventCreate(hipEvent_t* e,unsigned flags) { return e?error(hipEventCreateWithFlags(e,flags)):INVALID; } +int cuEventDestroy_v2(hipEvent_t e) { return error(hipEventDestroy(e)); } +int cuEventRecord(hipEvent_t e,hipStream_t s) { return error(hipEventRecord(e,s)); } +int cuEventQuery(hipEvent_t e) { return error(hipEventQuery(e)); } +int cuEventSynchronize(hipEvent_t e) { return error(hipEventSynchronize(e)); } +int cuEventElapsedTime(float* ms,hipEvent_t a,hipEvent_t b) { return ms?error(hipEventElapsedTime(ms,a,b)):INVALID; } +int cuStreamCreate(hipStream_t* s,unsigned flags) { return s?error(hipStreamCreateWithFlags(s,flags)):INVALID; } +int cuStreamDestroy_v2(hipStream_t s) { return error(hipStreamDestroy(s)); } +int cuStreamSynchronize(hipStream_t s) { return error(hipStreamSynchronize(s)); } +int cuStreamQuery(hipStream_t s) { return error(hipStreamQuery(s)); } +int cuStreamWaitEvent(hipStream_t s,hipEvent_t e,unsigned flags) { return error(hipStreamWaitEvent(s,e,flags)); } +// Unimplemented in ZLUDA, which returns NOT_SUPPORTED; the bridge falls back to its own records on that +// error (descriptors it stored at creation, linear-texture stand-ins), so these must fail the same way. +int cuDestroyExternalMemory(void*) { return UNSUPPORTED; } +int cuMipmappedArrayDestroy(void*) { return UNSUPPORTED; } +int cuGetErrorString(int e,const char** p) { + if(!p) return INVALID; + switch(e) { + case 0: *p="success"; break; case 1: *p="invalid value"; break; + case 200: *p="module outside the MicroCUDA pack, or a pack/native file that fails verification"; break; + case 801: *p="outside the MicroCUDA subset"; break; + default: *p="MicroCUDA error (see numeric CUDA driver code)"; + } return SUCCESS; +} +} + +// ---- arrays, 2D copies, texture and surface objects (ZLUDA array.rs, memory.rs, surf.rs) ---- + +// CUDA formats HIP's legacy descriptor lacks get equivalent backing storage; the original format is +// remembered per array for surface descriptors and packed copies. +static std::mutex resources_lock; +static std::unordered_map original_formats; +static uint32_t original_format(hipArray_t a) { + std::lock_guard lock(resources_lock); + auto it=original_formats.find(a); return it==original_formats.end()?0:it->second; +} +extern "C" int cuArrayCreate_v2(hipArray_t* handle,const HIP_ARRAY_DESCRIPTOR* desc) { + if(!handle||!desc) return INVALID; + int format,channels; + switch(static_cast(desc->Format)) { + case 80: if(desc->NumChannels!=4) return error(hipArrayCreate(handle,desc)); + format=32; channels=4; break; // UNORM 10:10:10:2 -> RGBA32F backing + case 192: format=1; channels=1; break; case 193: format=1; channels=2; break; // UNORM_INT8 + case 194: format=1; channels=4; break; + case 195: format=2; channels=1; break; case 196: format=2; channels=2; break; // UNORM_INT16 + case 197: format=2; channels=4; break; + case 198: format=8; channels=1; break; case 199: format=8; channels=2; break; // SNORM_INT8 + case 200: format=8; channels=4; break; + case 201: format=9; channels=1; break; case 202: format=9; channels=2; break; // SNORM_INT16 + case 203: format=9; channels=4; break; + default: return error(hipArrayCreate(handle,desc)); + } + HIP_ARRAY_DESCRIPTOR backing=*desc; + backing.Format=static_cast(format); backing.NumChannels=channels; + int r=error(hipArrayCreate(handle,&backing)); if(r) return r; + std::lock_guard lock(resources_lock); + original_formats[*handle]=static_cast(desc->Format); return SUCCESS; +} +extern "C" int cuArray3DCreate_v2(hipArray_t* handle,const HIP_ARRAY3D_DESCRIPTOR* desc) { + return handle&&desc?error(hipArray3DCreate(handle,desc)):INVALID; +} +extern "C" int cuArrayDestroy(hipArray_t a) { + int r=error(hipArrayDestroy(a)); if(r) return r; + std::lock_guard lock(resources_lock); original_formats.erase(a); return SUCCESS; +} +extern "C" int cuArrayGetDescriptor_v2(void*,hipArray_t) { return UNSUPPORTED; } // as ZLUDA (see above) + +static uint32_t packed_component(float v,uint32_t scale) { + float c=std::isnan(v)?0.f:std::clamp(v,0.f,1.f); + return static_cast(std::floor(c*static_cast(scale)+0.5f)); +} +// Packed 10:10:10:2 pixels between public 4-byte words and the RGBA32F backing array. +static int copy_2d_packed(const hip_Memcpy2D& m,bool src_packed,bool dst_packed) { + if(m.WidthInBytes%4||(src_packed&&m.srcXInBytes%4)||(dst_packed&&m.dstXInBytes%4)) return INVALID; + const size_t backing_width=m.WidthInBytes*4; + if(src_packed&&dst_packed) { + hip_Memcpy2D b=m; b.srcXInBytes*=4; b.dstXInBytes*=4; b.WidthInBytes=backing_width; + return error(hipMemcpyParam2D(&b)); + } + if((src_packed&&m.dstMemoryType==hipMemoryTypeArray)||(dst_packed&&m.srcMemoryType==hipMemoryTypeArray)) + return INVALID; + const size_t width=m.WidthInBytes; + std::vector backing(backing_width/4*m.Height); + std::vector packed(width*m.Height); + if(dst_packed) { + // external (host or device) packed words -> host staging + if(m.srcMemoryType==hipMemoryTypeHost) { + if(!m.srcHost||m.srcXInBytes+width>m.srcPitch) return INVALID; + for(size_t row=0;row(m.srcHost)+(m.srcY+row)*m.srcPitch+m.srcXInBytes,width); + } else if(m.srcMemoryType==hipMemoryTypeDevice||m.srcMemoryType==hipMemoryTypeUnified) { + hip_Memcpy2D s=m; s.dstMemoryType=hipMemoryTypeHost; s.dstHost=packed.data(); + s.dstXInBytes=0; s.dstY=0; s.dstPitch=width; + if(int r=error(hipMemcpyParam2D(&s))) return r; + } else return UNSUPPORTED; + for(size_t p=0;p(bits&1023)/1023.f; + backing[4*p+1]=static_cast((bits>>10)&1023)/1023.f; + backing[4*p+2]=static_cast((bits>>20)&1023)/1023.f; + backing[4*p+3]=static_cast(bits>>30)/3.f; + } + hip_Memcpy2D t=m; t.srcMemoryType=hipMemoryTypeHost; t.srcHost=backing.data(); + t.srcXInBytes=0; t.srcY=0; t.srcPitch=backing_width; t.dstXInBytes=m.dstXInBytes*4; t.WidthInBytes=backing_width; + return error(hipMemcpyParam2D(&t)); + } + hip_Memcpy2D t=m; t.srcXInBytes=m.srcXInBytes*4; t.dstMemoryType=hipMemoryTypeHost; t.dstHost=backing.data(); + t.dstXInBytes=0; t.dstY=0; t.dstPitch=backing_width; t.WidthInBytes=backing_width; + if(int r=error(hipMemcpyParam2D(&t))) return r; + for(size_t p=0;pm.dstPitch) return INVALID; + for(size_t row=0;row(m.dstHost)+(m.dstY+row)*m.dstPitch+m.dstXInBytes,packed.data()+row*width,width); + return SUCCESS; + } + if(m.dstMemoryType!=hipMemoryTypeDevice&&m.dstMemoryType!=hipMemoryTypeUnified) return UNSUPPORTED; + hip_Memcpy2D s=m; s.srcMemoryType=hipMemoryTypeHost; s.srcHost=packed.data(); + s.srcXInBytes=0; s.srcY=0; s.srcPitch=width; + return error(hipMemcpyParam2D(&s)); +} +// CUDA_MEMCPY2D has hip_Memcpy2D's layout, but CUDA numbers memory types HOST 1, DEVICE 2, ARRAY 3, +// UNIFIED 4 where HIP uses 1, 2, 10, 11. +static bool memory_type(int cuda,hipMemoryType& hip) { + switch(cuda) { + case 1: hip=hipMemoryTypeHost; return true; case 2: hip=hipMemoryTypeDevice; return true; + case 3: hip=hipMemoryTypeArray; return true; case 4: hip=hipMemoryTypeUnified; return true; + default: return false; + } +} +extern "C" int cuMemcpy2D_v2(const hip_Memcpy2D* cuda) { + if(!cuda) return INVALID; + hip_Memcpy2D m=*cuda; + if(!memory_type(static_cast(cuda->srcMemoryType),m.srcMemoryType)|| + !memory_type(static_cast(cuda->dstMemoryType),m.dstMemoryType)) return UNSUPPORTED; + if(m.dstMemoryType==hipMemoryTypeDevice||m.dstMemoryType==hipMemoryTypeUnified) + track_write(reinterpret_cast(m.dstDevice)+m.dstY*m.dstPitch+m.dstXInBytes, + m.Height?(m.Height-1)*m.dstPitch+m.WidthInBytes:0); + bool src_packed=m.srcMemoryType==hipMemoryTypeArray&&original_format(m.srcArray)==80; + bool dst_packed=m.dstMemoryType==hipMemoryTypeArray&&original_format(m.dstArray)==80; + if(!src_packed&&!dst_packed) return error(hipMemcpyParam2D(&m)); + return copy_2d_packed(m,src_packed,dst_packed); +} + +extern "C" int cuTexObjectCreate(hipTextureObject_t* object,const HIP_RESOURCE_DESC* resource, + const HIP_TEXTURE_DESC* texture,const HIP_RESOURCE_VIEW_DESC* view) { + return object&&resource&&texture?error(hipTexObjectCreate(object,resource,texture,view)):INVALID; +} +extern "C" int cuTexObjectDestroy(hipTextureObject_t object) { return error(hipTexObjectDestroy(object)); } +extern "C" int cuTexObjectGetResourceDesc(void*,hipTextureObject_t) { return UNSUPPORTED; } // as ZLUDA + +// A surface handle is device memory: a copy of HIP's 48-byte image object (so the handle is also a +// valid texture object for texel fetches), a zeroed 32-byte sampler slot, then the CUDA format word +// and padding the bridge's output redirect writes into (bytes 84..95). +constexpr size_t IMAGE_OBJECT_BYTES=48, SAMPLER_OBJECT_BYTES=32, + FORMAT_OFFSET=IMAGE_OBJECT_BYTES+SAMPLER_OBJECT_BYTES, SURFACE_HANDLE_BYTES=FORMAT_OFFSET+16; +struct Surface { hipSurfaceObject_t native; HIP_RESOURCE_DESC desc; }; +static std::unordered_map surfaces; +extern "C" int cuSurfObjectCreate(uint64_t* object,const HIP_RESOURCE_DESC* resource) { + if(!object||!resource) return INVALID; + uint32_t format=resource->resType==HIP_RESOURCE_TYPE_ARRAY?original_format(resource->res.array.hArray):0; + hipSurfaceObject_t native{}; + int r=error(hipCreateSurfaceObject(&native,reinterpret_cast(resource))); if(r) return r; + void* handle=nullptr; + if((r=error(hipMalloc(&handle,SURFACE_HANDLE_BYTES)))) { auto ignored=hipDestroySurfaceObject(native); (void)ignored; return r; } + uint32_t tail[(SURFACE_HANDLE_BYTES-IMAGE_OBJECT_BYTES)/4]{}; + tail[(FORMAT_OFFSET-IMAGE_OBJECT_BYTES)/4]=format; + r=error(hipMemcpyDtoD(handle,reinterpret_cast(native),IMAGE_OBJECT_BYTES)); + if(!r) r=error(hipMemcpyHtoD(static_cast(handle)+IMAGE_OBJECT_BYTES,tail,sizeof(tail))); + if(r) { auto a=hipFree(handle); auto b=hipDestroySurfaceObject(native); (void)a; (void)b; return r; } + std::lock_guard lock(resources_lock); + surfaces[reinterpret_cast(handle)]={native,*resource}; + *object=reinterpret_cast(handle); return SUCCESS; +} +extern "C" int cuSurfObjectDestroy(uint64_t object) { + Surface s; + { std::lock_guard lock(resources_lock); + auto it=surfaces.find(object); if(it==surfaces.end()) return INVALID; s=it->second; surfaces.erase(it); } + int native=error(hipDestroySurfaceObject(s.native)); + int freed=error(hipFree(reinterpret_cast(object))); + return native?native:freed; +} +extern "C" int cuSurfObjectGetResourceDesc(HIP_RESOURCE_DESC* desc,uint64_t object) { + if(!desc) return INVALID; + std::lock_guard lock(resources_lock); + auto it=surfaces.find(object); if(it==surfaces.end()) return INVALID; + *desc=it->second.desc; return SUCCESS; +} + +// ---- module pack ---- + +struct PackEntry { uint64_t fatbin_hash; size_t fatbin_bytes; uint64_t ptx_hash; size_t ptx_bytes; + std::string sha256, file; bool empty=false; }; +struct Pack { bool loaded=false; int status=NOT_FOUND; std::string dir, arch, switches; int fp8native=0; + std::vector entries; }; +static std::mutex pack_lock; +static Pack pack; + +// The code-generation switch set ZLUDA would key its cache with (module.rs zluda_build_version), +// minus settings that do not change generated code. Pack and process must agree. +static std::string process_switches() { + static const char* ignored[]={"D4R_ZLUDA_NATIVE_DIR","D4R_ZLUDA_DIR","D4R_ZLUDA_DUMP_DIR", + "D4R_ZLUDA_DUMP_PTX_DIR","D4R_ZLUDA_LIBCUDA","D4R_ZLUDA_WMMA_FP8_NATIVE","D4R_ZLUDA_CACHE_HOME"}; + std::vector out; + for(char** e=environ;*e;++e) { + std::string kv=*e; if(kv.rfind("D4R_ZLUDA_",0)) continue; + size_t eq=kv.find('='); if(eq==std::string::npos) continue; + std::string key=kv.substr(0,eq), value=kv.substr(eq+1); + if(std::any_of(std::begin(ignored),std::end(ignored),[&](const char* i) { return key==i; })) continue; + size_t b=value.find_first_not_of(" \t\r\n"), e2=value.find_last_not_of(" \t\r\n"); + out.push_back(key+"="+(b==std::string::npos?"":value.substr(b,e2-b+1))); + } + std::sort(out.begin(),out.end()); + std::string joined; for(auto& s:out) { if(!joined.empty()) joined+=','; joined+=s; } + return joined; +} +static int load_pack() { + std::lock_guard lock(pack_lock); + if(pack.loaded) return pack.status; + pack.loaded=true; + const char* dir=std::getenv("D4R_MICROCUDA_PACK"); + if(!dir||!*dir) { note("%s","D4R_MICROCUDA_PACK is not set; no module can be loaded"); return pack.status=NOT_FOUND; } + pack.dir=dir; + FILE* f=std::fopen((pack.dir+"/microcuda-pack.txt").c_str(),"r"); + if(!f) { + // a directory of packs, one per GPU target (/gfx1101/microcuda-pack.txt, ...): take this GPU's + hipDeviceProp_t props{}; int device=0; + if(hipGetDevice(&device)==hipSuccess&&hipGetDeviceProperties(&props,device)==hipSuccess) { + std::string arch=props.gcnArchName; arch=arch.substr(0,arch.find(':')); + pack.dir=std::string(dir)+"/"+arch; + f=std::fopen((pack.dir+"/microcuda-pack.txt").c_str(),"r"); + } + } + if(!f) { note("cannot read %s/microcuda-pack.txt",dir); return pack.status=NOT_FOUND; } + char line[1024]; + while(std::fgets(line,sizeof(line),f)) { + std::string s=line; while(!s.empty()&&(s.back()=='\n'||s.back()=='\r')) s.pop_back(); + if(s.empty()||s[0]=='#') continue; + if(s.rfind("arch ",0)==0) { pack.arch=s.substr(5); continue; } + if(s.rfind("switches",0)==0) { pack.switches=s.size()>9?s.substr(9):""; continue; } + if(s.rfind("fp8native ",0)==0) { pack.fp8native=std::atoi(s.c_str()+10); continue; } + PackEntry e; char sha[65], file[256]; unsigned long long fh,ph; size_t fb,pb; + if(std::sscanf(line,"empty %llx %zu",&fh,&fb)==2) { + e.fatbin_hash=fh; e.fatbin_bytes=fb; e.ptx_hash=0; e.ptx_bytes=SIZE_MAX; e.empty=true; + pack.entries.push_back(e); continue; + } + if(std::sscanf(line,"module %llx %zu %llx %zu %64s %255s",&fh,&fb,&ph,&pb,sha,file)!=6) continue; + e.fatbin_hash=fh; e.fatbin_bytes=fb; e.ptx_hash=ph; e.ptx_bytes=pb; e.sha256=sha; e.file=file; + pack.entries.push_back(e); + } + std::fclose(f); + std::string current=process_switches(); + if(current!=pack.switches) { + note("pack was compiled with ZLUDA switches '%s', this process has '%s'; refusing the pack", + pack.switches.c_str(),current.c_str()); + return pack.status=INVALID_IMAGE; + } + return pack.status=pack.entries.empty()?NOT_FOUND:SUCCESS; +} +// RDNA4 objects differ with native FP8 WMMA (ZLUDA keys that by effect: gfx12 with WMMA and NativeFp8 on). +static bool fp8_matches(const std::string& arch) { + auto on=[](const char* k) { const char* v=std::getenv(k); return v&&std::strcmp(v,"1")==0; }; + bool expected=arch.rfind("gfx12",0)==0&&on("D4R_ZLUDA_WMMA")&&on("D4R_ZLUDA_WMMA_FP8_NATIVE"); + return expected==(pack.fp8native!=0); +} +static bool read_file(const std::string& path,std::vector& bytes) { + FILE* file=std::fopen(path.c_str(),"rb"); if(!file) return false; + unsigned char buffer[65536]; size_t count; + while((count=std::fread(buffer,1,sizeof(buffer),file))) bytes.insert(bytes.end(),buffer,buffer+count); + bool failed=std::ferror(file); std::fclose(file); return !failed; +} +static bool digest_matches(const std::vector& data,const std::string& expected) { + unsigned char digest[SHA256_DIGEST_LENGTH]; SHA256(data.data(),data.size(),digest); + char hex[65]; for(size_t i=0;i<32;++i) std::sprintf(hex+2*i,"%02x",digest[i]); + return expected==hex; +} + +// ---- modules and functions (ZLUDA module.rs / function.rs) ---- + +struct Post { hipFunction_t function; uint32_t gx, gy, bx, by, bz; }; +struct Function { + hipFunction_t main=nullptr, prep=nullptr; + uint32_t prep_blocks=0, block_z=0, grid_x=0, prep_slots=0; + bool has_prep_key=false; uint32_t prep_key=0; + std::atomic prep_last{0}; + std::mutex seen_lock; std::vector seen; + std::vector post; + bool native=false; + std::atomic prep_runs{0}, prep_skips{0}; + // identity policy (d4r_prep_identity_at): prepared images per weights pointer, valid while the pointer's + // allocation generation and write epoch are unchanged and no allocation has been freed since the slot + // table was last reset. Only on one stream; another stream turns the policy off for this function. + uint32_t identity_at=0; + hipDeviceptr_t prep_keys=nullptr; size_t prep_keys_bytes=0; + std::mutex identity_lock; + struct Prepared { uint64_t key, generation, epoch; }; + std::vector prepared; + uint64_t prepared_free_generation=0; + hipStream_t identity_stream=nullptr; bool identity_stream_set=false, identity_off=false; +}; +struct NativeImage { hipModule_t module; std::vector bytes; }; +struct Module { + hipModule_t hip=nullptr; + std::vector image; // retained until HIP module unload + std::mutex lock; + std::map> functions; + std::vector> natives; +}; +static std::mutex modules_lock; +static std::vector modules; +static bool registered(Module* m) { return std::find(modules.begin(),modules.end(),m)!=modules.end(); } + +static bool global_u32(hipModule_t m,const char* name,uint32_t& value) { + hipDeviceptr_t ptr=nullptr; size_t bytes=0; + return hipModuleGetGlobal(&ptr,&bytes,m,name)==hipSuccess&&bytes==4&&hipMemcpyDtoH(&value,ptr,4)==hipSuccess; +} +// d4r native replacement served by the bridge as DIR/NAME.hsaco (only for PTX it verified); same +// metadata contract as ZLUDA's native_override. Returns false to use the pack's compile of NAME. +static bool native_override(Module* m,const std::string& name,Function& f) { + const char* dir=std::getenv("D4R_MICROCUDA_NATIVE_DIR"); if(!dir||!*dir) return false; + auto image=std::make_unique(); + std::string path=std::string(dir)+"/"+name+".hsaco"; + if(!read_file(path,image->bytes)||image->bytes.empty()) return false; + if(hipModuleLoadData(&image->module,image->bytes.data())!=hipSuccess) { note("failed to load %s",path.c_str()); return false; } + hipModule_t nm=image->module; + auto reject=[&](const char* why) { note("%s: %s; using the pack's kernel",path.c_str(),why); + auto ignored=hipModuleUnload(nm); (void)ignored; return false; }; + if(hipModuleGetFunction(&f.main,nm,name.c_str())!=hipSuccess) return reject("no kernel of that name"); + if(hipModuleGetFunction(&f.prep,nm,(name+"_prep").c_str())==hipSuccess) { + if(!global_u32(nm,"d4r_prep_blocks",f.prep_blocks)||!f.prep_blocks) return reject("prep kernel without d4r_prep_blocks"); + } else f.prep=nullptr; + global_u32(nm,"d4r_block_z",f.block_z); + global_u32(nm,"d4r_grid_x",f.grid_x); + // d4r_prep_key_at holds offset + 1; the older d4r_prep_key_offset uses 0 for "no key" + uint32_t at=0, offset=0; + if(global_u32(nm,"d4r_prep_key_at",at)&&at>0) { f.has_prep_key=true; f.prep_key=at-1; } + else if(global_u32(nm,"d4r_prep_key_offset",offset)&&offset) { f.has_prep_key=true; f.prep_key=offset; } + global_u32(nm,"d4r_prep_key_slots",f.prep_slots); + if(f.prep&&global_u32(nm,"d4r_prep_identity_at",f.identity_at)&&f.identity_at) { + // the slot table the prep claims into (keys only ever go from 0 to a pointer until reset) + if(hipModuleGetGlobal(&f.prep_keys,&f.prep_keys_bytes,nm,"g_prep_keys")!=hipSuccess|| + f.prep_keys_bytes!=(f.prep_slots+1)*sizeof(uint64_t)||!f.prep_slots) { + note("%s: d4r_prep_identity_at without a matching g_prep_keys slot table; preparing every launch",path.c_str()); + f.identity_at=0; + } + } + for(unsigned k=1;k<=8;++k) { + Post p{}; + if(hipModuleGetFunction(&p.function,nm,(name+"_post"+std::to_string(k)).c_str())!=hipSuccess) break; + auto field=[&](const char* s) { uint32_t v=0; global_u32(nm,("d4r_post"+std::to_string(k)+"_"+s).c_str(),v); return v; }; + p.gx=field("grid_x"); p.gy=field("grid_y"); + p.bx=field("block_x"); p.by=field("block_y"); p.bz=field("block_z"); + if(!p.bx) p.bx=32; + if(!p.by) p.by=1; + if(!p.bz) p.bz=1; + f.post.push_back(p); + } + f.native=true; + m->natives.push_back(std::move(image)); + return true; +} + +static int load_module(Module** out,const void* image) { + if(!out||!image) return INVALID; + *out=nullptr; + if(!initialized.load(std::memory_order_acquire)) return NOT_INIT; + if(int r=load_pack()) return r; + const auto* p=static_cast(image); + uint32_t magic; std::memcpy(&magic,p,4); + const PackEntry* entry=nullptr; + if(magic==0xba55ed50U) { + uint16_t version,header; uint64_t payload; + std::memcpy(&version,p+4,2); std::memcpy(&header,p+6,2); std::memcpy(&payload,p+8,8); + if(version!=1||header!=16||payload>256*1024*1024) return INVALID_IMAGE; + size_t n=header+payload; uint64_t h=0; bool hashed=false; + for(const auto& e:pack.entries) if(n==e.fatbin_bytes) { + if(!hashed) { h=fnv(p,n); hashed=true; } + if(h==e.fatbin_hash) { entry=&e; break; } + } + } else { + // PTX text: identity of the text without trailing NULs, as the bridge's native manifest + size_t n=strnlen(reinterpret_cast(p),256*1024*1024); uint64_t h=fnv(p,n); + for(const auto& e:pack.entries) if(n==e.ptx_bytes&&h==e.ptx_hash) { entry=&e; break; } + } + if(!entry) return INVALID_IMAGE; + hipDeviceProp_t props{}; int device; + if(int r=error(hipGetDevice(&device))) return r; + if(int r=error(hipGetDeviceProperties(&props,device))) return r; + std::string arch=props.gcnArchName; arch=arch.substr(0,arch.find(':')); + if(arch!=pack.arch) return 209; + if(!fp8_matches(arch)) { note("%s pack was built with NativeFp8=%s; the process setting differs",arch.c_str(), + pack.fp8native?"1":"0"); return INVALID_IMAGE; } + auto m=std::make_unique(); + if(entry->empty) { // fatbin without PTX: a module without functions, as ZLUDA loads it + std::lock_guard lock(modules_lock); modules.push_back(m.get()); *out=m.release(); return SUCCESS; + } + if(!read_file(pack.dir+"/modules/"+entry->file,m->image)||!digest_matches(m->image,entry->sha256)) { + note("%s/modules/%s is missing or fails its digest",pack.dir.c_str(),entry->file.c_str()); + return INVALID_IMAGE; + } + if(int r=error(hipModuleLoadData(&m->hip,m->image.data()))) return r; + std::lock_guard lock(modules_lock); modules.push_back(m.get()); *out=m.release(); return SUCCESS; +} +extern "C" int cuModuleLoadData(Module** out,const void* image) { return load_module(out,image); } +extern "C" int cuModuleLoadDataEx(Module** m,const void* image,unsigned,int*,void**) { + return load_module(m,image); // JIT options are irrelevant to precompiled objects (ZLUDA ignores them too) +} +extern "C" int cuModuleLoad(Module** m,const char* path) { + if(!m||!path) return INVALID; + std::vector bytes; if(!read_file(path,bytes)) return INVALID; + bytes.push_back(0); + return load_module(m,bytes.data()); +} +extern "C" int cuModuleGetFunction(Function** out,Module* m,const char* name) { + if(!out||!name) return INVALID; + *out=nullptr; + { std::lock_guard lock(modules_lock); if(!registered(m)) return INVALID_HANDLE; } + if(!m->hip) return NOT_FOUND; // empty module + std::lock_guard lock(m->lock); + auto& slot=m->functions[name]; + if(!slot) { + auto f=std::make_unique(); + if(!native_override(m,name,*f)) { + f=std::make_unique(); + if(int r=error(hipModuleGetFunction(&f->main,m->hip,name))) { m->functions.erase(name); return r; } + } + slot=std::move(f); + } + *out=slot.get(); return SUCCESS; +} +extern "C" int cuModuleUnload(Module* m) { + { std::lock_guard lock(modules_lock); + if(!registered(m)) return INVALID_HANDLE; + modules.erase(std::find(modules.begin(),modules.end(),m)); } + static const bool stats=std::getenv("D4R_MICROCUDA_STATS")!=nullptr; + if(stats) for(auto& [name,f]:m->functions) if(f->prep) + std::fprintf(stderr,"[d4r-microcuda] %s: prep runs %llu, skips %llu%s\n",name.c_str(), + (unsigned long long)f->prep_runs.load(),(unsigned long long)f->prep_skips.load(), + f->identity_at?(f->identity_off?" (identity reuse, turned off: several streams)":" (identity reuse)"):""); + int r=m->hip?error(hipModuleUnload(m->hip)):SUCCESS; + for(auto& n:m->natives) { int nr=error(hipModuleUnload(n->module)); if(!r) r=nr; } + delete m; return r; +} +extern "C" int d4rMicrocudaGetDispatchStats(Function* f,uint64_t* prep,uint64_t* skips) { + if(!f||!prep||!skips) return INVALID; + *prep=f->prep_runs.load(std::memory_order_relaxed); + *skips=f->prep_skips.load(std::memory_order_relaxed); return SUCCESS; +} + +// Prep under the identity policy; launch(h, grid..., block..., shared) submits on the caller's stream. +template +static int identity_prep(Function* f,const unsigned char* block,hipStream_t stream,Launch& launch) { + std::lock_guard lock(f->identity_lock); + auto prep=[&] { + int r=launch(f->prep,f->prep_blocks,1,1,128,1,1,0); + if(!r) f->prep_runs.fetch_add(1,std::memory_order_relaxed); + return r; + }; + if(f->identity_off||!block) return prep(); + if(!f->identity_stream_set) { f->identity_stream=stream; f->identity_stream_set=true; } + else if(f->identity_stream!=stream) { + // a second stream: slot resets could race with the other stream's reads. Prepare every launch from now on. + f->identity_off=true; f->prepared.clear(); return prep(); + } + uint64_t key=0, generation=0, epoch=0; + std::memcpy(&key,block+f->identity_at-1,sizeof(key)); + const uint64_t frees=free_generation.load(std::memory_order_acquire); + if(frees!=f->prepared_free_generation) { + // an allocation was freed (feature release/recreation): addresses may be recycled. Forget every image + // and clear the device slot table on this stream, so no key can match an old image. + if(int r=error(hipMemsetAsync(f->prep_keys,0,f->prep_keys_bytes,stream))) return r; + f->prepared.clear(); f->prepared_free_generation=frees; + } + if(!key||!identity_of(key,generation,epoch)) return prep(); // untracked weights: no reuse + for(auto& p:f->prepared) if(p.key==key) { + if(p.generation==generation&&p.epoch==epoch) { f->prep_skips.fetch_add(1,std::memory_order_relaxed); return SUCCESS; } + p.generation=generation; p.epoch=epoch; return prep(); // same slot (claimed by key), new contents + } + // a new key claims the next free slot; beyond the slots it shares the overflow slot and prepares every launch + if(f->prepared.size()prep_slots) f->prepared.push_back({key,generation,epoch}); + return prep(); +} + +// CUDA's packed-argument tokens are POINTER=1, SIZE=2, END=0; HIP's END is 3. +extern "C" int cuLaunchKernel(Function* f,unsigned gx,unsigned gy,unsigned gz, + unsigned bx,unsigned by,unsigned bz,unsigned shared,hipStream_t stream,void** params,void** extra) { + if(!f) return INVALID_HANDLE; + void* translated[8]{}; void** hip_extra=nullptr; const unsigned char* block=nullptr; + if(extra) { + size_t i=0, o=0; bool ended=false; + while(i<8&&o<8) { + uintptr_t token=reinterpret_cast(extra[i]); + if(token==0) { translated[o]=reinterpret_cast(uintptr_t{3}); ended=true; break; } + if(token!=1&&token!=2) return UNSUPPORTED; + if(i+1>=8||o+1>=8) return INVALID; + if(token==1) block=static_cast(extra[i+1]); + translated[o]=extra[i]; translated[o+1]=extra[i+1]; i+=2; o+=2; + } + if(!ended) return INVALID; + hip_extra=translated; + } else if(params) block=static_cast(params[0]); // single-struct kernels + auto launch=[&](hipFunction_t h,unsigned x,unsigned y,unsigned z,unsigned tx,unsigned ty,unsigned tz,unsigned sm) { + return error(hipModuleLaunchKernel(h,x,y,z,tx,ty,tz,sm,stream,params,hip_extra)); + }; + if(f->prep&&f->identity_at) { + if(int r=identity_prep(f,block,stream,launch)) return r; + } else if(f->prep) { + uint64_t key=0; + if(f->has_prep_key&&block) std::memcpy(&key,block+f->prep_key,sizeof(key)); + bool run=true; + if(key) { + if(!f->prep_slots) run=f->prep_last.load(std::memory_order_relaxed)!=key; + else { + // one prepared image per key while slots last; keys beyond them prepare every launch + std::lock_guard lock(f->seen_lock); + if(std::find(f->seen.begin(),f->seen.end(),key)!=f->seen.end()) run=false; + else if(f->seen.size()prep_slots) f->seen.push_back(key); + } + } + if(run) { + if(int r=launch(f->prep,f->prep_blocks,1,1,128,1,1,0)) return r; + f->prep_last.store(key,std::memory_order_relaxed); + f->prep_runs.fetch_add(1,std::memory_order_relaxed); + } else f->prep_skips.fetch_add(1,std::memory_order_relaxed); + } + if(f->block_z) bz=f->block_z; + if(f->grid_x) { gx=f->grid_x; gy=1; gz=1; } + if(int r=launch(f->main,gx,gy,gz,bx,by,bz,shared)) return r; + for(const auto& p:f->post) + if(int r=launch(p.function,p.gx?p.gx:gx,p.gy?p.gy:gy,(p.gx||p.gy)?1:gz,p.bx,p.by,p.bz,0)) return r; + return SUCCESS; +} + +extern "C" int cuGetProcAddress_v2(const char* name,void** out,int version,uint64_t flags,int* status) { + if(!out||!name) return INVALID; + *out=nullptr; if(status) *status=1; + // No per-thread default-stream ABI or unversioned 32-bit size_t APIs. + if(flags>1||version<3020||version>12080) return UNSUPPORTED; + struct Symbol { const char* name; void* address; }; +#define SYMBOL(n) {#n,reinterpret_cast(&n)} + static const Symbol symbols[]={ + SYMBOL(cuInit),SYMBOL(cuDeviceGetCount),SYMBOL(cuDeviceGet), + SYMBOL(cuCtxCreate_v2),SYMBOL(cuCtxDestroy_v2),SYMBOL(cuCtxPushCurrent_v2), + SYMBOL(cuCtxPopCurrent_v2),SYMBOL(cuCtxSetCurrent),SYMBOL(cuCtxGetCurrent), + SYMBOL(cuCtxGetDevice),SYMBOL(cuCtxSynchronize),SYMBOL(cuDriverGetVersion), + SYMBOL(cuDeviceGetAttribute),SYMBOL(cuDeviceGetName),SYMBOL(cuDeviceGetUuid), + SYMBOL(cuDeviceGetLuid),SYMBOL(cuDeviceTotalMem_v2),SYMBOL(cuDeviceComputeCapability), + SYMBOL(cuMemAlloc_v2),SYMBOL(cuMemFree_v2),SYMBOL(cuMemAllocHost_v2), + SYMBOL(cuMemHostAlloc),SYMBOL(cuMemFreeHost),SYMBOL(cuMemGetInfo_v2), + SYMBOL(cuMemcpyHtoD_v2),SYMBOL(cuMemcpyDtoH_v2),SYMBOL(cuMemcpyDtoD_v2), + SYMBOL(cuMemcpyHtoDAsync_v2),SYMBOL(cuMemcpyDtoHAsync_v2),SYMBOL(cuMemcpyDtoDAsync_v2), + SYMBOL(cuMemsetD8_v2),SYMBOL(cuMemcpy2D_v2),SYMBOL(cuEventCreate),SYMBOL(cuEventDestroy_v2), + SYMBOL(cuEventRecord),SYMBOL(cuEventQuery),SYMBOL(cuEventSynchronize),SYMBOL(cuEventElapsedTime), + SYMBOL(cuStreamCreate),SYMBOL(cuStreamDestroy_v2),SYMBOL(cuStreamSynchronize), + SYMBOL(cuStreamQuery),SYMBOL(cuStreamWaitEvent),SYMBOL(cuGetErrorString), + SYMBOL(cuArrayCreate_v2),SYMBOL(cuArray3DCreate_v2),SYMBOL(cuArrayDestroy), + SYMBOL(cuArrayGetDescriptor_v2),SYMBOL(cuMipmappedArrayDestroy),SYMBOL(cuDestroyExternalMemory), + SYMBOL(cuTexObjectCreate),SYMBOL(cuTexObjectDestroy),SYMBOL(cuTexObjectGetResourceDesc), + SYMBOL(cuSurfObjectCreate),SYMBOL(cuSurfObjectDestroy),SYMBOL(cuSurfObjectGetResourceDesc), + SYMBOL(cuModuleLoad),SYMBOL(cuModuleLoadData),SYMBOL(cuModuleLoadDataEx),SYMBOL(cuModuleGetFunction), + SYMBOL(cuModuleUnload),SYMBOL(cuLaunchKernel) + }; +#undef SYMBOL + for(const auto& s:symbols) { + size_t n=std::strlen(s.name); + bool match=std::strcmp(name,s.name)==0; + // GetProcAddress receives CUDA base names and selects the modern ABI. + if(n>3&&std::strcmp(s.name+n-3,"_v2")==0) + match=match||(std::strlen(name)==n-3&&std::strncmp(name,s.name,n-3)==0); + if(match) { *out=s.address; if(status) *status=0; return SUCCESS; } + } + return NOT_FOUND; +} +extern "C" int cuGetProcAddress(const char* name,void** out,int version,uint64_t flags) { + return cuGetProcAddress_v2(name,out,version,flags,nullptr); +} diff --git a/tools/microcuda/zluda_precompile.py b/tools/microcuda/zluda_precompile.py new file mode 100644 index 0000000..2bb508a --- /dev/null +++ b/tools/microcuda/zluda_precompile.py @@ -0,0 +1,95 @@ +#!/usr/bin/env python3 +"""Compile every DLSS fatbin with one ZLUDA build into a cache, for make_pack.py. + +LD_LIBRARY_PATH= zluda_precompile.py DLL ZLUDA_LIBCUDA CACHE_HOME SWITCHES [--jobs N] + +SWITCHES is the D4R_ZLUDA_ code-generation set the pack is for, as make_pack.py prints it, e.g. +'D4R_ZLUDA_IGNORE_DENORMAL=1,D4R_ZLUDA_IMPLICIT_MAX_BLOCK=256,D4R_ZLUDA_WMMA=1,D4R_ZLUDA_WMMA_FP8=1'. +Each fatbin is loaded with cuModuleLoadData in a worker process (ZLUDA keeps the result in +CACHE_HOME/zluda/ComputeCache/zluda2.db) and unloaded again; nothing is launched. Already cached +modules return at once, so an interrupted run can be resumed. Fatbins without PTX fail (ZLUDA has +nothing to compile) and are left out of the pack. Compiling is CPU heavy and slow for +the big network modules (a minute or more each): do not run it next to a benchmark. +""" +import argparse +import concurrent.futures +import ctypes as c +import os +from pathlib import Path +import sys +import time + +sys.path.insert(0, str(Path(__file__).resolve().parents[2] / 'kernels/tools')) +from extract_dlss_ptx import fatbins + + +def compile_range(libcuda, dll, indices): + lib = c.CDLL(libcuda) + for name, args in (('cuInit', [c.c_uint]), ('cuDeviceGet', [c.POINTER(c.c_int), c.c_int]), + ('cuCtxCreate_v2', [c.POINTER(c.c_void_p), c.c_uint, c.c_int]), + ('cuModuleLoadData', [c.POINTER(c.c_void_p), c.c_void_p]), + ('cuModuleUnload', [c.c_void_p])): + getattr(lib, name).argtypes = args + getattr(lib, name).restype = c.c_int + assert lib.cuInit(0) == 0 + device, context = c.c_int(), c.c_void_p() + assert lib.cuDeviceGet(c.byref(device), 0) == 0 + assert lib.cuCtxCreate_v2(c.byref(context), 0, device.value) == 0 + images = list(fatbins(Path(dll).read_bytes())) + results = [] + for i in indices: + image = c.create_string_buffer(images[i][1]) + module = c.c_void_p() + start = time.time() + r = lib.cuModuleLoadData(c.byref(module), image) + if r == 0: + lib.cuModuleUnload(module) + results.append((i, r, time.time() - start)) + return results + + +def main(): + p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + p.add_argument('dll', type=Path) + p.add_argument('libcuda', type=Path) + p.add_argument('cache_home', type=Path) + p.add_argument('switches') + # ZLUDA's cache is one SQLite file: concurrent writers lose inserts silently (SQLITE_BUSY), so default to 1 + p.add_argument('--jobs', type=int, default=1) + # one process per fatbin: a module the compiler aborts on (e.g. WMMA under D4R_ZLUDA_WAVE64=1) is skipped + # instead of ending the run + p.add_argument('--isolate', action='store_true') + a = p.parse_args() + for key in [k for k in os.environ if k.startswith('D4R_ZLUDA_')]: + del os.environ[key] + for item in filter(None, a.switches.split(',')): + key, value = item.split('=', 1) + os.environ[key] = value + a.cache_home.mkdir(parents=True, exist_ok=True) + os.environ['XDG_CACHE_HOME'] = str(a.cache_home.resolve()) + count = len(list(fatbins(a.dll.read_bytes()))) + # interleave so every worker gets a mix of small and large modules + chunks = [list(range(j, count, a.jobs)) for j in range(a.jobs)] + failed = 0 + if a.isolate: + for i in range(count): + with concurrent.futures.ProcessPoolExecutor(1) as pool: + try: + [(_, r, seconds)] = pool.submit(compile_range, str(a.libcuda), str(a.dll), [i]).result() + except concurrent.futures.process.BrokenProcessPool: + r, seconds = 'aborted', 0.0 + failed += r != 0 + print(f'fatbin {i:3}: {"ok" if r == 0 else f"CUDA error {r}" if r != "aborted" else "compiler aborted"}' + f' ({seconds:.1f} s)', flush=True) + else: + with concurrent.futures.ProcessPoolExecutor(a.jobs) as pool: + for future in concurrent.futures.as_completed( + [pool.submit(compile_range, str(a.libcuda), str(a.dll), chunk) for chunk in chunks]): + for i, r, seconds in future.result(): + failed += r != 0 + print(f'fatbin {i:3}: {"ok" if r == 0 else f"CUDA error {r}"} ({seconds:.1f} s)', flush=True) + print(f'{count - failed}/{count} fatbins compiled into {a.cache_home}/zluda/ComputeCache/zluda2.db') + + +if __name__ == '__main__': + main() diff --git a/tools/wine_nvcuda_bridge.c b/tools/wine_nvcuda_bridge.c index db9b33e..3002658 100644 --- a/tools/wine_nvcuda_bridge.c +++ b/tools/wine_nvcuda_bridge.c @@ -62,6 +62,7 @@ enum }; static void* cuda_library; +static int microcuda_k_requested; static pthread_once_t load_once = PTHREAD_ONCE_INIT; static pthread_mutex_t trace_lock = PTHREAD_MUTEX_INITIALIZER; static pthread_mutex_t array_descriptor_lock = PTHREAD_MUTEX_INITIALIZER; @@ -777,6 +778,61 @@ int WINAPI d4rOutputKernelNative(void) static void load_zluda(void) { + /* Explicit experimental backend. Never silently mix module/context/array + handles from MicroCUDA and ZLUDA, and never fall back after selection. */ + const char* backend = getenv("D4R_CUDA_BACKEND"); + if (backend != NULL && strcmp(backend, "microcuda-k") == 0) + { + microcuda_k_requested = 1; + backend = "zluda"; + tracef("microcuda-k selected: audited K enc0 via MicroCUDA; all remaining APIs/kernels via ZLUDA (partial integration)"); + } + if (backend != NULL && backend[0] != '\0' && strcmp(backend, "zluda") != 0) + { + if (strcmp(backend, "microcuda") != 0) + { + set_load_error("unknown D4R_CUDA_BACKEND: %s", backend); + return; + } + const char* library = getenv("D4R_MICROCUDA_LIBCUDA"); + if (library == NULL || library[0] == '\0') + { + set_load_error("microcuda requires D4R_MICROCUDA_LIBCUDA and D4R_MICROCUDA_PACK"); + return; + } + const char* prefer = getenv("D4R_PREFER_ACCURACY"); + if (prefer != NULL && strcmp(prefer, "1") == 0) + { + set_load_error("microcuda has no PreferAccuracy module pack"); + return; + } + char native_path[1024]; + expand_home(library, native_path, sizeof(native_path)); + init_library_dirs(); + preload_rocm(); + cuda_library = load_library_resolving(native_path, 0); + if (cuda_library != NULL) + { + load_error[0] = '\0'; + /* Native kernels are verified and served exactly as for ZLUDA (per-process directory, + linked when NGX loads a module whose PTX matches); MicroCUDA reads that directory. */ + const char* cache = getenv("D4R_ZLUDA_CACHE_HOME"); + char cache_home[1024] = {0}; + if (cache != NULL && cache[0] != '\0') + { + expand_home(cache, cache_home, sizeof(cache_home)); + make_directories(cache_home); + } + prepare_native_kernels(cache_home[0] != '\0' ? cache_home : NULL); + const char* served = getenv("D4R_ZLUDA_NATIVE_DIR"); + if (served != NULL && served[0] != '\0') + setenv("D4R_MICROCUDA_NATIVE_DIR", served, 1); + else + unsetenv("D4R_MICROCUDA_NATIVE_DIR"); + tracef("MicroCUDA backend selected: %s (modules from D4R_MICROCUDA_PACK; no ZLUDA fallback)", native_path); + } + return; + } /* A policy switch, applied before ZLUDA loads or reads its compilation/cache settings. This also covers native-off and missing/version-mismatched native kernels. */ const int accuracy = d4r_apply_accuracy_policy(); @@ -1061,7 +1117,7 @@ static void create_default_zluda_context(void) return; } - tracef("created ZLUDA context for device %d", device); + tracef("created CUDA context for device %d", device); } static void ensure_context(void) @@ -1171,6 +1227,148 @@ static int find_allocation(CUdeviceptr address, CUdeviceptr* base, size_t* bytes return found; } +/* Explicit partial integration for game validation. Keep CUDA/ZLUDA handles + visible to NGX; replace only the audited enc0 dispatch, using the same HIP + context/default stream and linear pointers. This is NOT a ZLUDA-free backend. */ +static pthread_once_t microcuda_k_once = PTHREAD_ONCE_INIT; +static pthread_mutex_t microcuda_k_lock = PTHREAD_MUTEX_INITIALIZER; +static CUMODULELOADDATA_FN microcuda_k_load; +static CUMODULEGETFUNCTION_FN microcuda_k_get; +static CUMODULEUNLOAD_FN microcuda_k_unload; +static CULAUNCHKERNEL_FN microcuda_k_launch; +typedef void(__attribute__((sysv_abi)) *MICROCUDA_RANGE_FN)(uint64_t, size_t); +typedef void(__attribute__((sysv_abi)) *MICROCUDA_FREE_FN)(uint64_t); +typedef CUresult(__attribute__((sysv_abi)) *MICROCUDA_STATS_FN)(CUfunction, uint64_t*, uint64_t*); +static MICROCUDA_RANGE_FN microcuda_k_track, microcuda_k_write; +static MICROCUDA_FREE_FN microcuda_k_forget; +static MICROCUDA_STATS_FN microcuda_k_stats; +static CUmodule microcuda_k_reference[32], microcuda_k_modules[32]; +static CUfunction microcuda_k_original[32], microcuda_k_functions[32]; +static unsigned int microcuda_k_highwater; +static uint64_t microcuda_k_launches; + +static void load_microcuda_k(void) +{ + const char* configured = getenv("D4R_MICROCUDA_LIBCUDA"); + if (configured == NULL || configured[0] == '\0') + { + tracef("microcuda-k unavailable: D4R_MICROCUDA_LIBCUDA required"); + return; + } + char path[1024]; + expand_home(configured, path, sizeof(path)); + void* library = load_library_resolving(path, 0); + if (library == NULL) return; + CUINIT_FN init = (CUINIT_FN)dlsym(library, "cuInit"); + microcuda_k_load = (CUMODULELOADDATA_FN)dlsym(library, "cuModuleLoadData"); + microcuda_k_get = (CUMODULEGETFUNCTION_FN)dlsym(library, "cuModuleGetFunction"); + microcuda_k_unload = (CUMODULEUNLOAD_FN)dlsym(library, "cuModuleUnload"); + microcuda_k_launch = (CULAUNCHKERNEL_FN)dlsym(library, "cuLaunchKernel"); + microcuda_k_track = (MICROCUDA_RANGE_FN)dlsym(library, "d4rMicrocudaTrackAllocation"); + microcuda_k_write = (MICROCUDA_RANGE_FN)dlsym(library, "d4rMicrocudaNotifyWrite"); + microcuda_k_forget = (MICROCUDA_FREE_FN)dlsym(library, "d4rMicrocudaForgetAllocation"); + microcuda_k_stats = (MICROCUDA_STATS_FN)dlsym(library, "d4rMicrocudaGetDispatchStats"); + if (microcuda_k_track) + { + pthread_mutex_lock(&instrumentation_lock); + for (AllocationRecord* a = allocations; a != NULL; a = a->next) + microcuda_k_track(a->base, a->bytes); + pthread_mutex_unlock(&instrumentation_lock); + } + if (!init || !microcuda_k_load || !microcuda_k_get || !microcuda_k_unload || !microcuda_k_launch || init(0)) + { + microcuda_k_load = NULL; + tracef("microcuda-k initialization failed; no native MicroCUDA dispatch available"); + } +} + +static void microcuda_k_note_module(CUmodule reference, const void* image) +{ + if (!microcuda_k_requested || native_source[0] == '\0') return; + /* The experimental dispatch must use the same selected accuracy/fast/native + set as the reference. A registry for another variant fails its SHA check. */ + setenv("D4R_MICROCUDA_NATIVE_DIR", native_source, 1); + pthread_once(µcuda_k_once, load_microcuda_k); + if (!microcuda_k_load) return; + CUmodule module = NULL; + const CUresult result = microcuda_k_load(&module, image); + if (result != CUDA_SUCCESS) return; /* explicit hybrid: unregistered modules use reference */ + pthread_mutex_lock(µcuda_k_lock); + unsigned int slot; + for (slot = 0; slot < 32; ++slot) if (microcuda_k_modules[slot] == NULL) break; + if (slot < 32) + { + microcuda_k_reference[slot] = reference; + microcuda_k_modules[slot] = module; + if (slot >= microcuda_k_highwater) + __atomic_store_n(µcuda_k_highwater, slot + 1, __ATOMIC_RELEASE); + tracef("microcuda-k registered reference=%p native=%p slot=%u", reference, module, slot); + } + else + { + microcuda_k_unload(module); + tracef("microcuda-k module table exhausted"); + } + pthread_mutex_unlock(µcuda_k_lock); +} + +static void microcuda_k_note_function(CUmodule reference, CUfunction original, const char* name) +{ + if (!microcuda_k_requested || strcmp(name, "dltss_pwin_enc0_layer") != 0) return; + pthread_mutex_lock(µcuda_k_lock); + int mapped = 0; + for (unsigned int slot = 0; slot < microcuda_k_highwater; ++slot) + if (microcuda_k_reference[slot] == reference && microcuda_k_modules[slot] != NULL) + { + CUfunction function = NULL; + CUresult result = microcuda_k_get(&function, microcuda_k_modules[slot], name); + if (result == CUDA_SUCCESS) + { + mapped = 1; + __atomic_store_n(µcuda_k_functions[slot], function, __ATOMIC_RELEASE); + __atomic_store_n(µcuda_k_original[slot], original, __ATOMIC_RELEASE); + tracef("microcuda-k dispatch mapped %s original=%p native=%p", name, original, function); + } + else tracef("microcuda-k function lookup failed %d", result); + } + if (!mapped) tracef("microcuda-k enc0 remains reference: no compatible module/selected native variant"); + pthread_mutex_unlock(µcuda_k_lock); +} + +static CUfunction microcuda_k_lookup(CUfunction original) +{ + if (!microcuda_k_requested) return NULL; + const unsigned int count = __atomic_load_n(µcuda_k_highwater, __ATOMIC_ACQUIRE); + for (unsigned int i = 0; i < count; ++i) + if (__atomic_load_n(µcuda_k_original[i], __ATOMIC_ACQUIRE) == original) + return __atomic_load_n(µcuda_k_functions[i], __ATOMIC_ACQUIRE); + return NULL; +} + +static void microcuda_k_remove_module(CUmodule reference) +{ + if (!microcuda_k_requested) return; + pthread_mutex_lock(µcuda_k_lock); + for (unsigned int slot = 0; slot < microcuda_k_highwater; ++slot) + if (microcuda_k_reference[slot] == reference && microcuda_k_modules[slot] != NULL) + { + if (microcuda_k_stats && microcuda_k_functions[slot]) + { + uint64_t prep = 0, skips = 0; + microcuda_k_stats(microcuda_k_functions[slot], &prep, &skips); + tracef("microcuda-k prep stats runs=%llu skips=%llu", (unsigned long long)prep, (unsigned long long)skips); + } + __atomic_store_n(µcuda_k_original[slot], NULL, __ATOMIC_RELEASE); + __atomic_store_n(µcuda_k_functions[slot], NULL, __ATOMIC_RELEASE); + CUresult result = microcuda_k_unload(microcuda_k_modules[slot]); + if (result != CUDA_SUCCESS) tracef("microcuda-k unload failed %d", result); + microcuda_k_modules[slot] = NULL; + microcuda_k_reference[slot] = NULL; + tracef("microcuda-k completed dispatches=%llu", (unsigned long long)microcuda_k_launches); + } + pthread_mutex_unlock(µcuda_k_lock); +} + static void remember_surface_object(CUsurfObject object, const void* resource) { uint32_t resource_type = UINT32_MAX; @@ -2119,6 +2317,7 @@ CUresult WINAPI cuModuleLoadData(CUmodule* module, const void* image) capture_ptx_if_present(image, "cuModuleLoadData"); image = replacement_image(image, "cuModuleLoadData"); CUresult result = function != NULL ? function(module, image) : missing("cuModuleLoadData"); + if (result == CUDA_SUCCESS) microcuda_k_note_module(*module, image); tracef("cuModuleLoadData image=%p result=%d module=%p", image, result, module != NULL ? *module : NULL); return result; } @@ -2131,6 +2330,7 @@ CUresult WINAPI cuModuleLoadDataEx(CUmodule* module, const void* image, unsigned capture_ptx_if_present(image, "cuModuleLoadDataEx"); CUresult result = function != NULL ? function(module, image, option_count, options, option_values) : missing("cuModuleLoadDataEx"); + if (result == CUDA_SUCCESS && option_count == 0) microcuda_k_note_module(*module, image); tracef("cuModuleLoadDataEx image=%p options=%u result=%d module=%p", image, option_count, result, module != NULL ? *module : NULL); return result; @@ -2149,6 +2349,7 @@ CUresult WINAPI cuModuleUnload(CUmodule module) { CUMODULEUNLOAD_FN function = (CUMODULEUNLOAD_FN)find_zluda_symbol("cuModuleUnload"); CUresult result = function != NULL ? function(module) : missing("cuModuleUnload"); + if (result == CUDA_SUCCESS) microcuda_k_remove_module(module); tracef("cuModuleUnload module=%p result=%d", module, result); return result; } @@ -2159,6 +2360,7 @@ CUresult WINAPI cuModuleGetFunction(CUfunction* function_out, CUmodule module, c CUresult result = function != NULL ? function(function_out, module, name) : missing("cuModuleGetFunction"); if (result == CUDA_SUCCESS && function_out != NULL) { + microcuda_k_note_function(module, *function_out, name); remember_function_name(*function_out, name); note_function_lookup(*function_out, name); } @@ -2184,10 +2386,25 @@ CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x, note_launch(function_handle); PendingKernelProfile* profile = function != NULL ? begin_kernel_profile(sequence, function_handle, stream) : NULL; - CUresult result = function != NULL - ? function(function_handle, grid_x, grid_y, grid_z, block_x, block_y, block_z, + CUfunction native_function = microcuda_k_lookup(function_handle); + if (native_function != NULL) + { + const uint64_t count = __atomic_add_fetch(µcuda_k_launches, 1, __ATOMIC_RELAXED); + if (count == 1 || count % 4096 == 0) + tracef("microcuda-k actual dispatch %llu original=%p native=%p", (unsigned long long)count, function_handle, native_function); + } + CULAUNCHKERNEL_FN selected = native_function != NULL ? microcuda_k_launch : function; + CUresult result = selected != NULL + ? selected(native_function != NULL ? native_function : function_handle, grid_x, grid_y, grid_z, block_x, block_y, block_z, shared_bytes, stream, kernel_params, extra) : missing("cuLaunchKernel"); + if (native_function != NULL && result == CUDA_SUCCESS && microcuda_k_stats && + (__atomic_load_n(µcuda_k_launches, __ATOMIC_RELAXED) % 4096) == 0) + { + uint64_t prep = 0, skips = 0; + microcuda_k_stats(native_function, &prep, &skips); + tracef("microcuda-k prep stats runs=%llu skips=%llu", (unsigned long long)prep, (unsigned long long)skips); + } end_kernel_profile(profile, result, stream); if (result != CUDA_SUCCESS || trace_verbose()) { @@ -2233,7 +2450,10 @@ CUresult WINAPI cuMemAlloc(CUdeviceptr* pointer, size_t bytes) CUMEMALLOC_FN function = (CUMEMALLOC_FN)find_zluda_symbol("cuMemAlloc_v2"); CUresult result = function != NULL ? function(pointer, bytes) : missing("cuMemAlloc_v2"); if (result == CUDA_SUCCESS && pointer != NULL) + { remember_allocation(*pointer, bytes); + if (microcuda_k_track) microcuda_k_track(*pointer, bytes); + } TRACE_CALL(result, "cuMemAlloc->v2 bytes=%zu result=%d device_ptr=0x%llx", bytes, result, (unsigned long long)(pointer != NULL ? *pointer : 0)); return result; @@ -2254,6 +2474,7 @@ CUresult WINAPI cuMemAllocHost(void** pointer, size_t bytes) CUresult WINAPI cuMemFree(CUdeviceptr pointer) { CUMEMFREE_FN function = (CUMEMFREE_FN)find_zluda_symbol("cuMemFree_v2"); + if (microcuda_k_forget) microcuda_k_forget(pointer); CUresult result = function != NULL ? function(pointer) : missing("cuMemFree_v2"); if (result == CUDA_SUCCESS) forget_allocation(pointer); @@ -2272,6 +2493,7 @@ CUresult WINAPI cuMemFreeHost(void* pointer) CUresult WINAPI cuMemcpy2D(const void* copy) { CUMEMCPY2D_FN function = (CUMEMCPY2D_FN)find_zluda_symbol("cuMemcpy2D_v2"); + if (microcuda_k_write) microcuda_k_write(0, SIZE_MAX); CUresult result = function != NULL ? function(copy) : missing("cuMemcpy2D_v2"); TRACE_CALL(result, "cuMemcpy2D->v2 descriptor=%p result=%d", copy, result); return result; @@ -2280,6 +2502,7 @@ CUresult WINAPI cuMemcpy2D(const void* copy) CUresult WINAPI cuMemcpyHtoDAsync(CUdeviceptr destination, const void* source, size_t bytes, CUstream stream) { CUMEMCPYHTODASYNC_FN function = (CUMEMCPYHTODASYNC_FN)find_zluda_symbol("cuMemcpyHtoDAsync_v2"); + if (microcuda_k_write) microcuda_k_write(destination, bytes); CUresult result = function != NULL ? function(destination, source, bytes, stream) : missing("cuMemcpyHtoDAsync_v2"); TRACE_CALL(result, "cuMemcpyHtoDAsync->v2 dst=0x%llx src=%p bytes=%zu stream=%p result=%d", (unsigned long long)destination, source, bytes, stream, result); @@ -2610,12 +2833,18 @@ CUresult WINAPI cuTexObjectGetResourceDesc(void* descriptor, CUtexObject object) /* Direct pass-throughs for driver APIs whose arguments are all integers or pointers, so the Windows x64 call forwards unchanged to ZLUDA (SysV). Unversioned names resolve to the versions the CUDA 12 driver returns for them. */ +static void microcuda_k_notify_generic_write(const char* target) +{ + if (microcuda_k_write && (strncmp(target, "cuMemcpyHtoD", 12) == 0 || strncmp(target, "cuMemcpyDtoD", 12) == 0 || strncmp(target, "cuMemset", 8) == 0)) + microcuda_k_write(0, SIZE_MAX); /* unsupported/general writer: invalidate all prepared weights */ +} typedef uintptr_t D4rArg; #define D4R_FORWARD(name, target, count, params, args) \ CUresult WINAPI name params \ { \ typedef CUresult(__attribute__((sysv_abi)) * function_type) params; \ function_type function = (function_type)find_zluda_symbol(#target); \ + microcuda_k_notify_generic_write(#target); \ CUresult result = function != NULL ? function args : missing(#target); \ TRACE_CALL(result, #name "->" #target " result=%d", result); \ return result; \