Struct CudaPipeline 

pub struct CudaPipeline<'a> { /* private fields */ }

Implementations

impl<'a> CudaPipeline<'a>

pub fn compile_cuda_c( context: &'a CudaComputeContext, kernel: &KernelSpec, schedule: Schedule, ) -> Result<Self, ForgeError>

Emit CUDA-C for the kernel and compile it to PTX via NVRTC (mirrors the HLSL -> DXC path), then load the resulting module.

pub fn compile_cuda_c_source( context: &'a CudaComputeContext, source: &str, entry_point: &str, storage_buffer_bindings: &[u32], ) -> Result<Self, ForgeError>

Compile a raw CUDA-C source string (entry point + storage-buffer bindings supplied directly) to PTX via NVRTC and load it. This is for kernels that have no portable-IR analogue — notably the nvcuda::wmma tensor-core GEMM, whose f16/f32 fragment API and fixed 16x16x16 shape cannot be expressed in WGSL/IR. storage_buffer_bindings lists the kernel’s pointer parameters in binding order (all treated as storage pointers; no by-value uniform).

pub fn from_ptx( context: &'a CudaComputeContext, ptx: &Ptx, entry_point: &str, spec: KernelSpec, ) -> Result<Self, ForgeError>

Load a pipeline from already-compiled PTX (no NVRTC). Used by the process-wide WMMA cache path so hot GEMM calls only pay load_module.

pub fn compile_cuda_c_source_cached( context: &'a CudaComputeContext, source: &str, entry_point: &str, storage_buffer_bindings: &[u32], ) -> Result<Self, ForgeError>

Compile (or reuse cached PTX for) raw CUDA-C and load — same as [compile_cuda_c_source] but shares the process NVRTC cache when source matches a previously compiled kernel body.

pub fn compile_ptx( context: &'a CudaComputeContext, ptx_source: &str, entry_point: &str, storage_buffer_bindings: &[u32], ) -> Result<Self, ForgeError>

Load a hand-emitted PTX module (from emit/ptx.rs) directly into the CUDA driver — no NVRTC compilation step. This is the PTX execution bridge: the emitter produces complete PTX text with .version, .target, .address_size, entry point, and full kernel body; the driver JITs it to the actual GPU ISA.

Shared-memory size is passed via LaunchConfig.shared_mem_bytes at dispatch time, not at compile time.

impl<'a> CudaPipeline<'a>

pub fn dispatch_async( &self, buffers: &[BufferView], schedule: &Schedule, element_count: usize, ) -> Result<(), ForgeError>

Launch without a host fence. Same-stream kernels stay ordered; the next read_buffer_* / synchronize is the completion barrier. Used by the P4 decode attention chain to avoid one PCIe-class fence per micro-kernel.

pub fn dispatch_ptx( &self, buffers: &[BufferView], grid: (u32, u32, u32), block: (u32, u32, u32), shared_mem_bytes: u32, ) -> Result<(), ForgeError>

Launch a PTX kernel with shared-memory size and a 3D grid/block config. Used by hand-emitted PTX kernels (RMSNorm, Q4K GEMV, WMMA GEMV, SDPA) that need shared_mem_bytes and multi-dimensional dispatch.

pub fn dispatch_async_sorted( &self, buffers: &[BufferView], schedule: &Schedule, element_count: usize, ) -> Result<(), ForgeError>

Fast-path async dispatch for pre-sorted buffer views.

Assumes buffers are already in ascending binding order (as the mega-pass always provides). Skips spec.buffers.clone() + sort + linear search — eliminating 2 Vec allocations and O(n²) search per dispatch.

pub fn dispatch_gpu_timed_ms_sorted( &self, buffers: &[BufferView], schedule: &Schedule, element_count: usize, ) -> Result<f32, ForgeError>

Measure one pre-sorted kernel launch with CUDA events.

This is a lab/profiling operation, not a decode hot-path primitive: creating and synchronizing timing events intentionally fences the stream. It remains useful when hardware performance counters are unavailable because the elapsed value is device time rather than host submission/synchronization wall time.

Trait Implementations

impl<'a> QualiaCompute for CudaPipeline<'a>

Available on crate feature cuda only.

fn dispatch( &self, buffers: &[BufferView], schedule: &Schedule, element_count: usize, ) -> Result<u64, ForgeError>

Dispatches a kernel using the supplied Schedule and buffer views.