Repository navigation
Conversation
A system-scope relaxed notification store can become visible to the host before earlier stream data writes. The notification dispatch's ending release cannot order data visibility before that store. Insert a system-release barrier before ordinary WRITE/update blits and batches containing writes. Capture the barrier into graphs so every replay releases earlier stream writes. Keep metadata slots aligned and avoid an additional barrier for WAIT-only batches. Add a 32/64-bit data-publication regression covering eager operations, native capture, captured batches, and explicit graphs. Only the host re-arms the response; a preceding GPU reset WRITE can conceal this bug. Both host input and the resulting GPU output use independent references. On gfx1201 with same-develop CLR/ROCr, the same regression binary fails all submission modes before this change and passes after it. Record the WRITE packet cost and the unrelated suite-order hipGetLastError failure in the PR test results.
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
There was a problem hiding this comment.
🟡 Changes recommended
Graph updates that change the captured packet layout can leave completion-signal patches on incorrect or removed packets.
1 open issue
What changed in this PR
Adds a leading system-release barrier so stream memop notifications publish earlier writes, extending the dispatch-scope fix from #12692.
Changes:
- Apply publication barriers to ordinary writes and write-containing batches.
- Capture barriers for graph replay while preserving WAIT-only behavior.
- Add 32/64-bit publication regression tests across five submission modes.
| File | Description |
|---|---|
| projects/hip-tests/catch/unit/stream/hipStreamMemOpPublication.cc | Adds publication regression coverage. |
| projects/hip-tests/catch/unit/stream/CMakeLists.txt | Builds the new test. |
| projects/hip-tests/catch/config/configs/unit/stream.yaml | Registers test metadata. |
| projects/clr/rocclr/platform/command.hpp | Tracks whether batches contain writes. |
| projects/clr/rocclr/device/rocm/rocvirtual.hpp | Declares the publication-barrier helper. |
| projects/clr/rocclr/device/rocm/rocvirtual.cpp | Emits or captures publication barriers. |
| projects/clr/hipamd/src/hip_stream_ops.cpp | Detects writes in eager batches. |
| projects/clr/hipamd/src/hip_graph_internal.hpp | Detects writes in graph batches. |
Review effort: Balanced
| hsa_barrier_and_packet_t packet{}; | ||
| packet.header = kBarrierPacketReleaseHeader; | ||
| std::memcpy(const_cast<uint8_t*>(cmd.getAqlPacket()), &packet, sizeof(packet)); |
| void VirtualGPU::dispatchMemoryPublicationBarrier(amd::Command& cmd) { | ||
| if (cmd.getPktCapturingState()) { | ||
| // This release must precede the notification on every graph launch. Issuing | ||
| // a barrier on the capture queue would only fence graph instantiation. | ||
| hsa_barrier_and_packet_t packet{}; | ||
| packet.header = kBarrierPacketReleaseHeader; | ||
| std::memcpy(const_cast<uint8_t*>(cmd.getAqlPacket()), &packet, sizeof(packet)); | ||
| if (metadata_preloader_.HasMetadataQueue()) { | ||
| // Keep the metadata and AQL packet arrays parallel, with an invalid slot | ||
| // for this non-kernel packet. | ||
| cmd.getMetadataPacket(); | ||
| } | ||
| } else { | ||
| // Also consumes pending cross-engine dependencies, including SDMA copies. | ||
| dispatchBarrierPacket(kBarrierPacketReleaseHeader); | ||
| } |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[IMPORTANT] Captured batch memop packet count now depends on params, which can break graph-exec param updates
When a batch memop is captured, the barrier is emitted only if the batch contains a write. A WAIT-only batch captures one packet ([kernel]). A batch with a write captures two ([release barrier, kernel]), plus an extra metadata slot. The patch does not show the instantiated graph's per-packet bookkeeping being rebuilt when an update changes this layout. Calling hipGraphExecBatchMemOpNodeSetParams() to switch a node between WAIT-only and write-containing can leave the bookkeeping indexed to the old layout. That bookkeeping covers segment completion/synchronization patches and metadata slots. The update can then land a completion signal on the leading barrier instead of the memop kernel, so a dependent segment may start before the memops finish. The reverse update can leave a stale patch pointing past the shorter packet list.
Suggested fix: Make the captured layout independent of params, for example by always emitting the release barrier (or a no-op barrier slot) for graph-captured batch memops. Alternatively, have the exec param-update path detect a layout change and rebuild or retarget the segment synchronization and metadata plan, or reject the update. Add a test that updates a batch memop exec node WAIT-only→write and write→WAIT-only with hipGraphExecBatchMemOpNodeSetParams(), adds a dependent node/segment after it, and checks ordering and data.
| // Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. | ||
| // | ||
| // Permission is hereby granted, free of charge, to any person obtaining a copy | ||
| // of this software and associated documentation files (the "Software"), to deal | ||
| // in the Software without restriction, including without limitation the rights | ||
| // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell | ||
| // copies of the Software, and to permit persons to whom the Software is | ||
| // furnished to do so, subject to the following conditions: | ||
| // | ||
| // The above copyright notice and this permission notice shall be included in | ||
| // all copies or substantial portions of the Software. | ||
| // | ||
| // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | ||
| // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | ||
| // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | ||
| // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | ||
| // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | ||
| // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN | ||
| // THE SOFTWARE. |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[IMPORTANT] New test file is missing the SPDX-License-Identifier header
The new source file has the full MIT license text but no SPDX-License-Identifier: MIT line, which new files are required to carry before merge.
Suggested fix: Add // SPDX-License-Identifier: MIT to the file header, following the convention used in other recent hip-tests sources.
| // The system-scope release at the end of the blit dispatch is too late: | ||
| // the host or a peer can observe the relaxed notification store first. | ||
| dispatchMemoryPublicationBarrier(cmd); |
There was a problem hiding this comment.
🤖 AI-assisted pre-review (for @mangupta)
[SUGGESTION] Every eager hipStreamWriteValue now issues an extra barrier packet even when nothing is pending
ROCCLR_COMMAND_STREAM_WRITE_VALUE now always calls dispatchMemoryPublicationBarrier, which in eager mode dispatches a separate system-release barrier packet. Signaling patterns that issue many back-to-back writes pay one extra AQL packet and fence per write. This includes cases where the previous packet already ended with a system-scope release and no stream writes have happened since.
Suggested fix: Consider skipping the barrier when the last packet dispatched on this queue already carried a system-scope release and no later work has been enqueued. Reuse the existing fence-state tracking rather than adding new state. Document the measured overhead in the PR.
A memop graph node can grow from one packet to two when it starts writing, or shrink when it becomes WAIT-only. Index-based packet replacement leaves completion patches on the leading barrier or pointing at the removed kernel, allowing dependent segments to run early or never complete. Use the existing full recapture path when the packet count changes. This clears pending batch references, rebuilds the sync plan and refreshes the flat and filtered packet buffers. Keep incremental updates for unchanged packet counts. Add 32/64-bit regressions that alternate WAIT-only and WAIT+WRITE nodes at a segment boundary with dependent consumers. Exercise executable setters and whole-graph updates with an earlier pending kernel update, checking actual producer data and the notification value on every replay. Validated on gfx1201: default, round-robin, DFS and small-queue regressions; existing graph update, segment ordering and node enable tests; strict publication probes and shared-IO protocol. The old candidate fails the same regression binary. A peer-device test skips on the single GPU.
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
|
@mangupta, the repeated CI failures have a configuration cause:
The refreshed snapshot at 2026-10-21 08:47 UTC contains 274 check-runs and still no I prepared a separate CI fix: #13884. It derives applicability from trusted producer workflows, retains covered/observed/unknown checks, and fetches all check-run pages. All 63 local unit/integration tests pass. A completely offline replay of this PR's saved metadata/checks returns 1 on the baseline and 0 on the candidate; no upstream check result was changed by that replay. Could the CI maintainers review and merge that fix into the trusted base, then start a new policy event for this PR using the updated base SHA? The bot checks out the base, so putting the fix only on this HIP branch or rerunning an event pinned to the old base would not apply it. This account does not have upstream Actions write access. The HIP/CLR head remains unchanged. |
|
Related downstream report and application-level workaround: FlashML-org/FreeToken#764. FreeToken's CPU+single-GPU hybrid MoE path exhibited stale CPU expert inputs during ROCm graph replay even though capture succeeded. That PR bypasses the memcpy + memop handoff with existing mapped-memory publish/combine kernels and adds graph-output checks before model capture. Its R9700 / ROCm 7.14.1 validation includes 72 passing tests and Qwen3-30B-A3B-FP8 hybrid Graph requests at 512 and 32K input lengths. This CLR PR addresses the underlying publication-ordering problem independently of FreeToken. The two changes operate at different layers: the FreeToken workaround protects unpatched runtimes; the CLR change fixes the notification's publication boundary. Full-model validation against this patched CLR is still pending. The remaining end-to-end check is to restore FreeToken's original memcpy + memop path, run its multi-layer replay regression, then validate model outputs and hybrid Graph performance. FreeToken's existing results should not be treated as model-level validation or speedup evidence for this runtime patch. |
AlexXAmd
left a comment
AlexXAmd
left a comment
There was a problem hiding this comment.
The PR states:
“On gfx1201, the CPU can observe a memop notification before preceding D2H copies are visible, read stale inputs, and produce incorrect results.”
This is a serious claim. If correct, it suggests a fundamental ordering/visibility issue on gfx1201. Before making a broad change in CLR, I think we need to establish the following:
Prove the issue with a minimal reproducer.
Please provide a simpler reproducer that isolates the D2H copy and memop notification ordering, without graph-specific complexity.
Identify the root cause.
If the stated behavior is confirmed, we need to understand why it occurs specifically on gfx1201. In particular, we should determine whether an underlying software or hardware component is failing to provide the ordering/visibility required by the relevant HSA/AQL semantics.
Determine the correct component for the fix.
Even assuming the claim is correct, it does not establish that CLR is the appropriate layer to fix it. We should identify the component responsible for the missing guarantee before adding a workaround here.
I am also concerned that the proposed change appears to affect all ASICs while the reported problem is gfx1201-specific. If this introduces additional synchronization or release operations, we need to understand the performance cost imposed on unaffected ASICs and why that cost is necessary.
Finally, I don't understand why the fix is limited to the graph path. The stated problem is broader than graph execution: if a memop notification can become CPU-visible before preceding D2H copies, the same ordering issue should be considered wherever this sequence can occur.
For these reasons, I don't think this PR is ready to proceed in its current form. We first need proof of the isolated issue, root-cause analysis, and identification of the correct component in which to fix it.
|
@AlexXAmd I have completed the isolated reproducer and additional HSA/AQL controls requested in your review. The evidence now identifies a specific missing operation in CLR's software sequence: publication to SYSTEM before the notification store becomes observable. Below are the source-level explanation, controls, and the proposed correction. 1. Graph-free proof and controlled CLR replacement The 85-line HIP/C++ reproducer below tests one 544-byte D2H copy followed by one default WRITE. A CPU observer acquires the notification and checks a snapshot before the main thread synchronizes the stream. Device input preparation finishes before the tested sequence; the host destination is poisoned. There is no graph, batch, GPU WAIT, reset WRITE, custom GPU kernel, or application computation. On one R9700/gfx1201, using the same binary, ROCr, COMGR, driver, and GPU, changing only the HIP/CLR DSO gives:
Neither run timed out. The current locked ROCm 10.1 foundation (HIP 7.16.26385) separately fails 9962/10000; that is a different runtime control, not part of the DSO-only comparison above. The develop libraries are local rebuilds reporting HIP 7.17.26422, not a release-package claim. I have not tested a newer develop revision in this follow-up. 2. The concrete CLR composition and why the existing scopes are insufficient The eager call chain is: The baseline WRITE branch calls The actual AQL trace for the failing minimal sequence is: The HSA header definitions distinguish three guarantees:
Consequently, the notification's ending SYSTEM release is too late to provide the required pre-store publication. This locates the insufficiency in the sequence CLR constructs from a copy and a low-level relaxed notification primitive. The primitive itself is allowed to be relaxed; its system scope is not a substitute for release ordering. Under the intended default WRITE publication contract, the layer composing those operations must establish the missing boundary. 3. Independent HSA controls distinguish a missing request from a violated guarantee I bypassed CLR's copy/memop submission and queued the tested copy and notification directly on an independent HSA queue. HIP is used only for allocation/input setup. The copy kernel and memory types are fixed, and the notification dispatch always has acquire/release SYSTEM. Only the indicated publication operation changes:
All groups completed without timeouts. The no-fence barrier changes timing but does not restore correctness. In the separately compiled control kernels, relaxed notification emits These results are consistent with HSA executing the requested weak/strong sequences. They do not require a hardware or ROCr guarantee violation to explain the failure. The demonstrated software defect is the absence of the publication request in the default-memop composition, not evidence that a requested SYSTEM release was ignored. 4. Proposed CLR fix, including eager and graph paths The proposed helper changes the sequence to:
CLR is a concrete fix point because it owns the API-to-command translation, dispatch scopes, pending dependencies, and graph packet replay. A release operation in the device helper is another implementation option; the controls show it works for this case, but it would require its own integration and dependency analysis. I am not claiming that an extra AQL barrier is the only correct or optimal implementation. 5. Scope, API contract, and cost The contract being implemented is default stream WRITE publishing earlier stream writes before its value can serve as a notification, consistent with the intent of #12692. HIP's public header describes execution after earlier stream commands, but is less explicit about CPU publication. I would appreciate confirmation of that intended contract; I am not using CUDA documentation as a substitute for a HIP specification. Only gfx1201 was tested; I cannot explain a gfx1201-only distinction without measurements on another ASIC. This is an absent software publication boundary, rather than evidence of a hardware-specific erratum, so the candidate currently supplies it on the ROCr path for all ASICs. That does not establish that every existing sequence needs an additional packet, or that all-ASIC cost is acceptable. Historical local GPU-event medians on gfx1201 were eager WRITE 2.889→5.988 µs/op, graph single WRITE node 2.554→3.833, graph batch of 128 0.331→0.342. Seven samples of 64000 operations were used, without exclusive hardware or locked clocks; these are cost estimates, not cross-ASIC qualification. Proving when an existing release can safely be reused remains necessary before eliding barriers. PAL test gating/CI failures, peer/NIC publication, metadata-queue hardware, and other ASICs remain open. The PR description has been revised to distinguish these limits from the demonstrated CLR sequence defect. 85-line graph-free reproducer: copy_notify_repro.cppSelect exactly one GPU. hipcc -O2 -pthread copy_notify_repro.cpp -o repro
HIP_VISIBLE_DEVICES=GPU-541f8c51daca3a7e EXPECT_PCI=0000:43:00.0 timeout 60s ./repro 10000// Graph-free publication reproducer. Build: hipcc -O2 -pthread copy_notify_repro.cpp -o repro
#include <hip/hip_runtime.h>
#include <array>
#include <atomic>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <fstream>
#include <iostream>
#include <thread>
#define HIP(call) do { hipError_t e = (call); if (e != hipSuccess) { \
std::fprintf(stderr, "%s: %s\n", #call, hipGetErrorString(e)); std::exit(2); } } while (0)
constexpr unsigned words = 136;
static uint32_t expected(unsigned epoch, unsigned index) {
return 0x80000001u ^ (epoch * 107u) ^ (index * 13007u);
}
int main(int argc, char** argv) {
const unsigned iterations = argc > 1 ? std::strtoul(argv[1], nullptr, 10) : 10000;
int count = 0, version = 0;
HIP(hipGetDeviceCount(&count));
if (count != 1 || !iterations) return 2;
HIP(hipSetDevice(0));
HIP(hipRuntimeGetVersion(&version));
char pci[32]{};
HIP(hipDeviceGetPCIBusId(pci, sizeof(pci), 0));
const char* wanted_pci = std::getenv("EXPECT_PCI");
if (wanted_pci && std::strcmp(wanted_pci, pci)) return 2;
std::printf("IDENTITY hip=%d pci=%s iterations=%u\n", version, pci, iterations);
std::ifstream maps("/proc/self/maps");
for (std::string line; std::getline(maps, line);)
if (line.find("libamdhip64") != std::string::npos || line.find("libhsa-runtime64") != std::string::npos)
std::cout << "MAP " << line << '\n';
std::cout.flush();
uint32_t *device = nullptr, *host = nullptr;
uint64_t* ready = nullptr;
HIP(hipMalloc(reinterpret_cast<void**>(&device), words * sizeof(*device)));
HIP(hipHostMalloc(reinterpret_cast<void**>(&host), words * sizeof(*host),
hipHostMallocMapped | hipHostMallocCoherent));
HIP(hipExtMallocWithFlags(reinterpret_cast<void**>(&ready), sizeof(*ready), hipMallocSignalMemory));
__atomic_store_n(ready, uint64_t{0}, __ATOMIC_RELEASE);
hipStream_t stream{};
HIP(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
std::atomic<unsigned> checked{0};
std::atomic<bool> stopped{false};
unsigned wrong_before_sync = 0, wrong_after_sync = 0;
std::thread observer([&] {
std::array<uint32_t, words> snapshot{};
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
while (__atomic_load_n(ready, __ATOMIC_ACQUIRE) != epoch) {
if (stopped.load()) return;
if (std::chrono::steady_clock::now() > deadline) { stopped = true; return; }
}
std::memcpy(snapshot.data(), host, sizeof(snapshot));
for (unsigned i = 0; i < words; ++i) {
if (snapshot[i] != expected(epoch, i)) { ++wrong_before_sync; break; }
}
checked.store(epoch, std::memory_order_release);
}
});
std::array<uint32_t, words> input{};
for (unsigned epoch = 1; epoch <= iterations && !stopped.load(); ++epoch) {
for (unsigned i = 0; i < words; ++i) input[i] = expected(epoch, i);
HIP(hipMemcpyAsync(device, input.data(), sizeof(input), hipMemcpyHostToDevice, stream));
HIP(hipStreamSynchronize(stream)); // Finish input setup before the tested sequence.
std::memset(host, 0xa5, sizeof(input));
HIP(hipMemcpyAsync(host, device, sizeof(input), hipMemcpyDeviceToHost, stream));
HIP(hipStreamWriteValue64(stream, ready, epoch, hipStreamWriteValueDefault));
while (checked.load(std::memory_order_acquire) != epoch && !stopped.load()) std::this_thread::yield();
if (stopped.load()) break;
HIP(hipStreamSynchronize(stream)); // CPU snapshot was already checked.
for (unsigned i = 0; i < words; ++i) {
if (host[i] != expected(epoch, i)) { ++wrong_after_sync; break; }
}
}
const bool timeout = checked.load() != iterations;
stopped = true;
observer.join();
std::printf("RESULT checked=%u before_sync_errors=%u after_sync_errors=%u timeout=%d\n",
checked.load(), wrong_before_sync, wrong_after_sync, int(timeout));
HIP(hipStreamDestroy(stream)); HIP(hipFree(ready)); HIP(hipHostFree(host)); HIP(hipFree(device));
return timeout ? 124 : (wrong_before_sync || wrong_after_sync ? 1 : 0);
}Direct HSA/AQL control sources and build/run instructionsThese are the exact control sources used above.
#include <hip/hip_runtime.h>
#include <cstdint>
extern "C" __global__ void raw_copy(const uint32_t* src, uint32_t* dst, uint32_t words) {
for (uint32_t i = 0; i < words; ++i) dst[i] = src[i];
}
extern "C" __global__ void raw_relaxed(uint64_t* flag, uint64_t epoch) {
__hip_atomic_store(flag, epoch, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
extern "C" __global__ void raw_release(uint64_t* flag, uint64_t epoch) {
__hip_atomic_store(flag, epoch, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
}
#include <hip/hip_runtime.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <atomic>
#include <chrono>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <fcntl.h>
#include <fstream>
#include <iostream>
#include <stdexcept>
#include <string>
#include <thread>
#include <unistd.h>
#include <vector>
static void hip_check(hipError_t s) { if (s) throw std::runtime_error(hipGetErrorString(s)); }
static void hsa_check(hsa_status_t s) {
if (s) { const char* msg = nullptr; hsa_status_string(s, &msg); throw std::runtime_error(msg ? msg : "HSA error"); }
}
static uint32_t pattern(unsigned epoch, size_t i) {
return 0x80000001u ^ (epoch * 107u) ^ (static_cast<uint32_t>(i) * 13007u);
}
static hsa_status_t find_agent(hsa_agent_t agent, void* p) {
hsa_device_type_t type{};
hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
if (type != HSA_DEVICE_TYPE_GPU) return HSA_STATUS_SUCCESS;
char uuid[32]{};
hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_UUID), uuid);
if (std::string(uuid) == "GPU-541f8c51daca3a7e") *static_cast<hsa_agent_t*>(p) = agent;
return HSA_STATUS_SUCCESS;
}
static hsa_status_t find_kernarg(hsa_region_t region, void* p) {
hsa_region_segment_t segment{};
hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment);
if (segment == HSA_REGION_SEGMENT_GLOBAL) {
uint32_t flags = 0;
hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) *static_cast<hsa_region_t*>(p) = region;
}
return HSA_STATUS_SUCCESS;
}
struct Kernel {
uint64_t object = 0;
uint32_t private_size = 0, group_size = 0, arg_size = 0;
void* args = nullptr;
};
static Kernel load_kernel(hsa_executable_t executable, hsa_agent_t agent, hsa_region_t region,
const char* name) {
hsa_executable_symbol_t symbol{};
hsa_check(hsa_executable_get_symbol_by_name(executable, name, &agent, &symbol));
Kernel k;
hsa_check(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &k.object));
hsa_check(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &k.private_size));
hsa_check(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &k.group_size));
hsa_check(hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &k.arg_size));
hsa_check(hsa_memory_allocate(region, k.arg_size, &k.args));
std::memset(k.args, 0, k.arg_size);
return k;
}
static uint16_t header(unsigned type, unsigned acquire, unsigned release) {
return type | (1u << HSA_PACKET_HEADER_BARRIER) |
(acquire << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(release << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
}
static void dispatch(hsa_queue_t* queue, uint64_t index, const Kernel& k, uint16_t hdr,
hsa_signal_t completion) {
hsa_kernel_dispatch_packet_t packet{};
packet.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
packet.workgroup_size_x = packet.workgroup_size_y = packet.workgroup_size_z = 1;
packet.grid_size_x = packet.grid_size_y = packet.grid_size_z = 1;
packet.kernel_object = k.object;
packet.kernarg_address = k.args;
packet.private_segment_size = k.private_size;
packet.group_segment_size = k.group_size;
packet.completion_signal = completion;
auto* target = static_cast<hsa_kernel_dispatch_packet_t*>(queue->base_address) + (index & (queue->size - 1));
std::memcpy(target, &packet, sizeof(packet));
__atomic_store_n(&target->header, hdr, __ATOMIC_RELEASE);
}
int main(int argc, char** argv) {
try {
if (argc < 3) return 2;
const std::string mode = argv[2];
const unsigned iterations = argc > 3 ? std::stoul(argv[3]) : 10000;
if (mode != "agent-relaxed" && mode != "system-relaxed" && mode != "agent-release" &&
mode != "agent-barrier" && mode != "agent-nofence") return 2;
hip_check(hipSetDevice(0));
int count = 0;
hip_check(hipGetDeviceCount(&count));
char pci[32]{};
hip_check(hipDeviceGetPCIBusId(pci, sizeof(pci), 0));
if (count != 1 || std::string(pci) != "0000:43:00.0") throw std::runtime_error("Wrong device selection");
hsa_check(hsa_init());
hsa_agent_t agent{};
hsa_check(hsa_iterate_agents(find_agent, &agent));
if (!agent.handle) throw std::runtime_error("Target HSA agent unavailable");
hsa_region_t region{};
hsa_check(hsa_agent_iterate_regions(agent, find_kernarg, ®ion));
if (!region.handle) throw std::runtime_error("No kernarg region");
int fd = open(argv[1], O_RDONLY);
if (fd < 0) throw std::runtime_error("Cannot open HSACO");
hsa_code_object_reader_t reader{};
hsa_check(hsa_code_object_reader_create_from_file(fd, &reader));
hsa_executable_t executable{};
hsa_check(hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &executable));
hsa_check(hsa_executable_load_agent_code_object(executable, agent, reader, nullptr, nullptr));
hsa_check(hsa_executable_freeze(executable, nullptr));
Kernel copy = load_kernel(executable, agent, region, "raw_copy.kd");
Kernel notify = load_kernel(executable, agent, region, mode == "agent-release" ? "raw_release.kd" : "raw_relaxed.kd");
hsa_queue_t* queue = nullptr;
hsa_check(hsa_queue_create(agent, 64, HSA_QUEUE_TYPE_SINGLE,
[](hsa_status_t e, hsa_queue_t*, void*) { std::fprintf(stderr, "QUEUE_ERROR %d\n", int(e)); std::_Exit(2); },
nullptr, UINT32_MAX, UINT32_MAX, &queue));
hsa_signal_t complete{};
hsa_check(hsa_signal_create(1, 0, nullptr, &complete));
constexpr size_t words = 136, bytes = words * 4;
uint32_t *device = nullptr, *host = nullptr;
uint64_t* ready = nullptr;
hip_check(hipMalloc(reinterpret_cast<void**>(&device), bytes));
hip_check(hipHostMalloc(reinterpret_cast<void**>(&host), bytes, hipHostMallocMapped | hipHostMallocCoherent));
hip_check(hipExtMallocWithFlags(reinterpret_cast<void**>(&ready), 8, hipMallocSignalMemory));
__atomic_store_n(ready, uint64_t{0}, __ATOMIC_RELEASE);
struct CopyArgs { uint32_t* src; uint32_t* dst; uint32_t words; } ca{device, host, words};
struct NotifyArgs { uint64_t* flag; uint64_t value; } na{ready, 0};
const size_t copy_arg_bytes = offsetof(CopyArgs, words) + sizeof(ca.words);
if (copy.arg_size < copy_arg_bytes || notify.arg_size < sizeof(na))
throw std::runtime_error("Invalid kernarg layout: copy=" + std::to_string(copy.arg_size) +
" notify=" + std::to_string(notify.arg_size));
std::memcpy(copy.args, &ca, copy_arg_bytes);
std::atomic<unsigned> checked{0};
std::atomic<bool> stop{false}, timeout{false};
unsigned bad = 0, bad_after = 0;
std::cout << "IDENTITY pci=" << pci << " mode=" << mode << " hsa_queue=" << queue << std::endl;
std::ifstream maps("/proc/self/maps");
for (std::string line; std::getline(maps, line);)
if (line.find(".so") != std::string::npos) std::cout << "MAP " << line << '\n';
std::thread observer([&] {
std::vector<uint32_t> snapshot(words);
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
while (__atomic_load_n(ready, __ATOMIC_ACQUIRE) != epoch) {
if (stop.load()) return;
if (std::chrono::steady_clock::now() > deadline) { timeout = true; return; }
}
std::memcpy(snapshot.data(), host, bytes);
for (size_t i = 0; i < words; ++i) if (snapshot[i] != pattern(epoch, i)) { ++bad; break; }
checked.store(epoch, std::memory_order_release);
}
});
try {
std::vector<uint32_t> input(words);
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
for (size_t i = 0; i < words; ++i) input[i] = pattern(epoch, i);
hip_check(hipMemcpy(device, input.data(), bytes, hipMemcpyHostToDevice));
std::memset(host, 0xa5, bytes);
na.value = epoch;
std::memcpy(notify.args, &na, sizeof(na));
hsa_signal_store_screlease(complete, 1);
const bool barrier = mode == "agent-barrier" || mode == "agent-nofence";
const unsigned packets = barrier ? 3 : 2;
uint64_t start = hsa_queue_add_write_index_relaxed(queue, packets);
while (start + packets - hsa_queue_load_read_index_scacquire(queue) > queue->size) std::this_thread::yield();
dispatch(queue, start, copy, header(HSA_PACKET_TYPE_KERNEL_DISPATCH, 2, mode == "system-relaxed" ? 2 : 1), {});
if (barrier) {
hsa_barrier_and_packet_t p{};
auto* dst = static_cast<hsa_barrier_and_packet_t*>(queue->base_address) + ((start + 1) & (queue->size - 1));
std::memcpy(dst, &p, sizeof(p));
__atomic_store_n(&dst->header, header(HSA_PACKET_TYPE_BARRIER_AND, 0, mode == "agent-barrier" ? 2 : 0), __ATOMIC_RELEASE);
}
dispatch(queue, start + packets - 1, notify, header(HSA_PACKET_TYPE_KERNEL_DISPATCH, 2, 2), complete);
hsa_signal_store_screlease(queue->doorbell_signal, start + packets - 1);
auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
while (checked.load(std::memory_order_acquire) != epoch && !timeout.load()) {
if (std::chrono::steady_clock::now() > deadline) timeout = true;
std::this_thread::yield();
}
if (timeout.load()) break;
while (hsa_signal_load_scacquire(complete) != 0) {
if (std::chrono::steady_clock::now() > deadline) { timeout = true; break; }
}
if (timeout.load()) break;
for (size_t i = 0; i < words; ++i) if (host[i] != pattern(epoch, i)) { ++bad_after; break; }
}
} catch (...) { stop = true; observer.join(); throw; }
stop = true;
observer.join();
std::printf("RESULT iterations=%u checked=%u bad=%u bad_after_completion=%u timeout=%d\n", iterations, checked.load(), bad, bad_after, int(timeout.load()));
if (timeout.load()) return 124;
hsa_check(hsa_queue_destroy(queue));
hsa_check(hsa_signal_destroy(complete));
hsa_check(hsa_memory_free(copy.args));
hsa_check(hsa_memory_free(notify.args));
hsa_check(hsa_executable_destroy(executable));
hsa_check(hsa_code_object_reader_destroy(reader));
close(fd);
hip_check(hipFree(ready)); hip_check(hipHostFree(host)); hip_check(hipFree(device));
hsa_check(hsa_shut_down());
return bad || bad_after ? 1 : 0;
} catch (const std::exception& e) { std::cerr << "ERROR " << e.what() << '\n'; return 2; }
}The following are the tested wheel-SDK build steps, with output paths made relative. The shim supplies an unversioned HIP link only inside the experiment directory; it does not modify the SDK. Use a matched compiler/device-library/ROCr installation. The bundler step is needed when sdk=/opt/venv/lib/python3.13/site-packages/_rocm_sdk_core
mkdir -p sdk-shim/lib
ln -s "$sdk/include" sdk-shim/include
ln -s "$sdk/lib/libamdhip64.so.7" sdk-shim/lib/libamdhip64.so
hipcc --genco -O2 --offload-arch=gfx1201 \
--rocm-device-lib-path="$sdk/lib/llvm/amdgcn/bitcode" \
--hip-path="$PWD/sdk-shim" raw_kernels.cpp -o raw-kernels.bundle
"$sdk/lib/llvm/bin/clang-offload-bundler" --unbundle --type=o \
--targets=hipv4-amdgcn-amd-amdhsa--gfx1201 \
--input=raw-kernels.bundle --output=raw-kernels.hsaco
g++ -O2 -std=c++17 -D__HIP_PLATFORM_AMD__ -I"$sdk/include" \
raw_aql.cpp -L"$sdk/lib" -l:libamdhip64.so.7 -l:libhsa-runtime64.so.1 \
-pthread -o raw-aql
for mode in agent-relaxed system-relaxed agent-release agent-barrier agent-nofence; do
HIP_VISIBLE_DEVICES=GPU-541f8c51daca3a7e timeout 60s ./raw-aql raw-kernels.hsaco "$mode" 10000
doneThe two negative controls return 1 for data errors; the three publication controls return 0. Exit 124/timeouts are failures. Recorded source, binary, and runtime SHA256 identitiesThese distinguish the two separately rebuilt HIP DSOs, the fixed ROCr control, and the current foundation controls. The candidate was built at source HEAD 538d0d4 with diff SHA256 2d05884e50868e3883924a4e44808517c6de4bd0a5ebb98d33476f59678dacf4; that diff was subsequently committed verbatim as e736720 (verified against the commit delta). Current PR head 1e6b4da adds only the test SPDX header. The original evidence ZIP is preserved locally with SHA256 dba967fe06d6949c2b4e52ccfebb08abb059f2b8da8549adea7bd4e5a519d7b7; the reproducible source and result tables are included in this comment so access to my local archive is not required. |
Add a graph-free D2H-to-WRITE regression for 32/64-bit notifications. Check poisoned host data on a CPU observer before stream synchronization, without a GPU WAIT, reset WRITE, result copy, or application GPU kernel. Bound notification waits and join the observer during exceptional teardown. Register eager publication separately from batch/captured modes so PAL retains its supported coverage. Gate batch tests by backend selection rather than treating single-WAIT capability as batch support. Require an explicit ROCr selection for ambiguous automatic backend choices. Replace the Windows-wide layout-test skip with the same backend gate and retain the classic-executor guard, allowing Windows ROCr segmented graphs.
|
@AlexXAmd Follow-up in commit
The same new test binary and matching ROCr were used for the baseline/candidate comparison on gfx1201:
Post-sync data is correct in both minimal-case libraries; no run timed out. All actual test-process HIP/ROCr mappings and SHA256 identities were recorded. These are correctness runs on shared hardware, with no new performance claim. Ten CPU selection checks, including the simulated Windows branch, and Windows YAML registration pass. PAL/auto/classic skip-only paths were exercised before HIP initialization; these checks are not Windows/PAL runtime validation. Actual Windows ROCr/PAL results remain pending CI, as do the API-contract and cross-ASIC/cost questions from the earlier root-cause reply. timeout 90s ./hipStreamMemOpPublication \
'Unit_hipStreamMemOp_D2HWritePublication*' --reporter compact
timeout 90s ./hipStreamMemOpPublication --reporter compact
GPU_ENABLE_PAL=0 timeout 90s ./hipGraphExecBatchMemOpNodeSetParams \
'Unit_hipGraphExecBatchMemOpNodeSetParams_LayoutChange*' --reporter compact |
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
|
@AlexXAmd I traced the caller/helper ownership to an existing CLR ordering boundary and a specific change that removed it. This follow-up adds evidence; the published code remains 1. The ordering was provided by CLR before calling the relaxed helper The device-library WRITE helper was introduced as a relaxed SYSTEM-scope store in 6918714d635c. When CLR switched to that helper, cbb8d82bdb2b retained this operation before the WRITE blit: // Ensure memory ordering preceding the write
dispatchBarrierPacket(kBarrierPacketReleaseHeader);#2304, commit 2. Restoring that one caller-side operation fixes the isolated counterexample On the fixed modern baseline
There were no timeouts. Actual process mappings show identical non-HIP library content hashes, and the same test binary was used. The regression source is already in this PR. The exact diagnostic diff and identities are below so this result does not depend on access to my workspace. The restored AQL trace is: This is a controlled restoration on a modern tree, not a full build/bisect of the 2025 parent, a complete replacement implementation, or new Graph validation of the published candidate. 3. The missing guarantee is before the observable store The original barrier releases completed producers before the notification kernel starts. The replacement upgrades the notification dispatch's scopes, but its SYSTEM release occurs after that kernel's active phase; the CPU can already have observed the relaxed flag store. The HSA packet definitions and HSA specification §2.9.2/§3.3.8 distinguish this completion release from acquire's guarantee for subsequent device loads. The two sequences therefore do not provide the same pre-notification publication boundary. The earlier direct-AQL controls and their source honor explicitly requested SYSTEM release and packet-completion observations. Weak ordinary-payload cases are diagnostics, not HSA conformance failures. Those controls support the timing explanation; the caller/helper history and controlled restoration above provide the additional ownership evidence. This does not require a firmware/driver violation or a compiler dropping a requested fence. 4. Batch/capture is a separate path, and cost still needs review The batch submission introduction did not include the ordinary WRITE's pre-release. #12692 later added dispatch scopes only. Native captured WRITE is represented as a batch node, so that path also needs the publication boundary. I am not attributing all batch/graph behavior to #2304 or proposing a blind whole-PR revert. This provides a concrete reason to review the correction in CLR: restore the caller-owned ordering that #2304 intended to preserve, and give the independent batch path the corresponding guarantee. Formal HIP wording and per-item batch semantics still deserve owner confirmation. The smallest production implementation, safe fence reuse, all-ASIC cost, and remaining platform CI are still open; the existing packet-layout/synchronization-plan fix remains in the published code. Exact diagnostic diff, command and runtime identitiesThe diagnostic source is diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
index df24a95a3e..a178d9c9bc 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -4357,9 +4357,8 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
}
}
} else if (type == ROCCLR_COMMAND_STREAM_WRITE_VALUE) {
- // Even though the blit kernel uses system scope atomic, we still need to add system scope on
- // the AQLPacket because atomics on kernels can bypass L2 cache on some hardware.
- addSystemScope();
+ // Diagnostic: restore the caller-owned publication boundary removed by #2304.
+ dispatchBarrierPacket(kBarrierPacketReleaseHeader);
bool result;
switch (flags) {The existing regression executable was run unchanged against the two private library prefixes: timeout 90s ./hipStreamMemOpPublication \
'Unit_hipStreamMemOp_D2HWritePublication*' --reporter compact -sBoth libraries use the same compiler/build options and dependency paths, with only source/build/install prefixes differing. They are local rebuilds reporting HIP 7.17.0-0000000 / 71700000, with git version stamping disabled; they are not release-package identities. Target: R9700/gfx1201, PCI SHA256 identities:
Locked foundation image: The baseline's two cases each failed the stale-snapshot assertion (exit 42); the restoration passed all 8 assertions in 2 cases (exit 0). Both completed every iteration, with zero post-sync errors and no timeout. |
| // HIP has no public batch capability query. Follow explicit backend selection | ||
| // and platform defaults; pin auto-selection to GPU_ENABLE_PAL=0 to run these | ||
| // ROCr-only tests rather than assuming that an automatic fallback supports them. | ||
| inline bool IsRocrMemOpBackendSelected() { |
There was a problem hiding this comment.
For API design, if HIP wants to expose certain feature to application, add an info/property into HIP API so the HIP application can query about the feature. HIP application or HIP test should not decide whether a certain feature is supported by a specific implementation.
| unsigned badAfterSync = 0; | ||
| for (unsigned epoch = 1; epoch <= iterations && !observer.timedOut.load(); ++epoch) { | ||
| for (unsigned i = 0; i < kCopyWords; ++i) input[i] = CopyPattern(epoch, i); | ||
| HIP_CHECK(hipMemcpyAsync(deviceInput.ptr(), input.data(), sizeof(input), hipMemcpyHostToDevice, |
There was a problem hiding this comment.
See my offline communication. This is not guaranteed behavior by HIP API document.
| } else { | ||
| HIP_CHECK(hipStreamWriteValue64(stream, ready, epoch, hipStreamWriteValueDefault)); | ||
| } | ||
| while (observer.checked.load(std::memory_order_acquire) != epoch && !observer.timedOut.load()) { |
There was a problem hiding this comment.
See my offline communication. This is not guaranteed behavior by HIP API document.
|
@AlexXAmd Following your suggested kernel-based test, I completed the CUDA comparison on an RTX 5090 (the actual tested model), with nvcc 13.0.88, libcudart 13.0.96, driver 615.71.09, and native The sequence is: The producer has 256 blocks × 256 threads. Payload and notification use separate mapped pinned allocations and their device aliases. There is no memcpy, application-side GPU fence, GPU WAIT, or reset WRITE. Payload is poisoned each round; ready uses increasing epochs. All six processes exited 0, completed every iteration, and had no skips or timeouts:
That is 5000 completed CUDA rounds with zero stale snapshots. For the corresponding 1 MiB mapped Eager kernel cases, both widths total 4000 rounds:
These are separate platform runs using each platform's mapped-memory API, not one common memory pool or a newer-develop/full-stack bisect. The HIP libraries are local rebuilds reporting HIP 7.17.0-0000000 / 71700000, not release-package identities. The CUDA test is Eager only; it does not validate CUDA Graph or D2H publication. This supplies the requested CUDA behavior comparison and supports discussing publication as a compatibility target. It does not itself settle the HIP API contract or prove a firmware/driver violation. Could we confirm whether the intended default WRITE contract is:
This is proposed wording for owner review. If this should be treated as a new guarantee, I can proceed on that basis with the corresponding documentation and tests. I also prepared a local one-runtime-file alternative: four release fences in CLR's WRITE blit variants and batch WRITE32/64 branches, before the relaxed stores. It retains existing AQL scopes/dependencies and adds no packets, so the Graph executor stays at baseline. Its exact diff is below. This alternative has not replaced the public PR head ( The folded sections include complete CUDA/HIP test source, exact commands, the minimal alternative diff, all six raw CUDA logs, and the original CUDA manifest. Source/binary hashes were independently recomputed against the supplied archive. CUDA DSO hashes are the test host's recorded values, consistent across cases; the DSO files themselves were not included in that archive. Reproduce: full CUDA source and commandsThe source below is the exact tested CUDA file. On a different GPU, adapt the target-name guard and compiler architecture while retaining the protocol. Link stub directories in the recorded command are link-only; actual runtime mappings below use the real driver. # Recorded CUDA target; select the corresponding UUID when reproducing elsewhere.
export CUDA_VISIBLE_DEVICES=GPU-17390cc0-bf3c-53e3-170c-21eb949de43b
export CUDA_LAUNCH_BLOCKING=0
# Exact recorded CUDA compilation:
/usr/local/cuda-13.0/bin/nvcc -std=c++17 -O2 -lineinfo -arch=sm_120 --cudart shared -Xcompiler -pthread /work/cuda_memop_publication.cu -o /work/results-20261027T050314.144422Z/cuda_memop_publication -L /usr/local/cuda-13.0/lib64/stubs -L /usr/local/cuda-13.0/targets/x86_64-linux/lib/stubs -lcuda
# Matching invocations (the original absolute output paths are in the manifest):
timeout 180s ./cuda_memop_publication 32 250 4096
timeout 180s ./cuda_memop_publication 32 2000 1048576
timeout 180s ./cuda_memop_publication 32 250 8388608
timeout 180s ./cuda_memop_publication 64 250 4096
timeout 180s ./cuda_memop_publication 64 2000 1048576
timeout 180s ./cuda_memop_publication 64 250 8388608
# HIP standalone build and matching Eager calls on the recorded gfx1201 target:
export HIP_VISIBLE_DEVICES=GPU-541f8c51daca3a7e
export GPU_STREAMOPS_CP_WAIT=0
hipcc -x hip -O2 -std=c++17 -pthread --offload-arch=gfx1201 -DHIP_MEMOP_PUBLICATION_STANDALONE=1 hipStreamMemOpKernelPublication.cc -o hip-kernel-publication
timeout 90s ./hip-kernel-publication 32 2000 1048576 eager mapped
timeout 90s ./hip-kernel-publication 64 2000 1048576 eager mapped// SPDX-License-Identifier: MIT
// Kernel -> default stream WRITE -> CPU snapshot, without an application fence.
#include <cuda.h>
#include <cuda_runtime.h>
#include <atomic>
#include <chrono>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <fstream>
#include <set>
#include <stdexcept>
#include <string>
#include <thread>
#include <vector>
#if CUDA_VERSION < 12000
#error "Use CUDA Toolkit 12 or newer for this handoff."
#endif
static void runtime_check(cudaError_t rc, const char* call) {
if (rc == cudaSuccess) return;
std::fprintf(stderr, "CUDA_RUNTIME_ERROR call=%s code=%d text=%s\n",
call, int(rc), cudaGetErrorString(rc));
std::fflush(nullptr);
std::_Exit(2);
}
static void driver_check(CUresult rc, const char* call, bool allow_skip = false) {
if (rc == CUDA_SUCCESS) return;
const char *name = nullptr, *text = nullptr;
cuGetErrorName(rc, &name);
cuGetErrorString(rc, &text);
const int exit_code = allow_skip && rc == CUDA_ERROR_NOT_SUPPORTED ? 77 : 2;
std::fprintf(stderr, "%s call=%s code=%d name=%s text=%s\n",
exit_code == 77 ? "SKIP" : "CUDA_DRIVER_ERROR", call, int(rc),
name ? name : "unknown", text ? text : "unknown");
std::fflush(nullptr);
std::_Exit(exit_code);
}
#define RT(call) runtime_check((call), #call)
#define DR(call) driver_check((call), #call)
static __host__ __device__ uint32_t pattern(uint32_t epoch, size_t i) {
return 0x80000001u ^ (epoch * 107u) ^ (uint32_t(i) * 13007u);
}
__global__ void writeDataShader(uint32_t* output, size_t words, uint32_t epoch) {
for (size_t i = size_t(blockIdx.x) * blockDim.x + threadIdx.x; i < words;
i += size_t(gridDim.x) * blockDim.x)
output[i] = pattern(epoch, i);
}
static void dump_dsos() {
std::ifstream maps("/proc/self/maps");
std::set<std::string> paths;
for (std::string line; std::getline(maps, line);) {
const auto pos = line.find('/');
if (pos != std::string::npos && line.find(".so", pos) != std::string::npos)
paths.insert(line.substr(pos));
}
for (const auto& path : paths) std::printf("DSO %s\n", path.c_str());
std::fflush(stdout);
}
template <class T>
static int run(unsigned iterations, size_t bytes) {
using Clock = std::chrono::steady_clock;
const size_t words = bytes / sizeof(uint32_t);
uint32_t *host = nullptr, *device = nullptr;
T *ready_host = nullptr, *ready_device = nullptr;
RT(cudaHostAlloc(reinterpret_cast<void**>(&host), bytes, cudaHostAllocMapped));
RT(cudaHostGetDevicePointer(reinterpret_cast<void**>(&device), host, 0));
// Separate allocation keeps notification and payload out of the same cache line.
RT(cudaHostAlloc(reinterpret_cast<void**>(&ready_host), 4096, cudaHostAllocMapped));
RT(cudaHostGetDevicePointer(reinterpret_cast<void**>(&ready_device), ready_host, 0));
if (reinterpret_cast<uintptr_t>(ready_host) % alignof(T)) {
std::fprintf(stderr, "SETUP_ERROR unaligned notification\n");
return 2;
}
__atomic_store_n(ready_host, T{0}, __ATOMIC_RELEASE);
cudaStream_t stream = nullptr;
RT(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
dump_dsos();
std::vector<uint32_t> snapshot(words);
std::atomic<unsigned> checked{0};
std::atomic<bool> started{false}, stop{false}, timed_out{false};
unsigned bad = 0, bad_after_sync = 0;
uint64_t bad_words = 0;
std::thread observer([&] {
started.store(true, std::memory_order_release);
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
const auto deadline = Clock::now() + std::chrono::seconds(10);
while (__atomic_load_n(ready_host, __ATOMIC_ACQUIRE) != T(epoch)) {
if (stop.load(std::memory_order_relaxed)) return;
if (Clock::now() > deadline) { timed_out.store(true); return; }
}
std::memcpy(snapshot.data(), host, bytes);
size_t wrong = 0, first = 0;
for (size_t i = 0; i < words; ++i) {
if (snapshot[i] != pattern(epoch, i)) {
if (wrong == 0) first = i;
++wrong;
}
}
if (wrong) {
++bad;
bad_words += wrong;
if (bad <= 5)
std::printf("MISMATCH epoch=%u wrong_words=%zu first_index=%zu "
"actual=%08x expected=%08x\n", epoch, wrong, first,
snapshot[first], pattern(epoch, first));
}
checked.store(epoch, std::memory_order_release);
}
});
while (!started.load(std::memory_order_acquire)) std::this_thread::yield();
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
std::memset(host, 0xa5, bytes);
writeDataShader<<<256, 256, 0, stream>>>(device, words, epoch);
RT(cudaGetLastError());
CUresult rc;
if constexpr (sizeof(T) == 4)
rc = cuStreamWriteValue32(reinterpret_cast<CUstream>(stream),
reinterpret_cast<CUdeviceptr>(ready_device),
cuuint32_t(epoch), 0);
else
rc = cuStreamWriteValue64(reinterpret_cast<CUstream>(stream),
reinterpret_cast<CUdeviceptr>(ready_device),
cuuint64_t(epoch), 0);
driver_check(rc, "cuStreamWriteValue(default)", true);
const auto deadline = Clock::now() + std::chrono::seconds(10);
while (checked.load(std::memory_order_acquire) != epoch && !timed_out.load()) {
if (Clock::now() > deadline) timed_out.store(true);
std::this_thread::yield();
}
if (timed_out.load()) break;
// CPU snapshot and validation already finished. This cannot repair that snapshot.
RT(cudaStreamSynchronize(stream));
for (size_t i = 0; i < words; ++i) {
if (host[i] != pattern(epoch, i)) { ++bad_after_sync; break; }
}
}
stop.store(true);
observer.join();
const unsigned completed = checked.load();
const int timeout = timed_out.load() || completed != iterations;
std::printf("RESULT {\"width\":%zu,\"bytes\":%zu,\"iterations\":%u,"
"\"checked\":%u,\"bad_before_sync\":%u,\"bad_words\":%llu,"
"\"bad_after_sync\":%u,\"timeout\":%d}\n",
sizeof(T) * 8, bytes, iterations, completed, bad,
static_cast<unsigned long long>(bad_words), bad_after_sync, timeout);
dump_dsos();
if (timeout) { std::fflush(nullptr); std::_Exit(124); }
RT(cudaStreamDestroy(stream));
RT(cudaFreeHost(ready_host));
RT(cudaFreeHost(host));
return bad || bad_after_sync ? 1 : 0;
}
int main(int argc, char** argv) {
std::setvbuf(stdout, nullptr, _IOLBF, 0);
try {
if (argc != 4) throw std::runtime_error("usage: binary WIDTH ITERATIONS BYTES");
const unsigned long width = std::stoul(argv[1]);
const unsigned long rounds = std::stoul(argv[2]);
const size_t bytes = std::stoull(argv[3]);
if ((width != 32 && width != 64) || !rounds || rounds > 1000000 ||
bytes < 4 || bytes % 4 || bytes > (size_t{1} << 30))
throw std::runtime_error("invalid width, iterations or byte size");
const char* blocking = std::getenv("CUDA_LAUNCH_BLOCKING");
if (blocking && *blocking && std::strcmp(blocking, "0"))
throw std::runtime_error("CUDA_LAUNCH_BLOCKING must be unset or 0");
DR(cuInit(0));
int count = 0;
DR(cuDeviceGetCount(&count));
if (count != 1) throw std::runtime_error("expose exactly one RTX 5090 with CUDA_VISIBLE_DEVICES");
RT(cudaSetDeviceFlags(cudaDeviceMapHost));
RT(cudaSetDevice(0));
cudaDeviceProp prop{};
RT(cudaGetDeviceProperties(&prop, 0));
if (!std::strstr(prop.name, "5090") || !prop.canMapHostMemory)
throw std::runtime_error("requires a 5090 with mapped host memory support");
CUdevice dev;
DR(cuDeviceGet(&dev, 0));
int can64 = 0, driver = 0, runtime = 0;
DR(cuDeviceGetAttribute(&can64, CU_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS, dev));
DR(cuDriverGetVersion(&driver));
RT(cudaRuntimeGetVersion(&runtime));
char pci[32]{};
RT(cudaDeviceGetPCIBusId(pci, sizeof(pci), 0));
std::printf("IDENTITY name=%s pci=%s sm=%d%d header=%d runtime=%d driver=%d "
"can64=%d mapped=%d host_native_atomics=%d flags=0 blocks=256 threads=256\n",
prop.name, pci, prop.major, prop.minor, CUDA_VERSION, runtime, driver,
can64, prop.canMapHostMemory, prop.hostNativeAtomicSupported);
dump_dsos();
if (width == 64 && !can64) {
std::fprintf(stderr, "SKIP 64-bit stream memops not supported\n");
return 77;
}
return width == 32 ? run<uint32_t>(unsigned(rounds), bytes)
: run<uint64_t>(unsigned(rounds), bytes);
} catch (const std::exception& error) {
std::fprintf(stderr, "SETUP_ERROR %s\n", error.what());
return 2;
}
}Full HIP kernel counterpart and runtime alternative (local, not pushed)HIP source: local candidate // Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
// SPDX-License-Identifier: MIT
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
// Opt-in contract investigation. API-owner confirmation of publication semantics
// and public batch/capture capability discovery precede normal CI registration.
// No OS or backend environment variable is interpreted as a feature query.
#include <hip/hip_runtime.h>
#include <atomic>
#include <chrono>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <fstream>
#include <stdexcept>
#include <string>
#include <thread>
#include <vector>
#ifndef HIP_MEMOP_PUBLICATION_STANDALONE
#include <hip_test_common.hh>
#endif
namespace {
void Check(hipError_t error, const char* call) {
if (error != hipSuccess) {
throw std::runtime_error(std::string(call) + ": " + hipGetErrorString(error));
}
}
#define MEMOP_CHECK(call) Check((call), #call)
enum class Submission { Eager, Batch, Capture, CaptureBatch, ExplicitGraph };
__host__ __device__ uint32_t Pattern(unsigned epoch, size_t i) {
return 0x80000001u ^ (epoch * 107u) ^ (static_cast<uint32_t>(i) * 13007u);
}
__global__ void WriteDataShader(uint32_t* output, size_t words, unsigned epoch) {
for (size_t i = size_t{blockIdx.x} * blockDim.x + threadIdx.x; i < words;
i += size_t{gridDim.x} * blockDim.x) {
output[i] = Pattern(epoch, i);
}
}
struct Buffer {
void* host = nullptr;
void* device = nullptr;
bool signal;
Buffer(size_t bytes, bool signalMemory) : signal(signalMemory) {
if (signal) {
MEMOP_CHECK(hipExtMallocWithFlags(&host, bytes, hipMallocSignalMemory));
device = host;
} else {
MEMOP_CHECK(hipHostMalloc(&host, bytes, hipHostMallocMapped | hipHostMallocCoherent));
const auto error = hipHostGetDevicePointer(&device, host, 0);
if (error != hipSuccess) {
static_cast<void>(hipHostFree(host));
host = nullptr;
Check(error, "hipHostGetDevicePointer");
}
}
}
~Buffer() {
if (host) static_cast<void>(signal ? hipFree(host) : hipHostFree(host));
}
Buffer(const Buffer&) = delete;
Buffer& operator=(const Buffer&) = delete;
};
struct Stream {
hipStream_t handle = nullptr;
Stream() { MEMOP_CHECK(hipStreamCreateWithFlags(&handle, hipStreamNonBlocking)); }
~Stream() { static_cast<void>(hipStreamDestroy(handle)); }
};
struct Graph {
hipGraph_t graph = nullptr;
hipGraphExec_t exec = nullptr;
~Graph() {
if (exec) static_cast<void>(hipGraphExecDestroy(exec));
if (graph) static_cast<void>(hipGraphDestroy(graph));
}
};
template <typename T> struct Observer {
std::atomic<unsigned> checked{0};
std::atomic<bool> started{false}, stop{false}, timedOut{false};
unsigned bad = 0;
uint64_t badWords = 0;
std::vector<uint32_t> snapshot;
std::thread worker;
Observer(T* ready, const uint32_t* host, size_t words, unsigned iterations)
: snapshot(words), worker([this, ready, host, words, iterations] {
started.store(true, std::memory_order_release);
for (unsigned epoch = 1; epoch <= iterations; ++epoch) {
const auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
while (__atomic_load_n(ready, __ATOMIC_ACQUIRE) != static_cast<T>(epoch)) {
if (stop.load(std::memory_order_relaxed)) return;
if (std::chrono::steady_clock::now() > deadline) {
timedOut = true;
return;
}
}
std::memcpy(snapshot.data(), host, words * sizeof(uint32_t));
size_t wrong = 0;
for (size_t i = 0; i < words; ++i) wrong += snapshot[i] != Pattern(epoch, i);
if (wrong) {
++bad;
badWords += wrong;
if (bad <= 5) std::printf("MISMATCH epoch=%u wrong_words=%zu\n", epoch, wrong);
}
checked.store(epoch, std::memory_order_release);
}
}) {}
void Join() {
stop = true;
if (worker.joinable()) worker.join();
}
~Observer() { Join(); }
};
struct Result {
unsigned checked, badBeforeSync, badAfterSync;
uint64_t badWords;
bool timeout;
};
template <typename T>
Result Run(Submission submission, unsigned iterations, size_t bytes, bool signalReady) {
const size_t words = bytes / sizeof(uint32_t);
Buffer payload(bytes, false), readyMemory(signalReady ? sizeof(uint64_t) : 4096, signalReady);
auto* host = static_cast<uint32_t*>(payload.host);
auto* device = static_cast<uint32_t*>(payload.device);
auto* ready = static_cast<T*>(readyMemory.host);
__atomic_store_n(ready, T{0}, __ATOMIC_RELEASE);
Stream stream;
Graph graph;
unsigned epoch = 1;
size_t kernelWords = words;
void* kernelArgs[] = {&device, &kernelWords, &epoch};
hipKernelNodeParams kernelParams{};
kernelParams.func = reinterpret_cast<void*>(WriteDataShader);
kernelParams.gridDim = dim3(256);
kernelParams.blockDim = dim3(256);
kernelParams.kernelParams = kernelArgs;
hipStreamBatchMemOpParams write{};
write.operation = sizeof(T) == 4 ? hipStreamMemOpWriteValue32 : hipStreamMemOpWriteValue64;
write.writeValue.address = readyMemory.device;
write.writeValue.value64 = 1;
hipBatchMemOpNodeParams batchParams{};
MEMOP_CHECK(hipCtxGetCurrent(&batchParams.ctx));
batchParams.count = 1;
batchParams.paramArray = &write;
hipGraphNode_t producer = nullptr, notification = nullptr;
const auto enqueueWrite = [&] {
if constexpr (sizeof(T) == 4) {
MEMOP_CHECK(hipStreamWriteValue32(stream.handle, readyMemory.device, epoch,
hipStreamWriteValueDefault));
} else {
MEMOP_CHECK(hipStreamWriteValue64(stream.handle, readyMemory.device, epoch,
hipStreamWriteValueDefault));
}
};
if (submission == Submission::Capture || submission == Submission::CaptureBatch) {
MEMOP_CHECK(hipStreamBeginCapture(stream.handle, hipStreamCaptureModeThreadLocal));
hipLaunchKernelGGL(WriteDataShader, dim3(256), dim3(256), 0, stream.handle, device, words,
epoch);
MEMOP_CHECK(hipGetLastError());
if (submission == Submission::CaptureBatch) {
MEMOP_CHECK(hipStreamBatchMemOp(stream.handle, 1, &write, 0));
} else {
enqueueWrite();
}
MEMOP_CHECK(hipStreamEndCapture(stream.handle, &graph.graph));
size_t count = 0;
MEMOP_CHECK(hipGraphGetNodes(graph.graph, nullptr, &count));
std::vector<hipGraphNode_t> nodes(count);
MEMOP_CHECK(hipGraphGetNodes(graph.graph, nodes.data(), &count));
for (const auto node : nodes) {
hipGraphNodeType type;
MEMOP_CHECK(hipGraphNodeGetType(node, &type));
if (type == hipGraphNodeTypeKernel && !producer) producer = node;
if (type == hipGraphNodeTypeBatchMemOp && !notification) notification = node;
}
if (count != 2 || !producer || !notification) {
throw std::runtime_error("capture did not record one producer and one notification");
}
} else if (submission == Submission::ExplicitGraph) {
MEMOP_CHECK(hipGraphCreate(&graph.graph, 0));
MEMOP_CHECK(hipGraphAddKernelNode(&producer, graph.graph, nullptr, 0, &kernelParams));
MEMOP_CHECK(hipGraphAddBatchMemOpNode(¬ification, graph.graph, &producer, 1, &batchParams));
}
if (graph.graph) MEMOP_CHECK(hipGraphInstantiate(&graph.exec, graph.graph, nullptr, nullptr, 0));
Observer<T> observer(ready, host, words, iterations);
while (!observer.started.load(std::memory_order_acquire)) std::this_thread::yield();
unsigned badAfterSync = 0;
for (epoch = 1; epoch <= iterations; ++epoch) {
std::memset(host, 0xa5, bytes);
write.writeValue.value64 = epoch;
if (graph.exec) {
MEMOP_CHECK(hipGraphExecKernelNodeSetParams(graph.exec, producer, &kernelParams));
MEMOP_CHECK(hipGraphExecBatchMemOpNodeSetParams(graph.exec, notification, &batchParams));
MEMOP_CHECK(hipGraphLaunch(graph.exec, stream.handle));
} else {
hipLaunchKernelGGL(WriteDataShader, dim3(256), dim3(256), 0, stream.handle, device, words,
epoch);
MEMOP_CHECK(hipGetLastError());
if (submission == Submission::Batch) {
MEMOP_CHECK(hipStreamBatchMemOp(stream.handle, 1, &write, 0));
} else {
enqueueWrite();
}
}
const auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(10);
while (observer.checked.load(std::memory_order_acquire) != epoch && !observer.timedOut.load()) {
if (std::chrono::steady_clock::now() > deadline) observer.timedOut = true;
std::this_thread::yield();
}
if (observer.timedOut.load()) break;
MEMOP_CHECK(hipStreamSynchronize(stream.handle)); // Snapshot is already checked.
for (size_t i = 0; i < words; ++i) {
if (host[i] != Pattern(epoch, i)) {
++badAfterSync;
break;
}
}
}
observer.Join();
const auto checked = observer.checked.load();
return {checked, observer.bad, badAfterSync, observer.badWords,
observer.timedOut.load() || checked != iterations};
}
} // namespace
#ifdef HIP_MEMOP_PUBLICATION_STANDALONE
int main(int argc, char** argv) {
std::setvbuf(stdout, nullptr, _IOLBF, 0);
try {
if (argc != 6) throw std::runtime_error("usage: WIDTH ITERATIONS BYTES MODE mapped|signal");
const auto width = std::stoul(argv[1]);
const auto iterations = std::stoul(argv[2]);
const auto bytes = std::stoull(argv[3]);
const std::string mode = argv[4], ready = argv[5];
if ((width != 32 && width != 64) || !iterations || iterations > 1000000 || bytes < 4 ||
bytes % 4 || bytes > (size_t{1} << 30) || (ready != "mapped" && ready != "signal")) {
throw std::runtime_error("invalid arguments");
}
Submission submission;
if (mode == "eager")
submission = Submission::Eager;
else if (mode == "batch")
submission = Submission::Batch;
else if (mode == "capture")
submission = Submission::Capture;
else if (mode == "capture-batch")
submission = Submission::CaptureBatch;
else if (mode == "graph")
submission = Submission::ExplicitGraph;
else
throw std::runtime_error("unknown submission mode");
const char* blocking = std::getenv("HIP_LAUNCH_BLOCKING");
if (blocking && blocking[0] && std::strcmp(blocking, "0")) {
throw std::runtime_error("HIP_LAUNCH_BLOCKING must be unset or 0");
}
int count = 0, version = 0;
MEMOP_CHECK(hipGetDeviceCount(&count));
if (count != 1) throw std::runtime_error("requires exactly one visible target GPU");
MEMOP_CHECK(hipSetDevice(0));
MEMOP_CHECK(hipRuntimeGetVersion(&version));
char pci[32]{};
MEMOP_CHECK(hipDeviceGetPCIBusId(pci, sizeof(pci), 0));
if (std::strcmp(pci, "0000:43:00.0")) throw std::runtime_error("unexpected target PCI");
hipDeviceProp_t prop{};
MEMOP_CHECK(hipGetDeviceProperties(&prop, 0));
std::printf("IDENTITY hip=%d pci=%s arch=%s width=%lu bytes=%llu mode=%s ready=%s\n", version,
pci, prop.gcnArchName, width, bytes, mode.c_str(), ready.c_str());
const auto result = width == 32
? Run<uint32_t>(submission, iterations, bytes, ready == "signal")
: Run<uint64_t>(submission, iterations, bytes, ready == "signal");
std::printf(
"RESULT {\"width\":%lu,\"bytes\":%llu,\"iterations\":%lu,\"checked\":%u,"
"\"bad_before_sync\":%u,\"bad_after_sync\":%u,\"bad_words\":%llu,"
"\"timeout\":%d}\n",
width, bytes, iterations, result.checked, result.badBeforeSync, result.badAfterSync,
static_cast<unsigned long long>(result.badWords), int(result.timeout));
if (result.timeout) return 124;
return result.badBeforeSync || result.badAfterSync ? 1 : 0;
} catch (const std::exception& error) {
std::fprintf(stderr, "ERROR %s\n", error.what());
return 2;
}
}
#else
// Hidden until the default publication contract is agreed. A matching explicit
// filter opts in; this does not claim a generic batch feature-discovery solution.
TEMPLATE_TEST_CASE("Unit_hipStreamMemOp_KernelPublication", "[.memop-publication-contract]",
uint32_t, uint64_t) {
const size_t bytes = GENERATE(size_t{4096}, size_t{1048576}, size_t{8388608});
const unsigned iterations = bytes == 1048576 ? 2000 : 250;
const auto result = Run<TestType>(Submission::Eager, iterations, bytes, false);
INFO("bytes=" << bytes << ", checked=" << result.checked << ", bad=" << result.badBeforeSync);
CHECK_FALSE(result.timeout);
CHECK(result.checked == iterations);
CHECK(result.badBeforeSync == 0);
CHECK(result.badAfterSync == 0);
}
TEMPLATE_TEST_CASE("Unit_hipStreamMemOp_KernelPublication_BatchGraph", "[.memop-batch-contract]",
uint32_t, uint64_t) {
const auto mode = GENERATE(Submission::Batch, Submission::Capture, Submission::CaptureBatch,
Submission::ExplicitGraph);
const auto result = Run<TestType>(mode, 64, 1048576, false);
INFO("mode=" << static_cast<unsigned>(mode) << ", bad=" << result.badBeforeSync);
CHECK_FALSE(result.timeout);
CHECK(result.checked == 64);
CHECK(result.badBeforeSync == 0);
CHECK(result.badAfterSync == 0);
}
#endifRuntime alternative relative to fixed baseline diff --git a/projects/clr/rocclr/device/blitcl.cpp b/projects/clr/rocclr/device/blitcl.cpp
index 4db6780484..e2e6fcf545 100644
--- a/projects/clr/rocclr/device/blitcl.cpp
+++ b/projects/clr/rocclr/device/blitcl.cpp
@@ -8,6 +8,9 @@ namespace amd::device {
#define BLIT_KERNELS(...) #__VA_ARGS__
+// Default WRITE helpers use a relaxed device-library store. Release prior writes
+// before that store; the dispatch release fence runs only at kernel completion.
+// An empty LLVM syncscope denotes full system scope, as in __threadfence_system.
const char* BlitLinearSourceCode = BLIT_KERNELS(
// Extern
extern void __amd_fillBufferAligned2D(__global uchar*, __global ushort*, __global uint*,
@@ -237,6 +240,7 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
(ulong)~0UL);
break;
case ROCCLR_STREAM_WRITE_VALUE_32:
+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
__amd_streamOpsWrite((__global uint*)param[i].writeValue.address, NULL,
(uint)param[i].writeValue.value);
break;
@@ -246,6 +250,7 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
(uint)param[i].waitValue.flags, (ulong)~0UL);
break;
case ROCCLR_STREAM_WRITE_VALUE_64:
+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
__amd_streamOpsWrite(NULL, (__global ulong*)param[i].writeValue.address,
param[i].writeValue.value64);
break;
@@ -258,6 +263,7 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
const char* HipExtraSourceCode = BLIT_KERNELS(
__kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong,
ulong value) {
+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
__amd_streamOpsWrite(ptrInt, ptrUlong, value);
}
@@ -286,6 +292,7 @@ const char* HipExtraSourceCode = BLIT_KERNELS(
const char* HipExtraSourceCodeNoGWS = BLIT_KERNELS(
__kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong,
ulong value) {
+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
__amd_streamOpsWrite(ptrInt, ptrUlong, value);
}Matched HIP DSO SHA256: base Raw CUDA environment and all six case logsnvcc-version.log nvidia-smi-L.log write32-4096bytes.log write32-1048576bytes.log write32-8388608bytes.log write64-4096bytes.log write64-1048576bytes.log write64-8388608bytes.log Original CUDA manifest and archive identityOriginal supplied archive SHA256: {
"created_utc": "20261027T050314.144422Z",
"platform": "Linux-7.0.0-31-generic-x86_64-with-glibc2.39",
"python": "3.12.3 (main, Aug 14 2025, 17:47:21) [GCC 13.3.0]",
"source_sha256": "3a90793a23a751ab5c36677f4578ba512f0d734ecd119ceb60191f7b6b090ac4",
"cases": [
{
"width": 32,
"bytes": 4096,
"iterations": 250,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write32-4096bytes.log",
"result": {
"width": 32,
"bytes": 4096,
"iterations": 250,
"checked": 250,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
},
{
"width": 32,
"bytes": 1048576,
"iterations": 2000,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write32-1048576bytes.log",
"result": {
"width": 32,
"bytes": 1048576,
"iterations": 2000,
"checked": 2000,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
},
{
"width": 32,
"bytes": 8388608,
"iterations": 250,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write32-8388608bytes.log",
"result": {
"width": 32,
"bytes": 8388608,
"iterations": 250,
"checked": 250,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
},
{
"width": 64,
"bytes": 4096,
"iterations": 250,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write64-4096bytes.log",
"result": {
"width": 64,
"bytes": 4096,
"iterations": 250,
"checked": 250,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
},
{
"width": 64,
"bytes": 1048576,
"iterations": 2000,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write64-1048576bytes.log",
"result": {
"width": 64,
"bytes": 1048576,
"iterations": 2000,
"checked": 2000,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
},
{
"width": 64,
"bytes": 8388608,
"iterations": 250,
"status": "PASS",
"exit_code": 0,
"outer_timeout": false,
"log": "write64-8388608bytes.log",
"result": {
"width": 64,
"bytes": 8388608,
"iterations": 250,
"checked": 250,
"bad_before_sync": 0,
"bad_words": 0,
"bad_after_sync": 0,
"timeout": 0
},
"loaded_dsos": [
{
"path": "/usr/lib/x86_64-linux-gnu/ld-linux-x86-64.so.2",
"sha256": "6222a16be7f2d458d6870efe6e715fc0c8d45766fb79cf7dcc3125538d703e28"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libc.so.6",
"sha256": "511f825ee075610ac9c0f7f91e2c13de2000d0f7b859f6461137e809a0a009d0"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libcuda.so.615.71.09",
"sha256": "85a8cc8cf00419ceb295ec62c481d05b4cdbdab10f25c5375d6d24bdbe00b4d9"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libdl.so.2",
"sha256": "feee7549138dc862300b9ba9b2b02502059c838a6c0b92d2ea4c0a357e8e0b45"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libgcc_s.so.1",
"sha256": "02f3f192bf5f79b811f1e34a650fea443d407408703448906585232383957f60"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libm.so.6",
"sha256": "f06f2ce1f1833df5f41cf13b6447ff07bea993ad9b27297d3428c2f70ab3f0e7"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libpthread.so.0",
"sha256": "4472cde19acc87a431ad1eefe9e86f8238dd1f26f8dac692cff7040474d9d085"
},
{
"path": "/usr/lib/x86_64-linux-gnu/librt.so.1",
"sha256": "4b3a59575a0812a466235bc745814ca5b13825ef5f8ad5680637872ac64cc7ca"
},
{
"path": "/usr/lib/x86_64-linux-gnu/libstdc++.so.6.0.33",
"sha256": "a68762c86d371e6041f03f03a33a78fa235809ae7d81c90185940c93c3535aed"
},
{
"path": "/usr/local/cuda-13.0/targets/x86_64-linux/lib/libcudart.so.13.0.96",
"sha256": "96c42e418cec19054186b9429c321603cc190bf26a18104e19408117a2a817b0"
}
],
"cuda_dso_identity_valid": true
}
],
"commands": [
{
"argv": [
"/usr/local/cuda-13.0/bin/nvcc",
"--version"
],
"log": "nvcc-version.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"nvidia-smi",
"-L"
],
"log": "nvidia-smi-L.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"nvidia-smi",
"-q"
],
"log": "nvidia-smi-q.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/usr/local/cuda-13.0/bin/nvcc",
"-std=c++17",
"-O2",
"-lineinfo",
"-arch=sm_120",
"--cudart",
"shared",
"-Xcompiler",
"-pthread",
"/work/cuda_memop_publication.cu",
"-o",
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"-L",
"/usr/local/cuda-13.0/lib64/stubs",
"-L",
"/usr/local/cuda-13.0/targets/x86_64-linux/lib/stubs",
"-lcuda"
],
"log": "build.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"32",
"250",
"4096"
],
"log": "write32-4096bytes.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"32",
"2000",
"1048576"
],
"log": "write32-1048576bytes.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"32",
"250",
"8388608"
],
"log": "write32-8388608bytes.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"64",
"250",
"4096"
],
"log": "write64-4096bytes.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"64",
"2000",
"1048576"
],
"log": "write64-1048576bytes.log",
"exit_code": 0,
"outer_timeout": false
},
{
"argv": [
"/work/results-20261027T050314.144422Z/cuda_memop_publication",
"64",
"250",
"8388608"
],
"log": "write64-8388608bytes.log",
"exit_code": 0,
"outer_timeout": false
}
],
"environment": {
"CUDA_VISIBLE_DEVICES": "GPU-17390cc0-bf3c-53e3-170c-21eb949de43b",
"CUDA_LAUNCH_BLOCKING": "0",
"LD_LIBRARY_PATH": "/usr/local/lib/python3.12/dist-packages/torch/lib:/usr/local/lib/python3.12/dist-packages/torch_tensorrt/lib:/usr/local/cuda/compat/lib:/usr/local/nvidia/lib:/usr/local/nvidia/lib64",
"CUDA_MODULE_LOADING": "LAZY",
"CUDA_FORCE_PTX_JIT": null,
"CUDA_DEVICE_MAX_CONNECTIONS": null,
"LD_PRELOAD": null,
"CUDA_INJECTION64_PATH": null
},
"nvcc": "/usr/local/cuda-13.0/bin/nvcc",
"nvcc_sha256": "b8e2347d9d7fcadf8812a27324d515de99d033d3e48611811ece7e66d7895f8a",
"binary_sha256": "700ef78e4cb772b82d5417824b94adc4e679b178fd7d3e563fcec0d7f37b0b51",
"effective_LD_LIBRARY_PATH": "/usr/local/cuda-13.0/lib64:/usr/local/cuda-13.0/targets/x86_64-linux/lib:/usr/local/lib/python3.12/dist-packages/torch/lib:/usr/local/lib/python3.12/dist-packages/torch_tensorrt/lib:/usr/local/cuda/compat/lib:/usr/local/nvidia/lib:/usr/local/nvidia/lib64",
"status": "PASS"
} |
|
@AlexXAmd Following the RTX 5090/kernel evidence, I completed application qualification of the separate, one-file pre-store release candidate. This is not validation of the currently submitted barrier/Graph implementation, whose head remains Contract and implementation are still separate review questions. One documentation reference may help: CUDA 13.0's On the same R9700/gfx1201, matched local CLR builds and identical ROCr/COMGR give these new no-reset-WRITE results:
No timeout or diagnostic unlock occurred. The application then passed eight Graph/Eager logits comparisons (NRMSE=0, matching argmax). Fifteen independent model services, including matched rebuilt controls, completed three rounds at 512/32K inputs; all 30 measured requests produced matching text within each round/context. These tests validate the tested candidate/configuration, not every ASIC or the missing HIP contract. There is no demonstrated stable model speedup. Release/raw and stock/event both averaged about 84.49 tok/s at 512. Long inter-token pauses occurred in stock/kernel, stock/event and release/raw; all slow rounds are retained and the pauses are not attributed to this patch. Kernel IO also fuses copies/accumulation and skips empty CPU tasks, so removing an event alone does not reproduce that application optimization. My proposed next code revision is the one-file CLR approach already shown above: four pre-store release sites in the ordinary GWS/NoGWS WRITE helpers and the batch WRITE32/64 cases. It adds no AQL packets, so it would remove the publication-packet-related Graph runtime changes while retaining the dependent WAIT-to-WRITE and WRITE-to-WAIT tests. WAIT-only batches acquire no new release. Backend-name inference would not be retained as public capability discovery; that API/error-handling question should be handled separately. Could we settle whether HIP should provide this stream-scoped pre-notification publication guarantee, including each default WRITE in a batch, and whether this CLR direction is appropriate? I am happy to proceed as a contract clarification/restoration or an enhancement. The remaining technical work is to establish the full producer-completion/acquire-to-notification chain, including other workgroups and SDMA, then qualify the final revision and cross-ASIC cost. A fence in one work-item is not, by itself, proof of that stream-wide chain. I am not claiming a firmware/driver violation or proposing to merge based only on single-GPU pass counts. Build identities and all model rounds
Qwen3-30B-A3B FP8, BF16 compute/KV, TP1, concurrency1, 20 GiB/4550-slot expert cache, 17 CPU workers, 256 output tokens; warmup is separate. The private event/control and release/raw comparison changes the whole solution (release plus event removal), not just the fence cost.
Maximum observed 32K token gaps in the affected measured runs: stock/kernel 2.053 s, stock/event 0.536 s, release/raw 0.877 s and 1.812 s. The correctness result is not a localization of these pauses. The exact commands, sources, DSO identities, raw requests and failure/cleanup records are retained in the local evidence bundle, SHA256 |
Review status (2026-10-27)
The default pre-notification publication contract remains under review. Head
65467398dfbelow still implements the barrier-packet approach; a separate one-file CLR release candidate has now passed local strict-publication and model correctness checks, without demonstrating a stable model throughput advantage. Those candidate results do not qualify the submitted head. See the updated evidence and proposed revision.The proposed revision would remove the publication-packet-related Graph runtime changes while retaining dependent update tests. Stream-wide ordering sufficiency, public capability discovery and cross-ASIC cost remain open; no firmware/driver violation is claimed.
Motivation
Restore publication before stream-WRITE notifications. CLR historically issued a SYSTEM-release barrier before its relaxed WRITE blit. #2304 replaced that barrier with dispatch scopes while intending to preserve ordering. On the tested R9700/gfx1201, a minimal
D2H → WRITE(default) → CPU acquire/readshows that these sequences are not equivalent: the notification can be observed before the copied data.The ownership trace and controlled restoration identify the caller-side change. Batch/capture has a separate history: its submission path lacked the ordinary WRITE's pre-release, and #12692 added dispatch scopes only. This PR addresses both paths; it does not attribute every graph/batch failure to #2304 or allege a firmware/HSA violation.
Technical Details
The device helper explicitly uses a system-scope relaxed store. The notification dispatch's SYSTEM release occurs after kernel completion, which is too late for a CPU observing a store inside that kernel.
BARRIER_ANDbefore ordinary WRITE/update operations and write-containing batches; preserve existing dispatch scopes and dependencies. WAIT-only batches add no publication barrier.The goal is to restore the existing pre-WRITE ordering and provide consistent batch publication. Formal default/batch contract wording, implementation cost, and safe fence reuse remain review points.
Issue Tracking
Eager ordering follows up #2304; batch dispatch scope follows up #12692, already present in the baseline. Targets
develop; release backports require separate validation.Test Plan
The automated minimal regression tests one D2H and one default WRITE, 2000 iterations per width. The destination is poisoned; CPU validation precedes stream synchronization. No graph, GPU WAIT, reset WRITE, result H2D, or application kernel is needed. Timeouts fail the test.
Broader tests validate CPU inputs and GPU outputs across eager, batch, native capture, captured batch, and explicit graphs. Update regressions cover both WAIT↔WRITE directions with dependent segments. Backend selection preserves supported PAL eager checks and Windows ROCr coverage while excluding unsupported batch modes.
timeout 90s ./hipStreamMemOpPublication --reporter compact timeout 90s ./hipGraphExecBatchMemOpNodeSetParams \ 'Unit_hipGraphExecBatchMemOpNodeSetParams_LayoutChange*' --reporter compactTest Result
Previously recorded validation of candidate
65467398dfon gfx1201: minimal publication 0/4000 errors; publication suite 58 assertions/6 cases with segmented and classic executors; layout updates 1276 assertions in each recorded configuration. Actual Windows/PAL and wider platform qualification remain pending; CI is not claimed fully passing.Separate ownership diagnostic, not a new Graph validation of this candidate: on fixed baseline
02d6cbc71114, restoring only the pre-#2304 eager WRITE release changes stale snapshots from 3149/4000 to 0/4000. Post-sync data is correct in both. The same test binary and identical non-HIP loaded-library hashes were verified. This is a modern-tree restoration, not a full historical-parent build/bisect.Matched local event-time medians for this PR's implementation were eager WRITE 2.943→6.758 µs/op, graph single-operation nodes 2.555→3.909, and graph batches of 128 0.370→0.379 per operation. Shared hardware/unlocked clocks limit these cost estimates; cross-ASIC cost remains unqualified. Peer/NIC, metadata-queue hardware, full-model validation, and release backports remain unverified. Captured custom WAIT masks are unchanged.
Submission Checklist