Repository navigation
Conversation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
c75ccdd to
e3545ce
Compare
There was a problem hiding this comment.
🟡 Changes recommended
One or more issues must be addressed before approval.
2 open issues
What changed in this PR
Extends rocjitsu’s local GPU peer-memory model for stock-driver KFD and VFIO traffic.
Changes:
- Adds fabric backing, freshness, access-admission and lifetime interfaces.
- Updates shader apertures, deferred transfers, fault reporting and queue teardown.
- Adds RCCL tooling, dependency patches and regression coverage.
| File | Description |
|---|---|
| emulation/rocjitsu/tools/rocjitsu/main.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/xcd_fanout_kfd_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/vmm/vfio_user_client.h | Updated as part of this pull request. |
| emulation/rocjitsu/tests/vmm/vfio_user_client.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/sdma_ring_consumer_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/sdma_queue_scheduler_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/scripts/test_run_rccl_pod_emulated_node.py | Updated as part of this pull request. |
| emulation/rocjitsu/tests/pci_bar_address_decoder_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/instruction_cache_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/guest/hip_owner_peer_refetch.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/gpu_vm_service_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/gpu_queue_registry_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/execution_plugin_memory_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/CMakeLists.txt | Updated as part of this pull request. |
| emulation/rocjitsu/tests/cdna5_sdma_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/canonical_line_freshness_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/cache_allocation_failure_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/tests/aql_packet_processor_test.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/run-rccl-pod-vfio.py | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/run-rccl-pod-emulated-node.py | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/qemu-10.1-xio3130-link-capability.patch | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/qemu-10.1-vfio-user-class-code-default.patch | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/qemu-10.1-pcie-atomic-capabilities.patch | Updated as part of this pull request. |
| emulation/rocjitsu/scripts/apply-qemu-pcie-fixes.sh | Updated as part of this pull request. |
| emulation/rocjitsu/patches/libvfio-user/0001-preserve-small-64-bit-bar-upper-address.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0011-trace-kfd-import-map.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0006-abort-failed-sdma-pde-job.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0005-finalize-failed-cpu-pde-prefix.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0004-publish-cpu-tlb-seq-after-hdp.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0003-retry-failed-compute-tlb-flush.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0002-abort-sdma-pt-clear-dependency.patch | Updated as part of this pull request. |
| emulation/rocjitsu/patches/amdgpu/7.1.3.31500000-2390945.24.04/0001-link-pt-after-clear.patch | Updated as part of this pull request. |
| emulation/rocjitsu/lib/simdojo/include/simdojo/components/sparse_memory.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/simdojo/include/simdojo/components/pci_device.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/simdojo/include/simdojo/components/pci_bar_address_decoder.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vmm/vfu/vfio_server.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vmm/vfu/CMakeLists.txt | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/virtual_machine.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/virtual_machine.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/topology_assembly.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/soc.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/rj_vm.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/rj_vm_impl.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/fabric_backing.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/fabric_backing_kfd.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/fabric_backing_kfd.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/CMakeLists.txt | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/canonical_line_freshness.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/canonical_backing_id.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/canonical_backing_id.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/xcd.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/translated_local_cache_policy.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/spi.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/sdma_ring_consumer.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/sdma_ring_consumer.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/sdma_queue_scheduler.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/sdma_packet_processor.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/raw_backing_write_gate.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/vram_store.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/vram_store.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/physical_memory_access.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/peer_alias_observations.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/nbif_hdp_block_model.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/nbif_hdp_block_model.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/mes_block_model.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/mes_block_model.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/memory_hub_block_model.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/memory_hub_block_model.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/ip_block_model.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/ip_block_model.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/graphics_block_model.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/graphics_block_model.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/gpu_pci_device_spec.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/gpu_generation_registry.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/gpu_generation_registry.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/pci/CMakeLists.txt | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/mes_engine.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/mes_engine.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_ingress.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/mem_state.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/legacy_page_table.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/legacy_address_space.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/l1_vector_cache.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/l1_scalar_cache.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/interrupt_sink.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/gpu_queue_registry.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/gpu_queue_registry.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/dispatch_entry.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_queue_binding_factory.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/completion_tracker.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/command_processor.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/circular_ring_reader.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/circular_ring_reader.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/aql/aql_packet_processor.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/aql/aql_packet_processor.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/remote_driver.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/remote_driver.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/private_gtt_copy_source_gate.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/guest_kfd.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/events.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/kmd/linux/events.cpp | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/tensor_dma.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/include/rocjitsu/vm/rj_vm.h | Updated as part of this pull request. |
| emulation/rocjitsu/lib/rocjitsu/include/rocjitsu/daemon/rj_daemon.h | Updated as part of this pull request. |
| emulation/rocjitsu/CMakeLists.txt | Updated as part of this pull request. |
| emulation/rocjitsu/cmake/rj_libvfio_user.cmake | Updated as part of this pull request. |
| emulation/rocjitsu/AGENTS.md | Updated as part of this pull request. |
Review effort: Balanced
| if (identity < vmid_apertures_.size() && vmid_apertures_[identity]) | ||
| return vmid_apertures_[identity]; |
There was a problem hiding this comment.
I think we should preserve VMID versus PASID identity through queue dispatch and make the aperture resolver's name and parameter reflect that distinction. mes_engine.cpp puts a VMID in Queue::process_id for direct PM4 queues and a PASID there for MES queues; Wavefront::set_process_id passes either integer to this lookup.
The collision matters when programmed banks differ; equal banks mask it. The existing VFIO test makes banks 3 and 4 equal before testing PASID 77. Please also test a MES PASID below 16 with divergent banks, separately from a direct VMID lookup. On the PCI path, an unresolved bank should produce an explicit failure or diagnostic: currently set_process_id silently retains the default apertures when this returns null, so refusing the fallback does not itself reject the dispatch.
There was a problem hiding this comment.
We should make an unresolved PCI aperture lookup visible rather than silently retaining empty defaults. The identity tag fixes the PASID/VMID collision at 10ba7dde5735, but divergent banks still make vmid_apertures() return null and Wavefront::set_process_id() continues without a diagnostic. On the PCI path, the CU fallback apertures are zero, so a dispatch cannot route FLAT LDS through its intended shared aperture. Please fault the unsupported dispatch or at least warn.
It would also help to put the identity-kind regression in the default test suite: the current case lives only in vmm/vfio_device_host_test.cpp, behind ROCJITSU_ENABLE_VFIO. Cover divergent-bank PASIDs separately from direct VMIDs and verify the queue identity tags.
Please update wavefront.h:262–266 too: process_id() is still described as a PASID analog even though direct PM4 passes a VMID. Document the namespace selected by kind and the setter's default of Pasid.
There was a problem hiding this comment.
We should move the aperture identity regression into the default test suite. ProgramsGfx1250ShaderAperturesFromGuestRegistersByVmid is still gated by ROCJITSU_ENABLE_VFIO, which defaults to OFF; no in-tree workflow enables it. The test uses no VFIO transport: copying its body and configuration helper into a scratch translation unit linked against the default build passed the divergent-bank PASID, direct-VMID, warning, and reset cases. tests/pci/gpu_pci_device_test.cpp would let ordinary builds exercise this fix.
The new warning addresses the silent fallback. Please also update wavefront.h:262, :265, and :1133: the integer is a VMID or PASID according to kind, whose default is Pasid, rather than always a “PASID analog.”
| public: | ||
| FillAuthority() = default; | ||
| FillAuthority(FillAuthority &&) noexcept = default; | ||
| FillAuthority &operator=(FillAuthority &&) noexcept = default; |
c20b2fc to
ee0ddea
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
I think we should move the new emulation/rocjitsu/AGENTS.md commit-message and signing policies into a separate change. Those rules affect future contributions across rocjitsu independently of peer fabric behavior; reviewing them separately would keep repository policy decisions out of this feature's implementation review.
We should follow CONTRIBUTING.md's class convention for the new non-trivial types with methods. Under emulation/rocjitsu/, the related examples are lib/rocjitsu/src/rocjitsu/vm/fabric_endpoint.h:68 (PreparedTransfer, with ownership and a conversion operator) and lib/rocjitsu/src/rocjitsu/vm/fabric_endpoint.cpp:25 / :75 (Request and State, with lifecycle methods). Please use class with explicit public sections and update their forward declarations together, so the endpoint types follow the project's existing distinction between behavior-bearing types and simple records.
| return VmAccessOutcome::Revoked; | ||
| VmAccessOutcome outcome = VmAccessOutcome::Complete; | ||
| cache_.for_each_valid_line([&](simdojo::CacheTag &tag, uint64_t line_addr, uint8_t *data) { | ||
| if (outcome != VmAccessOutcome::Complete || tag.vmid != vmid || |
There was a problem hiding this comment.
We should invalidate translated-local entries using their full cache key in both L2 completion flush and V$ completion invalidation. These entries are allocated with vmid == 0 and address_space_identity == access.cache_identity(), so the nonzero-VMID filter skips them; the subsequent two-argument invalidate also uses the old key.
A regression populated through a real translated load still returns 0x13579bdf after both completion operations succeed and backing changes to 0x55667788, with no second backing read. Clearing V$ alone still hits the retained L2 entry; clearing both caches refetches correctly. Please cover this path alongside the existing completion tests that seed numeric-VMID entries.
There was a problem hiding this comment.
We should include the real translated-load path in the completion fix and its regression test. At 10ba7dde5735, the rebuilt repro still caches 0x13579bdf, then returns it after V$ completion invalidation and L2Cache::flush_for_completion succeed and backing changes to 0x55667788. The backing read count stays at one. Clearing V$ alone still hits L2; full invalidation refetches correctly.
The existing completion test seeds numeric-VMID entries, so it misses the translated entries keyed by address-space identity.
| } | ||
| void note_potential_execution_effect() { | ||
| if (execution_effect_possible_) | ||
| execution_effect_possible_->store(true, std::memory_order_release); |
There was a problem hiding this comment.
We should avoid repeatedly storing true to the dispatch-wide effect marker after a wave has marked it. Every non-NOP/non-ENDPGM instruction reaches this release store, and waves on different CPU workers share the same atomic, creating avoidable cache-line contention.
A per-wave “already marked” bit would let later instructions skip both publication and the mnemonic checks in compute_unit.cpp:1866. Alternatively, guard the store with a load if the marker is only a monotonic flag. The standalone marker benchmark shows contention, but an end-to-end workload slowdown has not been measured; this request is about removing redundant work from the instruction issue path.
There was a problem hiding this comment.
It would be useful to avoid the repeated shared store on the current head too. At 10ba7dde5735, compute_unit.cpp:1867 still calls this for every instruction except s_nop and s_endpgm, while completion only needs the marker's monotonic false-to-true transition. A relaxed-load guard before the release store would remove the redundant writes; a per-wave flag could also skip the mnemonic checks. No end-to-end slowdown has been measured.
| rj_vm_destroy(daemon->vm); | ||
| bool trace_complete = true; | ||
| if (daemon->trace_fabric_effects) { | ||
| auto *fabric = daemon->vm->vm ? daemon->vm->vm->fabric_endpoint() : nullptr; |
There was a problem hiding this comment.
We should separate trace completeness from retryable teardown failure and make the retained daemon handle safe to stop again. After a trace overflow, teardown() can successfully destroy the VM and set daemon->vm = nullptr, yet return ROCJITSU_STATUS_ERROR; rj_daemon_stop then retains the handle. Retrying as the public header instructs dereferences daemon->vm in this trace block.
Please report dropped trace records separately from VM-close failure, guard access to a destroyed VM, and add a stop/retry regression with a forced trace drop.
| struct CacheTag { | ||
| uint64_t tag = 0; | ||
| uint64_t coherence_epoch = 0; ///< Controller-defined lazy-invalidation generation. | ||
| uint64_t coherence_epoch = 0; ///< Controller-defined lazy-invalidation generation. | ||
| uint64_t admission_generation = 0; ///< Address-space admission generation at fill. |
There was a problem hiding this comment.
Can we keep translated-line metadata in a controller-supplied extension or side table, rather than adding it to every generic cache tag? CacheTag grows from 40 to 96 bytes here, including backing addresses, PTE attributes, and publication epochs used by rocjitsu's translated-local policy. I$, scalar caches, and other users also carry that storage, and generic allocation/reset code must now maintain those fields.
The address-space lookup key belongs in the generic cache interface; the remaining policy metadata could stay with the controllers that use it. If the shared layout is intentional, please explain that tradeoff and check a cache-heavy workload. The size increase alone does not establish a proportional throughput regression.
| if (soc_ != nullptr) { | ||
| // CU callbacks retain only weak device-local state. A completed read may | ||
| // record its peer translation, but never serves a hit or fill. | ||
| const std::weak_ptr<PeerAliasObservations> weak_catalogue = peer_alias_observations_; |
There was a problem hiding this comment.
Can we remove this PCI peer-read catalogue until it has a consumer, or make it an opt-in diagnostic with a reader and test? The handler revalidates the translation and freshness, takes the shared VFIO grant locks, and scans up to 4096 observations, then always returns false. snapshot_for_test() has no callers in the library, tools, or tests, so successful peer reads pay to maintain state that nothing currently uses.
If the catalogue stays, peer_alias_observations.h should also include canonical_line_freshness.h directly for its CanonicalLineFreshness::ReadStamp interface instead of relying on gpu_vm.h to include it.
There was a problem hiding this comment.
Can we also remove the unused VmPeerReadOrigin payload left behind by the catalogue removal? gpu_vm.h:245 still describes an observer that validates its 19 fields, and vfio_device_host.cpp:457 heap-allocates that payload for each prepared peer read, retaining the grant gate. No field is read now: memory_pipeline.cpp:867 only checks whether the pointer is non-null, and the PCI constructor clears the peer-cache handlers.
If the presence check is still needed for other callers, represent that condition directly with a bool or enum. The diagnostic TraceOrigin has a separate consumer and can stay.
|
|
||
| def guest_kfd_maps(guest_log: Path, config_path: Path, imported: dict) -> dict: | ||
| guest = guest_log.read_text(encoding="utf-8", errors="replace") | ||
| if "guest: amdgpu.p4v_ipc_map_trace=Y" not in guest: |
There was a problem hiding this comment.
It would help to document the exact external artifacts and patch sets used for the reported RCCL gates, including how to build the driver module, QEMU, kernel, and guest image. The launcher accepts prebuilt inputs, while this four-rank trace checker requires amdgpu.p4v_ipc_map_trace=Y, introduced by 0011-trace-kfd-import-map.patch. The tree also carries eight behavioral amdgpu patches without an application/build recipe.
Please distinguish which patches are needed for traffic, which only enable diagnostics, and which were used in each reported eight-rank run. Clarify what “stock-driver” means with those modifications, and whether the numbering gaps represent intentionally omitted patches. This would make the validation reproducible without assuming that this four-rank checker is the eight-rank gate.
There was a problem hiding this comment.
Should we remove the legacy checker, or label its external prerequisites in its docstring? At this head, emulation/rocjitsu/scripts/check-vfio-fabric-trace.py:88 still requires amdgpu.p4v_ipc_map_trace=Y and the custom P4V KFD records, but the driver patch series has been removed. tests/guest/README.md now describes a stock module and says the diagnostic patches are outside the repository. The script should identify the required diagnostic driver and its four-rank scope so readers do not mistake it for a checker for the documented eight-rank zero-patch gate.
| /// @details A permit spans admission and the backing effect without holding a | ||
| /// page-table or allocation lock across callbacks that can reenter KFD. | ||
| class VmRawBackingWriteGate { | ||
| private: |
There was a problem hiding this comment.
It would be easier to read this gate if its public API came first, followed by the private Span/Entry helpers. That also follows CONTRIBUTING.md's public/protected/private order. Moving the helper definitions below Permit, Closed, and the admission methods lets readers find the lifetime contract before the implementation; inline member bodies can still refer to those later private declarations.
There was a problem hiding this comment.
I think we should still put the public gate API first at 10ba7dde5735. The private Span/Entry definitions and helpers occupy lines 23–51 before Permit and Closed; moving them below the admission methods exposes the lifetime contract first and follows CONTRIBUTING.md's access-section order.
|
|
||
| namespace rocjitsu::amdgpu { | ||
|
|
||
| struct MemoryIngressSpan { |
There was a problem hiding this comment.
Can we put MemoryIngressSpan alongside MemoryIngressPermit, MemoryIngressAdmission, and VmTransferIngress in gpu_vm.h? This header already includes gpu_vm.h and adds only the span, so it does not establish a separate dependency boundary. Keeping the ingress vocabulary together would make it easier to find; it would also avoid the current file brief describing a permit that is defined elsewhere.
There was a problem hiding this comment.
Can we keep MemoryIngressSpan with the other ingress types in gpu_vm.h? At 10ba7dde5735, memory_ingress.h still includes that entire header and adds only this span, while its file brief describes a permit defined elsewhere. Consolidating them would make the related API easier to find without losing a dependency boundary.
| // pages of one owner allocation. The requester uses the ordinary peer mapping; | ||
| // there is no guest-side control register or publication protocol. | ||
|
|
||
| #include <hip/hip_runtime.h> |
There was a problem hiding this comment.
We should provide a discoverable build-and-run path for this guest test. Neither hip_owner_peer_refetch.cpp nor the adjacent rccl_pod_init is referenced by an in-tree build or packaging rule, so a reader cannot tell how to put the probe in the guest image and reproduce it. An explicit guest-payload target or documented build/packaging commands would address this; it need not use the host HIP test harness.
There was a problem hiding this comment.
We should still add a reproducible guest build-and-run recipe for this probe at 10ba7dde5735. I could find no in-tree build, packaging, or documentation reference to hip_owner_peer_refetch.cpp or rccl_pod_init. A guest-payload target or documented compiler, image-installation, and invocation steps would make the test usable without requiring the host HIP harness.
| exit 1 | ||
| fi | ||
|
|
||
| class_patch=$script_dir/qemu-10.1-vfio-user-class-code-default.patch |
There was a problem hiding this comment.
It would help to collect these three QEMU diffs under patches/qemu/10.1/ and update this script's paths. The same PR places the amdgpu and libvfio-user series under patches/<project>/; keeping all vendor changes there gives reviewers one place to find the versioned dependencies while the application script stays in scripts/.
There was a problem hiding this comment.
It would help to keep the QEMU series alongside the other vendor patches. At 10ba7dde5735, this script still loads three qemu-10.1-*.patch files from scripts/, whereas amdgpu and libvfio-user changes live under patches/. Moving the diffs to patches/qemu/10.1/ and updating these paths would give readers one place to find the versioned dependency changes.
ee0ddea to
cb4058b
Compare
cb4058b to
0730b13
Compare
github-actions
Bot
added
the
documentation
10ba7dd to
6d1c61e
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
I think it would help to document the peer transport contract together with the shipped fixtures. Under emulation/rocjitsu/:
- Add
gfx1250_mi455x_helios_4gpu.jsonandgfx1250_mi455x_kmd_4node_2gpu_if.jsonto the pre-built table indocs/configuration.md:21: Helios supplies local ScaleUp routes; the IF fixture supplies local GPU pairs in four nodes, with NET between nodes. - Explain in the architecture overview that
FabricEndpointexecutes these routes while explicit manifests still suppress KFD hive discovery. Updatelib/rocjitsu/src/rocjitsu/kmd/linux/simulated_kfd.cpp:3138, whose comment says this suppression lasts only until routes can execute; they already execute in this PR. - Update
lib/rocjitsu/src/rocjitsu/config/system_graph.h:64, which describes only ScaleUp despite supporting InfinityFabric, and clarify that both currently use the same unit-latency transport. This documents the present model without promising hardware-specific timing.
A link from docs/qemu-vfio.md to the pod launcher and the existing gate/artifact documentation would also help readers find the supported peer path.
| public: | ||
| MemoryIngressPermit() = default; | ||
| explicit MemoryIngressPermit(std::function<void()> release, | ||
| std::optional<MemoryPtEpoch> pt_epoch = {}) |
There was a problem hiding this comment.
Can we use one representation for the optional epoch within MemoryIngressPermit? The constructor accepts std::optional<MemoryPtEpoch>, stores an empty MemoryPtEpoch for absence, and pt_epoch() reconstructs the optional using operator bool(). Keeping the optional throughout, or using the empty-epoch convention throughout, would remove the conversion in both directions and make the absence contract easier to follow.
| return -EINVAL; | ||
| const bool had_memfd = alloc.memfd >= 0; | ||
|
|
||
| if (alloc.memfd < 0 && alloc.host_ptr) { |
There was a problem hiding this comment.
We should restore the USERPTR -EPERM check before snapshot promotion and restore the original IPC-export test expectation. The base tree rejects this export through share_allocation_locked(), but this path copies the caller's pages to a memfd and remaps the exporting process's GPU view to that snapshot.
The scratch KfdIoctlTest at 10ba7dde5735 returns success from IPC_EXPORT_HANDLE; a subsequent CPU write of 0x2222 leaves the GPU host extent holding 0x1111. The renamed UserptrDmabufExportFailsAndIpcSnapshotLeavesCallerMemoryAlone test checks only the caller's CPU memory, so it misses the loss of shared USERPTR backing.
There was a problem hiding this comment.
We should restore the USERPTR rejection before snapshot promotion and cover the exporter's GPU view in the regression. At 6d1c61e, simulated_kfd.cpp:8247–8291 still copies caller memory into a memfd and redirects GPU pages with remap_page_host_ptrs, breaking their alias with later CPU writes. UserptrDmabufExportFailsAndIpcSnapshotLeavesCallerMemoryAlone only checks the caller's CPU memory, so its passing result does not cover the issue demonstrated above.
There was a problem hiding this comment.
I think we should reject USERPTR export before snapshot promotion and check the exporter's GPU view in the regression. Rebuilding the scratch case against 8ef6a6ffcb7c still returns success from IPC_EXPORT_HANDLE and moves that view off the caller's pages: after a CPU store of 0x2222, the GPU host extent still contains 0x1111. The existing test checks only the caller's CPU memory, so it misses this loss of shared backing.
| alloc_flags = alloc.flags; | ||
| alloc_gpu_id = alloc.gpu_id; | ||
| backing_owner = alloc.backing_owner; | ||
| const int public_fd = alloc.memfd >= 0 ? alloc.memfd : alloc.dmabuf_fd; |
There was a problem hiding this comment.
This would be safer if IPC export retained export_access_mode through the same descriptor helper as dma-buf export. An IPC import of an O_RDONLY public descriptor keeps a writable private memfd and records the original mode. Here, F_DUPFD_CLOEXEC duplicates that private fd without applying the recorded mode, so a second IPC import derives O_RDWR and loses the original public restriction.
Please share the descriptor-retention logic with share_allocation_locked() and routed export_dmabuf_ioctl, while leaving their PTE-policy differences at the call sites. Extend PrimeHandleToFdReturnsTheBuffer through a second IPC export/import: its current single-hop case does not exercise this path. This follows from the source; I have not run a second-hop repro.
There was a problem hiding this comment.
This still needs the descriptor-mode preservation on the IPC path. At 6d1c61e, simulated_kfd.cpp:8386 duplicates the private writable fd with F_DUPFD_CLOEXEC, while dma-buf exports at :7961 and :7996 reopen it with alloc.export_access_mode. A subsequent IPC import therefore reads O_RDWR from the re-exported fd and loses the original O_RDONLY restriction. Please share that descriptor-retention step and add a second IPC export/import hop to the regression; this remaining path was checked from source.
6d1c61e to
8ef6a6f
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
I think it would help to finish documenting the transport/discovery distinction raised above. The new text at emulation/rocjitsu/docs/configuration.md:47 still says hive suppression lasts until local routes can execute, but these routes already execute while explicit manifests continue to suppress KFD hive discovery. The same stale explanation remains in simulated_kfd.cpp:3145, rj_vm.cpp:100, and sysfs.h:136.
Please add the Helios and four-node IF fixtures to the configuration table and show the VM-owned FabricEndpoint in docs/vm-design.md:49. Its file brief (lib/rocjitsu/src/rocjitsu/vm/fabric_endpoint.h:5) also still says ScaleUp only, although local IF routes use the same unit-latency transport. These descriptions should distinguish emulator transport from guest discovery without implying hardware-specific IF timing.
We should apply CONTRIBUTING.md's class convention to the remaining behavior-bearing types as well as the endpoint types covered in the earlier comment. Under emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/, these include vm/fabric_endpoint.cpp:110 (WireMessage, with message initialization), vm/amdgpu/gpu_vm.h:182 (VmAccessAdmission, with a conversion operator), and vmm/vfu/vfio_target_grants.h:118 (PreparedSegment, with owned leases and copy()). Use explicit public sections and update matching forward declarations together; simple data records can retain their existing form.
| if (blockIdx.x != 0 || threadIdx.x != 0) | ||
| return; | ||
|
|
||
| const uint32_t first = *owner; |
There was a problem hiding this comment.
We should make the owner-data reads cacheable and keep the cached line resident until the peer write, then verify this probe fails when peer invalidation is disabled. With the documented gfx1250 hipcc -O2 build, all three owner reads emit flat_load_b32 ... scope:SCOPE_SYS, which the simulator maps to Mtype::UC. Both owner fences also emit global_inv scope:SCOPE_SYS, explicitly invalidating V$ and flushing L2 before the final read.
The value labelled hit therefore does not demonstrate a cache hit, and the probe can pass despite broken peer-driven invalidation. Separately ordered control flags and cacheable owner reads would make it test the cached-owner refetch claimed in tests/guest/README.md:172. This conclusion comes from the generated ISA and executor; the guest probe was not run.
There was a problem hiding this comment.
We should still use cacheable owner reads and independently ordered control flags, then check that disabling peer invalidation makes the probe fail. Recompiling the documented gfx1250 build at 8ef6a6ffcb7c emits flat_load_b32 ... scope:SCOPE_SYS for all three owner-data reads and global_inv scope:SCOPE_SYS at both owner fences. Those reads bypass the simulator caches, so a passing result does not exercise peer-driven eviction of a resident owner line. This was checked in generated ISA; the guest was not run.
There was a problem hiding this comment.
Please keep the ready/done handshake from evicting the owner data line, and add a negative control for the peer data write's invalidation. The cacheable loads fix the previous ISA issue, but ready and done still live in the peer-mapped owner allocation. Reading ready here enters FabricEndpoint::execute_effect, whose VFIO fallback performs whole-domain WritebackInvalidate, even for a different page. That read necessarily completes after the two owner loads and before the requester data write, so the write no longer encounters the resident owner line the probe is supposed to test. The later peer store to done can invalidate it too.
I reproduced this against the current build using FabricEndpointCoherenceTest: cache the old data, verify that a raw backing write alone remains stale, issue a real peer read of a ready flag 4096 bytes away, then perform the same raw data write with no invalidation. The owner reads the new value and the test passes. Thus this handshake can mask a missing invalidation on the data write. Use control storage that does not invoke the fabric's cache maintenance, or another handshake that preserves the tested line, and verify that suppressing data-write invalidation makes the probe fail. The scratch test exercised the endpoint directly; I did not run the guest.
There was a problem hiding this comment.
We should exercise the host-mapped control accesses in the negative control at emulation/rocjitsu/tests/fabric_endpoint_test.cpp:976. Moving the guest flags to hipHostMalloc addresses the previous peer-page handshake, but AControlHandshakeCannotHideMissingDataInvalidation only assigns stack integers to ready and done. It assumes those accesses bypass the fabric rather than testing that property.
A focused case should perform requester reads/writes through a system-memory mapping, check that fabric effect counts stay unchanged, and check that the seeded owner line remains resident before the data write. Then suppressing data invalidation can demonstrate that the probe fails. The current test establishes only that a raw backing write leaves a cached line stale; rename it to that narrower claim if it cannot exercise the control path.
|
|
||
|
|
||
| class RankCleanupTest(unittest.TestCase): | ||
| def test_exited_group_does_not_prevent_peer_cleanup(self): |
There was a problem hiding this comment.
It would be good to register this regression in tests/CMakeLists.txt alongside VfioGuestRunUnitTest and RcclScaleoutRunUnitTest. Direct execution passes its one cleanup-race test, but ctest -N -R 'RcclPod|RankCleanup' discovers zero tests, so ordinary CTest runs miss it. A Python -B entry would make the coverage effective.
| typedef struct rj_daemon_options_t { | ||
| uint32_t struct_size; ///< Set to sizeof(rj_daemon_options_t). | ||
| uint32_t flags; ///< Bitwise OR of RJ_DAEMON_OPTION_* values. | ||
| uint32_t fabric_trace_capacity; ///< Zero selects the default trace capacity. |
There was a problem hiding this comment.
We should document this field's accepted values in the public header and use a named constant for the shared limit/default. rj_daemon_start_with_options rejects capacities above 200000 and any nonzero capacity without RJ_DAEMON_OPTION_TRACE_FABRIC_EFFECTS; zero selects 200000 when tracing is enabled. The header currently mentions only an unspecified default, so C callers cannot discover the valid contract without reading the implementation or getting ROCJITSU_STATUS_INVALID_ARGUMENT.
8ef6a6f to
c46f6fc
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
Added one comment and some inline replies
| /// window awaiting its last bypass permit. Group cache maintenance invalidates | ||
| /// every local domain epoch, so a second closed dirty alias cannot be left | ||
| /// unresolved. | ||
| class GpuAliasIngressGate : public std::enable_shared_from_this<GpuAliasIngressGate> { |
There was a problem hiding this comment.
This would be easier to maintain if GpuAliasIngressGate's out-of-line methods lived in gpu_alias_ingress.cpp, registered in kmd/linux/CMakeLists.txt. The gate already has its own header and dedicated tests, but its implementation occupies simulated_kfd.cpp:149–803 alongside the ioctl implementation. Pairing the header with its implementation, as events and guest_kfd do here, would make the admission and drain logic easier to find and review independently of the large driver translation unit.
aef001a to
3f5260d
Compare
06adf17 to
6eceac0
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
Original anchor: emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/config/system_graph.cpp:391
A unique local path is dropped here when its ScaleUp and Infinity Fabric edges are mixed. SystemGraph::local_peer_routes documents missing, ambiguous, GPU-transit, remote-tainted, and overlong paths as fail-closed; mixed class is not in that list, and I did not find a regression for it next to InfinityFabricProjectsOnlyLocalGpuRoutes.
That skip is not required by today's transport. FabricEndpoint::wire never reads SystemPeerRoute.link_class -- fabric_endpoint.cpp:366 always installs latency 1 -- so a unique mixed path would execute the same way as the uniform Helios ScaleUp and four-node IF fixtures. The projector is refusing a route only so it can stamp one enum value.
A tray that uses IF for one hop and ScaleUp for another is a natural combination of the two fixtures this PR ships. It would currently get no FabricEndpoint route and look like a path-finding failure.
I would prefer accepting mixed local fabric paths so that follow-up does not have to change the projector. If fail-closed is intentional, please name mixed class in system_graph.h:109 and add a test that a unique mixed path grants no route. This is not asking for class-specific timing.
| /// Unlike the request reservation, this permit never survives a retry. | ||
| using PreEffectAdmission = std::function<amdgpu::MemoryIngressAdmission()>; | ||
|
|
||
| struct PreparedTransfer { |
There was a problem hiding this comment.
Nit: CONTRIBUTING.md says non-trivial types with methods must be class, not struct.
PreparedTransfer owns a unique deferred operation and exposes operator bool(), so it is a behavior-bearing type. The adjacent result type VmTranslationResult in gpu_vm.h already follows that rule: class with a public conversion operator. Please match that here, keep a public: section, and update any forward declarations together.
The same pattern is still used for Request and State in fabric_endpoint.cpp, which have lifecycle methods rather than being simple records.
| namespace rocjitsu { | ||
| using amdgpu::VmAccessOutcome; | ||
|
|
||
| struct FabricEndpoint::Request { |
There was a problem hiding this comment.
Nit: for consistency with CONTRIBUTING.md's class rule, Request should be a class.
It has cancel_queued() and cancel_before_effect(), so it is not a simple record. State in this file has the same issue (take_idle_reservation(), finish_reservation_release()). Nearby nested types already use class when they have methods (WireMessage, Node, Operation). Please switch these two and keep explicit public: sections so the endpoint internals match that distinction.
| /// every local domain epoch, so a second closed dirty alias cannot be left | ||
| /// unresolved. | ||
| class GpuAliasIngressGate : public std::enable_shared_from_this<GpuAliasIngressGate> { | ||
| class Batch; |
There was a problem hiding this comment.
Nit: please keep public: / protected: / private: ordering as required by CONTRIBUTING.md.
The class currently opens with an implicit private nested forward declaration of Batch before the first public: section. shared_ptr<Batch> only needs an incomplete type, so the forward declaration can live in the later private section with Drain, Batch, and Entry. That is the layout VmRawBackingWriteGate already uses: public API first, then private nested helpers.
| /// fresh again. A monotonic epoch also rejects stamps after maintenance clears | ||
| /// the table and later inserts the same canonical line. | ||
| class CanonicalLineFreshness { | ||
| struct State; |
There was a problem hiding this comment.
Nit: please keep public: first, then private:, as required by CONTRIBUTING.md.
State and Snapshot are declared in an implicit private section before the public API. Nested FillAuthority / ReadStamp constructors only mention those types through shared_ptr, so the forward declarations can move to the existing private section at the bottom of the class. That matches the public-then-private layout used by VmRawBackingWriteGate.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once |
There was a problem hiding this comment.
Nit: every other header in this directory starts with a /// @file / /// @brief file comment after the license (vfio_device_host.h, vfio_server.h, guest_region_map.h, bus_plan.h).
Please add the same boilerplate here so this new grant header is discoverable the same way as the rest of the vfio-user layer. CONTRIBUTING.md asks for @brief / @details on public APIs; the file brief is the established local form of that for this directory.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once |
There was a problem hiding this comment.
Nit: the other headers in simdojo/components/ carry a /// @file / /// @brief file comment (pci_device.h, sparse_memory.h, cache.h, memory_interface.h).
This new decoder is the only sibling without that boilerplate. Please add a file brief here so the BAR decoder is documented the same way as the rest of the component library. CONTRIBUTING.md asks for @brief / @details when documenting public APIs; this directory applies that at file scope.
| [[nodiscard]] std::optional<amdgpu::MemoryIngressPermit> | ||
| try_acquire_for_drained_owner(std::span<const Key> aliases, | ||
| const DrainAuthority &authority) noexcept; | ||
| /// Whether every key is an exact member of this active drained close. |
There was a problem hiding this comment.
Nit: the neighboring methods document their public contract with @brief / @details (try_acquire_for_drained_owner above, try_freeze_for_effect below).
This comment is already a full-sentence public-API note, but it uses a bare /// Whether... form. Please match the file's established doxygen tags so the drained-owner coverage rule is formatted like the rest of this header. CONTRIBUTING.md asks for @brief / @details when documenting public APIs.
| class Batch; | ||
|
|
||
| public: | ||
| struct Key { |
There was a problem hiding this comment.
Nit: can we match the class/struct convention used by the rest of this gate?
Key has operator<=>, so CONTRIBUTING.md wants it as a class with an explicit public section (see MemoryPtEpoch in gpu_vm.h, which is a class with comparison and conversion operators). The later nested Drain, Batch, and Entry types are the opposite: they are classes with only public data and no methods, which is how this file treats simple records elsewhere as struct (Span / Entry in raw_backing_write_gate.h). Please invert those choices so method-bearing types are class and data-only nested records stay struct.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once |
There was a problem hiding this comment.
Nit: the other headers in this directory, including the new gpu_alias_ingress.h, start with a /// @file / /// @brief comment after the license. legacy_gpu_vm.h is the local #pragma once example of that form.
Please add the same file brief here so this gate is documented like the rest of kmd/linux/. CONTRIBUTING.md asks for @brief / @details when documenting public APIs; that is the established header convention in this directory.
| auto *wave = cu->dispatch_wf_at(/*wf_id=*/0, /*wg_id=*/0, /*pc=*/0, | ||
| /*num_sgprs=*/104, /*num_vgprs=*/32); | ||
| ASSERT_NE(wave, nullptr); | ||
| wave->set_process_id(3, rocjitsu::amdgpu::ProcessIdentityKind::Vmid); |
There was a problem hiding this comment.
Can we extend this test to execute one FLAT access through the VMID bank, so it covers the per-wave FLAT LDS routing named in the PR description? It currently checks only the wave's aperture fields and src_shared_base. The routing consumers in compute_unit.cpp changed from the CU's shared_aperture_base_/shared_aperture_limit_ to wf.shared_aperture_base()/wf.shared_aperture_limit(): lane classification and LDS rebasing at :948-:982, the wait-scoreboard plan at :1223, and the debugger fault predicate at :2211.
No test executes FLAT while the wave's apertures differ from the CU defaults. This is the only test that programs a VMID bank, and it does not issue a memory instruction. The existing FlatLds tests program the CU apertures, which each wave inherits at dispatch. If any routing consumer reverted to the CU fields, those tests would still pass. On the PCI path, where the CU defaults are zero, RCCL's flat LDS pointers would then go to global memory.
A small case would set CU defaults that differ from bank 3 (as the ambiguous-PASID block below already does), dispatch with set_process_id(3, Vmid), and run a flat_store_b32/flat_load_b32 pair at kSharedBase + offset. It should check that the value lands in the wave's LDS, that DScnt rather than LOADcnt is charged, and that no global backing access occurs. I found this gap by reading the source and did not run a mutation test. The aperture encoding itself matches GC 12.1: SH_MEM_BASES.PRIVATE_BASE[15:10] gives address bits 63:58, SHARED_BASE[31:16] gives bits 63:48, and the shared window is 4 GiB. That agrees with the CDNA5 ISA aperture check, the gc_12_1_0_sh_mask.h field masks, and upstream gfx_v12_1.c, which programs every nonzero VMID with the same >> 58 / >> 48 bases.
There was a problem hiding this comment.
This is partly addressed at 6eceac0, and my earlier claim was too broad. GpuVmPipeline.DdsStepReportsUnsupportedInstructionAndPcBeforeFaulting and RoutedMemoryObservationTest.AFlatAccessToTheSharedApertureIsSeenAsTheLdsAccessItBecame now give the wave its own apertures while the CU defaults stay zero. Reverting the lane classification and LDS rebasing in route_memory_inst to the CU fields fails both tests.
The wait-scoreboard consumer is still not covered. I reverted only track_memory_wait (compute_unit.cpp:1223) to shared_aperture_base_/shared_aperture_limit_ in an out-of-tree build. That mutant passed every runnable rocjitsu_tests case: 7763 passed, and the 4 failures (two EffectiveConfigTest and two BinaryTranslatorE2E handoff tests) fail identically on the unmodified binary from the same directory.
On the PCI path, where the CU defaults are zero, that reversion would make the wait checker plan FLAT LDS lanes against LOADcnt/STOREcnt instead of DScnt. A small wait-checked FLAT-to-LDS case would pin it: give the wave a per-wave aperture with a zero CU default, and assert that the destination is tracked on DScnt, for example by checking that no premature-overwrite warning appears after s_wait_dscnt 0. This is a coverage request; the current code reads the correct fields.
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
Original anchor: emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/config/system_graph.cpp:391
A unique local path is still dropped here when its ScaleUp and Infinity Fabric edges are mixed. 6eceac0e renamed the Helios fixture to gfx1250_mi455x_4gpu_scaleup.json and left this projector unchanged. SystemGraph::local_peer_routes documents missing, ambiguous, GPU-transit, remote-tainted, and overlong paths as fail-closed; mixed class is not in that list. There is still no regression next to InfinityFabricProjectsOnlyLocalGpuRoutes, which only covers a homogeneous IF path.
That skip is not required by today's transport. FabricEndpoint::wire never reads SystemPeerRoute.link_class -- it always installs latency 1 -- so a unique mixed path would execute the same way as the uniform ScaleUp and four-node IF fixtures. The projector is refusing a route only so it can stamp one enum value. The switch-taint comment a few lines above still says "ScaleUp edges" even though adjacency already includes Infinity Fabric, which is the same leftover homogeneity assumption.
A tray that uses IF for one hop and ScaleUp for another is the natural combination of the two fixtures this PR ships. It would currently get no FabricEndpoint route and look like a path-finding failure. Stamping a Mixed enum would still throw the per-hop identity away; keeping the edge classes on the projected route would let today's unit-latency wire accept the path and leave a later class-specific hop as a fill-in.
I would prefer accepting mixed local fabric paths so that follow-up does not have to change the projector. If fail-closed is intentional, please name mixed class in system_graph.h next to the other refusals and add a test that a unique mixed path grants no route. This is not asking for class-specific timing.
| /// every local domain epoch, so a second closed dirty alias cannot be left | ||
| /// unresolved. | ||
| class GpuAliasIngressGate : public std::enable_shared_from_this<GpuAliasIngressGate> { | ||
| class Batch; |
There was a problem hiding this comment.
Nit: please keep public: / protected: / private: ordering as required by CONTRIBUTING.md.
The class currently opens with an implicit private nested forward declaration of Batch before the first public: section. shared_ptr<Batch> only needs an incomplete type, so the forward declaration can live in the later private section with Drain, Batch, and Entry. That is the layout VmRawBackingWriteGate already uses: public API first, then private nested helpers.
| class Batch; | ||
|
|
||
| public: | ||
| struct Key { |
There was a problem hiding this comment.
Nit: can we match the class/struct convention used by the rest of this gate?
Key has operator<=>, so CONTRIBUTING.md wants it as a class with an explicit public section (see MemoryPtEpoch in gpu_vm.h, which is a class with comparison and conversion operators). The later nested Drain, Batch, and Entry types are the opposite: they are classes with only public data and no methods, which is how this file treats simple records elsewhere as struct (Span / Entry in raw_backing_write_gate.h). Please invert those choices so method-bearing types are class and data-only nested records stay struct.
| [[nodiscard]] std::optional<amdgpu::MemoryIngressPermit> | ||
| try_acquire_for_drained_owner(std::span<const Key> aliases, | ||
| const DrainAuthority &authority) noexcept; | ||
| /// Whether every key is an exact member of this active drained close. |
There was a problem hiding this comment.
Nit: the neighboring methods document their public contract with @brief / @details (try_acquire_for_drained_owner above, try_freeze_for_effect below).
This comment is already a full-sentence public-API note, but it uses a bare /// Whether... form. Please match the file's established doxygen tags so the drained-owner coverage rule is formatted like the rest of this header. CONTRIBUTING.md asks for @brief / @details when documenting public APIs.
| /// fresh again. A monotonic epoch also rejects stamps after maintenance clears | ||
| /// the table and later inserts the same canonical line. | ||
| class CanonicalLineFreshness { | ||
| struct State; |
There was a problem hiding this comment.
Nit: please keep public: first, then private:, as required by CONTRIBUTING.md.
State and Snapshot are declared in an implicit private section before the public API. Nested FillAuthority / ReadStamp constructors only mention those types through shared_ptr, so the forward declarations can move to the existing private section at the bottom of the class. That matches the public-then-private layout used by VmRawBackingWriteGate.
The out-of-line State in canonical_line_freshness.cpp also has a constructor and invalidate_all_locked(), so CONTRIBUTING.md's class rule wants that definition (and this forward declaration) to be class with an explicit public section. Snapshot can stay a struct; it is data-only.
| /// Unlike the request reservation, this permit never survives a retry. | ||
| using PreEffectAdmission = std::function<amdgpu::MemoryIngressAdmission()>; | ||
|
|
||
| struct PreparedTransfer { |
There was a problem hiding this comment.
Nit: CONTRIBUTING.md says non-trivial types with methods must be class, not struct.
PreparedTransfer owns a unique deferred operation and exposes operator bool(), so it is a behavior-bearing type. The adjacent result type VmTranslationResult in gpu_vm.h already follows that rule: class with a public conversion operator. Please match that here, keep a public: section, and update any forward declarations together.
The same pattern is still used for Request and State in fabric_endpoint.cpp, which have lifecycle methods rather than being simple records. Nearby nested types already use class when they have methods (WireMessage, Node, Operation).
| bool has_pt_epoch_ = false; | ||
| }; | ||
|
|
||
| class MemoryIngressAdmission { |
There was a problem hiding this comment.
Nit: can we match the class/struct split used immediately below?
MemoryIngressAdmission has only public data and no methods, which is the same shape as MemoryIngressSpan. CONTRIBUTING.md requires class for types with methods; the local convention for data-only records is struct (MemoryIngressSpan here, Span / Entry in raw_backing_write_gate.h). MemoryIngressPermit and MemoryPtEpoch stay class because they have conversion operators. Please make this result a struct so the two adjacent record types read the same way.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once |
There was a problem hiding this comment.
Nit: other test support headers start with a /// @file / /// @brief comment after the license and before #pragma once (vfio_user_client.h, execution_plugin_test_util.h, amdgpu_elf_test_support.h).
This file puts a prose comment after the pragma and before the includes. Please use the same file-brief form so the TCP rendezvous helper is documented like the rest of tests/. CONTRIBUTING.md asks for @brief / @details when documenting public APIs; that is the established header convention here.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once |
There was a problem hiding this comment.
Nit: other test support headers start with a /// @file / /// @brief comment after the license (vfio_user_client.h, execution_plugin_test_util.h, elf_test_support.h).
Please add the same file brief here so this async VM helper is documented like the rest of tests/. CONTRIBUTING.md asks for @brief / @details when documenting public APIs; the file brief is the established local form.
| } | ||
| // The translated peer route cannot preserve local cache semantics for a | ||
| // second lane. Deny a cacheable peer effect before any lane is issued. | ||
| if (has_local_request && has_cacheable_peer_request) |
There was a problem hiding this comment.
Mixing a local lane with an uncached peer lane in one vector instruction makes the local lane bypass the requester's write-back L2. This check rejects only cacheable peers, but once peer_data_route is set, every request in the instruction goes through execute_translated_transfer. Under the KFD translator, supports_translated_local_cache() is false. A local load therefore reads backing directly, and a local store writes backing directly, while the L2 line for that address is left untouched.
I reproduced both directions at 6eceac0 with a scratch KfdPeerBindingTest (the gfx950 MI355X simulated-KFD fixture). The setup has a requester-local VRAM word whose requester L2 line holds a dirty 0x22222222 while backing holds 0x11111111. Owner VRAM is mapped to the requester without the cacheable opt-in, so its route is Deferred and its mtype is UC.
- Control:
flat_load_dwordwith EXEC=1 on the local address returns0x22222222. - The same load with lane 1 on the peer address returns
0x11111111in lane 0. Lane 1 gets the correct peer word, and the wave keeps running with no fault. - A mixed
flat_store_dwordof0x33333333writes both backings, but the next single-lane local load still returns0x22222222from L2. I infer that a later writeback of that dirty line would also revert the store; I did not run that step.
The scalar path has the same shape: classify_scalar_data_route sends the whole load down the deferred path when any dword is peer, so the local dwords skip the scalar cache and L2 as well.
A lane's architectural result should not depend on which other lanes are active. The smallest repair is probably to keep local requests on the legacy cache path and route only the peer requests through the deferred transfer. The other option is to write back and invalidate the affected local L2 lines before a translated local access. Extending the existing mixed-lane block in CuScalarAndVectorPeerReadsReuseRequesterCache with a UC peer lane and a dirty local line would cover both directions. Failing closed, as the cacheable case does, would avoid the stale result, but it would fault legal gather kernels whose pointers mix local and peer memory.
| @@ -283,6 +283,24 @@ function(rj_add_test) | |||
| list(APPEND ARG_ENVIRONMENT "${_runtime_dir_env}") | |||
| endif() | |||
|
|
|||
| if(RJ_ENABLE_TSAN) | |||
| # A two-GPU scoped cache operation holds more than 128 CU cache locks. | |||
There was a problem hiding this comment.
I think we should scope this TSan relaxation to the cases that need it, and treat the lock count behind it as a design limit of scoped maintenance rather than a sanitizer setting. As written, every rj_add_test and every discovered gtest in a TSan build runs with detect_deadlocks=0. That removes lock-order-inversion checking from the whole suite, including unrelated suites, to accommodate one new protocol.
The cause is new in this PR. acquire_scoped_alias_maintenance_group reaches acquire_all in device_cache_coherence.cpp:590. For every domain in the group, acquire_all takes the domain's atomic and coherence locks, then locks every registered L1ScalarCache, every L1VectorCache, every L2Cache, and every memory-side gate. It does this twice per operation: once for preparation and once for the returned leases. The base tree had no L1 registration, so the number of simultaneously held locks now grows with CUs times GPUs in the group, and the 8-GPU pod topology is larger than the two-GPU unit fixtures.
I reproduced the limit at this head with build-clang-23-tsan. Running KfdIoctlTest.FixedAllocationRemapPublishesDirtyAliasAndRetries with TSAN_OPTIONS=detect_deadlocks=1 hits TSan's n_all_locks_ < 128 CHECK (0x80, 0x80), and the process then hangs until my 600-second timeout. With detect_deadlocks=0 it passes in 156 ms.
Two narrower shapes would keep the rest of the suite covered:
- Put the multi-GPU scoped-maintenance cases in a separate gtest target or CTest label, and set
detect_deadlocks=0only there. The comment would then name the specific protocol that exceeds TSan's per-thread limit. - Bound the lock set by the number of domains rather than CUs.
L1VectorCache::load_translated_localalready takesacquire_local_cache_access(), which holds the domainmutex_shared, andacquire_alltakes that same mutex exclusively. If every L1 path that can hold a selected alias entered through that gate, the exclusive domain lock would exclude L1 access without sweeping each L1 mutex. I have not established whether the scalar and non-translated vector paths can adopt the gate without other ordering problems, so treat this as a direction to evaluate rather than a drop-in fix.
Separately, the installed CTest fragments written by cmake/rj_generate_installed_gtest_ctest.cmake set only ROCJITSU_CONFIG_DIR and ROCJITSU_KERNEL_DIR. An installed TSan run of these cases would therefore still abort on the CHECK.
There was a problem hiding this comment.
We should make the explicit TSan test registrations reject an empty GoogleTest selection, in both the build and installed trees. At 657e3ccc, all 79 listed names exist, but deleting or renaming one leaves a registration whose old --gtest_filter runs zero tests and exits 0. I checked that behavior with this build's GoogleTest, and these registrations have no passing-test assertion.
A small fix is to require exactly one passing test with PASS_REGULAR_EXPRESSION, forwarding it through rj_add_test and rj_install_ctest. Renamed tests may still be discovered under their new names, so this is about preventing stale explicit entries from reporting a false pass. Keep the protocol's detector-limit explanation near the list: new or renamed cases do not automatically inherit its override.
6eceac0 to
bd26684
Compare
bd26684 to
3267cd2
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
We should pass d.translated.transfer_token to the masked scalar-store write at emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.cpp:671 and add a deferred-owner regression. The three-argument overload rejects a deferred transfer as Malformed, so this path never posts the owner write.
The decoded CDNA4 s_buffer_store_dword repro uses base 0x100, bound 4, and an asynchronous peer range [0x100, 0x104). Its mask is 1; issuing it returns Malformed and leaves backing zero, while a plain s_store_dword control writes 0x11223344. This affects ordinary in-range buffer stores too. The existing masked-store retry test uses synchronous backing and misses the missing token.
We should correct the current commit's validation statement to gfx1250. For emulation/rocjitsu/tests/guest/hip_owner_peer_refetch.cpp:23, hipcc --offload-arch=gfx950 -O2 -c fails because global_load_b32 and s_wait_loadcnt are unsupported; the same compilation for gfx1250 succeeds. The README already specifies gfx1250, so the claim that this probe compiles for gfx950 overstates the validated architectures.
| : execute_translated_scalar_atomic_rmw<uint64_t>(d); | ||
| } | ||
| if (d.translated.mixed_local_peer_route) { | ||
| if (d.elem_size < sizeof(uint32_t)) { |
There was a problem hiding this comment.
Can we remove the sub-dword branch from the mixed scalar path and assert its dword-size precondition? mixed_local_peer_route is set only under d.elem_size >= sizeof(uint32_t) at line 552, so lines 569–583 cannot execute. Byte and halfword accesses already use the ordinary deferred path below. Removing this duplicate makes the per-dword routing contract clearer.
| throw util::ConfigError("fabric GPU and coherence domain counts differ"); | ||
| prepared_domains.reserve(socs_.size()); | ||
| for (size_t i = 0; i < socs_.size(); ++i) { | ||
| auto domain = socs_[i]->cache_coherence(); |
There was a problem hiding this comment.
It would help to document SystemCoherenceDomain in config/system_graph.h:56 as descriptive membership that grants no coherence or peer-access capability, matching the schema. install_peer_routes() builds runtime maintenance domains from each SoC here; it does not consume system.coherence_domains. Stating that distinction on the C++ type would prevent callers from mistaking the parsed node0_candidate grouping for active cache-coherence policy.
There was a problem hiding this comment.
It would help to carry the clarified SystemCoherenceDomain contract into the user-facing configuration docs too. At 657e3ccc, docs/configuration.md:20 lists the four-GPU ScaleUp fixture, and docs/sphinx/conceptual/json-configuration.md:54 describes system, but neither explains that system.coherence_domains is descriptive inventory only.
The fixture's node0_candidate membership is parsed and validated, while install_peer_routes() still builds runtime maintenance domains from each SoC. Please state that changing this grouping does not change runtime cache maintenance or grant peer access, so users do not mistake it for an active configuration control.
| namespace rocjitsu { | ||
| namespace { | ||
|
|
||
| struct LineKey { |
There was a problem hiding this comment.
We should make LineKey a class with an explicit public section, following CONTRIBUTING.md's rule for types with methods. It defines operator==; using the same form as the new GpuAliasIngressGate::Key keeps comparison keys consistent with that documented convention.
| class EffectAuthority { | ||
| public: | ||
| EffectAuthority() = default; | ||
| EffectAuthority(amdgpu::MemoryIngressAdmission admission) |
There was a problem hiding this comment.
This would be clearer with an explicit constructor: converting a MemoryIngressAdmission into an EffectAuthority should be visible at the call site. The VFIO callers already use direct initialization, so that usage remains supported. This is a small API safeguard, not a demonstrated runtime failure.
Stock-driver UC PCIe peer mappings need live grants, coherent backing, and VMID-specific apertures. Mixed local and peer instructions must retain requester cache behavior. Freed GPU ranges can alias later host buffers, and closing an SDMA queue must cancel its deferred effects. Model node-local effects through a VM-owned endpoint, scoped ingress leases, and keyed cache invalidation. Preserve IPC export modes and USERPTR peer refusal. Route local scalar words and vector requests through requester caches while deferring peer requests. Let new SVM access grants claim retired host spans; keep fixed BO views denied until munmap. Cancel in-flight SDMA work on force close, and exercise host-mapped control flags. Scope TSan's 128-lock detector workaround to 79 maintenance cases and four interposer cases in both CTest trees while retaining race detection. Initialize the Tensor DMA fixtures and give the controlled peer wave its LDS size. Keep pinned QEMU and libvfio-user transport patches; include no AMDGPU patches. Verify the full serial Release CTest suite (9,202 tests, zero failures), 79 scoped TSan cases, four interposer cases, and five normal-detector cases. Pass the focused ASan/UBSan regressions, including 50 SDMA cancellation repeats, and 10 Release FlashAttention repeats. Compile the HIP peer probe for gfx1250 and check the diff.
3267cd2 to
657e3cc
Compare
kuhar
left a comment
kuhar
left a comment
There was a problem hiding this comment.
Some comments, plus a few inline replies
| // ROCr locks a pageable host copy destination with ACCESS_IN_PLACE. | ||
| // A fresh SVM grant can reuse a freed USERPTR VA; flag-only SET_ATTR | ||
| // cannot, and the old tombstone stays denied until this admission. | ||
| proc.claim_gpu_range_locked(args->start_addr, args->size); |
There was a problem hiding this comment.
We should validate and preserve the GPU-specific SVM access policy before reopening a retired range. grants_host_access checks only the attribute type, and claim_gpu_range_locked() clears the denial without using the requested GPU ID. A later NO_ACCESS request updates the stored attributes but never republishes a denial.
Extending NewSvmHostGrantClaimsOnlyItsFreedUserptrSpan reproduces both cases: ACCESS_IN_PLACE for nonexistent GPU 0xdeadbeef succeeds and lets GPU 38144 read the freed USERPTR page; after a legitimate grant, NO_ACCESS for GPU 38144 also succeeds but leaves that page readable. Both negative tests fail while the existing partial-span test passes. Please keep grants and revocations scoped to the requested GPU, or reject unsupported requests without reopening access.
| /// One virtual span admitted before a local GPU memory effect. | ||
| struct MemoryIngressSpan { | ||
| uint64_t address = 0; | ||
| size_t size = 0; |
There was a problem hiding this comment.
It would help to document MemoryIngressAdmission's outcome/permit contract alongside these ingress types, using the project's @brief / @details form. In the transfer loop, outcome determines whether access may proceed, while any returned permit keeps the admission alive through the local effect; an empty permit is not itself a refusal. A short note on the result type would make that distinction discoverable without reading the transfer implementation.
| if (!matched) { | ||
| generation = range_generation; | ||
| // A canonical peer alias can change behind either GPU's caches. | ||
| // Until the fabric provider supplies coherence, its local path must |
There was a problem hiding this comment.
We should describe the uncached policy here as the production UC PCIe BAR contract, rather than waiting for a missing fabric coherence provider. FabricEndpoint now performs cache maintenance before backing access, and the controlled cacheable-peer tests use its selected-alias path; production peer aliases intentionally remain uncached.
The same stale explanation appears at simulated_kfd.cpp:8015 on IPC import. Updating both comments would distinguish today's mapping policy from the existing maintenance mechanism without suggesting that cacheable peer access should be enabled.
| // The worker can retain a QueueRecord after publishing service_active_=false. | ||
| // Cancel its deferred owner effect before detach reports force-close done. | ||
| if (force) | ||
| queue->ring_consumer_.cancel_in_flight(); |
There was a problem hiding this comment.
Can we add a deterministic regression for force-close cancellation while the worker still retains its QueueRecord? Removing this cancel_in_flight() call in a scratch build leaves all 541 SDMA/queue/close/detach tests passing. The full binary also has the same four unrelated direct-run failures as the unmodified binary.
VmTransferToken already cancels on destruction, so dropping the last record reference can hide the missing explicit cancellation. The test needs to retain the worker's record after service_active_ becomes false, detach with a deferred owner write pending, and then step the owner engine. It should check that no write lands and the ring consumer releases its pinned VM snapshot before that retained record is destroyed.
The gfx1250 peer path uses UC PCIe BAR mappings. Local pod traffic must reach modeled GPU memory through stock driver registers and standard VFIO DMA maps, with safe ordering and lifetime across deferred effects.
Add fabric endpoints, hub and HDP behavior, backing freshness, and fail-closed KFD and VFIO peer grants. Program shader apertures from GC register writes, route FLAT LDS accesses by wave bank, and preserve GuestKfd map progress on retry. Focused tests and both eight-rank RCCL gates pass.