Repository navigation
Unify Tensor backend services and integrate AMD HIP/ROCm - #3844
francescofugazzi wants to merge 68 commits into
Conversation
|
Thanks for the detailed review. This update addresses the common-dispatch requests for all four backends, including the existing CUDA/Vulkan/Metal branches in the shared dispatch files, rather than just replacing the HIP compile-time branches with runtime checks. Common services and registration
Remaining shared operation paths
Regression tests and architecture checksShared contracts cover the compiled backends, including worker first use, inventory without context creation, shutdown/reopen, target ordering, retained ownership, error recovery, indices, exact/empty allocation, conversion semantics and RNG reseeding. Normal sequences for counts 1, 2, 3 and 16 were also compared bit-for-bit before/after on Windows CUDA and Vulkan. The backend-neutrality check now explicitly guards the touched shared Tensor dispatch files against native SDK includes/calls, backend compile guards and operational backend choices. This is a structural check alongside the runtime tests, not a substitute for ownership/ordering validation. Public documentation now includes Latest local validation:
Clang-format passed for all 98 changed C/C++ files checked, and changed ranges in the four Objective-C++ files. Backend neutrality reports zero violations; error debt, CUDA log format and I/O discipline checks pass. The error-debt baseline change relocates the existing cuRAND check macro to its native source; it does not add a new allowance. Normalized hashes were checked across the isolated platform test sources, including new native files. Before/after performance evidenceThe comparison baseline is the previous PR head, The following are tiny 16-element operations, not training throughput. Windows CUDA uses nine alternating before/after process pairs on the same machine/SDK. Host code is /O2 on both sides: the baseline build's generated /Od flags were replaced with /O2; the updated build uses fastcompile OFF. Values are medians in microseconds.
The residual overhead is explicit: exact allocation costs approximately 0.10 us more per call, and fence recording approximately 10 ns more. Removing unnecessary disabled-profiler work did not remove that allocation overhead. The Float32-to-Bool improvement was consistent across the repeated series (roughly halving the cost); fused and normal timings varied between series, so I am not claiming a reproducible speedup for those cases or a training improvement. The updated macOS Release comparison uses three alternating pairs on the same M4 Pro. These are CPU submission medians in microseconds; the underlying measurements also retain completion timings.
The Mac runs show dispersion, particularly for queue execution, so these results should not be read as proof of either a small speedup or zero overhead. Numerical checks passed. I am reporting the measured costs rather than claiming the abstraction is free. Scope of this validation: these are targeted Core/contract builds and runtime checks, not a full portable rebuild or a complete application CI matrix. WSL was used to compile HIP and run CPU contracts; Linux GPU execution was not tested. Runtime coverage is limited to the hardware listed above. Benchmark scripts, assembly dumps and development reports remain outside the committed repository; regression tests and public documentation are committed. |
|
This follow-up adds two measured, local-first refinements to the common Tensor layer after the R4 restructuring:
The scope is shared host-side preparation and dispatch. These commits do not change native numerical kernels, rasterization/training algorithms, SDK versions, or the public backend interface/StorageRef layout. Changes retainedAllocation teardown. The common deleter now constructs a temporary StorageMeta only when the allocation actually has native metadata. Pointer-only allocations avoid that construction/destruction. The retained version keeps the existing deleter captures and owner structure, descriptor snapshot for native teardown, accounting and error cleanup. An earlier experiment with a separate pointer-only deleter was not retained. Fence dispatch. The TensorExecutionTarget(queue) constructor is defined beside the queue accessors, allowing those calls to inline within the same translation unit. Optimized assembly confirms that queue.record no longer needs the intermediate target/accessor calls. The backend-mismatch check and the native virtual fence operation remain. Same-target synchronization. sync_to_stream still validates the Tensor, but avoids native driver dispatch when the requested target is already the Tensor's home target. Different targets retain native synchronization. This is a common identity check, not a backend-specific fast path. Fused launch preparation. Host binding/output vectors reserve their known capacities before being populated, avoiding growth/reallocation. Binding validation, alias checks, dimensions, host parameters and generated arithmetic remain unchanged. Execution-to-home ordering for fused outputs now uses the common different-target check rather than a CUDA identity branch. The other native adapters retain their canonical home-target behavior. Contracts and gate. Added FusedOutputsRetainTheirQueueOrdering to the shared backend semantics tests. It uses two queues, two fused outputs and a CPU parameter, then consumes the second output on its own queue. Initializers are settled before the fused launch; there is no host barrier between that launch and the consuming queue operation. The test passed on CUDA, Vulkan, Metal and HIP. The GPU-expression case is explicitly skipped on CPU; required GPU cases still fail if their requested backend is unavailable. The neutrality gate now also guards tensor_fused.cpp, and the public target-ordering documentation describes the same-target no-op contract. Experiments discardedCaching the activated operations pointer locally through expression describe() did not improve warm preparation: its measured series was approximately 1.324 -> 1.335 us. That code was reverted before committing. The first allocation experiment with a separate deleter type was also replaced by the more conservative implementation described above. Experimental sources, scripts, assembly and reports remain outside the committed repository. Allocation/fence measurementsThis series compared the old PR head, the published R4 snapshot (6814e0f) and the allocation/fence refinement. Values are medians in microseconds from nine alternating runs per version. The table reports one matched series rather than combining values from different sessions.
Exact allocation improved by approximately 7.3% against the published snapshot in this series; the other retained series also show recovery, with different absolute values. Fence recording improved by approximately 3%, while retaining about 7 ns overhead against the old head in this series. This does not establish that every abstraction overhead is gone. Isolated CPU preparation measurementsTo separate host preparation from GPU launch variability, additional benchmarks measured same-target synchronization and warm preparation of an already compiled fused expression. Windows fixes only the benchmark's main thread to logical CPU 2; it does not change application or system scheduling settings. Each process takes seven samples, with nine alternating comparisons per version. Same-target sync uses 262144 calls/sample; warm preparation uses 8192. Both Windows sides use optimized /O2 host code. macOS uses optimized Release code without forced CPU affinity.
For reference, the old CUDA head's warm preparation was 1.306470 us in the same final Windows preparation series. The updated path was therefore also lower than that old baseline in this measurement. These are host-side costs, not kernel execution or training iteration times. GPU-inclusive timings and uncertaintyLonger, single-case comparisons also measured 16-element operations including submission of the GPU work: 2048 calls/sample, seven samples/process, nine alternating comparisons per version, with the benchmark main thread fixed to the same CPU. The final series was:
These GPU-inclusive results remain inconclusive for the incremental change: the updated version was faster in only 4/9 matched pairs for each case, and exploratory paired bootstrap intervals include zero broadly. Other retained series reverse some differences. The larger Bool improvement against the old head remains evident; neither its kernel nor the add kernel was changed in these two commits. I am not presenting the small fused median difference as a demonstrated GPU speedup, nor claiming that every GPU-inclusive timing is unchanged. The Mac investigation also retained initially unfavorable Vulkan-add series. Comparing the same saved libraries with the two changes separated subsequently reversed that difference. All those series are preserved in the local report; none was dropped to select a favorable result. Background filesystem activity was possible on Windows, so the measurements are not described as perfectly noise-free. Validation of the retained code
Clang-format, backend neutrality, error-debt, CUDA log-format, I/O discipline and diff checks passed. The neutrality check reports zero violations, with no additional error-debt allowance. Normalized hashes of 836 Core/test/gate files match the isolated Mac and WSL sources. The new multi-output ordering case ran successfully on all four GPU backends. Linux GPU runtime execution was not tested under WSL. These were targeted incremental builds and tests. There was no full portable rebuild, no new long training comparison, and no claim of a complete application CI matrix simulation. The measured CPU improvements should not be interpreted as measured training throughput gains. Only production changes, regression tests and public documentation are committed. |
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Thanks, this is the refactor we asked for. Common services for all four backends, one driver interface, an indexed ops table with the per-thread cache, generated registration, and the neutrality check covering the dispatch files. The shape is right now. What's left is two build/test breaks, a few backend special cases that survived in shared code, one HIP lifetime bug and a Vulkan dispatch cost against dev.
I built 956363e on Linux (CUDA 13.3 + Vulkan, Release, RTX 4080). TensorOperationContracts.{CPU,CUDA,Vulkan} and TensorQueueContracts pass, 4/4.
Breaks
- The build fails at
src/training/rasterization/gsplat/Common.cpp:core/gpu_backend_fwd.hpp:29now includes the generatedcore/gpu_backend_config.hpp, which lives in${CMAKE_BINARY_DIR}/include.gsplat_backend_lfsaddssrc/core/includeby hand and never linkslfs_backend_configuration, so it can't find the header. Same failure in the Ubuntu Debug job. Linkinglfs_backend_configurationthere fixes it; please check any other target that addssrc/core/includeby hand. - macOS fails at test registration: "Duplicate parameterized test name 'MetalGather'".
tests/test_tensor_vulkan_index.cpp:835names every backend that isn't CUDA or Vulkan "Metal", andkGpuBackendsnow has HIP in every build. Usegpu_backend_name()liketests/test_memory_pressure.cpp:374does.
Backend special cases still in shared code
UnavailableDriver(backend_registry.cpp:16-84) carries CUDA behaviour through its constructor arguments. The uncompiled CUDA slot reports device 0, returnsgetCurrentCUDAStream()as its target and reports renderer storage access. The Metal and HIP slots also get"cuda.direct"as their default direct label. The fallback driver should be neutral. If no-CUDA builds still needselected_cuda_device() == 0(dev returns 0 there), keep that in the CUDA compatibility function, not in the shared fallback. The twostd::runtime_errorthrows in the same class should beTensorErrorlikethrow_unavailable().Tensor::storage_memory_summary(tensor.cpp:347-353) asks each driver for its supplementary accounting kind and then checks forHipOwnedto printhip_owned. Map the kind to its label and counter, so this loop doesn't name a backend.TensorCompletionAccess(backend/tensor_completion.hpp,tensor_backend.cpp:553-590) has one factory per backend in shared code, each with its ownsettle_on_releaseand timeline handling. One factory that takes the backend and the native record (target, value, timeline point, settle policy) moves that policy into the backends that own it.hip_driver.cpp:34labels HIP direct allocations"cuda.direct", whilehip_backend_ops.cpp:612usesAllocationLabel::HipDirect. Take the string from the label.- Smaller:
memory_pressure.cpp:49,174fall back to CUDA for any domain thatdomain_backend()doesn't map. Map the CUDA domain explicitly instead of making it the default.
HIP after shutdown on another thread
shutdown_gpu_backend() only clears the calling thread's activation cache. A worker that already cached HIP ops never goes through activate() again, so it never calls ensure_initialized(). allocate() recovers because it calls ensure_initialized() itself, but sync() (hip_backend_ops.cpp:372), synchronize_device and device_barrier don't, and throw "HIP Tensor context is closed" after a shutdown/reopen. Calling ensure_initialized() there (one acquire load on the warm path) fixes it. Please add a contract test: a worker caches HIP ops, the main thread shuts down and reopens, then the worker synchronizes.
Performance against dev
Your tables compare against earlier heads of this PR. I measured against the merge base (eecd2cc) instead: same machine, alternating base/head processes, both Release, GPU otherwise idle. Medians in ms, 7 pairs:
| Case | dev | 956363e | Change |
|---|---|---|---|
| CUDA eager_binary_dispatch_tight_loop | 12.800 | 12.869 | +0.5% |
| CUDA masking_ops_tight_loop | 1.550 | 1.561 | +0.7% |
| CUDA dispatch_overhead_small (gui bench) | 0.8589 | 0.8668 | +0.9% |
| CUDA broadcast_binary_add | 0.7290 | 0.7286 | -0.1% |
| CUDA sort_middle_axis_64x1024x64 | 0.6647 | 0.6657 | +0.2% |
| Vulkan dispatch_overhead_small (gui bench) | 0.5661 | 0.5796 | +2.4% |
CUDA is within run-to-run noise now. At 6814e0f the eager loop was +1.6 to +2.6% in two runs, so 956363e fixed that. Vulkan small dispatch is slower than dev in all three runs I did (+1.6%, +1.8%, +2.4%), with little overlap between the sample ranges. The bandwidth-bound GUI cases (RGB conversion, fused pointwise, mask reduce, non-contiguous copy) are within ±1.5% on both backends. Please profile bench_tensor_gui --backend vulkan --case dispatch_overhead_small against dev and get it back to parity. I'd start with the per-op paths that went from direct calls to the driver interface (stream/target binding, the scope and availability checks).
Requesting changes for the two breaks, the shared-code special cases and the HIP thread bug. The Vulkan number has to reach parity before merge too.
Follow-up: review corrections and fresh comparisons against upstream devPushed Build and test registration
Shared contracts and native policy
HIP lifetime regressionHIP synchronization now calls Dispatch changes and allocation cleanup
Validation on the updated content
GPU runtime behavior on native Linux was not exercised: WSL here supports the CPU checks, not the required native Vulkan GPU comparison. These local checks do not substitute for the complete CI runs triggered by this push. Performance methodologyBoth dev and updated head are rerun for each new comparison, in alternating processes on the same machine. Historical dev timings are not reused as the current baseline. Windows timing builds are Release The original Requested small-dispatch comparisonAll values are head versus the upstream dev merge-base. Negative percentages mean lower time. Rows are separate repeated series and should not be combined as if they used identical sampling protocols.
The Windows Vulkan 1001-sample confirmation had three constituent series at -4.99%, -1.54% and -1.20%. Other Vulkan repetitions still showed dispersion: one 101-sample series was +3.32%, and another three-series check produced -4.65%, -3.99% and +5.05%. The complete-suite measurements below are also retained, including +2.31% for macOS Vulkan small dispatch. I am not selecting only the favorable runs or claiming that every measurement improved. CUDA small-dispatch repetitions changed sign as well: two normal 1001-sample series gave -4.08% and +4.78%, and a separate nine-pair series gave +2.46%. To measure variability within an unchanged build, the final control used eight alternating ABBA/BAAB blocks, 32 process runs total, 1001 samples each, where A=dev and B=head. The block-geometric-mean difference was +1.27%, with a block-bootstrap 95% interval of [-4.25%, +6.25%]. Repeats of the same dev build varied by up to 12.92%; the same head by up to 11.44%. This interval includes zero but is too broad to prove strict performance equivalence. CUDA profiling cross-checkTwo CUDA Nsight captures, in reversed dev/head order, showed the same 28,416 kernel launches, 213 device synchronizations, launch geometry and register count. Kernel median was 800 ns. Host time outside traced CUDA APIs was lower in head in both captures (approximately 0.503→0.493 ms and 0.503→0.487 ms per measured dispatch batch), while time inside launch APIs varied. This locates the observed difference in those traces but does not establish that submission cadence and driver time are independent of the refactoring. A separate process-local API-count probe checked five runtime APIs not normally recorded by the installed profiler. Dev/head counts matched exactly: Complete Windows GUI benchmark comparison: 7 alternating pairs, 21 samples/process
The small positive CUDA mask/copy results were repeated in isolated 101-sample runs: CUDA launch-overhead suite: 7 alternating pairs, 101 samples/process
Separate nine-pair diagnostic-affinity repetitions gave sort -0.02%, broadcast +0.16%, masking -4.03%, eager dispatch -16.41% and alias +5.52% (about 1 microsecond absolute for alias). The large eager improvement is not claimed as an acquired code speedup; it and the opposite-sign normal results remain part of the variability evidence. Complete macOS GUI benchmark comparison: 5 alternating pairs, 21 samples/process
Some fused-pointwise measurements are bimodal in both versions, so the large percentage differences are not treated as demonstrated improvements. Small-dispatch follow-up measurements with 101 and 1001 samples are reported separately above. Review status and remaining confirmationThe functional review findings have been addressed and the relevant local contracts/build consumers pass. The focused Windows/macOS Vulkan results support returning dispatch overhead to the dev range, but they do not certify parity on the reviewer's native Linux RTX 4080/CUDA 13.3 setup. That original hardware comparison remains the final confirmation for the hard performance requirement. The CUDA measurements likewise justify reporting variability, not asserting a guaranteed speedup or absolute absence of regressions. Only production sources, permanent regression/CI tests, gate tooling and public backend-contract/build documentation are committed. Test reports, raw timings, diagnostic wrappers and experimental scripts remain outside the branch. The existing PR description and its earlier measurement tables are preserved. |
francescofugazzi
changed the title
Ubuntu Debug: keep backend drivers valid through process exitFixed the teardown regression in The failed job had built successfully; three standalone Python checks then aborted with Registered drivers, including unavailable-backend placeholders, now have process lifetime, matching the existing activated-ops lifetime. These are bounded control objects; native resource cleanup remains explicit through backend shutdown. Operation dispatch and numerical kernels are unchanged by this correction. Added a fresh-process CPU regression executable that registers its exit callback before the first driver lookup and uses normal process exit. GoogleTest discovery and the existing contract runner's explicit exit path could otherwise mask this ordering. The executable is built with Recompiled only the registry and driver objects without optimization in an isolated Linux library: the previous code reproduced the abort for all four backend slots, and the corrected code passed. The permanent exit test also failed against the old library and passed against the new one. Additional local checks passed:
The original three full-module Python checks were not rebuilt locally; the new CI run provides their direct confirmation. No tests were disabled or relaxed. |
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Thanks, this round fixes everything from my last review except the Vulkan number.
I rebuilt 74ab187 on Linux (CUDA 13.3 + Vulkan, Release, RTX 4080):
- gsplat and fastgs build again, macOS registers the index tests, and Ubuntu Debug passes. Windows Release only fails the UNC rename test that every Windows job hits today (fix in #3873).
TensorOperationContracts.{CPU,CUDA,Vulkan,Runner}andTensorQueueContractspass 5/5,tensor_backend_shutdown_contractspasses, and the backend gtests (semantics including the new cached-worker shutdown test on CUDA and Vulkan, memory pressure, index bounds) pass 71/71.- The shared-code special cases are gone: the fallback driver is neutral, the summary maps kinds through a table, there is one completion factory, and HIP takes its label from the allocation label.
Performance against dev
Same machine, dev merge base eecd2cc rebuilt for this run, 9 alternating base/head process pairs, GPU otherwise idle. Medians in ms:
| Case | dev | 74ab187 | Change |
|---|---|---|---|
| Vulkan dispatch_overhead_small (gui bench) | 0.5699 | 0.5841 | +2.5% |
| CUDA dispatch_overhead_small (gui bench) | 0.8573 | 0.8674 | +1.2% |
| CUDA eager_binary_dispatch_tight_loop | 12.871 | 12.879 | +0.1% |
| CUDA masking_ops_tight_loop | 1.565 | 1.564 | -0.1% |
| CUDA broadcast_binary_add | 0.7255 | 0.7253 | -0.0% |
| CUDA sort_middle_axis_64x1024x64 | 0.6639 | 0.6642 | +0.0% |
CUDA is within noise; the gui-bench CUDA case flips sign between pairs. Vulkan small dispatch is still slower. Head lost 8 of the 9 pairs (+1.6, +2.1, +2.4, +2.7, -0.8, +5.6, +2.8, +1.7, +2.7%), and the head median equals the slowest dev sample. That matches 956363e (+1.6, +1.8, +2.4% in three runs), so the r6 changes didn't move it on Linux.
One lead: the inlining fix only reaches MSVC. With GCC, Tensor::adopt_native_tensor_allocation is an exported member of Tensor, so under -fPIC it stays out of line, and both call sites in tensor_unified_ops.cpp call it through the PLT. Making it a non-exported helper in that file would let every compiler inline it, and the #if defined(_MSC_VER) block goes away too. I can't say that is all of it: the gap is about 55 ns per op (256 adds per sample), and I couldn't profile on this machine. Since you can't run native Linux Vulkan, push a candidate and I'll rerun this comparison.
Requesting changes for Vulkan parity only. The rest is ready.
|
Updated in commit 20eda28. Allocation adoption
The GCC Release object and shared library confirm that the adoption helper is no longer dynamically exported and neither call site uses the PLT. GCC still emits two direct local calls, so this is not a claim of full inlining. Ownership, allocation accounting, exception rollback, descriptors, execution targets and backend lookup on the releasing thread retain the same behavior. The commit changes only the two production files involved in moving the helper. Current performance measurements against devBaseline: dev merge base For The time columns below are medians of the process results across the three series. Each percentage is the geometric mean of the matched head/dev ratios within that series; it is therefore not calculated from the two pooled median columns. Negative percentages mean lower elapsed time.
Vulkan is close to dev in these Windows and macOS runs. CUDA and especially Metal show more variation between series; these results do not establish a speedup. The local results do not establish native Linux Vulkan parity, which remains to be verified. Validation
|
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Thanks for moving the adoption helper into the file. That part works: with GCC it is now inlined and there is no PLT call left. The Vulkan gap is smaller, but not gone.
I rebuilt 20eda28 on Linux (CUDA 13.3 + Vulkan, Release, RTX 4080).
Tests
- Contracts 5/5,
tensor_backend_shutdown_contractsand the 71 backend gtests pass. - CI: Windows only fails the UNC rename test that every Windows job hits right now.
Performance against dev
Dev merge base eecd2cc vs head, 15 alternating base/head process pairs, machine otherwise idle. Medians in ms:
| Case | dev | 20eda28 | Change | Head slower |
|---|---|---|---|---|
| Vulkan dispatch_overhead_small (gui bench) | 0.5685 | 0.5780 | +1.7% | 11 of 15 pairs |
| CUDA deferred_alias_copy_materialize | 0.01164 | 0.01238 | +6.4% | 13 of 15 pairs |
| CUDA dispatch_overhead_small (gui bench) | 0.8678 | 0.8661 | -0.2% | 5 of 15 |
| CUDA eager_binary_dispatch_tight_loop | 12.889 | 12.861 | -0.2% | |
| CUDA masking_ops_tight_loop | 1.565 | 1.554 | -0.7% | |
| CUDA broadcast_binary_add | 0.7242 | 0.7241 | -0.0% | |
| CUDA sort_middle_axis_64x1024x64 | 0.6643 | 0.6643 | -0.0% |
- Vulkan small dispatch went from +2.5% at r6 to +1.7%.
- CUDA deferred alias copy. I should have raised this one at r6: it was +2.5% there and I left it out of the table. Here it is consistently slower, by about 0.7 µs per iteration.
One cause I can point to
internal::backend_ops() now reads the thread_local std::array activated cache on every call. liblfs_core.so is built with -fPIC, so that read is a __tls_get_addr@plt call each time:
lea 0x90cc7a(%rip),%rdi # <$tlsld>
call __tls_get_addr$plt
lea 0x190(%rax),%rbp
mov 0x0(%rbp,%rbx,8),%rax
On dev, Vulkan returned a function-local static and never touched TLS; only CUDA used a thread_local. That fits the pattern: Vulkan got slower and CUDA did not.
As an experiment, I marked the cache [[gnu::tls_model("initial-exec")]]. The lookup became a single %fs: load, and Vulkan went from +1.7% to +1.3% (head still slower in 12 of 15 pairs). So TLS is part of it, but about two thirds of the gap comes from somewhere else in the driver path.
What I'd suggest
- The TLS lookup. One way is a process-wide cache for drivers whose activation is not per thread, with a per-thread cache only where the driver asks for one. That would be a driver capability, not a backend check in shared code. Another is an initial-exec cache, though that is GCC/Clang only and uses the static TLS reserve when the Python module loads the library.
- The remaining Vulkan cost and the deferred alias case. Please compare the per-op path against dev for extra virtual calls or atomics between the facade and the native dispatch.
Push a candidate and I'll run the same comparison again.
Requesting changes for these two performance gaps only. Everything else is ready.
|
Follow-up to the latest performance review, in commit 10b22f3. The candidate addresses the process/TLS lookup and deferred-alias overhead, and includes a small additional Vulkan CPU-side optimization. The comparisons below rerun dev merge base Activation cache and deferred aliases
Additional reductions of repeated workThe cache change alone did not remove the entire Vulkan gap, so the per-operation path was examined further:
No kernel mathematics, shader code, random-number algorithm or trainer behavior was changed by this follow-up. CPU affinity is used only in one benchmark control; the application has no new affinity setting or forced scheduling policy. Performance against devWindows: RTX 5060 Ti, driver 617.42, optimized Release ( For each case and mode, two series of four alternating ABBA/BAAB blocks give 16 processes per side, with 1,001 timed samples per process. Each process reports its median. Table times are geometric means of those process medians; percentages are geometric means of the paired block ratios. All blocks, including unfavorable results and outliers, are retained. The baseline was rerun during each comparison. No local compilation ran alongside the Windows measurements. Windows, normal scheduling — no forced affinity:
The Vulkan series were -5.33% and -4.79%; the alias series were -0.90% and -0.08%. Windows scheduling control — benchmark main thread fixed to logical CPU 2 on both sides:
The Vulkan series were -8.27% and -4.96%. Alias results were -0.16% and +0.08%. CUDA dispatch varied between -3.31% and +3.44% across the two series. The CUDA aggregate results are consistent with parity in these measurements, not evidence of a meaningful speedup. macOS, the same dev/candidate comparison and repetition scheme:
Mac Vulkan stayed close to dev in both series (-0.43% / -0.26%). Metal was much more variable: series +4.03% / -1.65%, with individual blocks ranging from -16.10% to +19.59%. These results do not establish a Metal speedup or an absolute parity guarantee. The Windows Vulkan improvement repeated in all eight blocks under both scheduling modes. These are small-operation microbenchmarks, not full-training measurements, and they do not establish native Linux NVIDIA performance. Regression and additional checksSix new parameterized contracts cover:
Final validation:
The Linux checks exercise compilation and CPU contracts; they are not GPU performance validation. GPU-specific cases skipped in CPU runs are not counted as evidence of GPU correctness. Additional local diagnostics and checks:
The commit contains production code, regression tests and public contract documentation. Benchmark probes, raw measurements and development reports remain outside the repository changes. This candidate keeps the review's dev baseline unchanged; the native Linux NVIDIA comparison is the remaining performance verification. |
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Thanks, the process-scoped activation cache works. Vulkan small dispatch is now faster than dev.
I rebuilt 10b22f3 on Linux (CUDA 13.3 + Vulkan, Release, RTX 4080).
Tests
- Contracts 5/5,
tensor_backend_shutdown_contractsand the 71 backend gtests pass.
Performance against dev
Same setup as last round: dev merge base eecd2cc vs head, 15 alternating base/head process pairs, machine otherwise idle. Medians in ms:
| Case | dev | 10b22f3 | Change | Head slower |
|---|---|---|---|---|
| Vulkan dispatch_overhead_small (gui bench) | 0.5723 | 0.5119 | -10.5% | 0 of 15 pairs |
| CUDA deferred_alias_copy_materialize | 0.01173 | 0.01214 | +3.5% | 14 of 15 pairs |
| CUDA dispatch_overhead_small (gui bench) | 0.8808 | 0.8771 | -0.4% | 4 of 15 |
| CUDA eager_binary_dispatch_tight_loop | 13.005 | 12.914 | -0.7% | 1 of 15 |
| CUDA masking_ops_tight_loop | 1.600 | 1.588 | -0.8% | 4 of 15 |
| CUDA broadcast_binary_add | 0.7291 | 0.7294 | +0.0% | 10 of 15 |
| CUDA sort_middle_axis_64x1024x64 | 0.6641 | 0.6642 | +0.0% | 8 of 15 |
What is left
- Vulkan: the gap from r7 (+1.7%) is gone.
- CUDA deferred alias copy: it is still slower, about 0.4 µs per iteration (+6.4% at r7). Head is slower in 14 of 15 pairs, so this is not noise.
The case runs three lazy elementwise ops, one materialization, then data_ptr() on four aliases of the result (bench_tensor_launch_overhead.cpp:152). Every CUDA case that only dispatches is now at or below dev. So the remaining cost sits in the deferred path, Tensor::materialize_deferred_slow and the alias data_ptr() that follows it, not in dispatch. Please compare that path against dev for work that is new per materialization. Candidates are the TensorTargetBinding for the home stream, canonical_home_target, validate_materialized and the storage adoption.
Push a candidate and I'll run the same comparison again.
Requesting changes for this one gap only. Everything else is ready.
|
Pushed candidate 73a7125 for the remaining CUDA Scoped change
This removes two redundant reference-count increments and their matching temporary decrements per publication. GCC disassembly confirms the two increments disappear from this path. Validation, gate locking, stream ordering and backend checks are unchanged. Dispatch, activation caching, Vulkan-specific code, kernels and trainers are unchanged by this commit. For the exact benchmark on Windows, the source, deferred result, materialized result and current CUDA stream were all null, so the guarded home-target binding and canonicalization calls were not executed. The added backend-parameterized regression test destroys the first materialized handle while a sibling remains deferred, then verifies the sibling's storage, values and placement. It also checks that a view remains usable after that sibling is released. Validation
The other-backend runs were regression checks of the shared materialization helper; there is no additional HIP or other backend optimization in this commit. CUDA measurements against devThe reference remains Times below are geometric means of process medians; changes use paired geometric ratios.
Across the 30 WSL pairs, the aggregate change is -3.33%, but the median paired change is only -0.70% and individual pairs range from approximately -35.7% to +25.1%. Windows is effectively unchanged at the precision of these measurements. CUDA ran successfully under WSL, providing additional functional and GCC coverage. However, the timing variability is too large to establish that the approximately 0.4 microsecond native-Linux gap is closed, or to attribute the aggregate WSL improvement to this change alone. These results do not invalidate the review's consistent regression measurement on the RTX 4080 with CUDA 13.3. Parity on that native-Linux setup remains to be verified. Only the implementation and regression test are committed; benchmark experiments and raw reports remain local. |
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Built 73a7125 and ran the same comparison as last round (dev merge base eecd2cc vs head, 15 alternating base/head process pairs, idle machine, RTX 4080). Contracts 5/5, tensor_backend_shutdown_contracts and the 71 backend gtests pass. Vulkan small dispatch stays fixed at -10.3%, 0 of 15 pairs slower.
One thing is still open, the same one as before:
-
CUDA
deferred_alias_copy_materializeis 0.01149 -> 0.01180 ms (+2.7%), head slower in 15 of 15 pairs. The move ofdata_owner_/storage_meta_helped a bit (it was +3.5% at r8), but the gap is still there. Every dispatch-only CUDA case is within noise (dispatch_overhead_small +0.7%, eager_binary +0.6%, masking -0.1%, broadcast -0.0%, sort +0.1%, each 8-11 of 15 pairs), so the cost is per materialization.Diffing
Tensor::materialize_deferred_slowagainst dev leaves two changes on that path besides the moves: the home-stream binding is now built for every GPU tensor with a stream (dev only did it for CUDA when no current stream was set and the stream wasn't retired), and every CUDA stream now goes throughbackend_driver(backend).canonical_home_target(...), where dev left CUDA streams untouched. Both are a driver lookup plus a virtual call per materialization. Could you keep dev's early-outs there, or give CUDA's identity case a fast path? If neither moves it, the remaining suspects arevalidate_materializedand the aliasdata_ptr()path after it.
Push a candidate and I'll rerun it.
|
Addressed the remaining deferred-materialization paths in dd2a575. The CUDA stream policy now preserves the original early-outs: an explicit current stream takes precedence, and a retired home is not bound. CUDA home canonicalization is an identity fast path. When an implicit CUDA binding is actually needed, saving, setting and restoring the current stream also use the native CUDA operations directly, avoiding the registry lookup and virtual binding/restoration calls. These compatibility paths stay inside the private tensor backend layer; other backends retain their existing driver policy. I also reduced the publication work paid by deferred aliases:
The materialized-result validity/deferred checks, backend identity assertion, stream retirement checks and raw-pointer validation remain enabled. There are no changes to GPU kernels, trainers, dispatch activation caches, Vulkan pipeline keys, public API or driver vtables. ValidationAdded regression coverage for explicit-current CUDA stream precedence, retired CUDA homes, binding/restoration after a materializer exception and successful retry, sibling retry after failure, and concurrent alias publication with exactly one expression execution. Existing alias lifetime and queue coverage remains enabled.
Comparison with the requested dev baselineBaseline remains eecd2cc, rerun alongside the candidate. These are comparisons against dev, not against an intermediate PR revision. The probes were rebuilt against each side's actual headers, and performance runs were sequential after compilation and functional tests finished. Complete
The Windows median block change is +0.006%. Both available environments are close to parity with mixed pair/block results; I would not interpret these small aggregate differences as a reliable complete-workload speedup. To isolate publication cost, a separate diagnostic materializes and synchronizes the producer before timing const
Each diagnostic uses 8 processes per side in ABBA/BAAB order and 1001 retained samples per process. An additional complete-workload check with a non-null CUDA home was also at parity with dev (-0.068% aggregate, slower in 3/8 blocks). I investigated the other suggested suspects as well. The native Linux/RTX 4080 comparison remains to be verified. The earlier consistent +2.7% measurement was on 73a7125, before these local changes; Windows and WSL results do not establish that the native result is resolved. This candidate removes the identified CUDA stream-policy indirections and measurably reduces alias publication work while preserving the validation contracts. |
MrNeRF
left a comment
MrNeRF
left a comment
There was a problem hiding this comment.
Built dd2a575 and reran the same 15-pair comparison against the base snapshot.
- Contracts 5/5, the exit runner and the 71 tensor tests pass.
- Vulkan
dispatch_overhead_small: −10.5%, faster in 15/15 pairs. - CUDA
deferred_alias_copy_materialize: still +2.86%, slower in 13/15 pairs (r9 was +2.72%, 15/15). The other CUDA cases are within noise.
- dd2a575 does not move the number, so please revert it. It also puts
#if LFS_HAS_CUDAandbackend == GpuBackend::CUDAbranches into the shared materializer path, and we want backends to go through the driver interface without special cases in shared code. The CUDA fast path was my suggestion last round; that was the wrong direction. - With the gate held across both phases and
publishedtaken by reference,materialize_deferred_slownow does less work than dev on the alias path. So the remaining ~0.3 µs per run is somewhere else in that benchmark: building the lazyadd/mul/subexpression, planning it, or the fused launch. Timing those three phases separately on base and head should show which one.dispatch_overhead_smalland the eager loop are flat, which points at the lazy-node or planner side rather than the launch. - dev moved today (master sync). Merging it gives one conflict, in
tests/test_collision_jobs.cpp.
Summary
This PR gives CUDA, Vulkan, Metal and AMD HIP/ROCm one shared Tensor dispatch and resource-ownership contract, and connects the opt-in HIP implementation to the public Tensor facade.
Shared Tensor code describes operations and owns opaque native implementations. Backend drivers implement device inventory and selection, activation, execution targets, transfers, queues, completion, timing, memory-pressure policy and protocol capabilities. Native SDK calls and backend-specific decisions stay in the corresponding implementation files.
This is the next Tensor-focused contribution from the work originally collected in #3759, following the native operators merged in #3809. The scope includes the existing CUDA/Vulkan/Metal integration paths as well as HIP. HIP training, application preferences and portable distribution profiles remain separate work.
Common backend architecture
CUDA = 0,Vulkan = 1,Metal = 2,HIP = 3in every build. CMake generates compiled-backend lists and registration references from the enabled options. The shared registry does not include SDK headers or construct concrete backends.StorageReflayout does not grow.Unsupported timestamps, raw imports or protocol exports are reported explicitly through the interface. Public compatibility entry points retain their native contracts, including the no-CUDA compatibility default. The touched shared dispatch files are covered by the backend-neutrality gate.
AMD HIP/ROCm integration
LFS_TENSOR_HIPremains OFF by default, with opt-in support on Windows and Linux. macOS rejects that option. Standard build defaults and backend preference order remain unchanged; compiling HIP does not automatically prefer it.hip. HIP DLPack export is explicitly unavailable and directs callers to a CPU copy; it does not invoke CUDA pointer queries on HIP memory.Build instructions, the ROCm SDK 10.1.0 development reference, architecture selection and capability limits are documented in the HIP backend guide. Compiled wave32/wave64 architecture targets are not a claim of hardware qualification on every corresponding GPU.
Ordering, numerical contracts and failure handling
The common contract covers Int32/Int64 indices, selected counts, overlapping indexed writes and native conversion scratch. Clamp/wrap normalization happens before narrowing wide indices, with corresponding changes in the existing CUDA, Vulkan and Metal index implementations. Bool conversion preserves nonzero semantics, including fractional values and NaN. Random operations retain native seed/offset and odd-count behavior; explicit CUDA generator compatibility remains native.
Completions use one factory driven by native target/value/timeline and release policy. Readbacks retain their source and destination, capture execution faults at their queue position, and check them before publishing bytes. Empty readbacks and zero-byte completions still carry error snapshots. Repeated polling or consumption of a later context error cannot turn an earlier failed snapshot into success.
Queue destruction distinguishes checked execution errors after a successful drain from failure of native synchronization. The former releases the queue lease; the latter retains live ownership for safe recovery. HIP shutdown trims cached buffers before closing the context, retains resources whose cleanup failed, and supports retry. A worker with cached HIP ops reinitializes before synchronization after another thread shuts the backend down.
Vulkan expression-cache reuse also validates the push-constant layout. Incompatible cached artifacts are rejected with a diagnostic and recompiled; failed fresh compilation is not accepted as a usable null kernel.
These changes affect shared integration and selected native Tensor semantics. This description does not claim that every existing backend source is untouched. Trainer and rendering algorithm integration is outside this PR.
Dispatch overhead and comparison with dev
Host-side refinements avoid no-op binding dispatch, same-target synchronization calls, unnecessary allocation-metadata construction and repeated vector growth during fused preparation. Allocation adoption uses scoped rollback so MSVC can inline the helper; the normal deleter still resolves the backend on the releasing thread. Optimized objects were inspected alongside runtime measurements.
The current baseline is upstream dev merge-base
eecd2cc187652f00a9bcaa3692d438329df63b61, rebuilt from unmodified sources. Each new comparison reruns both dev and head in alternating processes on the same hardware. Windows uses matching Release/O2 /Ob2builds with FAST_COMPILE off, CUDA 12.8, RTX 5060 Ti and driver 617.42; macOS uses matching Release builds on an M4 Pro. These are Tensor microbenchmarks, not training-throughput measurements.Representative repeated
dispatch_overhead_smallcomparisons:Unfavorable repetitions are retained too: a Windows Vulkan series was +3.32%, a separate check reached +5.05%, and the complete macOS suite measured Vulkan small dispatch at +2.31%. CUDA repetitions changed sign. In the final CUDA control, eight ABBA/BAAB blocks (32 process runs, 1001 samples each) gave a block-geometric-mean difference of +1.27%, with a 95% block-bootstrap interval of [-4.25%, +6.25%]. The unchanged dev executable itself varied by up to 12.92% between repeated runs. This establishes measurement variability, not strict equivalence or a guaranteed speedup.
CUDA profiling found matching kernel/synchronization counts and launch parameters; separate API-count checks also matched. The attempted Vulkan captures did not contain API events and are not used to attribute driver costs. Parity on the reviewer's native Linux RTX 4080/CUDA 13.3 setup still requires confirmation.
The latest detailed follow-up preserves complete Windows/macOS GUI tables, CUDA launch-overhead tables, profiling details and limitations. Earlier common-dispatch measurements and host-preparation measurements remain available with their original intermediate-PR baselines; they are not substituted for the current dev comparison.
Validation
Shared permanent contracts cover compiled backends, worker first use, inventory without context creation, shutdown/reopen, scope restoration, queue/fused-output ordering, retained ownership, indices, empty/exact allocations, conversion semantics, RNG behavior and fault recovery. Python tests exercise identity, capabilities and CPU DLPack. Separate instrumented runtime tests cover cleanup failures without adding fault hooks to production modules.
Latest relevant local checks on head
359524b76:The generated-header consumers and actual indexing-test registration were also checked directly. Earlier native-module, host-loader and numerical-operator suites are recorded in the linked follow-ups; their earlier results are not represented as fresh execution of every suite on this head.
The existing Ubuntu Release workflow now additionally compiles the main-tree HIP adapter and executes CPU/runner/Python contracts without a GPU, reusing its dependencies. The native portable compile/export checks remain separate. The standard Ubuntu builds still run when the HIP path-filter job fails; HIP-specific steps require the filter output. Required GPU tests fail explicitly when the requested backend is unavailable.
Local checks are targeted builds and runtime tests, not a complete portable rebuild or certification of every full application CI job. Native Linux GPU execution and all untested GPU architectures remain outside the available hardware coverage. Current CI status is shown by the PR checks.
Public documentation includes Tensor backend contracts and the HIP build guide. Only production code, permanent tests, gate tooling and public documentation are committed; development reports, raw benchmark data and experimental scripts remain outside the branch.