Re-exports§
pub use coopmat::matmul_tc_wgsl;pub use coopmat::matmul_tc_wgsl_tiled;pub use coopmat::MATMUL_TC_TILED_ENTRY;pub use cuda_c::emit_cuda_c;pub use cuda_graph::emit_graph_cuda_c;pub use cuda_graph::graph_cuda_entry;pub use cuda_graph::CudaCLowerer;pub use df64::GEMM_DF64_ENTRY;pub use df64::GEMM_DF64_WGSL;pub use dxc::compile_hlsl_to_spirv;pub use dxc_cache::clear_dxc_cache;pub use dxc_cache::compile_hlsl_to_spirv_cached;pub use dxc_cache::dxc_cache_len;pub use graph_hlsl::conv2d_hlsl;pub use graph_hlsl::emit_graph_hlsl;pub use graph_hlsl::HlslLowerer;pub use graph_msl::conv2d_msl;pub use graph_msl::emit_graph_msl;pub use graph_msl::MslLowerer;pub use hlsl::emit_hlsl;pub use msl::emit_msl;pub use ptx::emit_ptx;pub use spirv::decode_spirv_words;pub use spirv::emit_spirv;pub use spirv::emit_spirv_patched;pub use spirv::patch_spirv_workgroup_size;pub use wgsl::emit_graph_wgsl;pub use wgsl::emit_wgsl;
Modules§
- coopmat
- Cooperative-matrix (tensor-core) WGSL emission (plan §18).
- cuda_c
- CUDA-C emission, compiled to PTX by NVRTC at runtime (mirrors the HLSL -> DXC -> DXIL path). Storage buffers become pointer parameters in binding order; a uniform block is passed by value as the last parameter. This is what the native CUDA backend executes for the differential oracle.
- cuda_
c_ fused - Fused CUDA-C kernels for the mega-pass — combine multiple kernel launches into one to reduce launch overhead and intermediate global memory traffic.
- cuda_
graph - CUDA-C lowering of the compute-graph IR — the second backend in one pass (plan §4,
Phase 5). The same
ComputeGraphthatemit_graph_wgsllowers to WGSL is lowered here to CUDA-C by aCudaCLowerer, with no per-kernel.idbranches — both backends are justLowererimpls walked by the onelower_graphdriver. - df64
- Double-single (
df64) emulated double-precision GEMM in raw WGSL. - dxc
- dxc_
cache - DXC SPIR-V binary cache — avoids re-running DXC on identical HLSL source.
- graph_
hlsl - HLSL lowering of the portable compute-graph nodes (plan §4 / P7) — the fourth backend
the one
lower_graphdriver walks.HlslLowererimplements the sameLowerertrait, emitting HLSL compute shaders for the portable native kit (Elementwise/Reduce/Broadcast) with the same binding ABI + math as the WGSLgraph_opskernels. - graph_
msl - MSL (Metal) lowering of the portable compute-graph nodes (plan §4 / P7) — the third
backend the one
lower_graphdriver walks, after WGSL and CUDA-C.MslLowererimplements the sameLowerertrait, emitting Metal for the portable native kit (Elementwise/Reduce/Broadcast) — same binding ABI + math as the WGSLgraph_opskernels, so a portable graph lowers to Metal with no per-id branch. - hlsl
- hlsl_
wave - HLSL wave-intrinsic decode kernel emitters.
- msl
- ptx
- spirv
- SPIR-V emission target.
- wgsl