From 66a9edb8985c92374cd010a981ec6a54ad8a90cd Mon Sep 17 00:00:00 2001 From: Tommi Niemi Date: Mon, 6 Apr 2026 20:38:10 +0700 Subject: [PATCH] Initial extraction from isis: direct /dev/kfd GPU compute AMD RDNA3 GPU driver via raw KFD ioctls. No ROCm, no OpenCL. - Device discovery and topology queries - VRAM/GTT/Userptr memory with shared address space - PM4 compute queue and kernel dispatch - Hand-written RDNA3 ASM kernels: matmul (3100 GFLOP/s), matvec, superlinear - Tile-safe buffer padding for OOB protection - Event-based interrupt wait (no CPU polling) Co-Authored-By: Claude Opus 4.6 (1M context) --- Cargo.lock | 16 + Cargo.toml | 12 + src/compute.rs | 67 + src/device.rs | 1260 +++++++++++++++++ src/dispatch.rs | 413 ++++++ src/ioctl.rs | 330 +++++ src/kernels/addr_dump.co | Bin 0 -> 3216 bytes src/kernels/addr_dump.s | 149 ++ src/kernels/coop_test.co | Bin 0 -> 2696 bytes src/kernels/coop_test.s | 92 ++ src/kernels/lds_test.co | Bin 0 -> 2640 bytes src/kernels/lds_test.s | 72 + src/kernels/matmul.co | Bin 0 -> 3632 bytes src/kernels/matmul.s | 279 ++++ src/kernels/matmul_blocked.co | Bin 0 -> 4360 bytes src/kernels/matmul_blocked.s | 338 +++++ src/kernels/matmul_dbg.co | Bin 0 -> 3656 bytes src/kernels/matmul_dbg.s | 282 ++++ src/kernels/matmul_small.co | Bin 0 -> 4232 bytes src/kernels/matmul_small.s | 310 ++++ src/kernels/matvec.cl | 106 ++ src/kernels/matvec.co | Bin 0 -> 16928 bytes src/kernels/matvec.s | 138 ++ src/kernels/matvec_asm.co | Bin 0 -> 3040 bytes src/kernels/superlinear.co | Bin 0 -> 3208 bytes src/kernels/superlinear.s | 180 +++ src/kernels/test_store.co | Bin 0 -> 2032 bytes src/kernels/test_store.s | 41 + src/lib.rs | 25 + src/memory.rs | 328 +++++ src/queue.rs | 398 ++++++ target/.rustc_info.json | 1 + target/CACHEDIR.TAG | 3 + target/debug/.cargo-lock | 0 .../kfd-2447e9870f0e8752/invoked.timestamp | 1 + .../kfd-2447e9870f0e8752/output-lib-kfd | 50 + .../kfd-c036a4411ecb62c5/dep-lib-kfd | Bin 0 -> 307 bytes .../kfd-c036a4411ecb62c5/invoked.timestamp | 1 + .../.fingerprint/kfd-c036a4411ecb62c5/lib-kfd | 1 + .../kfd-c036a4411ecb62c5/lib-kfd.json | 1 + .../kfd-c036a4411ecb62c5/output-lib-kfd | 4 + .../libc-184a2b4f65324410/dep-lib-libc | Bin 0 -> 14 bytes .../libc-184a2b4f65324410/invoked.timestamp | 1 + .../libc-184a2b4f65324410/lib-libc | 1 + .../libc-184a2b4f65324410/lib-libc.json | 1 + .../build-script-build-script-build | 1 + .../build-script-build-script-build.json | 1 + .../dep-build-script-build-script-build | Bin 0 -> 14 bytes .../libc-6578899ebd06082d/invoked.timestamp | 1 + .../run-build-script-build-script-build | 1 + .../run-build-script-build-script-build.json | 1 + .../libc-6578899ebd06082d/build-script-build | Bin 0 -> 4139544 bytes .../build_script_build-6578899ebd06082d | Bin 0 -> 4139544 bytes .../build_script_build-6578899ebd06082d.d | 5 + .../libc-af64a192d6d29af3/invoked.timestamp | 1 + .../debug/build/libc-af64a192d6d29af3/output | 25 + .../build/libc-af64a192d6d29af3/root-output | 1 + .../debug/build/libc-af64a192d6d29af3/stderr | 0 target/debug/deps/kfd-2447e9870f0e8752.d | 16 + target/debug/deps/kfd-c036a4411ecb62c5.d | 16 + target/debug/deps/libc-184a2b4f65324410.d | 44 + .../debug/deps/libkfd-c036a4411ecb62c5.rmeta | Bin 0 -> 232298 bytes .../debug/deps/liblibc-184a2b4f65324410.rmeta | Bin 0 -> 3496652 bytes .../dep-graph.bin | Bin 0 -> 2714056 bytes .../metadata.rmeta | Bin 0 -> 232298 bytes .../query-cache.bin | Bin 0 -> 924301 bytes .../work-products.bin | Bin 0 -> 100 bytes .../s-hhcthcl33v-0whmljb.lock | 0 .../dep-graph.part.bin | Bin 0 -> 2753571 bytes .../s-hhcth4cuny-11e5e6s.lock | 0 tests/kfd_bench.rs | 160 +++ tests/kfd_gpu.rs | 198 +++ tests/kfd_matmul_bench.rs | 167 +++ tests/kfd_matmul_bench_small.rs | 97 ++ tests/kfd_remu.rs | 222 +++ 75 files changed, 5858 insertions(+) create mode 100644 Cargo.lock create mode 100644 Cargo.toml create mode 100644 src/compute.rs create mode 100644 src/device.rs create mode 100644 src/dispatch.rs create mode 100644 src/ioctl.rs create mode 100755 src/kernels/addr_dump.co create mode 100644 src/kernels/addr_dump.s create mode 100755 src/kernels/coop_test.co create mode 100644 src/kernels/coop_test.s create mode 100755 src/kernels/lds_test.co create mode 100644 src/kernels/lds_test.s create mode 100755 src/kernels/matmul.co create mode 100644 src/kernels/matmul.s create mode 100755 src/kernels/matmul_blocked.co create mode 100644 src/kernels/matmul_blocked.s create mode 100755 src/kernels/matmul_dbg.co create mode 100644 src/kernels/matmul_dbg.s create mode 100755 src/kernels/matmul_small.co create mode 100644 src/kernels/matmul_small.s create mode 100644 src/kernels/matvec.cl create mode 100755 src/kernels/matvec.co create mode 100644 src/kernels/matvec.s create mode 100755 src/kernels/matvec_asm.co create mode 100755 src/kernels/superlinear.co create mode 100644 src/kernels/superlinear.s create mode 100755 src/kernels/test_store.co create mode 100644 src/kernels/test_store.s create mode 100644 src/lib.rs create mode 100644 src/memory.rs create mode 100644 src/queue.rs create mode 100644 target/.rustc_info.json create mode 100644 target/CACHEDIR.TAG create mode 100644 target/debug/.cargo-lock create mode 100644 target/debug/.fingerprint/kfd-2447e9870f0e8752/invoked.timestamp create mode 100644 target/debug/.fingerprint/kfd-2447e9870f0e8752/output-lib-kfd create mode 100644 target/debug/.fingerprint/kfd-c036a4411ecb62c5/dep-lib-kfd create mode 100644 target/debug/.fingerprint/kfd-c036a4411ecb62c5/invoked.timestamp create mode 100644 target/debug/.fingerprint/kfd-c036a4411ecb62c5/lib-kfd create mode 100644 target/debug/.fingerprint/kfd-c036a4411ecb62c5/lib-kfd.json create mode 100644 target/debug/.fingerprint/kfd-c036a4411ecb62c5/output-lib-kfd create mode 100644 target/debug/.fingerprint/libc-184a2b4f65324410/dep-lib-libc create mode 100644 target/debug/.fingerprint/libc-184a2b4f65324410/invoked.timestamp create mode 100644 target/debug/.fingerprint/libc-184a2b4f65324410/lib-libc create mode 100644 target/debug/.fingerprint/libc-184a2b4f65324410/lib-libc.json create mode 100644 target/debug/.fingerprint/libc-6578899ebd06082d/build-script-build-script-build create mode 100644 target/debug/.fingerprint/libc-6578899ebd06082d/build-script-build-script-build.json create mode 100644 target/debug/.fingerprint/libc-6578899ebd06082d/dep-build-script-build-script-build create mode 100644 target/debug/.fingerprint/libc-6578899ebd06082d/invoked.timestamp create mode 100644 target/debug/.fingerprint/libc-af64a192d6d29af3/run-build-script-build-script-build create mode 100644 target/debug/.fingerprint/libc-af64a192d6d29af3/run-build-script-build-script-build.json create mode 100755 target/debug/build/libc-6578899ebd06082d/build-script-build create mode 100755 target/debug/build/libc-6578899ebd06082d/build_script_build-6578899ebd06082d create mode 100644 target/debug/build/libc-6578899ebd06082d/build_script_build-6578899ebd06082d.d create mode 100644 target/debug/build/libc-af64a192d6d29af3/invoked.timestamp create mode 100644 target/debug/build/libc-af64a192d6d29af3/output create mode 100644 target/debug/build/libc-af64a192d6d29af3/root-output create mode 100644 target/debug/build/libc-af64a192d6d29af3/stderr create mode 100644 target/debug/deps/kfd-2447e9870f0e8752.d create mode 100644 target/debug/deps/kfd-c036a4411ecb62c5.d create mode 100644 target/debug/deps/libc-184a2b4f65324410.d create mode 100644 target/debug/deps/libkfd-c036a4411ecb62c5.rmeta create mode 100644 target/debug/deps/liblibc-184a2b4f65324410.rmeta create mode 100644 target/debug/incremental/kfd-2a7rsuw39z8ru/s-hhcthcl33v-0whmljb-0nanvm67qve0kwcedetfc5y1i/dep-graph.bin create mode 100644 target/debug/incremental/kfd-2a7rsuw39z8ru/s-hhcthcl33v-0whmljb-0nanvm67qve0kwcedetfc5y1i/metadata.rmeta create mode 100644 target/debug/incremental/kfd-2a7rsuw39z8ru/s-hhcthcl33v-0whmljb-0nanvm67qve0kwcedetfc5y1i/query-cache.bin create mode 100644 target/debug/incremental/kfd-2a7rsuw39z8ru/s-hhcthcl33v-0whmljb-0nanvm67qve0kwcedetfc5y1i/work-products.bin create mode 100644 target/debug/incremental/kfd-2a7rsuw39z8ru/s-hhcthcl33v-0whmljb.lock create mode 100644 target/debug/incremental/kfd-3vh5csval5g5u/s-hhcth4cuny-11e5e6s-working/dep-graph.part.bin create mode 100644 target/debug/incremental/kfd-3vh5csval5g5u/s-hhcth4cuny-11e5e6s.lock create mode 100644 tests/kfd_bench.rs create mode 100644 tests/kfd_gpu.rs create mode 100644 tests/kfd_matmul_bench.rs create mode 100644 tests/kfd_matmul_bench_small.rs create mode 100644 tests/kfd_remu.rs diff --git a/Cargo.lock b/Cargo.lock new file mode 100644 index 0000000..9353f29 --- /dev/null +++ b/Cargo.lock @@ -0,0 +1,16 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "kfd" +version = "0.1.0" +dependencies = [ + "libc", +] + +[[package]] +name = "libc" +version = "0.2.184" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "48f5d2a454e16a5ea0f4ced81bd44e4cfc7bd3a507b61887c99fd3538b28e4af" diff --git a/Cargo.toml b/Cargo.toml new file mode 100644 index 0000000..1fb2938 --- /dev/null +++ b/Cargo.toml @@ -0,0 +1,12 @@ +[package] +name = "kfd" +version = "0.1.0" +edition = "2021" +description = "Direct /dev/kfd GPU compute for AMD RDNA3 — no ROCm, no OpenCL, just ioctl" +license = "MIT OR Apache-2.0" +repository = "https://git.rotko.net/tommi/kfd" + +[dependencies] +libc = "0.2" + +[dev-dependencies] diff --git a/src/compute.rs b/src/compute.rs new file mode 100644 index 0000000..6d13fb9 --- /dev/null +++ b/src/compute.rs @@ -0,0 +1,67 @@ +//! compute dispatch primitives. +//! +//! GpuFuture: poll gpu signal for async completion. +//! ComputeState: what the ctm feels about its compute (interoception). +//! +//! the ctm doesn't know about gpus. it just calls the harness +//! and feels whether compute is fast or slow. + +use crate::memory::GpuBuffer; +use std::time::Instant; + +// ─── future ───────────────────────────────────────────────── + +/// async gpu completion. poll or block. +pub struct GpuFuture { + signal_ptr: *const u32, + expected: u32, + submitted_at: Instant, +} + +unsafe impl Send for GpuFuture {} + +impl GpuFuture { + pub fn new(signal: &GpuBuffer, expected: u32) -> Self { + GpuFuture { + signal_ptr: signal.cpu_ptr as *const u32, + expected, + submitted_at: Instant::now(), + } + } + + /// true if gpu finished. + pub fn poll(&self) -> bool { + let val = unsafe { self.signal_ptr.read_volatile() }; + val >= self.expected + } + + /// block until done. returns elapsed, or None on timeout. + pub fn wait(&self, timeout_us: u64) -> Option { + let start = Instant::now(); + loop { + if self.poll() { + return Some(self.submitted_at.elapsed()); + } + if start.elapsed().as_micros() as u64 > timeout_us { + return None; + } + std::hint::spin_loop(); + } + } +} + +// ─── interoception ────────────────────────────────────────── + +/// what the ctm feels about its compute. +/// read-only from organism side. harness updates it. +#[derive(Debug, Clone, Default)] +pub struct ComputeState { + /// rolling average dispatch latency (microseconds) + pub latency_us: f32, + /// true when gpu dispatch is in flight + pub gpu_busy: bool, + /// total dispatches + pub dispatch_count: u64, + /// faults (should be 0) + pub fault_count: u32, +} diff --git a/src/device.rs b/src/device.rs new file mode 100644 index 0000000..5080ff8 --- /dev/null +++ b/src/device.rs @@ -0,0 +1,1260 @@ +//! Heterogeneous compute via KFD (Kernel Fusion Driver). +//! +//! KFD manages both CPU and GPU as HSA compute agents. +//! One driver, unified memory, automatic dispatch. +//! +//! Node 0: CPU (32 cores, AVX-512, DDR5) +//! Node 1: GPU (64 SIMDs, wave32, 8GB VRAM) +//! +//! Small ops → CPU (lower latency, AVX-512) +//! Large matmuls → GPU (higher throughput, 64 SIMDs) +//! Shared address space — zero-copy between CPU and GPU. +//! +//! Reference: tinygrad/tinygrad/runtime/ops_amd.py + +use crate::memory::{GpuAllocator, GpuBuffer}; +use crate::queue::ComputeQueue; +use crate::dispatch::{CodeObject, GpuProgram, KernArgs, KernelEntry}; +use std::os::unix::io::RawFd; + +/// compiled rdna3 kernels (gfx1102), hand-written asm. +/// test_store: y[i] = 42.0 (minimal proof kernel) +const TEST_STORE_CO: &[u8] = include_bytes!("kernels/test_store.co"); +/// matvec: y = W*x + b (raw asm, no hidden opencl args) +const MATVEC_CO: &[u8] = include_bytes!("kernels/matvec_asm.co"); +/// matmul_blocked: TM=128 TT4x4 — best for M >= 1536 +const MATMUL_BLOCKED_CO: &[u8] = include_bytes!("kernels/matmul_blocked.co"); +/// matmul_small: TM=32 TT2x2 PLR — best for M < 1536 +const MATMUL_SMALL_CO: &[u8] = include_bytes!("kernels/matmul_small.co"); +/// superlinear_fwd: per-neuron batched matvec for NLM +/// Y[n*O+o] = B[n*O+o] + sum_k W[n*O*K+o*K+k] * X[n*K+k] +const SUPERLINEAR_CO: &[u8] = include_bytes!("kernels/superlinear.co"); + +// ─── Node properties ──────────────────────────────────────── + +/// Properties of a KFD compute node (CPU or GPU). +#[derive(Debug, Default, Clone)] +pub struct NodeProperties { + pub node_id: u32, + pub gpu_id: u32, + pub cpu_cores_count: u32, + pub simd_count: u32, + pub wave_front_size: u32, + pub gfx_target_version: u32, + pub max_waves_per_simd: u32, + pub array_count: u32, + pub cu_per_simd_array: u32, + pub drm_render_minor: u32, + pub cwsr_size: u32, + pub max_engine_clk: u32, + // Memory + pub mem_size_bytes: u64, + pub mem_width: u32, + pub mem_clk_max: u32, +} + +impl NodeProperties { + pub fn is_cpu(&self) -> bool { self.cpu_cores_count > 0 } + pub fn is_gpu(&self) -> bool { self.simd_count > 0 } + + /// Read properties from sysfs for a given node. + fn read(node_id: u32) -> std::io::Result { + let base = format!("/sys/devices/virtual/kfd/kfd/topology/nodes/{}", node_id); + let props_path = format!("{}/properties", base); + let gpu_id_path = format!("{}/gpu_id", base); + let mem_path = format!("{}/mem_banks/0/properties", base); + + let gpu_id: u32 = std::fs::read_to_string(&gpu_id_path) + .unwrap_or_default().trim().parse().unwrap_or(0); + + let props_text = std::fs::read_to_string(&props_path)?; + let mut p = NodeProperties { node_id, gpu_id, ..Default::default() }; + + for line in props_text.lines() { + let mut parts = line.split_whitespace(); + let key = parts.next().unwrap_or(""); + let val: u64 = parts.next().and_then(|v| v.parse().ok()).unwrap_or(0); + match key { + "cpu_cores_count" => p.cpu_cores_count = val as u32, + "simd_count" => p.simd_count = val as u32, + "wave_front_size" => p.wave_front_size = val as u32, + "gfx_target_version" => p.gfx_target_version = val as u32, + "cwsr_size" => p.cwsr_size = val as u32, + "max_waves_per_simd" => p.max_waves_per_simd = val as u32, + "array_count" => p.array_count = val as u32, + "cu_per_simd_array" => p.cu_per_simd_array = val as u32, + "drm_render_minor" => p.drm_render_minor = val as u32, + "max_engine_clk_ccompute" => p.max_engine_clk = val as u32, + _ => {} + } + } + + // Read memory bank + if let Ok(mem_text) = std::fs::read_to_string(&mem_path) { + for line in mem_text.lines() { + let mut parts = line.split_whitespace(); + let key = parts.next().unwrap_or(""); + let val: u64 = parts.next().and_then(|v| v.parse().ok()).unwrap_or(0); + match key { + "size_in_bytes" => p.mem_size_bytes = val, + "width" => p.mem_width = val as u32, + "mem_clk_max" => p.mem_clk_max = val as u32, + _ => {} + } + } + } + + Ok(p) + } +} + +// ─── Device specs & status ────────────────────────────────── + +/// Static device specifications. Readable without opening the device +/// (from sysfs topology). Use `DeviceSpecs::probe()` for zero-cost discovery. +#[derive(Debug, Clone)] +pub struct DeviceSpecs { + /// CPU node properties. + pub cpu: NodeProperties, + /// GPU node properties. + pub gpu: NodeProperties, + /// Total VRAM in bytes. + pub vram_total: u64, + /// VRAM bus width in bits (e.g. 256). + pub vram_width: u32, + /// VRAM max clock in MHz. + pub vram_clock_mhz: u32, + /// Peak VRAM bandwidth in GB/s (computed: width * clock * 2 / 8 / 1000). + pub vram_bandwidth_gbps: f32, + /// Total system RAM in bytes. + pub ram_total: u64, + /// GPU gfx target version (e.g. 0x110002 for gfx1102). + pub gfx_version: u32, + /// GPU compute units (simd_count * cu_per_simd_array, or simd_count if no CU info). + pub compute_units: u32, +} + +impl DeviceSpecs { + /// Probe device specs from sysfs. Does NOT open /dev/kfd — safe to call + /// even without permissions, just needs sysfs readable. + pub fn probe() -> std::io::Result { + let mut cpu = None; + let mut gpu = None; + for node in 0..16 { + match NodeProperties::read(node) { + Ok(p) if p.is_gpu() && gpu.is_none() => gpu = Some(p), + Ok(p) if p.is_cpu() && cpu.is_none() => cpu = Some(p), + _ => {} + } + } + let cpu = cpu.ok_or_else(|| std::io::Error::new( + std::io::ErrorKind::NotFound, "no CPU node in KFD topology"))?; + let gpu = gpu.ok_or_else(|| std::io::Error::new( + std::io::ErrorKind::NotFound, "no GPU node in KFD topology"))?; + + let vram_total = gpu.mem_size_bytes; + let vram_width = gpu.mem_width; + let vram_clock = gpu.mem_clk_max; + // DDR effective: clock * 2 (double data rate) + // Bandwidth = width_bits / 8 * effective_clock * 1e6 / 1e9 + let bw = (vram_width as f32 / 8.0) * (vram_clock as f32 * 2.0) / 1000.0; + + Ok(DeviceSpecs { + vram_total, + vram_width, + vram_clock_mhz: vram_clock, + vram_bandwidth_gbps: bw, + ram_total: cpu.mem_size_bytes, + gfx_version: gpu.gfx_target_version, + compute_units: if gpu.cu_per_simd_array > 0 { + gpu.simd_count * gpu.cu_per_simd_array + } else { + gpu.simd_count + }, + cpu, + gpu, + }) + } + + /// Human-readable summary. + pub fn summary(&self) -> String { + format!( + "GPU: gfx{:x}, {} CUs, {} GB VRAM (GDDR6-{} MHz, {}-bit, {:.0} GB/s)\n\ + CPU: {} cores, {} GB RAM", + self.gfx_version, + self.compute_units, + self.vram_total / (1024 * 1024 * 1024), + self.vram_clock_mhz, + self.vram_width, + self.vram_bandwidth_gbps, + self.cpu.cpu_cores_count, + self.ram_total / (1024 * 1024 * 1024), + ) + } + + /// Can this model's weights fit in VRAM? + /// `weight_bytes`: total weight size in bytes. + /// Returns (fits, headroom_bytes). + pub fn weights_fit(&self, weight_bytes: u64) -> (bool, i64) { + let headroom = self.vram_total as i64 - weight_bytes as i64; + (headroom > 0, headroom) + } +} + +/// Live device status. Requires an open KFD file descriptor. +#[derive(Debug, Clone)] +pub struct DeviceStatus { + /// Currently available (free) VRAM in bytes. + pub vram_available: u64, + /// Total VRAM in bytes (from specs). + pub vram_total: u64, + /// VRAM used = total - available. + pub vram_used: u64, + /// Current shader clock in MHz (from sysfs, None if unreadable). + pub sclk_mhz: Option, +} + +impl DeviceStatus { + /// Utilization ratio: 0.0 = empty, 1.0 = full. + pub fn vram_utilization(&self) -> f32 { + if self.vram_total == 0 { return 0.0; } + self.vram_used as f32 / self.vram_total as f32 + } + + /// Human-readable summary. + pub fn summary(&self) -> String { + let used_mb = self.vram_used / (1024 * 1024); + let total_mb = self.vram_total / (1024 * 1024); + let avail_mb = self.vram_available / (1024 * 1024); + let clk = self.sclk_mhz.map(|m| format!("{} MHz", m)) + .unwrap_or_else(|| "unknown".into()); + format!( + "VRAM: {} MB / {} MB ({} MB free, {:.0}% used), sclk: {}", + used_mb, total_mb, avail_mb, + self.vram_utilization() * 100.0, + clk, + ) + } +} + +// ─── Unified HSA device ───────────────────────────────────── + +/// Unified heterogeneous compute device via KFD. +/// +/// Manages both CPU and GPU nodes. Allocates shared memory visible +/// to both. Dispatches compute to the appropriate agent based on +/// problem size. +pub struct HsaDevice { + pub kfd_fd: RawFd, + pub drm_fd: RawFd, + pub cpu: NodeProperties, + pub gpu: NodeProperties, + pub alloc: GpuAllocator, + /// GPU resources — these fields drop BEFORE _fd_guard (declaration order). + _code_objects: Vec, + pub kernels: Option>, + pub queue: ComputeQueue, + pub signal: GpuBuffer, + pub event_page: GpuBuffer, + pub scratch: GpuBuffer, + signal_value: u32, + pub vram_available: u64, + pub queue_event_id: u32, + pub queue_event_mailbox_ptr: u64, + pub gpu_dispatch_threshold: usize, + /// MUST be last field: closes kfd_fd/drm_fd after all GpuBuffers are dropped. + _fd_guard: FdGuard, +} + +impl HsaDevice { + /// Open the HSA device. Discovers CPU + GPU from KFD topology. + pub fn open() -> std::io::Result { + let kfd_fd = unsafe { libc::open(b"/dev/kfd\0".as_ptr() as _, libc::O_RDWR) }; + if kfd_fd < 0 { + return Err(std::io::Error::new(std::io::ErrorKind::NotFound, + "/dev/kfd not available")); + } + + let (major, minor) = crate::ioctl::get_version(kfd_fd)?; + + // Discover all nodes + let mut cpu = None; + let mut gpu = None; + for node in 0..16 { + match NodeProperties::read(node) { + Ok(p) if p.is_gpu() && gpu.is_none() => gpu = Some(p), + Ok(p) if p.is_cpu() && cpu.is_none() => cpu = Some(p), + _ => {} + } + } + + let cpu = cpu.ok_or_else(|| std::io::Error::new( + std::io::ErrorKind::NotFound, "no CPU node in KFD topology"))?; + let gpu = gpu.ok_or_else(|| std::io::Error::new( + std::io::ErrorKind::NotFound, "no GPU node in KFD topology"))?; + + #[cfg(debug_assertions)] + eprintln!(" KFD {}.{} — HSA topology:", major, minor); + #[cfg(debug_assertions)] + eprintln!(" CPU: {} cores, {} MHz, {} GB RAM (DDR{}-{}, {}-bit)", + cpu.cpu_cores_count, cpu.max_engine_clk, + cpu.mem_size_bytes / (1024 * 1024 * 1024), + if cpu.mem_clk_max > 2400 { "5" } else { "4" }, + cpu.mem_clk_max, cpu.mem_width); + #[cfg(debug_assertions)] + eprintln!(" GPU: gfx{:x}, {} SIMDs, wave{}, {} GB VRAM (GDDR6-{}, {}-bit)", + gpu.gfx_target_version, gpu.simd_count, gpu.wave_front_size, + gpu.mem_size_bytes / (1024 * 1024 * 1024), + gpu.mem_clk_max, gpu.mem_width); + + // Open DRM render node for GPU + let drm_path = format!("/dev/dri/renderD{}\0", gpu.drm_render_minor); + let drm_fd = unsafe { libc::open(drm_path.as_ptr() as _, libc::O_RDWR) }; + if drm_fd < 0 { + unsafe { libc::close(kfd_fd); } + return Err(std::io::Error::last_os_error()); + } + + // Acquire VM for GPU + crate::ioctl::acquire_vm(kfd_fd, drm_fd, gpu.gpu_id) + .map_err(|e| { eprintln!(" acquire_vm failed: {}", e); e })?; + + // Runtime enable (required before queue creation on KFD >= 1.14) + if major > 1 || (major == 1 && minor >= 14) { + crate::ioctl::runtime_enable(kfd_fd) + .map_err(|e| { eprintln!(" runtime_enable failed: {}", e); e })?; + } + + let vram_available = crate::ioctl::available_memory(kfd_fd, gpu.gpu_id).unwrap_or(0); + + let alloc = GpuAllocator { + kfd_fd, + drm_fd, + gpu_id: gpu.gpu_id, + }; + + // Allocate event page (GTT, 0x8000 bytes) — must exist before CREATE_EVENT + let event_page = alloc.alloc_gtt(0x8000)?; + + // Create events before queue (KFD requirement) + // First event with event_page_offset=handle allocates the event page + let _first_event = crate::ioctl::create_event(kfd_fd, crate::ioctl::EVENT_SIGNAL, 0, + event_page.handle)?; + // Queue event (auto_reset=1) — used for interrupt signaling + let queue_event = crate::ioctl::create_event(kfd_fd, crate::ioctl::EVENT_SIGNAL, 1, 0)?; + let _mem_event = crate::ioctl::create_event(kfd_fd, crate::ioctl::EVENT_MEMORY, 0, 0)?; + let _hw_event = crate::ioctl::create_event(kfd_fd, crate::ioctl::EVENT_HW_EXCEPTION, 0, 0)?; + + // Compute mailbox pointer: event_page_va + event_slot_index * 8 + let queue_event_mailbox_ptr = event_page.va_addr + + queue_event.event_slot_index as u64 * 8; + + let queue = ComputeQueue::new(&alloc) + .map_err(|e| { eprintln!(" queue creation failed: {}", e); e })?; + // Scratch buffer for compute dispatch (2MB, VRAM) + let scratch = alloc.alloc_vram(2 * 1024 * 1024)?; + + // Signal buffer must be GPU-writable (flags 0xF0000004 matching tinygrad) + let signal = alloc.alloc_userptr_public(4096)?; + signal.write(0, &[0u8; 8]); + + // load compiled gpu kernels (raw asm) + // each code object owns its GpuBuffer — store them all to keep code in VRAM + let mut code_objects: Vec = Vec::new(); + for co_bytes in [TEST_STORE_CO, MATVEC_CO, MATMUL_BLOCKED_CO, MATMUL_SMALL_CO, SUPERLINEAR_CO] { + match CodeObject::load(&alloc, co_bytes) { + Ok(co) => code_objects.push(co), + Err(e) => eprintln!(" warning: kernel load failed: {}", e), + } + } + // merge kernel entries (they reference code_addr which stays valid while code_objects lives) + let mut kernel_map = std::collections::HashMap::new(); + for co in &code_objects { + for (name, entry) in &co.kernels { + kernel_map.insert(name.clone(), entry.clone()); + } + } + let kernels = if kernel_map.is_empty() { None } else { + #[cfg(debug_assertions)] + eprintln!(" kernels: {:?}", kernel_map.keys().collect::>()); + Some(kernel_map) + }; + + // GPU dispatch threshold: ~64K FLOPs + // Below this, CPU AVX-512 is faster due to dispatch overhead + let gpu_dispatch_threshold = 64 * 1024; + + // lock GPU clocks to max for consistent performance + let sysfs_device = format!("/sys/class/drm/renderD{}/device", gpu.drm_render_minor); + let clocks_locked = FdGuard::lock_clocks(&sysfs_device); + #[cfg(debug_assertions)] + if clocks_locked { + eprintln!(" gpu clocks: locked to max"); + } else { + eprintln!(" gpu clocks: auto (no sysfs write access)"); + } + + Ok(HsaDevice { + kfd_fd, + drm_fd, + cpu, + gpu, + alloc, + _code_objects: code_objects, + kernels, + queue, + signal, + event_page, + scratch, + signal_value: 0, + vram_available, + queue_event_id: queue_event.event_id, + queue_event_mailbox_ptr, + gpu_dispatch_threshold, + _fd_guard: FdGuard { + kfd_fd, drm_fd, + sysfs_device: if clocks_locked { Some(sysfs_device) } else { None }, + }, + }) + } + + /// Allocate a VRAM buffer and upload f32 data. + pub fn upload_f32(&self, data: &[f32]) -> std::io::Result { + // +4096: extra page for RDNA3 global_load_b128 cache line prefetch + // at allocation boundary (kernel prefetch reads 1 tile past last load) + let size = ((data.len() * 4 + 4096 + 4095) & !4095) as u64; + let buf = self.alloc.alloc_vram(size)?; + buf.write_f32(0, data); + Ok(buf) + } + + /// Upload a row-major matrix as column-major (transposed) in VRAM. + /// Input: data[rows][cols] row-major. Output: buf[cols][rows] in VRAM. + /// This gives coalesced access when threads read consecutive row indices + /// from the same column — critical for RDNA3's narrow 128-bit memory bus. + pub fn upload_f32_col_major(&self, data: &[f32], rows: usize, cols: usize) + -> std::io::Result + { + assert_eq!(data.len(), rows * cols); + let size = ((rows * cols * 4 + 4096 + 4095) & !4095) as u64; + let buf = self.alloc.alloc_vram(size)?; + let dst = unsafe { + std::slice::from_raw_parts_mut(buf.cpu_ptr as *mut f32, rows * cols) + }; + for r in 0..rows { + for c in 0..cols { + dst[c * rows + r] = data[r * cols + c]; + } + } + Ok(buf) + } + + /// Allocate a VRAM buffer of given size (bytes), zeroed. + pub fn alloc_output(&self, size_bytes: usize) -> std::io::Result { + let size = ((size_bytes + 4095) & !4095) as u64; + let buf = self.alloc.alloc_vram(size)?; + unsafe { std::ptr::write_bytes(buf.cpu_ptr, 0, size as usize); } + Ok(buf) + } + + /// dispatch a kernel and return a future (non-blocking). + pub fn dispatch_async(&mut self, name: &str, + kernargs: &GpuBuffer, + grid: [u32; 3], block: [u32; 3]) -> Option { + let entry = match self.kernels.as_ref().and_then(|m| m.get(name)) { + Some(e) => e.clone(), + None => return None, + }; + self.signal_value += 1; + self.queue.dispatch_lds( + entry.code_addr, + entry.desc.pgm_rsrc1, entry.desc.pgm_rsrc2, entry.desc.pgm_rsrc3, + kernargs.va_addr, self.scratch.va_addr, + grid, block, + entry.desc.group_segment_size, + ); + self.queue.signal(self.signal.va_addr, self.signal_value, + self.queue_event_mailbox_ptr, self.queue_event_id); + self.queue.submit(); + Some(crate::compute::GpuFuture::new(&self.signal, self.signal_value)) + } + + /// Queue a dispatch without signaling or submitting. + /// Call `submit_wait()` after all dispatches are queued. + pub fn dispatch_enqueue(&mut self, name: &str, + kernargs: &GpuBuffer, + grid: [u32; 3], block: [u32; 3]) -> bool { + let entry = match self.kernels.as_ref().and_then(|m| m.get(name)) { + Some(e) => e.clone(), + None => return false, + }; + self.queue.dispatch_lds( + entry.code_addr, + entry.desc.pgm_rsrc1, entry.desc.pgm_rsrc2, entry.desc.pgm_rsrc3, + kernargs.va_addr, self.scratch.va_addr, + grid, block, + entry.desc.group_segment_size, + ); + // CS_PARTIAL_FLUSH between dispatches (same as after dispatch in queue.rs) + true + } + + /// Signal + submit + wait via kfd event interrupt. + /// Call after dispatch_enqueue batch. Blocks until all dispatches complete. + pub fn submit_wait(&mut self, timeout_ms: u32) -> bool { + self.signal_value += 1; + self.queue.signal(self.signal.va_addr, self.signal_value, + self.queue_event_mailbox_ptr, self.queue_event_id); + self.queue.submit(); + self.wait_gpu(timeout_ms) + } + + /// Submit all queued packets and return a future for the last signal. + pub fn submit(&mut self) -> crate::compute::GpuFuture { + self.signal_value += 1; + self.queue.signal(self.signal.va_addr, self.signal_value, + self.queue_event_mailbox_ptr, self.queue_event_id); + self.queue.submit(); + crate::compute::GpuFuture::new(&self.signal, self.signal_value) + } + + /// Wait for the last signal via kfd event interrupt. + /// The GPU's RELEASE_MEM packet fires an interrupt when the signal + /// is written. The kfd driver wakes us — zero PCIe round-trips. + pub fn wait_gpu(&mut self, timeout_ms: u32) -> bool { + // fast path: check if already done (single volatile read) + let sig_ptr = self.signal.cpu_ptr as *const u32; + let val = unsafe { sig_ptr.read_volatile() }; + if val >= self.signal_value { return true; } + + let mut event = crate::ioctl::EventData::default(); + event.event_id = self.queue_event_id; + match crate::ioctl::wait_events(self.kfd_fd, &mut [event], false, timeout_ms) { + Ok(_) => { + let v = unsafe { sig_ptr.read_volatile() }; + if v < self.signal_value { + // event fired but signal not ready — brief spin + for _ in 0..10_000_000 { + let v2 = unsafe { sig_ptr.read_volatile() }; + if v2 >= self.signal_value { return true; } + std::hint::spin_loop(); + } + let v3 = unsafe { sig_ptr.read_volatile() }; + eprintln!("wait_gpu: event ok but signal={} < {}", v3, self.signal_value); + false + } else { + true + } + } + Err(e) => { + let val = unsafe { sig_ptr.read_volatile() }; + eprintln!("wait_gpu: err={} signal={} expected={}", e, val, self.signal_value); + val >= self.signal_value + } + } + } + + /// dispatch a kernel by name and wait for completion (blocking). + pub fn dispatch_kernel(&mut self, name: &str, + kernargs: &GpuBuffer, + grid: [u32; 3], block: [u32; 3]) -> bool { + let entry = match self.kernels.as_ref().and_then(|m| m.get(name)) { + Some(e) => e.clone(), + None => return false, + }; + self.signal_value += 1; + + // Record put position before dispatch for ring dump + let put_before = self.queue.put; + + self.queue.dispatch_lds( + entry.code_addr, + entry.desc.pgm_rsrc1, + entry.desc.pgm_rsrc2, + entry.desc.pgm_rsrc3, + kernargs.va_addr, + self.scratch.va_addr, + grid, block, + entry.desc.group_segment_size, + ); + self.queue.signal(self.signal.va_addr, self.signal_value, + self.queue_event_mailbox_ptr, self.queue_event_id); + + // Dump ring contents before submit (debug) + #[cfg(debug_assertions)] + { + let ring = self.queue.ring.as_slice::(); + let n = (self.queue.put - put_before) as usize; + eprintln!(" dispatch PM4 ({} DWORDs, code=0x{:x} kargs=0x{:x}):", + n, entry.code_addr, kernargs.va_addr); + for i in 0..n { + let idx = (put_before as usize + i) % (self.queue.ring.size as usize / 4); + eprintln!(" [{:3}] 0x{:08x}", i, ring[idx]); + } + } + + self.queue.submit(); + self.wait_gpu(10_000) + } + + /// Enqueue Y = X * W^T + B, selecting the optimal kernel for the shape. + /// W is row-major [M×K], X is row-major [N×K], Y is [N×M]. + /// Uses matmul_blocked (TM=128) for large M, matmul_small (TM=32) for small M. + /// + /// SAFETY: M, K, N must be padded to tile boundaries BEFORE calling this. + /// - M must be a multiple of 128 (matmul_blocked) or 32 (matmul_small) + /// - K must be a multiple of 8 + /// - N must be a multiple of 32 + /// - All buffers (W, B, X, Y) must be allocated for the padded dimensions. + /// Violating this will cause GPU OOB access → device hang → Xorg crash. + /// Use `dispatch_matmul_safe` for automatic padding. + pub fn dispatch_matmul_enqueue(&mut self, + w_row: &GpuBuffer, w_col: &GpuBuffer, + b: &GpuBuffer, x: &GpuBuffer, y: &GpuBuffer, + m: u32, k: u32, n: u32) -> bool { + // Guard: reject dimensions that would cause OOB + if k % 8 != 0 || n % 32 != 0 { + eprintln!("dispatch_matmul_enqueue: REJECTING unaligned dims M={} K={} N={}", m, k, n); + return false; + } + if m >= 1536 && m % 128 != 0 { + eprintln!("dispatch_matmul_enqueue: REJECTING M={} (not multiple of 128 for blocked)", m); + return false; + } + if m < 1536 && m % 32 != 0 { + eprintln!("dispatch_matmul_enqueue: REJECTING M={} (not multiple of 32 for small)", m); + return false; + } + + let (kernel, nwg) = if m >= 1536 { + ("matmul_blocked", ((m + 127) / 128) * ((n + 31) / 32)) + } else { + ("matmul_small", ((m + 31) / 32) * ((n + 31) / 32)) + }; + + let w = if m >= 1536 { w_row } else { w_col }; + + let mut args = KernArgs::new(); + args.push_ptr(w); args.push_ptr(b); + args.push_ptr(x); args.push_ptr(y); + args.push_u32(m); args.push_u32(k); args.push_u32(n); + let args_buf = match args.upload(&self.alloc) { + Ok(buf) => buf, + Err(_) => return false, + }; + + self.dispatch_enqueue(kernel, &args_buf, [nwg, 1, 1], [256, 1, 1]) + } + + /// Enqueue a SuperLinear forward: Y[n*O+o] = B[n*O+o] + sum_k W[..]*X[..]. + /// N=n_neurons, O=out_per, K=in_per. All buffers sized for N*O*K etc. + /// Returns false if kernel not loaded. + pub fn dispatch_superlinear_enqueue(&mut self, + w: &GpuBuffer, b: &GpuBuffer, + x: &GpuBuffer, y: &GpuBuffer, + n_neurons: u32, out_per: u32, in_per: u32) -> bool { + let total = n_neurons * out_per; + let nwg = (total + 255) / 256; + + let mut args = KernArgs::new(); + args.push_ptr(w); args.push_ptr(b); + args.push_ptr(x); args.push_ptr(y); + args.push_u32(n_neurons); args.push_u32(out_per); args.push_u32(in_per); + let args_buf = match args.upload(&self.alloc) { + Ok(buf) => buf, + Err(_) => return false, + }; + + self.dispatch_enqueue("superlinear_fwd", &args_buf, [nwg, 1, 1], [256, 1, 1]) + } + + /// Dispatch a GpuProgram directly and wait for completion. + pub fn dispatch_sync(&mut self, program: &GpuProgram, + kernargs: &GpuBuffer, + grid: [u32; 3], block: [u32; 3]) { + self.signal_value += 1; + program.dispatch(&mut self.queue, kernargs, self.scratch.va_addr, + grid, block, &self.signal, self.signal_value, + self.queue_event_mailbox_ptr, self.queue_event_id); + if !self.wait_gpu(10_000) { + eprintln!("WARNING: GPU dispatch timed out after 10s"); + } + } + + /// Load a code object into GPU memory. + pub fn load_program(&self, code_object: &[u8]) -> std::io::Result { + GpuProgram::load(&self.alloc, code_object) + } + + /// read current GPU shader clock in mhz from sysfs. + pub fn current_sclk_mhz(&self) -> Option { + let path = format!("/sys/class/drm/renderD{}/device/pp_dpm_sclk", + self.gpu.drm_render_minor); + let text = std::fs::read_to_string(&path).ok()?; + // active level marked with * + text.lines() + .find(|l| l.contains('*')) + .and_then(|l| { + // "2: 2330Mhz *" → extract 2330 + l.split(':').nth(1) + .and_then(|s| s.trim().split('M').next()) + .and_then(|s| s.trim().parse().ok()) + }) + } + + /// Should this operation go to GPU? Based on estimated FLOPs. + pub fn should_use_gpu(&self, flops: usize) -> bool { + flops >= self.gpu_dispatch_threshold + } + + /// Get static device specifications (from the already-read topology). + pub fn specs(&self) -> DeviceSpecs { + let vram_total = self.gpu.mem_size_bytes; + let vram_width = self.gpu.mem_width; + let vram_clock = self.gpu.mem_clk_max; + let bw = (vram_width as f32 / 8.0) * (vram_clock as f32 * 2.0) / 1000.0; + + DeviceSpecs { + vram_total, + vram_width, + vram_clock_mhz: vram_clock, + vram_bandwidth_gbps: bw, + ram_total: self.cpu.mem_size_bytes, + gfx_version: self.gpu.gfx_target_version, + compute_units: if self.gpu.cu_per_simd_array > 0 { + self.gpu.simd_count * self.gpu.cu_per_simd_array + } else { + self.gpu.simd_count + }, + cpu: self.cpu.clone(), + gpu: self.gpu.clone(), + } + } + + /// Query live device status (VRAM usage, clock speed). + pub fn status(&self) -> DeviceStatus { + let vram_available = crate::ioctl::available_memory(self.kfd_fd, self.gpu.gpu_id) + .unwrap_or(self.vram_available); + let vram_total = self.gpu.mem_size_bytes; + + DeviceStatus { + vram_available, + vram_total, + vram_used: vram_total.saturating_sub(vram_available), + sclk_mhz: self.current_sclk_mhz(), + } + } +} + +/// Guard that closes fds on drop. Placed as LAST field in HsaDevice +/// so it drops after all GpuBuffers (which need the fds for cleanup). +struct FdGuard { + kfd_fd: RawFd, + drm_fd: RawFd, + /// sysfs device path for clock control (e.g. /sys/class/drm/renderD128/device) + sysfs_device: Option, +} + +impl FdGuard { + /// Lock GPU to peak performance (max clocks, no DVFS). + fn lock_clocks(sysfs_device: &str) -> bool { + let perf_path = format!("{}/power_dpm_force_performance_level", sysfs_device); + // profile_peak forces max clocks immediately — no DVFS ramp needed + std::fs::write(&perf_path, "profile_peak").is_ok() + } + + fn unlock_clocks(sysfs_device: &str) { + let perf_path = format!("{}/power_dpm_force_performance_level", sysfs_device); + let _ = std::fs::write(&perf_path, "auto"); + } +} + +impl Drop for FdGuard { + fn drop(&mut self) { + if let Some(ref path) = self.sysfs_device { + FdGuard::unlock_clocks(path); + } + unsafe { + libc::close(self.drm_fd); + libc::close(self.kfd_fd); + } + } +} + +/// Check if KFD is available on this system. +pub fn is_available() -> bool { + std::path::Path::new("/dev/kfd").exists() +} + +// ─── Unified ComputeBackend ───────────────────────────────── + +// ─── Backend integration ─────────────────────────────────── +// HsaBackend and trait impls (ComputeBackend, TransformerOps) +// live in the downstream crate (e.g., isis) that depends on kfd. +// This crate only exposes HsaDevice for raw GPU access. + +// Backend trait implementations (ComputeBackend, TransformerOps) +// are defined in downstream crates that depend on `kfd`. +// See isis::host::device::kfd for the full integration. + +// ─── Tests ────────────────────────────────────────────────── + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn test_hsa_topology() { + if !is_available() { + eprintln!(" SKIP: /dev/kfd not available"); + return; + } + let dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { eprintln!(" Cannot open: {}", e); return; } + }; + + assert!(dev.cpu.is_cpu()); + assert!(dev.gpu.is_gpu()); + eprintln!(" CPU: {} cores, {} GB", dev.cpu.cpu_cores_count, + dev.cpu.mem_size_bytes / (1024*1024*1024)); + eprintln!(" GPU: {} SIMDs, {} GB VRAM", + dev.gpu.simd_count, dev.gpu.mem_size_bytes / (1024*1024*1024)); + } + + #[test] + fn test_hsa_alloc_vram() { + if !is_available() { return; } + let dev = match HsaDevice::open() { + Ok(d) => d, + Err(_) => return, + }; + + let buf = dev.alloc.alloc_vram(4096).expect("VRAM alloc failed"); + let data: Vec = (0..256).map(|i| i as f32).collect(); + buf.write_f32(0, &data); + let readback = buf.read_f32(0, 256); + assert_eq!(data, readback); + eprintln!(" VRAM write/read: OK"); + } + + #[test] + fn test_hsa_alloc_gtt() { + if !is_available() { return; } + let dev = match HsaDevice::open() { + Ok(d) => d, + Err(_) => return, + }; + + let buf = dev.alloc.alloc_gtt(4096).expect("GTT alloc failed"); + buf.write(0, b"hello hsa"); + assert_eq!(&buf.read(0, 9), b"hello hsa"); + eprintln!(" GTT write/read: OK"); + } + + #[test] + fn test_hsa_queue() { + if !is_available() { return; } + let dev = match HsaDevice::open() { + Ok(d) => d, + Err(_) => return, + }; + eprintln!(" Queue ID: {}", dev.queue.queue_id); + } + + #[test] + fn test_hsa_gpu_store42() { + // end-to-end: dispatch test_store kernel, verify 42.0 in output + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { println!(" SKIP: {}", e); return; }, + }; + if dev.kernels.is_none() { println!(" SKIP: no kernels"); return; } + + let n = 8u32; + + // verify code object in VRAM + { + let co = &dev._code_objects[0]; // test_store + let code_va_offset = 0x1300usize; // .text VA in the image + let code_bytes = co.buf.read(code_va_offset, 16); + println!(" code at va_offset 0x1300: {}", code_bytes.iter().map(|b| format!("{:02x}", b)).collect::>().join(" ")); + let is_zero = code_bytes.iter().all(|&b| b == 0); + println!(" code is zero: {}", is_zero); + } + + // allocate output buffer in VRAM (gpu can write, cpu reads via BAR) + let y_buf = dev.alloc.alloc_vram(4096).unwrap(); + // write sentinel to verify it gets overwritten + let sentinel = vec![-1.0f32; n as usize]; + y_buf.write_f32(0, &sentinel); + + // kernargs: just y pointer (8 bytes) + let mut args = KernArgs::new(); + args.push_ptr(&y_buf); + // pad to kernarg_size (8 bytes for test_store) + let args_buf = args.upload(&dev.alloc).unwrap(); + + // dispatch + let ok = dev.dispatch_kernel("test_store", &args_buf, [n, 1, 1], [n, 1, 1]); + assert!(ok, "test_store dispatch timed out"); + + // read back via cpu_ptr (direct BAR read, no DMA) + let y = y_buf.read_f32(0, n as usize); + println!(" gpu output: {:?}", y); + for i in 0..n as usize { + assert!((y[i] - 42.0).abs() < 0.01, + "y[{}] = {} expected 42.0", i, y[i]); + } + println!(" test_store: CORRECT — gpu wrote 42.0 to {} elements", n); + } + + #[test] + #[ignore] + fn test_hsa_signal_only() { + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { println!(" OPEN FAILED: {}", e); return; }, + }; + + // Just send release_mem to write a signal — no compute dispatch + dev.signal.write(0, &[0u8; 4]); // zero the signal + dev.signal_value = 42; + + println!(" signal_addr: 0x{:x}", dev.signal.va_addr); + println!(" mailbox_ptr: 0x{:x}", dev.queue_event_mailbox_ptr); + println!(" event_id: {}", dev.queue_event_id); + println!(" ring_addr: 0x{:x}", dev.queue.ring.va_addr); + println!(" put before: {}", dev.queue.put); + + dev.queue.signal(dev.signal.va_addr, 42, + dev.queue_event_mailbox_ptr, dev.queue_event_id); + + println!(" put after signal: {}", dev.queue.put); + + // Dump ring contents before submit + let ring_u32 = dev.queue.ring.as_slice::(); + println!(" ring contents (first {} DWORDs):", dev.queue.put); + for i in 0..dev.queue.put as usize { + println!(" [{:3}] 0x{:08x}", i, ring_u32[i]); + } + + dev.queue.submit(); + + let signal_ptr = dev.signal.cpu_ptr as *const u32; + let ok = ComputeQueue::poll_signal(signal_ptr, 42, 2_000_000); + let val = unsafe { signal_ptr.read_volatile() }; + println!(" signal_only: ok={} val={}", ok, val); + assert!(ok, "signal-only dispatch timed out (val={})", val); + } + + #[test] + fn test_hsa_backend_matvec() { + if !is_available() { return; } + let backend = match HsaBackend::new() { + Ok(b) => b, + Err(_) => return, + }; + + // Small matvec → should go to CPU (AVX-512) + let weight = vec![1.0f32, 2.0, 3.0, 4.0]; // 2x2 + let bias = vec![0.5, 0.5]; + let x = vec![1.0, 1.0]; + let mut y = vec![0.0; 2]; + backend.matvec(&weight, &bias, &x, &mut y, 2, 2); + assert!((y[0] - 3.5).abs() < 1e-5); // 1*1 + 2*1 + 0.5 + assert!((y[1] - 7.5).abs() < 1e-5); // 3*1 + 4*1 + 0.5 + eprintln!(" HsaBackend matvec (CPU): OK"); + } + + #[test] + fn test_hsa_kernels_loaded() { + if !is_available() { return; } + let dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { println!(" OPEN FAILED: {}", e); return; }, + }; + + let km = dev.kernels.as_ref().expect("kernels not loaded"); + assert!(km.contains_key("test_store"), "test_store kernel missing"); + assert!(km.contains_key("matvec"), "matvec kernel missing"); + + let mv = &km["matvec"]; + println!(" matvec: rsrc1=0x{:08x} rsrc2=0x{:08x} code=0x{:x} karg={}", + mv.desc.pgm_rsrc1, mv.desc.pgm_rsrc2, mv.code_addr, mv.desc.kernarg_size); + let ts = &km["test_store"]; + println!(" test_store: rsrc1=0x{:08x} code=0x{:x} karg={}", + ts.desc.pgm_rsrc1, ts.code_addr, ts.desc.kernarg_size); + println!(" {} kernels: {:?}", km.len(), km.keys().collect::>()); + } + + #[test] + fn test_hsa_dispatch_dryrun() { + // Dump PM4 packets WITHOUT submitting — safe, no GPU crash + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { println!(" OPEN FAILED: {}", e); return; }, + }; + if dev.kernels.is_none() { println!(" NO KERNELS"); return; } + + let out_dim = 64u32; + let in_dim = 64u32; + let weight = vec![0.0f32; (out_dim * in_dim) as usize]; + let bias = vec![0.0f32; out_dim as usize]; + let x = vec![0.0f32; in_dim as usize]; + + let w_buf = dev.upload_f32(&weight).unwrap(); + let b_buf = dev.upload_f32(&bias).unwrap(); + let x_buf = dev.upload_f32(&x).unwrap(); + let y_buf = dev.alloc_output(out_dim as usize * 4).unwrap(); + + let grid = [out_dim, 1, 1]; + let block = [64.min(out_dim), 1, 1]; + let kernarg_size = dev.kernels.as_ref().unwrap() + ["matvec"].desc.kernarg_size; + + let mut args = KernArgs::new(); + args.push_ptr(&w_buf); + args.push_ptr(&b_buf); + args.push_ptr(&x_buf); + args.push_ptr(&y_buf); + args.push_u32(out_dim); + args.push_u32(in_dim); + args.fill_hidden_args(grid, block, kernarg_size); + let args_buf = args.upload(&dev.alloc).unwrap(); + + let entry = dev.kernels.as_ref().unwrap()["matvec"].clone(); + + // Build packets but DON'T submit + dev.queue.dispatch( + entry.code_addr, + entry.desc.pgm_rsrc1, + entry.desc.pgm_rsrc2, + entry.desc.pgm_rsrc3, + args_buf.va_addr, + dev.scratch.va_addr, + grid, block, + ); + + // Dump + let ring = dev.queue.ring.as_slice::(); + let n = dev.queue.put as usize; + println!(" Our PM4 dispatch ({} DWORDs):", n); + for i in 0..n { + println!(" [{:3}] 0x{:08x}", i, ring[i]); + } + println!(" code_addr: 0x{:x}", entry.code_addr); + println!(" kernargs: 0x{:x}", args_buf.va_addr); + println!(" rsrc1: 0x{:08x} rsrc2: 0x{:08x} rsrc3: 0x{:08x}", + entry.desc.pgm_rsrc1, entry.desc.pgm_rsrc2, entry.desc.pgm_rsrc3); + + // Don't submit — just verify packet structure + assert!(n > 0, "no packets generated"); + } + + #[test] + fn test_hsa_gpu_matvec_dispatch() { + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(e) => { println!(" SKIP: {}", e); return; }, + }; + if dev.kernels.is_none() { println!(" SKIP: no kernels"); return; } + + let out_dim = 4u32; + let in_dim = 4u32; + + // W = identity matrix + let mut weight = vec![0.0f32; (out_dim * in_dim) as usize]; + for i in 0..out_dim as usize { + weight[i * in_dim as usize + i] = 1.0; + } + let bias = vec![0.5f32; out_dim as usize]; + let x: Vec = (1..=in_dim).map(|i| i as f32).collect(); // [1,2,3,4] + + let w_buf = dev.upload_f32(&weight).unwrap(); + let b_buf = dev.upload_f32(&bias).unwrap(); + let x_buf = dev.upload_f32(&x).unwrap(); + let y_buf = dev.alloc.alloc_userptr_public(4096).unwrap(); + + // raw asm kernel: 40 bytes kernargs, no hidden args + let grid = [out_dim, 1, 1]; + let block = [out_dim, 1, 1]; + + let mut args = KernArgs::new(); + args.push_ptr(&w_buf); + args.push_ptr(&b_buf); + args.push_ptr(&x_buf); + args.push_ptr(&y_buf); + args.push_u32(out_dim); + args.push_u32(in_dim); + let args_buf = args.upload(&dev.alloc).unwrap(); + + // Dispatch + let ok = dev.dispatch_kernel("matvec", &args_buf, grid, block); + assert!(ok, "GPU matvec dispatch timed out"); + + // Read back and verify + let y = y_buf.read_f32(0, out_dim as usize); + println!(" gpu matvec result: {:?}", y); + // identity * [1,2,3,4] + 0.5 = [1.5, 2.5, 3.5, 4.5] + let expected: Vec = (1..=out_dim).map(|i| i as f32 + 0.5).collect(); + for i in 0..out_dim as usize { + assert!((y[i] - expected[i]).abs() < 1e-3, + "y[{}] = {} expected {}", i, y[i], expected[i]); + } + println!(" gpu matvec: CORRECT ({}x{})", out_dim, in_dim); + } + + #[test] + #[ignore] // DISABLED: PM4 packets crash GPU + fn test_hsa_gpu_matvec_large() { + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(_) => return, + }; + if dev.kernels.is_none() { return; } + + // CTM-sized: 1024x1024 matvec + let out_dim = 1024u32; + let in_dim = 1024u32; + let n = (out_dim * in_dim) as usize; + + // Random-ish weights + let weight: Vec = (0..n).map(|i| ((i * 7 + 3) % 100) as f32 / 100.0 - 0.5).collect(); + let bias = vec![0.1f32; out_dim as usize]; + let x: Vec = (0..in_dim).map(|i| ((i * 13 + 5) % 100) as f32 / 100.0).collect(); + + // CPU reference + let mut y_cpu = vec![0.0f32; out_dim as usize]; + for row in 0..out_dim as usize { + let mut sum = bias[row]; + for col in 0..in_dim as usize { + sum += weight[row * in_dim as usize + col] * x[col]; + } + y_cpu[row] = sum; + } + + // GPU + let w_buf = dev.upload_f32(&weight).unwrap(); + let b_buf = dev.upload_f32(&bias).unwrap(); + let x_buf = dev.upload_f32(&x).unwrap(); + let y_buf = dev.alloc_output(out_dim as usize * 4).unwrap(); + + let mut args = KernArgs::new(); + args.push_ptr(&w_buf); + args.push_ptr(&b_buf); + args.push_ptr(&x_buf); + args.push_ptr(&y_buf); + args.push_u32(out_dim); + args.push_u32(in_dim); + let args_buf = args.upload(&dev.alloc).unwrap(); + + let ok = dev.dispatch_kernel("matvec", &args_buf, + [out_dim, 1, 1], [256, 1, 1]); + assert!(ok, "GPU matvec dispatch timed out"); + + let y_gpu = y_buf.read_f32(0, out_dim as usize); + + // Verify + let mut max_err = 0.0f32; + for i in 0..out_dim as usize { + let err = (y_gpu[i] - y_cpu[i]).abs(); + max_err = max_err.max(err); + } + eprintln!(" GPU matvec 1024x1024: max_err={:.6}", max_err); + assert!(max_err < 0.1, "GPU matvec too inaccurate: max_err={}", max_err); + } + + #[test] + #[ignore] // DISABLED: PM4 packets crash GPU + fn test_hsa_performance_comparison() { + if !is_available() { return; } + let mut dev = match HsaDevice::open() { + Ok(d) => d, + Err(_) => return, + }; + if dev.kernels.is_none() { return; } + + use std::time::Instant; + + // Simple CPU matvec for comparison + fn cpu_matvec(w: &[f32], b: &[f32], x: &[f32], y: &mut [f32], out: usize, inp: usize) { + for r in 0..out { + let mut s = b[r]; + for c in 0..inp { s += w[r * inp + c] * x[c]; } + y[r] = s; + } + } + + eprintln!("\n ┌─────────────────────────────────────────────────┐"); + eprintln!(" │ HSA Performance: CPU (AVX-512) vs GPU (RDNA3) │"); + eprintln!(" ├──────────┬──────────┬──────────┬────────────────┤"); + eprintln!(" │ Size │ CPU (μs) │ GPU (μs) │ Winner │"); + eprintln!(" ├──────────┼──────────┼──────────┼────────────────┤"); + + for &(out_dim, in_dim) in &[ + (64, 64), (128, 128), (256, 256), (512, 512), + (1024, 1024), (2048, 2048), (4096, 1024), + ] { + let n = out_dim * in_dim; + let weight: Vec = (0..n).map(|i| (i % 100) as f32 / 100.0 - 0.5).collect(); + let bias = vec![0.1f32; out_dim]; + let x: Vec = (0..in_dim).map(|i| (i % 100) as f32 / 100.0).collect(); + let mut y = vec![0.0f32; out_dim]; + + // CPU benchmark (10 iterations, take median) + let mut cpu_times = Vec::new(); + for _ in 0..10 { + let t = Instant::now(); + cpu_matvec(&weight, &bias, &x, &mut y, out_dim, in_dim); + cpu_times.push(t.elapsed().as_micros() as u64); + } + cpu_times.sort(); + let cpu_us = cpu_times[5]; // median + + // GPU benchmark + let w_buf = dev.upload_f32(&weight).unwrap(); + let b_buf = dev.upload_f32(&bias).unwrap(); + let x_buf = dev.upload_f32(&x).unwrap(); + let y_buf = dev.alloc_output(out_dim * 4).unwrap(); + + let mut args = KernArgs::new(); + args.push_ptr(&w_buf); + args.push_ptr(&b_buf); + args.push_ptr(&x_buf); + args.push_ptr(&y_buf); + args.push_u32(out_dim as u32); + args.push_u32(in_dim as u32); + let args_buf = args.upload(&dev.alloc).unwrap(); + + let block_size = 256.min(out_dim as u32); + + // Warmup + dev.dispatch_kernel("matvec", &args_buf, + [out_dim as u32, 1, 1], [block_size, 1, 1]); + + let mut gpu_times = Vec::new(); + for _ in 0..10 { + // Zero signal for fresh measurement + dev.signal.write(0, &(dev.signal_value + 1).to_le_bytes()[..4]); + let t = Instant::now(); + dev.dispatch_kernel("matvec", &args_buf, + [out_dim as u32, 1, 1], [block_size, 1, 1]); + gpu_times.push(t.elapsed().as_micros() as u64); + } + gpu_times.sort(); + let gpu_us = gpu_times[5]; // median + + let winner = if gpu_us < cpu_us { "GPU ←" } else { "CPU" }; + let flops = out_dim * in_dim * 2; + eprintln!(" │ {:4}x{:<4}│ {:>8} │ {:>8} │ {:>14} │", + out_dim, in_dim, cpu_us, gpu_us, winner); + } + eprintln!(" └──────────┴──────────┴──────────┴────────────────┘"); + } +} diff --git a/src/dispatch.rs b/src/dispatch.rs new file mode 100644 index 0000000..4b84696 --- /dev/null +++ b/src/dispatch.rs @@ -0,0 +1,413 @@ +//! Kernel dispatch: load code objects, set up kernel arguments, run. +//! +//! AMD code objects are ELF files containing multiple kernels. +//! Each kernel has a `.kd` descriptor in `.rodata` and ISA code in `.text`. +//! Symbol table maps kernel names to their descriptors. + +use crate::memory::{GpuAllocator, GpuBuffer}; +use crate::queue::ComputeQueue; +use std::collections::HashMap; + +/// Parsed kernel descriptor from an AMD code object. +#[derive(Debug, Clone)] +pub struct KernelDescriptor { + pub pgm_rsrc1: u32, + pub pgm_rsrc2: u32, + pub pgm_rsrc3: u32, + pub kernel_code_entry_offset: i64, + pub group_segment_size: u32, + pub private_segment_size: u32, + pub kernarg_size: u32, +} + +/// A single kernel within a loaded code object. +#[derive(Debug, Clone)] +pub struct KernelEntry { + pub name: String, + pub desc: KernelDescriptor, + pub code_addr: u64, +} + +/// A loaded code object in GPU memory, containing one or more kernels. +pub struct CodeObject { + pub buf: GpuBuffer, + pub kernels: HashMap, +} + +impl CodeObject { + /// Parse only (no GPU upload) — for diagnostics. + pub fn parse_only(elf_data: &[u8]) -> usize { + match parse_elf_kernels(elf_data, 0) { + Some(v) => v.len(), + None => 0, + } + } + + /// Load a code object ELF into GPU memory and parse all kernels. + /// + /// The ELF is loaded as a flat VA-indexed image: we allocate enough space + /// for the highest VA in the ELF, then copy each LOAD segment to its VA offset. + /// This way VA addresses in the code work correctly as offsets from gpu_base. + pub fn load(alloc: &GpuAllocator, elf_data: &[u8]) -> std::io::Result { + // Build VA-indexed image (like tinygrad's elf_loader) + let image = build_va_image(elf_data) + .ok_or_else(|| std::io::Error::new(std::io::ErrorKind::InvalidData, + "failed to parse ELF"))?; + + let size = ((image.len() + 4095) & !4095) as u64; + // VRAM with PUBLIC — code lives in GPU memory, CPU-accessible via BAR for upload + let buf = alloc.alloc_vram(size)?; + buf.write(0, &image); + + let kernels = parse_elf_kernels(elf_data, buf.va_addr) + .ok_or_else(|| std::io::Error::new(std::io::ErrorKind::InvalidData, + "failed to parse code object ELF"))?; + + let mut map = HashMap::new(); + for k in kernels { + map.insert(k.name.clone(), k); + } + + Ok(CodeObject { buf, kernels: map }) + } + + /// Get a kernel by name. + pub fn kernel(&self, name: &str) -> Option<&KernelEntry> { + self.kernels.get(name) + } +} + +/// A ready-to-dispatch GPU program (single kernel from a code object). +pub struct GpuProgram { + pub code_buf: GpuBuffer, + pub desc: KernelDescriptor, + pub code_addr: u64, +} + +impl GpuProgram { + /// Load a single-kernel code object. + pub fn load(alloc: &GpuAllocator, code_object: &[u8]) -> std::io::Result { + let co = CodeObject::load(alloc, code_object)?; + let first = co.kernels.values().next() + .ok_or_else(|| std::io::Error::new(std::io::ErrorKind::InvalidData, + "no kernels in code object"))?; + Ok(GpuProgram { + desc: first.desc.clone(), + code_addr: first.code_addr, + code_buf: co.buf, + }) + } + + /// Dispatch and signal. + pub fn dispatch(&self, queue: &mut ComputeQueue, + kernargs: &GpuBuffer, scratch_addr: u64, + grid: [u32; 3], block: [u32; 3], + signal: &GpuBuffer, signal_value: u32, + event_mailbox_ptr: u64, event_id: u32) { + queue.dispatch_lds( + self.code_addr, + self.desc.pgm_rsrc1, + self.desc.pgm_rsrc2, + self.desc.pgm_rsrc3, + kernargs.va_addr, + scratch_addr, + grid, block, + self.desc.group_segment_size, + ); + queue.signal(signal.va_addr, signal_value, event_mailbox_ptr, event_id); + queue.submit(); + } +} + +// ─── ELF parsing ──────────────────────────────────────────── + +/// ELF64 symbol types +const STT_OBJECT: u8 = 1; +const STT_FUNC: u8 = 2; + +/// Parse all kernels from an ELF code object. +/// Each kernel has a `.kd` OBJECT symbol in .rodata (the descriptor) +/// and a `` FUNC symbol in .text (the code entry). +fn parse_elf_kernels(elf: &[u8], gpu_base: u64) -> Option> { + if elf.len() < 64 || &elf[0..4] != b"\x7fELF" { return None; } + + let e_shoff = r64(elf, 40) as usize; + let e_shentsize = r16(elf, 58) as usize; + let e_shnum = r16(elf, 60) as usize; + let e_shstrndx = r16(elf, 62) as usize; + + // Get section name string table + let shstrtab_sh = e_shoff + e_shstrndx * e_shentsize; + let shstrtab_off = r64(elf, shstrtab_sh + 24) as usize; + + // Find section headers by name + let mut symtab_off = 0usize; + let mut symtab_size = 0usize; + let mut symtab_entsize = 0usize; + let mut symtab_link = 0usize; + let mut rodata_off = 0usize; + let mut rodata_addr = 0u64; + + for i in 0..e_shnum { + let sh = e_shoff + i * e_shentsize; + if sh + e_shentsize > elf.len() { break; } + + let sh_name_idx = r32(elf, sh) as usize; + let sh_type = r32(elf, sh + 4); + let sh_addr = r64(elf, sh + 16); + let sh_offset = r64(elf, sh + 24) as usize; + let sh_size = r64(elf, sh + 32) as usize; + + let name = read_str(elf, shstrtab_off + sh_name_idx); + + if sh_type == 2 { // SHT_SYMTAB + symtab_off = sh_offset; + symtab_size = sh_size; + symtab_entsize = r64(elf, sh + 56) as usize; + symtab_link = r32(elf, sh + 40) as usize; + } + if name == ".rodata" { + rodata_off = sh_offset; + rodata_addr = sh_addr; + } + } + + if symtab_entsize == 0 { return None; } + + // Get strtab + let strtab_sh = e_shoff + symtab_link * e_shentsize; + let strtab_off = r64(elf, strtab_sh + 24) as usize; + + // Parse symbols: collect .kd descriptors and FUNC addresses + let mut kd_map: HashMap = HashMap::new(); // name -> (vaddr, file_offset) + let mut func_map: HashMap = HashMap::new(); // name -> vaddr + + let n_syms = symtab_size / symtab_entsize; + for i in 0..n_syms { + let sym = symtab_off + i * symtab_entsize; + if sym + symtab_entsize > elf.len() { break; } + + let st_name = r32(elf, sym) as usize; + let st_info = elf[sym + 4]; + let st_value = r64(elf, sym + 8); + let _st_size = r64(elf, sym + 16); + let st_type = st_info & 0xf; + + let name = read_str(elf, strtab_off + st_name); + + if st_type == STT_OBJECT && name.ends_with(".kd") { + let kernel_name = name[..name.len() - 3].to_string(); + // .kd symbol value is the VA in the ELF. Since we upload VA-indexed image, + // the file_off in the original ELF for reading the descriptor is rodata_off + (va - rodata_addr) + // but the GPU offset is just the VA itself (since image is VA-indexed). + let elf_file_off = rodata_off + (st_value - rodata_addr) as usize; + kd_map.insert(kernel_name, (st_value, elf_file_off)); + } else if st_type == STT_FUNC && !name.starts_with("__clang") { + func_map.insert(name.to_string(), st_value); + } + } + + // Match descriptors with code addresses + let mut kernels = Vec::new(); + for (name, (kd_vaddr, kd_file_off)) in &kd_map { + if kd_file_off + 64 > elf.len() { continue; } + let kd = &elf[*kd_file_off..]; + + // AMDHSA kernel descriptor layout (verified against tinygrad): + // +00: group_segment_fixed_size (u32) + // +04: private_segment_fixed_size (u32) + // +08: kernarg_size (u32) + // +0c: reserved (u32) + // +10: kernel_code_entry_byte_offset (i64) + // +18: reserved (20 bytes) + // +2c: compute_pgm_rsrc3 (u32) + // +30: compute_pgm_rsrc1 (u32) + // +34: compute_pgm_rsrc2 (u32) + // +38: kernel_code_properties (u16) + kernarg_preload (u16) + let desc = KernelDescriptor { + group_segment_size: u32::from_le_bytes(kd[0x00..0x04].try_into().unwrap()), + private_segment_size: u32::from_le_bytes(kd[0x04..0x08].try_into().unwrap()), + kernarg_size: u32::from_le_bytes(kd[0x08..0x0c].try_into().unwrap()), + kernel_code_entry_offset: i64::from_le_bytes(kd[0x10..0x18].try_into().unwrap()), + pgm_rsrc3: u32::from_le_bytes(kd[0x2c..0x30].try_into().unwrap()), + pgm_rsrc1: u32::from_le_bytes(kd[0x30..0x34].try_into().unwrap()), + pgm_rsrc2: u32::from_le_bytes(kd[0x34..0x38].try_into().unwrap()), + }; + + // Code address: since we upload VA-indexed image, gpu_base + kd_vaddr = descriptor addr. + // kernel_code_entry_offset is relative to the descriptor's VA. + let kd_gpu_addr = gpu_base + *kd_vaddr; + let code_addr = (kd_gpu_addr as i64 + desc.kernel_code_entry_offset) as u64; + + eprintln!(" kernel '{}': kd_va=0x{:x} file_off=0x{:x} entry_off=0x{:x} code=0x{:x}", + name, kd_vaddr, kd_file_off, desc.kernel_code_entry_offset, code_addr); + + kernels.push(KernelEntry { + name: name.clone(), + desc, + code_addr, + }); + } + + Some(kernels) +} + +/// Build a VA-indexed image from an ELF: allocate buffer sized to max VA, +/// copy each LOAD segment to its VA offset. This way gpu_base + VA = correct GPU address. +fn build_va_image(elf: &[u8]) -> Option> { + if elf.len() < 64 || &elf[0..4] != b"\x7fELF" { return None; } + + let e_phoff = r64(elf, 32) as usize; + let e_phentsize = r16(elf, 54) as usize; + let e_phnum = r16(elf, 56) as usize; + + // Find max VA across all LOAD segments + let mut max_va: usize = 0; + for i in 0..e_phnum { + let ph = e_phoff + i * e_phentsize; + if ph + e_phentsize > elf.len() { break; } + let p_type = r32(elf, ph); + if p_type != 1 { continue; } // PT_LOAD = 1 + let p_vaddr = r64(elf, ph + 16) as usize; + let p_memsz = r64(elf, ph + 40) as usize; + max_va = max_va.max(p_vaddr + p_memsz); + } + + if max_va == 0 { return None; } + let mut image = vec![0u8; max_va]; + + // Copy each LOAD segment to its VA position + for i in 0..e_phnum { + let ph = e_phoff + i * e_phentsize; + if ph + e_phentsize > elf.len() { break; } + let p_type = r32(elf, ph); + if p_type != 1 { continue; } + let p_offset = r64(elf, ph + 8) as usize; + let p_vaddr = r64(elf, ph + 16) as usize; + let p_filesz = r64(elf, ph + 32) as usize; + let end = (p_offset + p_filesz).min(elf.len()); + let copy_len = end - p_offset; + if p_vaddr + copy_len <= image.len() { + image[p_vaddr..p_vaddr + copy_len].copy_from_slice(&elf[p_offset..p_offset + copy_len]); + } + } + + Some(image) +} + +fn r16(data: &[u8], off: usize) -> u16 { + u16::from_le_bytes(data[off..off+2].try_into().unwrap()) +} +fn r32(data: &[u8], off: usize) -> u32 { + u32::from_le_bytes(data[off..off+4].try_into().unwrap()) +} +fn r64(data: &[u8], off: usize) -> u64 { + u64::from_le_bytes(data[off..off+8].try_into().unwrap()) +} + +fn read_str(data: &[u8], off: usize) -> &str { + let end = data[off..].iter().position(|&b| b == 0).unwrap_or(0); + std::str::from_utf8(&data[off..off + end]).unwrap_or("") +} + +// ─── Kernel argument builder ──────────────────────────────── + +/// Helper to pack kernel arguments into a GPU buffer. +pub struct KernArgs { + data: Vec, +} + +impl KernArgs { + pub fn new() -> Self { + KernArgs { data: Vec::with_capacity(256) } + } + + pub fn push_ptr(&mut self, buf: &GpuBuffer) { + self.align(8); + self.data.extend_from_slice(&buf.va_addr.to_le_bytes()); + } + + pub fn push_u64(&mut self, val: u64) { + self.align(8); + self.data.extend_from_slice(&val.to_le_bytes()); + } + + pub fn push_u32(&mut self, val: u32) { + self.align(4); + self.data.extend_from_slice(&val.to_le_bytes()); + } + + pub fn push_f32(&mut self, val: f32) { + self.align(4); + self.data.extend_from_slice(&val.to_le_bytes()); + } + + fn align(&mut self, alignment: usize) { + let pad = (alignment - (self.data.len() % alignment)) % alignment; + self.data.extend(std::iter::repeat(0u8).take(pad)); + } + + /// Fill OpenCL hidden arguments based on dispatch grid/block sizes. + /// Must be called after all explicit args are pushed but before upload. + /// Standard hidden arg layout (from clang OpenCL compiler): + /// +0x00 from end of explicit: hidden_block_count_x/y/z (3x u32) + /// +0x0c: hidden_group_size_x/y/z (3x u16) + /// +0x12: hidden_remainder_x/y/z (3x u16) + /// +0x18: (padding to 0x28 from end of explicit) + /// +0x28: hidden_global_offset_x/y/z (3x u64) + /// +0x40: hidden_grid_dims (u16) + pub fn fill_hidden_args(&mut self, grid: [u32; 3], block: [u32; 3], kernarg_size: u32) { + // Pad to the expected kernarg_size + while self.data.len() < kernarg_size as usize { + self.data.push(0); + } + + // Block counts = grid / block (number of workgroups per dimension) + let block_count = [ + if block[0] > 0 { grid[0] / block[0] } else { 0 }, + if block[1] > 0 { grid[1] / block[1] } else { 0 }, + if block[2] > 0 { grid[2] / block[2] } else { 0 }, + ]; + let remainder = [ + if block[0] > 0 { (grid[0] % block[0]) as u16 } else { 0 }, + if block[1] > 0 { (grid[1] % block[1]) as u16 } else { 0 }, + if block[2] > 0 { (grid[2] % block[2]) as u16 } else { 0 }, + ]; + let grid_dims: u16 = if grid[2] > 1 { 3 } else if grid[1] > 1 { 2 } else { 1 }; + + // Write at standard offsets (relative to byte 0x28 in our matvec case) + // These offsets come from the .note metadata: hidden_block_count_x at 0x28 + let base = 0x28usize; // first hidden arg offset for our kernels + if base + 12 <= self.data.len() { + self.data[base..base+4].copy_from_slice(&block_count[0].to_le_bytes()); + self.data[base+4..base+8].copy_from_slice(&block_count[1].to_le_bytes()); + self.data[base+8..base+12].copy_from_slice(&block_count[2].to_le_bytes()); + } + let gs = base + 12; // hidden_group_size + if gs + 6 <= self.data.len() { + self.data[gs..gs+2].copy_from_slice(&(block[0] as u16).to_le_bytes()); + self.data[gs+2..gs+4].copy_from_slice(&(block[1] as u16).to_le_bytes()); + self.data[gs+4..gs+6].copy_from_slice(&(block[2] as u16).to_le_bytes()); + } + let rm = gs + 6; // hidden_remainder + if rm + 6 <= self.data.len() { + self.data[rm..rm+2].copy_from_slice(&remainder[0].to_le_bytes()); + self.data[rm+2..rm+4].copy_from_slice(&remainder[1].to_le_bytes()); + self.data[rm+4..rm+6].copy_from_slice(&remainder[2].to_le_bytes()); + } + // hidden_global_offset at 0x50 + // (all zeros — we don't use global offsets) + // hidden_grid_dims at 0x68 + if 0x68 + 2 <= self.data.len() { + self.data[0x68..0x6a].copy_from_slice(&grid_dims.to_le_bytes()); + } + } + + pub fn upload(&self, alloc: &GpuAllocator) -> std::io::Result { + let size = ((self.data.len().max(256) + 4095) & !4095) as u64; + // Kernargs must be GPU-readable (PUBLIC flag) + let buf = alloc.alloc_userptr_public(size)?; + buf.write(0, &self.data); + Ok(buf) + } +} diff --git a/src/ioctl.rs b/src/ioctl.rs new file mode 100644 index 0000000..cc9d4a8 --- /dev/null +++ b/src/ioctl.rs @@ -0,0 +1,330 @@ +//! Raw KFD ioctl bindings for AMD GPU access. +//! +//! Structs match /usr/include/linux/kfd_ioctl.h exactly. +//! No dependencies beyond libc. + +use std::os::unix::io::RawFd; + +// ─── ioctl direction bits ─────────────────────────────────── + +const IOC_NONE: u32 = 0; +const IOC_WRITE: u32 = 1; +const IOC_READ: u32 = 2; +const IOC_READWRITE: u32 = 3; + +const AMDKFD_IOCTL_BASE: u32 = b'K' as u32; + +const fn ioc(dir: u32, nr: u32, size: u32) -> u64 { + ((dir as u64) << 30) | ((size as u64) << 16) | ((AMDKFD_IOCTL_BASE as u64) << 8) | (nr as u64) +} + +const fn ior(nr: u32) -> u64 { ioc(IOC_READ, nr, std::mem::size_of::() as u32) } +const fn iow(nr: u32) -> u64 { ioc(IOC_WRITE, nr, std::mem::size_of::() as u32) } +const fn iowr(nr: u32) -> u64 { ioc(IOC_READWRITE, nr, std::mem::size_of::() as u32) } + +unsafe fn kfd_ioctl(fd: RawFd, request: u64, arg: &mut T) -> std::io::Result<()> { + let ret = libc::ioctl(fd, request as libc::c_ulong, arg as *mut T); + if ret < 0 { + Err(std::io::Error::last_os_error()) + } else { + Ok(()) + } +} + +// ─── Memory allocation flags ──────────────────────────────── + +pub const ALLOC_MEM_FLAGS_VRAM: u32 = 1 << 0; +pub const ALLOC_MEM_FLAGS_GTT: u32 = 1 << 1; +pub const ALLOC_MEM_FLAGS_USERPTR: u32 = 1 << 2; +pub const ALLOC_MEM_FLAGS_DOORBELL: u32 = 1 << 3; +pub const ALLOC_MEM_FLAGS_MMIO_REMAP: u32 = 1 << 4; +pub const ALLOC_MEM_FLAGS_WRITABLE: u32 = 1 << 31; +pub const ALLOC_MEM_FLAGS_EXECUTABLE: u32 = 1 << 30; +pub const ALLOC_MEM_FLAGS_PUBLIC: u32 = 1 << 29; +pub const ALLOC_MEM_FLAGS_NO_SUBSTITUTE: u32 = 1 << 28; +pub const ALLOC_MEM_FLAGS_AQL_QUEUE_MEM: u32 = 1 << 27; +pub const ALLOC_MEM_FLAGS_COHERENT: u32 = 1 << 26; +pub const ALLOC_MEM_FLAGS_UNCACHED: u32 = 1 << 25; + +// ─── Queue types ──────────────────────────────────────────── + +pub const QUEUE_TYPE_COMPUTE: u32 = 0; +pub const QUEUE_TYPE_SDMA: u32 = 1; +pub const QUEUE_TYPE_COMPUTE_AQL: u32 = 2; + +pub const MAX_QUEUE_PERCENTAGE: u32 = 100; +pub const MAX_QUEUE_PRIORITY: u32 = 15; + +// ─── Event types ──────────────────────────────────────────── + +pub const EVENT_SIGNAL: u32 = 0; +pub const EVENT_HW_EXCEPTION: u32 = 3; +pub const EVENT_MEMORY: u32 = 8; + +// ─── Doorbell mmap ────────────────────────────────────────── + +pub const MMAP_TYPE_DOORBELL: u64 = 0x3 << 62; + +// ─── Structs ──────────────────────────────────────────────── + +#[repr(C)] +#[derive(Debug, Default)] +pub struct GetVersionArgs { + pub major_version: u32, + pub minor_version: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct AcquireVmArgs { + pub drm_fd: u32, + pub gpu_id: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct CreateQueueArgs { + pub ring_base_address: u64, + pub write_pointer_address: u64, + pub read_pointer_address: u64, + pub doorbell_offset: u64, + pub ring_size: u32, + pub gpu_id: u32, + pub queue_type: u32, + pub queue_percentage: u32, + pub queue_priority: u32, + pub queue_id: u32, + pub eop_buffer_address: u64, + pub eop_buffer_size: u64, + pub ctx_save_restore_address: u64, + pub ctx_save_restore_size: u32, + pub ctl_stack_size: u32, + pub sdma_engine_id: u32, + pub pad: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct DestroyQueueArgs { + pub queue_id: u32, + pub pad: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct AllocMemoryArgs { + pub va_addr: u64, + pub size: u64, + pub handle: u64, + pub mmap_offset: u64, + pub gpu_id: u32, + pub flags: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct FreeMemoryArgs { + pub handle: u64, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct MapMemoryArgs { + pub handle: u64, + pub device_ids_array_ptr: u64, + pub n_devices: u32, + pub n_success: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct UnmapMemoryArgs { + pub handle: u64, + pub device_ids_array_ptr: u64, + pub n_devices: u32, + pub n_success: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct CreateEventArgs { + pub event_page_offset: u64, + pub event_trigger_data: u32, + pub event_type: u32, + pub auto_reset: u32, + pub node_id: u32, + pub event_id: u32, + pub event_slot_index: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct DestroyEventArgs { + pub event_id: u32, + pub pad: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct SetEventArgs { + pub event_id: u32, + pub pad: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct WaitEventsArgs { + pub events_ptr: u64, + pub num_events: u32, + pub wait_for_all: u32, + pub timeout: u32, + pub wait_result: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct RuntimeEnableArgs { + pub r_debug: u64, + pub mode_mask: u32, + pub capabilities_mask: u32, +} + +#[repr(C)] +#[derive(Debug, Default)] +pub struct GetAvailableMemoryArgs { + pub available: u64, + pub gpu_id: u32, + pub pad: u32, +} + +// kfd_event_data — used in wait_events array +// Union of memory_exception / hw_exception / signal_event, followed by ext + event_id + pad +// We only need the signal case, which is simplest +#[repr(C)] +#[derive(Debug, Default, Clone, Copy)] +pub struct EventData { + // Union — signal_event_data is just u64 (last_event_age) + // memory_exception_data is the largest at 32 bytes + // We use raw bytes for the union + pub union_data: [u8; 32], + pub kfd_event_data_ext: u64, + pub event_id: u32, + pub pad: u32, +} + +// ─── ioctl numbers ────────────────────────────────────────── + +const IOC_GET_VERSION: u64 = ior::(0x01); +const IOC_CREATE_QUEUE: u64 = iowr::(0x02); +const IOC_DESTROY_QUEUE: u64 = iowr::(0x03); +const IOC_CREATE_EVENT: u64 = iowr::(0x08); +const IOC_DESTROY_EVENT: u64 = iow::(0x09); +const IOC_SET_EVENT: u64 = iow::(0x0A); +const IOC_WAIT_EVENTS: u64 = iowr::(0x0C); +const IOC_ACQUIRE_VM: u64 = iow::(0x15); +const IOC_ALLOC_MEMORY: u64 = iowr::(0x16); +const IOC_FREE_MEMORY: u64 = iow::(0x17); +const IOC_MAP_MEMORY: u64 = iowr::(0x18); +const IOC_UNMAP_MEMORY: u64 = iowr::(0x19); +const IOC_AVAILABLE_MEMORY: u64 = iowr::(0x23); +const IOC_RUNTIME_ENABLE: u64 = iowr::(0x25); + +// ─── Typed wrappers ───────────────────────────────────────── + +pub fn get_version(fd: RawFd) -> std::io::Result<(u32, u32)> { + let mut args = GetVersionArgs::default(); + unsafe { kfd_ioctl(fd, IOC_GET_VERSION, &mut args)?; } + Ok((args.major_version, args.minor_version)) +} + +pub fn acquire_vm(fd: RawFd, drm_fd: RawFd, gpu_id: u32) -> std::io::Result<()> { + let mut args = AcquireVmArgs { drm_fd: drm_fd as u32, gpu_id }; + unsafe { kfd_ioctl(fd, IOC_ACQUIRE_VM, &mut args) } +} + +pub fn runtime_enable(fd: RawFd) -> std::io::Result { + let mut args = RuntimeEnableArgs::default(); + unsafe { kfd_ioctl(fd, IOC_RUNTIME_ENABLE, &mut args)?; } + #[cfg(debug_assertions)] + eprintln!(" runtime_enable: caps=0x{:x}", args.capabilities_mask); + Ok(args.capabilities_mask) +} + +pub fn alloc_memory(fd: RawFd, va_addr: u64, size: u64, gpu_id: u32, + flags: u32, mmap_offset: u64) -> std::io::Result { + let mut args = AllocMemoryArgs { + va_addr, size, gpu_id, flags, mmap_offset, handle: 0, + }; + unsafe { kfd_ioctl(fd, IOC_ALLOC_MEMORY, &mut args)?; } + Ok(args) +} + +pub fn free_memory(fd: RawFd, handle: u64) -> std::io::Result<()> { + let mut args = FreeMemoryArgs { handle }; + unsafe { kfd_ioctl(fd, IOC_FREE_MEMORY, &mut args) } +} + +pub fn map_memory(fd: RawFd, handle: u64, gpu_ids: &[u32]) -> std::io::Result { + let mut args = MapMemoryArgs { + handle, + device_ids_array_ptr: gpu_ids.as_ptr() as u64, + n_devices: gpu_ids.len() as u32, + n_success: 0, + }; + unsafe { kfd_ioctl(fd, IOC_MAP_MEMORY, &mut args)?; } + Ok(args.n_success) +} + +pub fn unmap_memory(fd: RawFd, handle: u64, gpu_ids: &[u32]) -> std::io::Result<()> { + let mut args = UnmapMemoryArgs { + handle, + device_ids_array_ptr: gpu_ids.as_ptr() as u64, + n_devices: gpu_ids.len() as u32, + n_success: 0, + }; + unsafe { kfd_ioctl(fd, IOC_UNMAP_MEMORY, &mut args)?; } + Ok(()) +} + +pub fn create_queue(fd: RawFd, args: &mut CreateQueueArgs) -> std::io::Result<()> { + unsafe { kfd_ioctl(fd, IOC_CREATE_QUEUE, args) } +} + +pub fn destroy_queue(fd: RawFd, queue_id: u32) -> std::io::Result<()> { + let mut args = DestroyQueueArgs { queue_id, pad: 0 }; + unsafe { kfd_ioctl(fd, IOC_DESTROY_QUEUE, &mut args) } +} + +pub fn create_event(fd: RawFd, event_type: u32, auto_reset: u32, + event_page_offset: u64) -> std::io::Result { + let mut args = CreateEventArgs { + event_page_offset, event_type, auto_reset, + ..Default::default() + }; + unsafe { kfd_ioctl(fd, IOC_CREATE_EVENT, &mut args)?; } + Ok(args) +} + +pub fn set_event(fd: RawFd, event_id: u32) -> std::io::Result<()> { + let mut args = SetEventArgs { event_id, pad: 0 }; + unsafe { kfd_ioctl(fd, IOC_SET_EVENT, &mut args) } +} + +pub fn wait_events(fd: RawFd, events: &mut [EventData], + wait_for_all: bool, timeout_ms: u32) -> std::io::Result { + let mut args = WaitEventsArgs { + events_ptr: events.as_mut_ptr() as u64, + num_events: events.len() as u32, + wait_for_all: wait_for_all as u32, + timeout: timeout_ms, + wait_result: 0, + }; + unsafe { kfd_ioctl(fd, IOC_WAIT_EVENTS, &mut args)?; } + Ok(args.wait_result) +} + +pub fn available_memory(fd: RawFd, gpu_id: u32) -> std::io::Result { + let mut args = GetAvailableMemoryArgs { gpu_id, available: 0, pad: 0 }; + unsafe { kfd_ioctl(fd, IOC_AVAILABLE_MEMORY, &mut args)?; } + Ok(args.available) +} diff --git a/src/kernels/addr_dump.co b/src/kernels/addr_dump.co new file mode 100755 index 0000000000000000000000000000000000000000..e00a79e29d7188fed989e060699f0e5dc1d43554 GIT binary patch literal 3216 zcmcgu-)mcS6hHZ~?xtzdB&M|v!3c$wGOn@c7$TBkwJx-x0|j5a-ZnRB+T`AF?@ibG zFth&94Z+nR`Zy7`Dbr05MG>0)1A-5NFMHUF;FC{L3(D#__nvc?cB`ut{6g~ie9t+b z@6Yo+$y?7&J|7N+3K8&I1a3iyIYEJWD&~_3&vpzH^dzAVdaxh$Y`(wXL!i6Cg9$3^ zCo}1GYpnymM-WnepLCv!CZ;y%gWx*PC)lkc$@)P*e*N%vZ16A*Ss>3BM*E(0tQYyo zU)&PyZ?E(5vf*Xa$<}ztJps+-ohL1r+73L+eL*>UIt~|oImF{r;}b8v0NQm_m2AUtG^e4@7?&$GYg(^eQLS>>F2N=5V04uWFoF5wP-tDK;OPaU$AZmC|Q6v8r z^_^O=t!6`Wj7r@wU9DVQG)fwsG<@S-5#CMU)N9p>seFBxox=0qs&#!)E7x>aJ7wAP zJS=*C2;tqCdfoVp?pG@{>y7W#hFx9IT_YH}Yu>A>Q)$@RjMX&Vfh(%BRG+qLp9flX zzJ#X|L>#)Vp3)bLvTfn&X(2gDum!_*s+M^nB-x5txV_Q4$V4zp<(;!Tj+n-mS>QTR zU^~Uj=a6R$V}IDQb;dtD^VM&cC#cp*y`$K>CJ#0O0mChMqwV&UXU7e$8@t6t;QFmq zbSm-ycMJTV>|yv|qXp3pPPYl6Xb0=pW(yIwu?gX5XNdZUIqFBCGlF>V?{n?mi-}fZ zEwR>u_cpR{expCp+m65|?f%61R`N`;_1D(N?TZjTx(3m){(+cqc1^SWq zF;5{+F;64UGEX2MVxC3*F!LN?NWpAE3@Z>v-w6A9kso2+hdj?b9UUz`!F&MMI)H0E zCjH20@oDC?XMiu;Qv1(xeo&rM!+Y==r#obs^s-(2zu(@n$n+!YCCxfQI9_0~YGRdt zDmOVfo*NreA6Li3b8qzK!M)4xr$o^^`Y-PkRMT<|P)kcDRu51sW>cNjomt>-vb#2@ zwpG$y9aPs?biHW3UY$X+QL`QowmDv_&CBhI4QXFlF}$#)pNX>I(q z7mrV!cvkkYU!-%O-_)R{2w{}pUzum;wUCzXTx605f-kU;uAk7RYs(V7 z%}j6Hp4?mXl(II-U#M6Q2erf!AH&{)aKgj*p?!%ysPCNb z1@v{dLU#xGWB4JxyrbZEgAKv|LeWuhS9=mUh05=P8|xm}ZVk>)#(5R}a(3OJL#%&;jR}sE<;cCfj`m>f980js ZfwF&b$-L5^MRG6wudIJm0vUJe{{fr<+N1yg literal 0 HcmV?d00001 diff --git a/src/kernels/addr_dump.s b/src/kernels/addr_dump.s new file mode 100644 index 0000000..49e5867 --- /dev/null +++ b/src/kernels/addr_dump.s @@ -0,0 +1,149 @@ +// Debug: dump computed addresses to Y buffer +// Y[lid*8 + 0] = v9 (W offset) +// Y[lid*8 + 1] = v11 (X offset) +// Y[lid*8 + 2] = s2 (W ptr lo) +// Y[lid*8 + 3] = s3 (W ptr hi) +// Y[lid*8 + 4] = s6 (X ptr lo) +// Y[lid*8 + 5] = s7 (X ptr hi) +// Y[lid*8 + 6] = s11 (K) +// Y[lid*8 + 7] = exec_lo + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1102" +.amdhsa_code_object_version 5 + +.text +.globl addr_dump +.p2align 8 +.type addr_dump, @function +addr_dump: + s_mov_b32 s20, s2 + + s_load_b64 s[2:3], s[0:1], 0x00 + s_load_b64 s[4:5], s[0:1], 0x08 + s_load_b64 s[6:7], s[0:1], 0x10 + s_load_b64 s[8:9], s[0:1], 0x18 + s_load_b64 s[10:11], s[0:1], 0x20 + s_load_b32 s12, s[0:1], 0x28 + s_waitcnt lgkmcnt(0) + + s_add_u32 s13, s10, 31 + s_lshr_b32 s13, s13, 5 + s_mov_b32 s14, 0 + s_mov_b32 s15, s20 +.Ldiv: + s_cmp_lt_u32 s15, s13 + s_cbranch_scc1 .Ldiv_done + s_sub_u32 s15, s15, s13 + s_add_u32 s14, s14, 1 + s_branch .Ldiv +.Ldiv_done: + + v_and_b32 v1, 31, v0 + v_lshrrev_b32 v2, 5, v0 + s_lshl_b32 s16, s15, 5 + s_lshl_b32 s17, s14, 3 + v_add_nc_u32 v3, s16, v1 + v_add_nc_u32 v4, s17, v2 + + // bias load (same as matmul) + v_lshlrev_b32 v5, 2, v3 + v_mov_b32 v6, 0 + v_cmp_lt_u32 vcc_lo, v3, s10 + s_and_saveexec_b32 s18, vcc_lo + global_load_b32 v6, v5, s[4:5] + s_mov_b32 exec_lo, s18 + s_waitcnt vmcnt(0) + + // precompute (same as matmul) + v_lshrrev_b32 v7, 3, v0 + v_and_b32 v8, 7, v0 + v_lshlrev_b32 v8, 2, v8 + v_add_nc_u32 v9, s16, v7 + v_mul_lo_u32 v9, v9, s11 + v_add_nc_u32 v9, v9, v8 + v_lshlrev_b32 v9, 2, v9 + + v_lshlrev_b32 v10, 4, v0 + + v_add_nc_u32 v11, s17, v2 + v_mul_lo_u32 v11, v11, s11 + v_add_nc_u32 v11, v11, v1 + v_lshlrev_b32 v11, 2, v11 + + // dump 16 values per thread: Y[lid*16..lid*16+15] + v_lshlrev_b32 v20, 6, v0 // lid * 64 (16 dwords * 4 bytes) + + global_store_b32 v20, v9, s[8:9] offset:0 // [0] W offset + global_store_b32 v20, v11, s[8:9] offset:4 // [1] X offset + v_mov_b32 v21, s14 + global_store_b32 v20, v21, s[8:9] offset:8 // [2] s14 (wg_n) + v_mov_b32 v21, s15 + global_store_b32 v20, v21, s[8:9] offset:12 // [3] s15 (wg_m) + v_mov_b32 v21, s16 + global_store_b32 v20, v21, s[8:9] offset:16 // [4] s16 (wg_m*32) + v_mov_b32 v21, s17 + global_store_b32 v20, v21, s[8:9] offset:20 // [5] s17 (wg_n*8) + v_mov_b32 v21, s13 + global_store_b32 v20, v21, s[8:9] offset:24 // [6] s13 (num_wg_m) + v_mov_b32 v21, s20 + global_store_b32 v20, v21, s[8:9] offset:28 // [7] s20 (wg_id) + global_store_b32 v20, v1, s[8:9] offset:32 // [8] v1 (thread_m) + global_store_b32 v20, v2, s[8:9] offset:36 // [9] v2 (thread_n) + v_mov_b32 v21, s11 + global_store_b32 v20, v21, s[8:9] offset:40 // [10] s11 (K) + v_mov_b32 v21, s10 + global_store_b32 v20, v21, s[8:9] offset:44 // [11] s10 (M) + v_mov_b32 v21, s12 + global_store_b32 v20, v21, s[8:9] offset:48 // [12] s12 (N) + // intermediate: v11 before lshlrev = (s17+v2)*s11 + v1 + // let's compute it fresh + v_add_nc_u32 v21, s17, v2 + global_store_b32 v20, v21, s[8:9] offset:52 // [13] s17+v2 + v_mul_lo_u32 v21, v21, s11 + global_store_b32 v20, v21, s[8:9] offset:56 // [14] (s17+v2)*K + v_add_nc_u32 v21, v21, v1 + global_store_b32 v20, v21, s[8:9] offset:60 // [15] (s17+v2)*K + v1 + s_waitcnt vmcnt(0) + s_endpgm + +.rodata +.p2align 6 +.amdhsa_kernel addr_dump + .amdhsa_group_segment_fixed_size 0 + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 48 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_next_free_vgpr 26 + .amdhsa_next_free_sgpr 21 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_wavefront_size32 1 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_ieee_mode 1 +.end_amdhsa_kernel + +.amdgpu_metadata +--- +amdhsa.version: [ 1, 2 ] +amdhsa.kernels: + - .name: addr_dump + .symbol: addr_dump.kd + .kernarg_segment_size: 48 + .group_segment_fixed_size: 0 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .sgpr_count: 21 + .vgpr_count: 26 + .max_flat_workgroup_size: 256 + .args: + - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 16, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 24, .value_kind: global_buffer, .address_space: global } + - { .size: 4, .offset: 32, .value_kind: by_value } + - { .size: 4, .offset: 36, .value_kind: by_value } + - { .size: 4, .offset: 40, .value_kind: by_value } +... +.end_amdgpu_metadata diff --git a/src/kernels/coop_test.co b/src/kernels/coop_test.co new file mode 100755 index 0000000000000000000000000000000000000000..5823f7b75af8e7f435ec96bcc47a0efa499f404d GIT binary patch literal 2696 zcmbtV&2Jl35FdXeT}%q5gsS9%k*K(!M{q@jI3QO|DWC*F1xQ@FJbRz*ZS39E?mA6S zrB<%1w&Dch&Z!hsXt{tu2&D(Ea6+O|d&q%G{R0a4fCw|YGaJ`Q+f+Ph{hK#4zc=q= zW_EN=`juV3ezopjx5d3+%|Wy2gv`rn_t$n@;e{ zG<=sF9#5T_b;oUa*0-DVMA-Aa>gr{?+0n6m+z%IdVOlDOx8mup`Lf~rfgPJD{!s;? zyQE{2Exi#LE~%&$gm%O4d9ig-MJwI9-}x%zsznD+oOzV5sN;IcG(#WvPaCPo(=C}W za((ZUoa6z%?^c7-VpjpPK7G6@g1$!%MT>C|fvJJ+~w@5ar+m9+sHEk4&NDBi~D5OT%aSi8ehh`7XX z$3N#+KZE@KGf=FahRNz#u=e*YYjy&j`@Qfe_MnGKZ}dy0S%9NACLllgcj5f%0H$$% z;(7@SnD_uct)9*AM%}e*c%M_G>%$DrkKNcgRRrPye1l)wV+e0?Bc}VOA?K2I7xTWQ zO^N-_f7gBTpJh7(ydwe`N5;JCVRgS;nV+AlRIBO*RUI*IAtrJ17D&Gomj}%sP@W$f zP|k{nxeco2^;BC&ZQx(h$04ZDcXX^l#b!B9ymi-YpxSgo-wv<~xn2uY!|zh*0@X>Z z>xgY+r&5X=8MYx&uGw?%zj+l!U+5*L~qj9tb|ee6b6k)$k+=w6@qKZv-psD0bPu;8SRqnkQGuQ$g7$ zFSDsHJc|Z8m#i;+N3bSD#_{>{s3(5*-zoi!3unN*49Iz8yofS!x%{#`m;dj8eWecN z{>rR>1(9^ZAMQ;JLnc^KuduP`K0_My4#U^oO-bP_7{u=Y27DvW${(pIaUjP6A literal 0 HcmV?d00001 diff --git a/src/kernels/coop_test.s b/src/kernels/coop_test.s new file mode 100644 index 0000000..4379ab3 --- /dev/null +++ b/src/kernels/coop_test.s @@ -0,0 +1,92 @@ +// test cooperative W tile load to LDS +// each thread: load 4 floats from W, store to LDS, barrier, read own row back +// kernargs: W(u64) Y(u64) M(u32) K(u32) +// dispatch: 1 WG of 256 threads, computes for first 32x8 tile + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1102" +.amdhsa_code_object_version 5 + +.text +.globl coop_test +.p2align 8 +.type coop_test, @function +coop_test: + s_load_b64 s[2:3], s[0:1], 0x00 // W + s_load_b64 s[4:5], s[0:1], 0x08 // Y + s_load_b64 s[6:7], s[0:1], 0x10 // M, K (s6=M, s7=K) + s_waitcnt lgkmcnt(0) + + // v0 = lid + v_and_b32 v1, 31, v0 // thread_m = lid & 31 + v_lshrrev_b32 v2, 5, v0 // thread_n = lid >> 5 + + // W coop load: tile_row = lid/8, tile_col = (lid%8)*4 + v_lshrrev_b32 v3, 3, v0 // tile_row + v_and_b32 v4, 7, v0 // lid & 7 + v_lshlrev_b32 v4, 2, v4 // tile_col = (lid&7)*4 + + // W global offset = (tile_row * K + tile_col) * 4 + v_mul_lo_u32 v5, v3, s7 // tile_row * K + v_add_nc_u32 v5, v5, v4 // + tile_col + v_lshlrev_b32 v5, 2, v5 // * 4 bytes + + // global load 4 floats + global_load_b128 v[6:9], v5, s[2:3] + s_waitcnt vmcnt(0) + + // LDS store offset = lid * 16 + v_lshlrev_b32 v10, 4, v0 + ds_store_b128 v10, v[6:9] + s_waitcnt lgkmcnt(0) + s_barrier + + // Now read back: thread (thread_m, thread_n) reads W[thread_m][0] from LDS + // LDS offset = thread_m * 32 * 4 = thread_m * 128 + v_lshlrev_b32 v11, 7, v1 // thread_m * 128 + ds_load_b32 v12, v11 // W[thread_m][0] + s_waitcnt lgkmcnt(0) + + // Store to Y[lid] = v12 + v_lshlrev_b32 v13, 2, v0 // lid * 4 + global_store_b32 v13, v12, s[4:5] + s_waitcnt vmcnt(0) + s_endpgm + +.rodata +.p2align 6 +.amdhsa_kernel coop_test + .amdhsa_group_segment_fixed_size 5120 + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 24 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_next_free_vgpr 16 + .amdhsa_next_free_sgpr 8 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_wavefront_size32 1 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_ieee_mode 1 +.end_amdhsa_kernel + +.amdgpu_metadata +--- +amdhsa.version: [ 1, 2 ] +amdhsa.kernels: + - .name: coop_test + .symbol: coop_test.kd + .kernarg_segment_size: 24 + .group_segment_fixed_size: 5120 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .sgpr_count: 8 + .vgpr_count: 16 + .max_flat_workgroup_size: 256 + .args: + - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global } + - { .size: 4, .offset: 16, .value_kind: by_value } + - { .size: 4, .offset: 20, .value_kind: by_value } +... +.end_amdgpu_metadata diff --git a/src/kernels/lds_test.co b/src/kernels/lds_test.co new file mode 100755 index 0000000000000000000000000000000000000000..4fdcade8799e37bea577283ed371f77dd84842db GIT binary patch literal 2640 zcmd5-&yN&E6n^u=VIm@+(SR4{z=m6a8#Kg&4sjJsW{nyo%Y~+PyJn`@p6*F^&$4(B zv%!P$1SgMP5R#362Mxp+Jb2(@!o_3{9QGjpfxiyM@xAW%wZpQ8)svNUf3MzqU%jeV zud2?!-g%>1skCcUv>L5bMP-4u%6A)ua5UE)rY2jSq$fz!^;}oj+Xa-o?{qL=Q|+P_ z-aoar8iBN{@D$!-BIZBFXQabT@p#cA;3QAxI;Xt3slpLBRNmFcRnT}~ojLmUdqgYO?7;7* z)-1Xta$$P$lgq{?z4YRRYsUJ1;?mShhjzhz-V5V-8(uMSx0||*E}JxX-!)gwk_|`B zTMQ!qi^9d5A9cGfxn+7uJQ{kb>kV9#dEH>y`JRw|sMFVZJPuTvZ5Z^T=2s8t^C&mX zz%F~;&}QD*I9b$zWvB{IB(ei{HS|+2b7}UC879G!&0IP1-+^h>q`hI{EySZJt6wqc zxxsuKu7YY7eO^xK(7bNW+9lUbVxGNBqlH{8xg-tZ=u!oV=Y($3WE-BUan@0{>LX>h zTPjUArmcGM$M;;WYtmiSbw5fv_rE;(shnM$UvZh7fp7_nDFt4Y=%r~q)Gbfq5$*c= z1@>qv?-9=bd-I1>|CuNG`?fzmzP&=#gH)@3Ky+wLwo!CM-`=j#)Q?oF(74ffY0SV_=|v;c zw`rgBFZwJYlf=HwESbzLXSuf>1PiQoVG?^ozEy&#M`j@&$h|;v3Nt&;*k^Vj6_70BI?6L7v2gBC7du(^ z>v$Z4kNpSkQ$FnAtD5$*AY;RhdI7f75q8WU_&nRC#9&8#0#TnjOuh!lg8TWCb%Brj z1=wcWWNS}$SU$j?@XI$4_oTkt<+~1^I8U?}Sr#si3F;f~Fy3h$uFU@~9pe>7$tn1E zrV7IyB^<^Yu^Q>@3_ c)isBgdGBou zW3-knNJu0S5`Tmc`4%GZi$-H?`3Fq=K_dzLz%M5LN-}kcgnG`q_dL2nWeGoclh@C= z=gU3!+ z`cnEYO73f+qE5}_^Sb$%0Rm4tK|fIV%UIp**;I6|JmaC~J0kuJnUaHH-6_t-Iqsa0^!yt57x7X=AQpufM2R zN6J%1>9R0XW^=eAt~hy3Ijqg=dDFm|(-cyiT=TkV6^+U%4`(W7;KPm1Mcjv3%I93F z&^a5xlbMlqA|v^v&)){51F(rC8RAW>-w{qsZ(E^SM;|*T*{UWY0XLvk!j8_M=%*$7 zhO8SMT5i`1*X}H0FTe-TmHa=17fv^7;9o}XE8y`j4S64X!v9;@@b^@PwQ;+40yLMfqx{r6kCiO z9axNcd$I%mK(^k~m#%sHKRyJzfAOscIFvfo0N8hX9r%0VaKG`l?|AblL`IH5YYrjwRBgxz#9Qirgc&R(fQ*iGjH zsYBRJrz&*_yXnvx1>xCWgDnr#aP>k+ha??88?T+iT;RP<;GX-R#yw9YP9##N8@+J8 z8O1&HV;)3%hqLwJVLVEQZby3pAok=suD`g02|ttPCPW=PfdBO}T-WrQVEB5^^$3*ReT*3D+N%djWS(TKv#pUO(D*pwV(vT z#5hdsfr++(o-Da-l~)~ochAo8-LG+6Pl~ipwDw&4T|zSWeZ+n`(}Qd& z_t>8Ek?~DwNNGsEc;7issl5EvY8hI_!D=jL)FF z{qE!Q!pP@E4&%hYG~sjjQC#fby?;CIVbZ0Raeg2CY(yqKV+u4Ma+uF-W zlPbkPjq{%0*;T<7{}GVp^%Lg@V29+feR7L1NCpM|;Z98SJKV!KvWbU9YE@%=&f{Hdd6_Y5y;{ CyEc&k literal 0 HcmV?d00001 diff --git a/src/kernels/matmul.s b/src/kernels/matmul.s new file mode 100644 index 0000000..c1808f0 --- /dev/null +++ b/src/kernels/matmul.s @@ -0,0 +1,279 @@ +// rdna3 LDS-tiled matmul: Y = X * W^T + B +// +// Each workgroup computes a 32x8 tile of Y. +// 256 threads/WG: thread_m = lid%32, thread_n = lid/32 +// Tiles K in chunks of TK=32. Each tile iteration: +// 1. Cooperative load W[32][32] and X[8][32] to LDS +// 2. barrier +// 3. Each thread reads its W row and X row from LDS, does 32 FMAs +// 4. barrier +// +// LDS layout: W at 0 (4096B), X at 4096 (1024B), total 5120B +// +// Cooperative load assignment (W: 1024 elts / 256 threads = 4 each): +// tile_row = lid/8, tile_col = (lid%8)*4 +// global_load_b128 loads W[tile_row][tile_col..+3] +// ds_store_b128 at LDS offset lid*16 +// +// Cooperative load assignment (X: 256 elts / 256 threads = 1 each): +// x_row = lid/32, x_col = lid%32 +// global_load_b32 loads X[x_row][x_col] +// ds_store_b32 at LDS offset 4096 + lid*4 +// +// SGPR: s[0:1]=kernarg, s2=TGID_X +// kernargs (48B): W(u64) B(u64) X(u64) Y(u64) M(u32) K(u32) N(u32) +// dispatch: grid.x = ceil(M/32)*ceil(N/8)*256, block.x = 256 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1102" +.amdhsa_code_object_version 5 + +.set TM, 32 +.set TN, 8 +.set TK, 32 +.set LDS_X, 4096 // W uses 0..4095, X uses 4096..5119 +.set LDS_SZ, 5120 + +.text +.globl matmul +.p2align 8 +.type matmul, @function +matmul: + s_mov_b32 s20, s2 // save wg_id + + s_load_b64 s[2:3], s[0:1], 0x00 // W + s_load_b64 s[4:5], s[0:1], 0x08 // B + s_load_b64 s[6:7], s[0:1], 0x10 // X + s_load_b64 s[8:9], s[0:1], 0x18 // Y + s_load_b64 s[10:11], s[0:1], 0x20 // M, K + s_load_b32 s12, s[0:1], 0x28 // N + s_waitcnt lgkmcnt(0) + + // num_wg_m = ceil(M/32) + s_add_u32 s13, s10, TM - 1 + s_lshr_b32 s13, s13, 5 + + // wg_m = wg_id % num_wg_m, wg_n = wg_id / num_wg_m + s_mov_b32 s14, 0 + s_mov_b32 s15, s20 +.Ldiv: + s_cmp_lt_u32 s15, s13 + s_cbranch_scc1 .Ldiv_done + s_sub_u32 s15, s15, s13 + s_add_u32 s14, s14, 1 + s_branch .Ldiv +.Ldiv_done: + // s15 = wg_m, s14 = wg_n + + // thread decomp + v_and_b32 v1, 31, v0 // thread_m = lid & 31 + v_lshrrev_b32 v2, 5, v0 // thread_n = lid >> 5 + + // global output coords + s_lshl_b32 s16, s15, 5 // wg_m * 32 + s_lshl_b32 s17, s14, 3 // wg_n * 8 + v_add_nc_u32 v3, s16, v1 // global_m + v_add_nc_u32 v4, s17, v2 // global_n + + // load bias into accumulator + v_lshlrev_b32 v5, 2, v3 // global_m * 4 + v_mov_b32 v6, 0 + v_cmp_lt_u32 vcc_lo, v3, s10 + s_and_saveexec_b32 s18, vcc_lo + global_load_b32 v6, v5, s[4:5] + s_mov_b32 exec_lo, s18 + s_waitcnt vmcnt(0) + + // ======== Precompute cooperative load offsets ======== + + // W coop: tile_row = lid/8, tile_col = (lid%8)*4 + v_lshrrev_b32 v7, 3, v0 // tile_row = lid >> 3 + v_and_b32 v8, 7, v0 // lid & 7 + v_lshlrev_b32 v8, 2, v8 // tile_col = (lid & 7) * 4 + + // W global byte offset for k_tile=0: + // ((wg_m*32 + tile_row) * K + tile_col) * 4 + v_add_nc_u32 v9, s16, v7 // wg_m*32 + tile_row + v_mul_lo_u32 v9, v9, s11 // * K + v_add_nc_u32 v9, v9, v8 // + tile_col + v_lshlrev_b32 v9, 2, v9 // * 4 bytes + // v9 = W global load voffset (running, += TK*4 each iter) + + // W LDS store offset = lid * 16 (4 floats * 4 bytes) + v_lshlrev_b32 v10, 4, v0 + + // X coop: x_row = lid/32 = v2, x_col = lid%32 = v1 + // X global byte offset for k_tile=0: + // ((wg_n*8 + x_row) * K + x_col) * 4 + v_add_nc_u32 v11, s17, v2 // wg_n*8 + lid/32 + v_mul_lo_u32 v11, v11, s11 // * K + v_add_nc_u32 v11, v11, v1 // + lid%32 + v_lshlrev_b32 v11, 2, v11 // * 4 bytes + // v11 = X global load voffset (running, += TK*4 each iter) + + // X LDS store offset = LDS_X + lid * 4 + v_lshlrev_b32 v12, 2, v0 + v_add_nc_u32 v12, LDS_X, v12 + + // Compute-phase LDS read bases + v_lshlrev_b32 v13, 7, v1 // W: thread_m * 128 + v_lshlrev_b32 v14, 7, v2 + v_add_nc_u32 v14, LDS_X, v14 // X: LDS_X + thread_n * 128 + + // ======== Tile loop over K ======== + s_mov_b32 s18, 0 // k_tile = 0 + +.Ltile_loop: + s_cmp_ge_u32 s18, s11 // k_tile >= K? + s_cbranch_scc1 .Ltile_done + + // Phase 1: cooperative global → LDS + global_load_b128 v[15:18], v9, s[2:3] // W: 4 consecutive floats + global_load_b32 v19, v11, s[6:7] // X: 1 float + s_waitcnt vmcnt(0) + + ds_store_b128 v10, v[15:18] // W → LDS + ds_store_b32 v12, v19 // X → LDS + s_waitcnt lgkmcnt(0) + s_barrier + + // Phase 2: compute 32 FMAs from LDS, unrolled 4x per block (8 blocks) + + // block 0: tk=0..3 + ds_load_b128 v[15:18], v13 offset:0 + ds_load_b128 v[20:23], v14 offset:0 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 1: tk=4..7 + ds_load_b128 v[15:18], v13 offset:16 + ds_load_b128 v[20:23], v14 offset:16 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 2: tk=8..11 + ds_load_b128 v[15:18], v13 offset:32 + ds_load_b128 v[20:23], v14 offset:32 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 3: tk=12..15 + ds_load_b128 v[15:18], v13 offset:48 + ds_load_b128 v[20:23], v14 offset:48 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 4: tk=16..19 + ds_load_b128 v[15:18], v13 offset:64 + ds_load_b128 v[20:23], v14 offset:64 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 5: tk=20..23 + ds_load_b128 v[15:18], v13 offset:80 + ds_load_b128 v[20:23], v14 offset:80 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 6: tk=24..27 + ds_load_b128 v[15:18], v13 offset:96 + ds_load_b128 v[20:23], v14 offset:96 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 7: tk=28..31 + ds_load_b128 v[15:18], v13 offset:112 + ds_load_b128 v[20:23], v14 offset:112 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + s_barrier + + // advance offsets to next K tile + v_add_nc_u32 v9, v9, TK * 4 // W += 128 bytes + v_add_nc_u32 v11, v11, TK * 4 // X += 128 bytes + s_add_u32 s18, s18, TK + s_branch .Ltile_loop + +.Ltile_done: + // store Y[global_n][global_m] with bounds check + v_cmp_lt_u32 vcc_lo, v3, s10 // global_m < M + v_cmp_lt_u32 s19, v4, s12 // global_n < N + s_and_b32 s19, vcc_lo, s19 + s_and_saveexec_b32 s20, s19 + s_cbranch_execz .Ldone + + v_mul_lo_u32 v15, v4, s10 // global_n * M + v_add_nc_u32 v15, v15, v3 // + global_m + v_lshlrev_b32 v15, 2, v15 // * 4 bytes + global_store_b32 v15, v6, s[8:9] + s_waitcnt vmcnt(0) + +.Ldone: + s_endpgm + +// ======== Kernel descriptor ======== +.rodata +.p2align 6 +.amdhsa_kernel matmul + .amdhsa_group_segment_fixed_size LDS_SZ + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 48 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_next_free_vgpr 24 + .amdhsa_next_free_sgpr 21 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_wavefront_size32 1 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_ieee_mode 1 +.end_amdhsa_kernel + +.amdgpu_metadata +--- +amdhsa.version: [ 1, 2 ] +amdhsa.kernels: + - .name: matmul + .symbol: matmul.kd + .kernarg_segment_size: 48 + .group_segment_fixed_size: 5120 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .sgpr_count: 21 + .vgpr_count: 24 + .max_flat_workgroup_size: 256 + .args: + - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 16, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 24, .value_kind: global_buffer, .address_space: global } + - { .size: 4, .offset: 32, .value_kind: by_value } + - { .size: 4, .offset: 36, .value_kind: by_value } + - { .size: 4, .offset: 40, .value_kind: by_value } +... +.end_amdgpu_metadata diff --git a/src/kernels/matmul_blocked.co b/src/kernels/matmul_blocked.co new file mode 100755 index 0000000000000000000000000000000000000000..f69f6a7fe4d2977cc54945c8ab61c9ad3f67c6df GIT binary patch literal 4360 zcmb_f|8G-O6hHm`Y`d;2V^*B;t(!QBUQ(U()<;#ou1I=6e+Lg6^P1_Cd zhf&HFASR%MM1RweAS6W4A2b@n#*hAD;tv{07rG5Z>cX30K=yOZ=ru9iReOb~IiGo&y zOHv{AfhJs)X4UkJ7M)II<6jyo(aD)aLd!qKf$*dbyie-D|5yhnDO%A`wXTl#%H*+V zHTIfT*SfU2|4Zw8DVf)2a?yg8%xKwSG?6;0#iMMee`kk3s7IrTv|5ZF(eu;XF>__Yua{Lb+O>>Y%*>>tlW9FRt;K(ka{1J(TGVR(A50Hd zr9v{7kH++wY;p5tscM!8xC$7pV_v*W4zHgfxIOG>d)e!i%TAvxZ1RRrf-iUqS_fKuZDAMcxX{M6xh-67Ym?o`J9nF) z?DxxM-xk^ZtUp}w`@`&wnKy5d=dJB>rME7@7WQz!b}5~sTNQDvTj?0< zQ37MVO6QoQh!I)ohy;~Dq)+LL3@GA6Na>guR00#bl+KAE#Q`t;>cBW$h%Q9mX4H|z79+n%= z*m8yHXwo;HZIttOkK{wFW5!vjF6RT!AwKH8$E)s$Iay{Lf{KbFOE1f*nkMw(I! zQ^0?jyqE_mugBG*3Q|!!TC8fTnN$qLS~{;sb6Dl3vPqC)dWMx;z)G88F|W?%d7M)fHd z)c%iFt+_Jw5&MytD%GL2Jm42nj_cC3Li=ar834;4CL{YC&!Zf-|A;n<`)Q0nf^sIq zi?yzi{lPk^q02<%N5MYr-D=ez`TPP`WcS5WNVI-BKZj68YhP{uv)kY(ej1~BiJ!vG RqwI%`3kvp-BF9Gie*yV}?2rHe literal 0 HcmV?d00001 diff --git a/src/kernels/matmul_blocked.s b/src/kernels/matmul_blocked.s new file mode 100644 index 0000000..00e0a63 --- /dev/null +++ b/src/kernels/matmul_blocked.s @@ -0,0 +1,338 @@ +// rdna3 register-blocked LDS matmul: Y = X * W^T + B +// +// 4x4 output tile per thread: each thread computes 16 elements. +// TM=128, TN=32, TK=8, 256 threads/WG. +// threads_m = 128/4 = 32, threads_n = 32/4 = 8 +// thread_m = lid % 32, thread_n = lid / 32 +// +// LDS layout (TRANSPOSED for vectorized reads): +// W: W[TK][TM] at offset 0 — 8*128*4 = 4096 bytes +// X: X[TK][TN] at offset 4096 — 8*32*4 = 1024 bytes +// Total: 5120 bytes +// +// Cooperative load (per thread, 4 W elements + 1 X element): +// W: 128*8 = 1024 elts / 256 threads = 4 per thread +// tile_row = lid/2, tile_col = (lid%2)*4 +// global: W[wg_m*128 + tile_row][k_tile + tile_col .. +3] +// LDS (transposed): W[tile_col+i][tile_row] for i=0..3 +// = 4 scattered ds_store_b32 (stride = TM*4 = 512 bytes) +// X: 32*8 = 256 elts / 256 threads = 1 per thread +// x_row = lid/8, x_col = lid%8 +// global: X[wg_n*32 + x_row][k_tile + x_col] +// LDS (transposed): X[x_col][x_row] at LDS_X + x_col*TN*4 + x_row*4 +// = 1 ds_store_b32 +// +// Compute phase (per thread, per tk step): +// W: ds_load_b128 reads W[tk][thread_m*4 .. thread_m*4+3] (contiguous after transpose) +// X: ds_load_b128 reads X[tk][thread_n*4 .. thread_n*4+3] (contiguous after transpose) +// 16 FMAs: outer product of 4 W values * 4 X values +// 2 LDS reads per 16 FMAs = 8x better than non-blocked kernel +// +// dispatch: grid.x = ceil(M/128)*ceil(N/32), block.x = 256 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1102" +.amdhsa_code_object_version 5 + +.set TM, 128 +.set TN, 32 +.set TK, 8 +.set BM, 4 // block size per thread in M +.set BN, 4 // block size per thread in N +.set LDS_X, 4096 // W: TK*TM*4 = 8*128*4 = 4096 +.set LDS_SZ, 5120 // + X: TK*TN*4 = 8*32*4 = 1024 + +.text +.globl matmul_blocked +.p2align 8 +.type matmul_blocked, @function +matmul_blocked: + s_mov_b32 s20, s2 // save wg_id + + s_load_b64 s[2:3], s[0:1], 0x00 // W + s_load_b64 s[4:5], s[0:1], 0x08 // B + s_load_b64 s[6:7], s[0:1], 0x10 // X + s_load_b64 s[8:9], s[0:1], 0x18 // Y + s_load_b64 s[10:11], s[0:1], 0x20 // M, K + s_load_b32 s12, s[0:1], 0x28 // N + s_waitcnt lgkmcnt(0) + + // num_wg_m = ceil(M/128) + s_add_u32 s13, s10, TM - 1 + s_lshr_b32 s13, s13, 7 // /128 + + // wg_m = wg_id % num_wg_m, wg_n = wg_id / num_wg_m + s_mov_b32 s14, 0 + s_mov_b32 s15, s20 +.Ldiv: + s_cmp_lt_u32 s15, s13 + s_cbranch_scc1 .Ldiv_done + s_sub_u32 s15, s15, s13 + s_add_u32 s14, s14, 1 + s_branch .Ldiv +.Ldiv_done: + // s15 = wg_m, s14 = wg_n + + // thread decomp: 32 threads in M, 8 in N + v_and_b32 v1, 31, v0 // thread_m = lid & 31 + v_lshrrev_b32 v2, 5, v0 // thread_n = lid >> 5 + + // global output coords (base of 4x4 block) + s_lshl_b32 s16, s15, 7 // wg_m * 128 + s_lshl_b32 s17, s14, 5 // wg_n * 32 + v_lshlrev_b32 v3, 2, v1 // thread_m * 4 = base_m offset + v_add_nc_u32 v3, s16, v3 // global_m_base = wg_m*128 + thread_m*4 + v_lshlrev_b32 v4, 2, v2 // thread_n * 4 = base_n offset + v_add_nc_u32 v4, s17, v4 // global_n_base = wg_n*32 + thread_n*4 + + // ======== Initialize 16 accumulators with bias ======== + // acc[i][j] for i=0..3, j=0..3 in v16..v31 + // acc[i][j] = B[global_m_base + i] + // Load 4 bias values + v_mov_b32 v16, 0 + v_mov_b32 v17, 0 + v_mov_b32 v18, 0 + v_mov_b32 v19, 0 + v_mov_b32 v20, 0 + v_mov_b32 v21, 0 + v_mov_b32 v22, 0 + v_mov_b32 v23, 0 + v_mov_b32 v24, 0 + v_mov_b32 v25, 0 + v_mov_b32 v26, 0 + v_mov_b32 v27, 0 + v_mov_b32 v28, 0 + v_mov_b32 v29, 0 + v_mov_b32 v30, 0 + v_mov_b32 v31, 0 + + // load B[global_m_base+0..3] + v_lshlrev_b32 v5, 2, v3 // global_m_base * 4 + v_cmp_lt_u32 vcc_lo, v3, s10 // bounds check + s_and_saveexec_b32 s18, vcc_lo + global_load_b128 v[16:19], v5, s[4:5] // B[m+0..3] → acc[0..3][0] + s_mov_b32 exec_lo, s18 + s_waitcnt vmcnt(0) + // Copy bias to all 4 N columns: acc[i][j] = B[m+i] for j=0..3 + v_mov_b32 v20, v16 // acc[0][1] = B[m+0] + v_mov_b32 v24, v16 // acc[0][2] + v_mov_b32 v28, v16 // acc[0][3] + v_mov_b32 v21, v17 // acc[1][1] = B[m+1] + v_mov_b32 v25, v17 // acc[1][2] + v_mov_b32 v29, v17 // acc[1][3] + v_mov_b32 v22, v18 // acc[2][1] + v_mov_b32 v26, v18 // acc[2][2] + v_mov_b32 v30, v18 // acc[2][3] + v_mov_b32 v23, v19 // acc[3][1] + v_mov_b32 v27, v19 // acc[3][2] + v_mov_b32 v31, v19 // acc[3][3] + + // ======== Precompute cooperative load offsets ======== + + // W coop: tile_row = lid/2 (0..127), tile_col = (lid%2)*4 (0 or 4) + v_lshrrev_b32 v5, 1, v0 // tile_row = lid >> 1 + v_and_b32 v6, 1, v0 // lid & 1 + v_lshlrev_b32 v6, 2, v6 // tile_col = (lid&1)*4 + + // W global byte offset for k_tile=0: + // ((wg_m*128 + tile_row) * K + tile_col) * 4 + v_add_nc_u32 v7, s16, v5 // wg_m*128 + tile_row + v_mul_lo_u32 v7, v7, s11 // * K + v_add_nc_u32 v7, v7, v6 // + tile_col + v_lshlrev_b32 v7, 2, v7 // * 4 bytes + // v7 = W global load voffset (running, += TK*4 each iter) + + // W LDS store offsets (transposed): W[tile_col+i][tile_row] + // base = tile_col * TM * 4 + tile_row * 4 + v_mul_lo_u32 v8, v6, TM // tile_col * 128 + v_lshlrev_b32 v8, 2, v8 // * 4 bytes + v_lshlrev_b32 v9, 2, v5 // tile_row * 4 + v_add_nc_u32 v8, v8, v9 // base LDS offset for W store + // stride between consecutive tile_col values = TM*4 = 512 + + // X coop: x_row = lid/8 (0..31), x_col = lid%8 (0..7) + v_lshrrev_b32 v9, 3, v0 // x_row = lid >> 3 + v_and_b32 v10, 7, v0 // x_col = lid & 7 + + // X global byte offset for k_tile=0: + // ((wg_n*32 + x_row) * K + x_col) * 4 + v_add_nc_u32 v11, s17, v9 // wg_n*32 + x_row + v_mul_lo_u32 v11, v11, s11 // * K + v_add_nc_u32 v11, v11, v10 // + x_col + v_lshlrev_b32 v11, 2, v11 // * 4 bytes + // v11 = X global load voffset (running, += TK*4 each iter) + + // X LDS store offset (transposed): X[x_col][x_row] + // = LDS_X + x_col * TN * 4 + x_row * 4 + v_mul_lo_u32 v12, v10, TN // x_col * 32 + v_lshlrev_b32 v12, 2, v12 // * 4 + v_lshlrev_b32 v13, 2, v9 // x_row * 4 + v_add_nc_u32 v12, v12, v13 + v_add_nc_u32 v12, LDS_X, v12 // + LDS_X + + // Compute-phase LDS read bases (transposed layout) + // W[tk][thread_m*4+0..3]: base = tk * TM * 4 + thread_m * 4 * 4 + // = tk * 512 + thread_m * 16 + // We use offset for tk, so base = thread_m * 16 + v_lshlrev_b32 v13, 4, v1 // thread_m * 16 + + // X[tk][thread_n*4+0..3]: base = LDS_X + tk * TN * 4 + thread_n * 4 * 4 + // = LDS_X + tk * 128 + thread_n * 16 + v_lshlrev_b32 v14, 4, v2 // thread_n * 16 + v_add_nc_u32 v14, LDS_X, v14 // + LDS_X + + // ======== Tile loop over K (with prefetch) ======== + // Prefetch hides global memory latency by issuing the next tile's + // load during the current tile's compute phase. + // v[40:43] = prefetch W, v44 = prefetch X + // v[32:35] = LDS W read, v[36:39] = LDS X read + s_mov_b32 s18, 0 // k_tile = 0 + + // Prologue: load first tile into prefetch regs + global_load_b128 v[40:43], v7, s[2:3] // W tile 0 + global_load_b32 v44, v11, s[6:7] // X tile 0 + +.Ltile_loop: + // Wait for current tile's global data (first iter: prologue, then: prefetch) + s_waitcnt vmcnt(0) + + // Store prefetched data to LDS (transposed) + ds_store_b32 v8, v40 // W[tile_col+0][tile_row] + ds_store_b32 v8, v41 offset:512 // W[tile_col+1][tile_row] + ds_store_b32 v8, v42 offset:1024 // W[tile_col+2][tile_row] + ds_store_b32 v8, v43 offset:1536 // W[tile_col+3][tile_row] + ds_store_b32 v12, v44 // X[x_col][x_row] + s_waitcnt lgkmcnt(0) + s_barrier + + // Compute from LDS — 8 k-steps × 16 FMAs = 128 FMAs + // Prefetch issued AFTER first LDS reads to not stall the barrier→compute path + + // tk=0 + ds_load_b128 v[32:35], v13 offset:0 + ds_load_b128 v[36:39], v14 offset:0 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v16, v32, v36 + v_fmac_f32 v17, v33, v36 + v_fmac_f32 v18, v34, v36 + v_fmac_f32 v19, v35, v36 + v_fmac_f32 v20, v32, v37 + v_fmac_f32 v21, v33, v37 + v_fmac_f32 v22, v34, v37 + v_fmac_f32 v23, v35, v37 + v_fmac_f32 v24, v32, v38 + v_fmac_f32 v25, v33, v38 + v_fmac_f32 v26, v34, v38 + v_fmac_f32 v27, v35, v38 + v_fmac_f32 v28, v32, v39 + v_fmac_f32 v29, v33, v39 + v_fmac_f32 v30, v34, v39 + v_fmac_f32 v31, v35, v39 + + // Issue prefetch after first tk step — overlaps with tk=1..7 compute + v_add_nc_u32 v7, v7, TK * 4 // W global += 32 bytes + v_add_nc_u32 v11, v11, TK * 4 // X global += 32 bytes + s_add_u32 s18, s18, TK + global_load_b128 v[40:43], v7, s[2:3] // prefetch next W + global_load_b32 v44, v11, s[6:7] // prefetch next X + + // tk=1..7 + .irp TK_OFF, 512, 1024, 1536, 2048, 2560, 3072, 3584 + ds_load_b128 v[32:35], v13 offset:\TK_OFF + ds_load_b128 v[36:39], v14 offset:(\TK_OFF / 4) + s_waitcnt lgkmcnt(0) + v_fmac_f32 v16, v32, v36 + v_fmac_f32 v17, v33, v36 + v_fmac_f32 v18, v34, v36 + v_fmac_f32 v19, v35, v36 + v_fmac_f32 v20, v32, v37 + v_fmac_f32 v21, v33, v37 + v_fmac_f32 v22, v34, v37 + v_fmac_f32 v23, v35, v37 + v_fmac_f32 v24, v32, v38 + v_fmac_f32 v25, v33, v38 + v_fmac_f32 v26, v34, v38 + v_fmac_f32 v27, v35, v38 + v_fmac_f32 v28, v32, v39 + v_fmac_f32 v29, v33, v39 + v_fmac_f32 v30, v34, v39 + v_fmac_f32 v31, v35, v39 + .endr + + s_barrier + + // Loop if prefetched tile is valid + s_cmp_lt_u32 s18, s11 // s18 < K? + s_cbranch_scc1 .Ltile_loop + +.Ltile_done: + // ======== Store 16 output elements ======== + // Y[global_n_base+j][global_m_base+i] for i=0..3, j=0..3 + // Y layout: row-major, Y[n][m], stride = M + // offset = (global_n_base + j) * M + (global_m_base + i) + + // Compute base offset: global_n_base * M + global_m_base + v_mul_lo_u32 v5, v4, s10 // global_n_base * M + v_add_nc_u32 v5, v5, v3 // + global_m_base + v_lshlrev_b32 v5, 2, v5 // * 4 bytes + + // Store row j=0: acc[0..3][0] = v16..v19 + global_store_b128 v5, v[16:19], s[8:9] + + // Row j=1: offset += M*4 + s_lshl_b32 s19, s10, 2 // M * 4 + v_add_nc_u32 v5, v5, s19 + global_store_b128 v5, v[20:23], s[8:9] + + // Row j=2 + v_add_nc_u32 v5, v5, s19 + global_store_b128 v5, v[24:27], s[8:9] + + // Row j=3 + v_add_nc_u32 v5, v5, s19 + global_store_b128 v5, v[28:31], s[8:9] + + s_waitcnt vmcnt(0) + s_endpgm + +// ======== Kernel descriptor ======== +.rodata +.p2align 6 +.amdhsa_kernel matmul_blocked + .amdhsa_group_segment_fixed_size LDS_SZ + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 48 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_next_free_vgpr 48 + .amdhsa_next_free_sgpr 21 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_wavefront_size32 1 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_ieee_mode 1 +.end_amdhsa_kernel + +.amdgpu_metadata +--- +amdhsa.version: [ 1, 2 ] +amdhsa.kernels: + - .name: matmul_blocked + .symbol: matmul_blocked.kd + .kernarg_segment_size: 48 + .group_segment_fixed_size: 5120 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .sgpr_count: 21 + .vgpr_count: 48 + .max_flat_workgroup_size: 256 + .args: + - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 16, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 24, .value_kind: global_buffer, .address_space: global } + - { .size: 4, .offset: 32, .value_kind: by_value } + - { .size: 4, .offset: 36, .value_kind: by_value } + - { .size: 4, .offset: 40, .value_kind: by_value } +... +.end_amdgpu_metadata diff --git a/src/kernels/matmul_dbg.co b/src/kernels/matmul_dbg.co new file mode 100755 index 0000000000000000000000000000000000000000..c41313ce2958a2a436bd71f12d57dbfe1c6683a2 GIT binary patch literal 3656 zcmbtWUrbwN6hHklXlZFnVXQ6$>w?On-c_R0n8kDvT~=Y48zX9rmkYfug|@f3_rky! zRttz`#wl6ii^<3moN?Lai$-HCJo@6BMw9JfFTQUJ8WNY)bMF04SAhy~Uvls7eCPbm z_xF6~`_8=F_lncuh`YcpF1QH}Arr)f{8N)H9Ib7}|1<}kDgsT(bL9fb5|uLnbcLokd1;eqJELiYQswQlBA_lhHApHWMtk^ zJ(s0PB|D+Yw77!YZC?8P}4azZ_266l$w-jrGIXNZ>>$j-Km<|eXoayb2iRM0b%im6uHzTY)mmW*^kmq)aTyt(_5 zWK8FVwd|Gq&eC`iSEfpdT$4^KlWI!Wa0WD#Ciz1K2g`UnCbws$=lB4*CRno=ZpSpME(4JF#mN+L8VdLPe^Ll@_ zV2c81@o@g+w{2;&88JaDio8?L&k5IV&Y{u62e=+_hk7TRFBieRhV*rCxYsa$trrn- z80+A4uSKYi7^S)k)-Y|ZH{ShlrRhR*v3b6Ez6c+e+u&R|*xa<@f{QD`=5xisdx7HL zm8F#n;5;}F?w;U$Xf80@G8b|-BwE~_M5&=ER&@Gj#^C5LuH67*(Zw>ri5t7X-4KR* zyEI{aBDHMu@0^x)gp8U(TX1*8DjgkGy$Fltu+O{l7LM_bvvcN-Gm76n{JQXaYzE-n>rFvgv}bwA{TO0B9q9aKC>r^n`z~bjh-(z!178E_t1?o5NzyBlsY@{7>wNrHqb5lHLksV_Y!_0;@w9TEEtzo5yu-2 z2|+wl;1OZ}^IkXi}LbR{1H`>!9JuUUHWoxwZ0Nk!bH<;~y zZddp{z-Ie-v6*j=`*wxjE6lgk0qRAbVEuC>I6c7c7uMr(fUWZ3KGbiOKlh`4t2_x( zUNcpYlGAxi?I5M|6Vj++jDq-+-qbwogintZx-hLwBYICi-I=u15P|BbYd zv^Ui}&w@aHH*p6XM_eagicNOocx4Go#UIadxJlZpYdU|(|yPOe4ZHjJc(wkIj~IlIDRxQ_UC(-F^;y{60fbGd;mZ4 zufBcxJ{OAWd!GG>TvMJ!nW%hc^SR*ncR_IZ*z@|?=L_IPw3CkgLv%2(guFK1st(Mq zZ|ENSQB68bVE-!YPw$c1!v2hdXymo4<#luaLWy^g6{kFYj-R(t&ttFkkBEfG^Jg3P S%k|MZ{vQbcI1>eK^#2FUpG8Oj literal 0 HcmV?d00001 diff --git a/src/kernels/matmul_dbg.s b/src/kernels/matmul_dbg.s new file mode 100644 index 0000000..ae0173f --- /dev/null +++ b/src/kernels/matmul_dbg.s @@ -0,0 +1,282 @@ +// rdna3 LDS-tiled matmul: Y = X * W^T + B +// +// Each workgroup computes a 32x8 tile of Y. +// 256 threads/WG: thread_m = lid%32, thread_n = lid/32 +// Tiles K in chunks of TK=32. Each tile iteration: +// 1. Cooperative load W[32][32] and X[8][32] to LDS +// 2. barrier +// 3. Each thread reads its W row and X row from LDS, does 32 FMAs +// 4. barrier +// +// LDS layout: W at 0 (4096B), X at 4096 (1024B), total 5120B +// +// Cooperative load assignment (W: 1024 elts / 256 threads = 4 each): +// tile_row = lid/8, tile_col = (lid%8)*4 +// global_load_b128 loads W[tile_row][tile_col..+3] +// ds_store_b128 at LDS offset lid*16 +// +// Cooperative load assignment (X: 256 elts / 256 threads = 1 each): +// x_row = lid/32, x_col = lid%32 +// global_load_b32 loads X[x_row][x_col] +// ds_store_b32 at LDS offset 4096 + lid*4 +// +// SGPR: s[0:1]=kernarg, s2=TGID_X +// kernargs (48B): W(u64) B(u64) X(u64) Y(u64) M(u32) K(u32) N(u32) +// dispatch: grid.x = ceil(M/32)*ceil(N/8)*256, block.x = 256 + +.amdgcn_target "amdgcn-amd-amdhsa--gfx1102" +.amdhsa_code_object_version 5 + +.set TM, 32 +.set TN, 8 +.set TK, 32 +.set LDS_X, 4096 // W uses 0..4095, X uses 4096..5119 +.set LDS_SZ, 5120 + +.text +.globl matmul_dbg +.p2align 8 +.type matmul, @function +matmul_dbg: + s_mov_b32 s20, s2 // save wg_id + + s_load_b64 s[2:3], s[0:1], 0x00 // W + s_load_b64 s[4:5], s[0:1], 0x08 // B + s_load_b64 s[6:7], s[0:1], 0x10 // X + s_load_b64 s[8:9], s[0:1], 0x18 // Y + s_load_b64 s[10:11], s[0:1], 0x20 // M, K + s_load_b32 s12, s[0:1], 0x28 // N + s_waitcnt lgkmcnt(0) + + // num_wg_m = ceil(M/32) + s_add_u32 s13, s10, TM - 1 + s_lshr_b32 s13, s13, 5 + + // wg_m = wg_id % num_wg_m, wg_n = wg_id / num_wg_m + s_mov_b32 s14, 0 + s_mov_b32 s15, s20 +.Ldiv: + s_cmp_lt_u32 s15, s13 + s_cbranch_scc1 .Ldiv_done + s_sub_u32 s15, s15, s13 + s_add_u32 s14, s14, 1 + s_branch .Ldiv +.Ldiv_done: + // s15 = wg_m, s14 = wg_n + + // thread decomp + v_and_b32 v1, 31, v0 // thread_m = lid & 31 + v_lshrrev_b32 v2, 5, v0 // thread_n = lid >> 5 + + // global output coords + s_lshl_b32 s16, s15, 5 // wg_m * 32 + s_lshl_b32 s17, s14, 3 // wg_n * 8 + v_add_nc_u32 v3, s16, v1 // global_m + v_add_nc_u32 v4, s17, v2 // global_n + + // load bias into accumulator + v_lshlrev_b32 v5, 2, v3 // global_m * 4 + v_mov_b32 v6, 0 + v_cmp_lt_u32 vcc_lo, v3, s10 + s_and_saveexec_b32 s18, vcc_lo + global_load_b32 v6, v5, s[4:5] + s_mov_b32 exec_lo, s18 + // s_waitcnt vmcnt(0) -- no global loads + + // ======== Precompute cooperative load offsets ======== + + // W coop: tile_row = lid/8, tile_col = (lid%8)*4 + v_lshrrev_b32 v7, 3, v0 // tile_row = lid >> 3 + v_and_b32 v8, 7, v0 // lid & 7 + v_lshlrev_b32 v8, 2, v8 // tile_col = (lid & 7) * 4 + + // W global byte offset for k_tile=0: + // ((wg_m*32 + tile_row) * K + tile_col) * 4 + v_add_nc_u32 v9, s16, v7 // wg_m*32 + tile_row + v_mul_lo_u32 v9, v9, s11 // * K + v_add_nc_u32 v9, v9, v8 // + tile_col + v_lshlrev_b32 v9, 2, v9 // * 4 bytes + // v9 = W global load voffset (running, += TK*4 each iter) + + // W LDS store offset = lid * 16 (4 floats * 4 bytes) + v_lshlrev_b32 v10, 4, v0 + + // X coop: x_row = lid/32 = v2, x_col = lid%32 = v1 + // X global byte offset for k_tile=0: + // ((wg_n*8 + x_row) * K + x_col) * 4 + v_add_nc_u32 v11, s17, v2 // wg_n*8 + lid/32 + v_mul_lo_u32 v11, v11, s11 // * K + v_add_nc_u32 v11, v11, v1 // + lid%32 + v_lshlrev_b32 v11, 2, v11 // * 4 bytes + // v11 = X global load voffset (running, += TK*4 each iter) + + // X LDS store offset = LDS_X + lid * 4 + v_lshlrev_b32 v12, 2, v0 + v_add_nc_u32 v12, LDS_X, v12 + + // Compute-phase LDS read bases + v_lshlrev_b32 v13, 7, v1 // W: thread_m * 128 + v_lshlrev_b32 v14, 7, v2 + v_add_nc_u32 v14, LDS_X, v14 // X: LDS_X + thread_n * 128 + + // ======== Tile loop over K ======== + s_mov_b32 s18, 0 // k_tile = 0 + +.Ltile_loop: + s_cmp_ge_u32 s18, s11 // k_tile >= K? + s_cbranch_scc1 .Ltile_done + + // Phase 1: cooperative global → LDS + v_mov_b32 v15, 1.0 + v_mov_b32 v16, 1.0 + v_mov_b32 v17, 1.0 + v_mov_b32 v18, 1.0 + v_mov_b32 v19, 1.0 + // s_waitcnt vmcnt(0) -- no global loads + + ds_store_b128 v10, v[15:18] // W → LDS + ds_store_b32 v12, v19 // X → LDS + s_waitcnt lgkmcnt(0) + s_barrier + + // Phase 2: compute 32 FMAs from LDS, unrolled 4x per block (8 blocks) + + // block 0: tk=0..3 + ds_load_b128 v[15:18], v13 offset:0 + ds_load_b128 v[20:23], v14 offset:0 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 1: tk=4..7 + ds_load_b128 v[15:18], v13 offset:16 + ds_load_b128 v[20:23], v14 offset:16 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 2: tk=8..11 + ds_load_b128 v[15:18], v13 offset:32 + ds_load_b128 v[20:23], v14 offset:32 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 3: tk=12..15 + ds_load_b128 v[15:18], v13 offset:48 + ds_load_b128 v[20:23], v14 offset:48 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 4: tk=16..19 + ds_load_b128 v[15:18], v13 offset:64 + ds_load_b128 v[20:23], v14 offset:64 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 5: tk=20..23 + ds_load_b128 v[15:18], v13 offset:80 + ds_load_b128 v[20:23], v14 offset:80 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 6: tk=24..27 + ds_load_b128 v[15:18], v13 offset:96 + ds_load_b128 v[20:23], v14 offset:96 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + // block 7: tk=28..31 + ds_load_b128 v[15:18], v13 offset:112 + ds_load_b128 v[20:23], v14 offset:112 + s_waitcnt lgkmcnt(0) + v_fmac_f32 v6, v15, v20 + v_fmac_f32 v6, v16, v21 + v_fmac_f32 v6, v17, v22 + v_fmac_f32 v6, v18, v23 + + s_barrier + + // advance offsets to next K tile + v_add_nc_u32 v9, v9, TK * 4 // W += 128 bytes + v_add_nc_u32 v11, v11, TK * 4 // X += 128 bytes + s_add_u32 s18, s18, TK + s_branch .Ltile_loop + +.Ltile_done: + // store Y[global_n][global_m] with bounds check + v_cmp_lt_u32 vcc_lo, v3, s10 // global_m < M + v_cmp_lt_u32 s19, v4, s12 // global_n < N + s_and_b32 s19, vcc_lo, s19 + s_and_saveexec_b32 s20, s19 + s_cbranch_execz .Ldone + + v_mul_lo_u32 v15, v4, s10 // global_n * M + v_add_nc_u32 v15, v15, v3 // + global_m + v_lshlrev_b32 v15, 2, v15 // * 4 bytes + global_store_b32 v15, v6, s[8:9] + // s_waitcnt vmcnt(0) -- no global loads + +.Ldone: + s_endpgm + +// ======== Kernel descriptor ======== +.rodata +.p2align 6 +.amdhsa_kernel matmul_dbg + .amdhsa_group_segment_fixed_size LDS_SZ + .amdhsa_private_segment_fixed_size 0 + .amdhsa_kernarg_size 48 + .amdhsa_user_sgpr_kernarg_segment_ptr 1 + .amdhsa_system_sgpr_workgroup_id_x 1 + .amdhsa_next_free_vgpr 24 + .amdhsa_next_free_sgpr 21 + .amdhsa_float_denorm_mode_32 3 + .amdhsa_float_denorm_mode_16_64 3 + .amdhsa_wavefront_size32 1 + .amdhsa_system_vgpr_workitem_id 0 + .amdhsa_ieee_mode 1 +.end_amdhsa_kernel + +.amdgpu_metadata +--- +amdhsa.version: [ 1, 2 ] +amdhsa.kernels: + - .name: matmul_dbg + .symbol: matmul_dbg.kd + .kernarg_segment_size: 48 + .group_segment_fixed_size: 5120 + .private_segment_fixed_size: 0 + .kernarg_segment_align: 8 + .wavefront_size: 32 + .sgpr_count: 21 + .vgpr_count: 24 + .max_flat_workgroup_size: 256 + .args: + - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 16, .value_kind: global_buffer, .address_space: global } + - { .size: 8, .offset: 24, .value_kind: global_buffer, .address_space: global } + - { .size: 4, .offset: 32, .value_kind: by_value } + - { .size: 4, .offset: 36, .value_kind: by_value } + - { .size: 4, .offset: 40, .value_kind: by_value } +... +.end_amdgpu_metadata diff --git a/src/kernels/matmul_small.co b/src/kernels/matmul_small.co new file mode 100755 index 0000000000000000000000000000000000000000..67c4736d206bef6881d1ec35d9b49fc708f18e59 GIT binary patch literal 4232 zcmeHK-EUM?5TE^OZj05BT68KvI=N!aB3(%(mSYTZ8e*lkEP^ znVH|5GxwZ1ckam-`d+l#Y&{MT%>lQ;Ci?_EvafhVXHQwy4GL@lS^6=&|KxvX~ER0@QfTY%Ett?X!;iFRK(r=y?b6h z3{bdiq@shCp$(bYw3)QtfA@lBWaHLHlb1Ck8qJ!P6}E8imclI&gvU2p3RI5hb=RnGSj(mEOFe7hUrQFsKD3dO~T!&sEuSIo>s0dk`qC{ zrlpMI;aJkhg^y*jL((z1vf=KfjgcSU^)-ZS)BGR*qpN%&fPpBhH)3{ z&e;|kBdRoZz-&8W?Y(!3?sJ~JXUa2`hxeufaC*AI<1RYjlVXGCblyAe&Hq_CUpx)= z?UUf_Y?$&*dZ!vE$L%$`ucjtA?o{+fMG5BZ>R8^s5r1v?+cZ|FaqCB6-%pO!07uoc z(*U6v7dU->xIg`eizr7ZozUe~ST6pZ|+6{;vZ?NF?>I%ts z+d%R7f(4&XcTY7)k~ngQxO{$RoH*WEa-L3{VAK(G)9=%qj)NSEn^ zy-A;Nw&V^WggG(7*p} zf8M{Uu6Pum6WC^Mt)kP6SPy`Yd7)qOp>8PKv=XAtIvZRqHg%|>=8lb_rmbzE=B=AT zfHJb}*)9ALl0UL--KmDw?%Wt^-PIOayKBz&4MX@PT+VNPT=*qi&hPn;gir8^I+HU&i=IehVE|e@J)#X7 z)*#4#lIOCZWiwGDXMmP7kLSwPMk*0Ov6;+f!oyf?Cem@xBAFDGK0rlA&KN)(M7HP& zghTysXg?f!3Ho{uhL4p055BhGrM+N4Z}4#MtNVBN?|YtOd>EwtqCKb-Fx-*IbvMsZ zeawhUXP9Yxt{L}AA{}R%i?@x>djraTFO@RP=X1?S`$UcP%okFCHP|8FtmiynRHY&7 zdHsx!qMmChGUdEup36Yqm`Q$0QOqk`6t7vO@%{KE^R z!j?qDp@xM(<=K@5@-NTx%Q1b%(^=#7^Z9ulWxRH2T81O|!llOXi_zS?JRZc(Qv4f| M|129y= out_dim) return; + + float sum = b[row]; + __global const float* w_row = W + row * in_dim; + + // Vectorized inner loop + uint i = 0; + for (; i + 4 <= in_dim; i += 4) { + sum += w_row[i] * x[i]; + sum += w_row[i+1] * x[i+1]; + sum += w_row[i+2] * x[i+2]; + sum += w_row[i+3] * x[i+3]; + } + for (; i < in_dim; i++) { + sum += w_row[i] * x[i]; + } + + y[row] = sum; +} + +// Fused synapse: matvec → GLU → SiLU → layer_norm +// W is [out_dim*2 x in_dim], produces out_dim outputs +// scratch is [out_dim*2] temp space +__kernel void synapse_fused( + __global const float* W, + __global const float* b, + __global const float* x, + __global float* output, + __global float* scratch, + uint out_dim, + uint in_dim +) { + uint row = get_global_id(0); + uint total_rows = out_dim * 2; + if (row >= total_rows) return; + + // Step 1: matvec into scratch + float sum = b[row]; + __global const float* w_row = W + row * in_dim; + for (uint i = 0; i < in_dim; i++) { + sum += w_row[i] * x[i]; + } + scratch[row] = sum; + + barrier(CLK_GLOBAL_MEM_FENCE); + + // Only first out_dim threads continue for GLU + SiLU + if (row >= out_dim) return; + + // Step 2: GLU — output[i] = scratch[i] * sigmoid(scratch[i + out_dim]) + float val = scratch[row]; + float gate = 1.0f / (1.0f + exp(-scratch[row + out_dim])); + float glu_out = val * gate; + + // Step 3: SiLU — x * sigmoid(x) + float silu_out = glu_out / (1.0f + exp(-glu_out)); + + output[row] = silu_out; +} + +// GLU activation +__kernel void glu( + __global const float* input, + __global float* output, + uint half_dim +) { + uint i = get_global_id(0); + if (i >= half_dim) return; + float gate = 1.0f / (1.0f + exp(-input[i + half_dim])); + output[i] = input[i] * gate; +} + +// SiLU (swish) in-place +__kernel void silu( + __global float* x, + uint n +) { + uint i = get_global_id(0); + if (i >= n) return; + float v = x[i]; + x[i] = v / (1.0f + exp(-v)); +} + +// Elementwise add: y[i] += alpha * x[i] +__kernel void axpy( + __global float* y, + __global const float* x, + float alpha, + uint n +) { + uint i = get_global_id(0); + if (i >= n) return; + y[i] += alpha * x[i]; +} diff --git a/src/kernels/matvec.co b/src/kernels/matvec.co new file mode 100755 index 0000000000000000000000000000000000000000..ec087cfaeb102c6055806218cdb98d94b91a48a3 GIT binary patch literal 16928 zcmeHOZ;)Hnb-(Yvr>D1)c2{dJo+a2ucA3UdEY^QmY;ZloxN!9tl>8R&K)@>?MT~S^ktGZ-SStq#w-L8~% zwzi?{Rdro@xbmpN5sA$6nVII8W)X?fkh2mie|BKp^6d&*7h3y>4-X?xl% zsV&hS8BAA0dFC^ctV`|tAjq%bSq$afJGO3nU?*dpdc5$b|H>8} zi>LBqg$r!)m)?k{Mn>}KLLphmrH0e5jE!f9Qsc+tnL=_nJ6R}x`_*`MbhMBzvSaZ= z=8I|PM7+2^mrhQmCep8rj%QQFTV9LrO^r{blY26gBd<&ONoHLKfNrKG&N|9V$uWF$SA z92(CK@4+}vO%{{;Y7)1$DY3sM(ch-T7i$vhr3Cj{^|rEVET5gqB`MfEMh&zl;M%mI zJptFGJADM+QYPio6DdS@I`6S)Q(OG5L7Us+ckOvV@;hD^y`xEuR@lj$|eZ@5a?grpG4Ih_KPjzVwI~SN4toBVbUnpdO~iGh>s^ zH=5`a2JaiIj>pGSlVekRL+O13046|dB30a*9{y%Lm(T1?71K2fJm$R?FO22#a$+nxDJz9sW_(iS;Kmb+LxChG;+5kRm0+T(c8@s3cW2aH2LKe7i1I~FDT(cj32ztn<4!0)kna! zDd7WRJ_4>uTYLoO%Me@J;&%=DTwDCEJ?*2tJ9LlAHrmK4u3ZZ4+mLZ>%lY)>T4sp; ztCt~OF954^g>-Th$USmZ62+T!DnD7i2wuH>5xutgVgZ13R497TYQY^v4IraZxVqo?l0JeJ6F`gzWv4Uulr0cALMdpuu3Njwcwg`i%$#lL7mDx zbqzH0y9UME((l@%u(}Nha)<8ovW+%mT)Wy|bi1}G7_RI4wS>C9s|R&|IW|6ZRR9jZ z8%+I@z}kxozz$PkaYeABFzYg5M}yTuAmYYuuQYN&5OE{7U1)RTRzX``-1?x1Yt(X| zN$Vrv+N7Xxp%z?|R{FHi3?r^Ze#qI(?;7L>oXz~MJsVW4cn5BwJ$U<2<=WLgRJpb( zsH$`BwS=AM)q|av3z@5gJFk+aPX;{p)mu;XTm$H7e~G5B>B=^o8lpU?cuf4f>ka+p z-y}6odBj@(lKJT}%m?0l0Mwo1pDx4PqXmE5b$#_KybOCewJ*2-s;t9aB~70U4+EbN z4_{ZuO-1~?n=bw>YGcEbt5MJg=TNIxkB#nIv!;KoK9_QTy_JVrKHqWB7MVERxxWU(G4{^PrP{)}Wux7e%`<+~7)xZ|FuT+Fd6J52*fbLY4T-F-cFz zf)!B`{Y*>}60%^=q7-&2!g^)TQAJ;Ww`9Ck(aC*uOXrDn#AvYacX};$B!YVuV{8>; z<)hJ98I}qO`l2&OI%Zw|!tZCI%f52=_+y8{uk3hs`le@&-+$Nj-}!iM@IZiKRY*iqt=kd{hU(jaMeG+ds3SwR~)z_FGF`+b5ho zQNf^GT+i60GP5o|fRb_RB5b`>=5{Ig@!^V9HNRY~GWI-svSLjC;s@O}v-mVS&pBr$ zcJOqGeeaxOarQ7fR$(l15f)!!Y}5IJpmBEV`3{sFbD_#ng?v=w77RRZ8aYl6oHk6{zfEO^Rjt6X{Bg3g0a%>}_1xQ^oM$%+HM zaO8wxg+k>(IDC*Fa%LENL3X$j>2e02hFoXV8SGl*43duS zUT5$})uR>O(NVbpc2JwTEh}U}Z{Q$6tZamhgWO}IscfV^a5fzPmdN&~uN?;CGkk`4 z%IJqR2R-G$LEaG(cCe5$_)+zx3IngL{9=GN2*3Q?JdFPCzQC+$xc5!mAFdb}yE*6w z9$m69cEH+8)y=&FRe%=xmsbrjZZt;Jzc9uw!dF=#j2&CTS(z^>anQy1T?cD1erqvC zois+h9Noj%fnGx2sgL{yISypMx$FdHjOe2y#uWUsi=KJ>X6SeZIz!-R=zWCz5u`Dj zIbCLRrwvxFjUA2OyWnL9zcj}V1!lS-69Rxosn2~UjK!#1sMC1#5-mh+K60U(O#|W| zxv-c`FJ^OWCa~zaX@2#dt<R=vP#&SN zESaSvCD>JB|8Oq6n6q`Tz0CVcG}h#=Zt?^CM1AQ-Uyg!C{fVMK0EOxSnj4F*<+w;01{p#=#{*{4$ zn+UTE>*7Nw2?uE|(Hyd?fqLuR-O7fH8?c+vCRqOrlaQXH7T4Z<7~^UnnKmAO-*cr+;C3}3XISjJ+_G%4F+etkI1{*ylJp7O1|Ab$QU3yL8n7t( zY5qKTxKg#8=v^mHq1-h6{XhE0Eh~XR0k-%DRcp~4(XDNfQq>9#CUeK$B5bib9mfhg zw%rpdy+!r=oYITkhbk6Zdfp1dUyj48*1&XM{^Ylaze>r8l@I(BHm;x?G~<-`0jp%~ zbV~2xdD*hIlRPJR%ya6q8L*L{ZjpUvhiwOX9ATGv zpDpY%q1WtyUAA4bYYT^N)X6S>AXGvum1v%ZFa|%hZJXw57-K-bFk!EGz!vtNguQ{F zE$!{FqGU4ZGpzvz30lnbd95gn_|3gv<3+~ z|HF%&z($gD&_CkuK>tF^+&Llr-8?5o)SL({)L+8&4&)I%C&*{i z2h1@h=LD~dI}ryW_G}S*#x@6Y!miB;qk}s}klR*|En?5OkBitd5Fen!|GXA^VouQF zjphXBVon%{KjQ!|@tu~O6P$|}BYAGhIl)XO=LA!>F*#p3*~PS7wwx0zbk+zmPIJNt zaNF!bd?SY6rm^4Fipk!SGQ$ihp5-W79MUz3}BZduo!Ez zdj;-;zKHNg#ltLj1}G2Ecu;K6I+x;t*0=`M!A~)P*f2TK70w%M+7K~fTQWv$ISwyD zkHAekBw~aEH_h-U@0_yqV#V_*uK+r=i!aj}Vuyi3?*6E<1B7R{vqgI^+?iJXFg zDq%fI_(ZtB1wOaJUH4il2a%0Iy6@tQ_6Mw+vmRgq?GKC|jS14nv^SUyg$BL=Kakx< zSYlaiK6a^??~=Z@DdxL$%tuAc2RolSYxFQ%#C#{Wot{YZTzIGMH{z?nAH@4o75^JhQcRw~z5>5-T(YWG{rbmY z*qX`M$Gk03%oD%fN0^9}0TKI_hx#0R`ToRYvwFUOy2AH~Om4PXdG0 z`h@&M_@I47emV&qUSBo(sonT%jK}}$x?cYi{?q!O@SoQIg#Q$?cGa$r{|MsW!ukT~ zvj;vQtO&3kq(<5U%%1L}*f%8Z>)0pXX{;w)r#1gk8Xw38fsbN651Mx^4~!SskG+7{ zuK|Ox*V}!;G>AUm{y*Z0K^H@V0z>-*hT`|o^NcW&=AZ{d1Hh*1V82+y17oRwvysSv zSmOa-$!5w=W~ahV&;(XeUFInSu3V4Ne6|9%uo+3sgGVnSFMj&@W$}UjZTt)p0A2JO zI$eFTkndXIK5?TY=+~n@RL}qW7u77?CYDPAe#j8lDOFe1TCqJdi66!9jBVe(HMVwb zd`)}}W8(Mnm^#$aRJ@E7YkJTh?;pz)W2-Ssam_hY{8oY_a;$xFj@m;CHE~dVTUd*$2K0A^srdYg~-d7ad^bQWl>G6Cv znZuipnaMF0mv2hO>5Uw`J(QW4nkcaNaCU;;T4MC7Pcb!wau-y}J6Y;eLeGmt^6e(| zR?y_sL^1`Xnvd%A+K;ZoV0tY`$l$#wdNIiRD4!ml60bEioGdmqS9hcsueI7t7LC`F z{cLFFa804Z$|eJdSZh{SL}OI(YeLgq#o~WlO_~L@rn**>9ie}I`6(A{^uX0F&PAKF z3u;Y!ttPaKnxFRC16TY1`$Oa8)`#!Ed&h69Nvr1Ul@tbqDA!)# z8U5d