From 6af8546530c8a15272581014d09758516721ce50 Mon Sep 17 00:00:00 2001 From: diego Date: Tue, 27 Oct 2026 11:49:10 -0300 Subject: [PATCH] shim: take the cubin CRC off the application's thread, which deadlocked it (#211) The MODULE_LOADED callback computed the cubin CRC with cuptiGetCubinCrc, which takes the adapter's one CUPTI guard. Under lazy module loading the callback fires from inside cuLaunchKernel, on an application thread that holds a driver lock. The timer thread held the guard across cuptiActivityFlushAll, waiting for a CUPTI worker that was waiting for that driver lock. Three threads, a permanent deadlock of the profiled process: 1 run in 8 on a PyTorch workload with only the shim injected, and 4 captures in 7 empty with the agent attached. CubinQueue now has three steps instead of two. capture() copies and parks the copy as pending. It takes no CRC function, so an application thread inside a vendor callback waits on nothing but the queue's leaf mutex. settle() runs on the timer thread every tick, attached or not: it computes the CRC over the copy, de-duplicates, fires gpu_module_load_v1, and hands the entry to drain(). It holds the mutex across neither the CRC nor the record. The record carries the callback's module_id and load_ns, so its timestamp is still the load time, not the time of the tick. The pending list is bounded separately from the offer queue, so a full, retained queue still lets a new module reach settle() and its record. Past the pending bounds a module is lost whole and counted as cubin_pending_full, which the report line now prints alongside cubin_pending_depth. check-cubin-defer gains a mode 6: passing a CRC function to capture() must not compile. That makes the property structural, the way the copy deadline already is. Verified on an RTX 3090: 30 of 30 shim-only runs complete (the old shim hung 1 in 8), and 10 of 10 captures with the agent attached are full, about 79k records each. Module counters are identical to the old shim's on the same workload. Holding the mutex across the CRC in settle() fails the new queue test. --- shim/Makefile | 13 +- shim/core/cubin_defer_test.cc | 25 +++- shim/core/cubinqueue.cc | 158 +++++++++++++------- shim/core/cubinqueue.h | 83 +++++++++-- shim/core/cubinqueue_test.cc | 273 +++++++++++++++++++++++++++++++--- shim/nvidia/cupti_adapter.cc | 65 +++++--- shim/stub/stub.cc | 30 ++-- 7 files changed, 517 insertions(+), 130 deletions(-) diff --git a/shim/Makefile b/shim/Makefile index daa0bf8..60e1ee5 100644 --- a/shim/Makefile +++ b/shim/Makefile @@ -329,6 +329,10 @@ nvidia/testdata/cuda_concurrent: nvidia/testdata/cuda_concurrent.cu # that the enforcement works, by trying to violate it five ways and requiring # every one to be rejected by the compiler. # +# Mode 6 guards the opposite direction (issue #211): capture() takes no CRC +# function, because the CRC is a guarded CUPTI call and the callback can run +# on an application thread holding a driver lock the guard's holder needs. +# # Mode 0 is the control and it MUST compile. Without it the whole check would # pass just as happily on a file that fails to build for an unrelated reason - # a typo, a renamed header - and would then be proving nothing, which is the @@ -344,17 +348,18 @@ check-cubin-defer: core/cubin_defer_test.cc core/cubinqueue.h echo "check-cubin-defer: the COMPLIANT capture does not compile; this check would"; \ echo " otherwise pass for the wrong reason. Fix core/cubinqueue.h."; \ exit 1; } - @rc=0; for m in 1 2 3 4 5; do \ + @rc=0; for m in 1 2 3 4 5 6; do \ if $(CXX) -std=c++17 -Wall -Werror -I core -fsyntax-only \ -DPERFAGENT_CUBIN_DEFER=$$m core/cubin_defer_test.cc 2>/dev/null; then \ - echo "check-cubin-defer: deferral mode $$m COMPILED. The vendor's pointer can escape"; \ + echo "check-cubin-defer: mode $$m COMPILED. Modes 1-5: the vendor's pointer can escape"; \ echo " the callback, so the copy can be deferred past the buffer's"; \ echo " lifetime and a late reader gets silently WRONG bytes (spec"; \ - echo " 6.3 finding 2). See core/cubin_defer_test.cc mode $$m."; \ + echo " 6.3 finding 2). Mode 6: the callback can take the CRC, a guarded"; \ + echo " CUPTI call, which deadlocks (issue #211). See mode $$m."; \ rc=1; \ fi; \ done; \ - [ $$rc -eq 0 ] && echo "check-cubin-defer: OK - the compliant capture compiles, all 5 deferrals do not"; \ + [ $$rc -eq 0 ] && echo "check-cubin-defer: OK - the compliant capture compiles, all 6 violations do not"; \ exit $$rc # Issue #99's structural assertion, run as a BUILD step for the same reason diff --git a/shim/core/cubin_defer_test.cc b/shim/core/cubin_defer_test.cc index db735d2..64d6be9 100644 --- a/shim/core/cubin_defer_test.cc +++ b/shim/core/cubin_defer_test.cc @@ -17,10 +17,18 @@ // 3 move the view into a container // 4 read the borrowed pointer straight out of the view // 5 assign one view over another +// 6 compute the CRC inside capture (issue #211) // // Each of 1-5 is a real thing an implementer would reach for while "just // deferring the copy to the drain thread"; each is a compile error, and the // Makefile prints which. +// +// Mode 6 is the opposite temptation: doing MORE in the callback. The CRC is +// a vendor call behind the CUPTI guard, and the callback can run on an +// application thread inside cuLaunchKernel, holding a driver lock that the +// guard's holder is waiting for. capture() therefore takes no CRC function, +// and putting one back is a compile error rather than a deadlock found a +// release later. #include "cubinqueue.h" #include @@ -34,12 +42,14 @@ namespace { +#if PERFAGENT_CUBIN_DEFER == 6 uint64_t fake_crc(const void *bytes, size_t len) { (void)bytes; return len ? 0x1234u : 0u; } -void fake_captured(void *, uint64_t, const void *, size_t) {} +void fake_captured(void *, uint64_t, uint64_t, uint64_t, const void *, size_t) {} +#endif } // namespace @@ -49,7 +59,7 @@ void fake_captured(void *, uint64_t, const void *, size_t) {} // capture, and dies there. The copy is taken before the function returns. bool compliant(perfagent::CubinQueue &q, const void *vendor_bytes, size_t len) { const perfagent::CubinView view(vendor_bytes, len); - return q.capture(view, fake_crc, fake_captured, nullptr); + return q.capture(view, /*module_id=*/1, /*load_ns=*/0); } #elif PERFAGENT_CUBIN_DEFER == 1 @@ -100,6 +110,15 @@ void defer_by_reassigning(perfagent::CubinView &slot, const void *vendor_bytes, slot = perfagent::CubinView(vendor_bytes, len); } +#elif PERFAGENT_CUBIN_DEFER == 6 + +// "The CRC is cheap, take it with the copy like it used to be." It is a +// vendor call behind the CUPTI guard, and that is issue #211. +bool crc_in_the_callback(perfagent::CubinQueue &q, const void *vendor_bytes, size_t len) { + const perfagent::CubinView view(vendor_bytes, len); + return q.capture(view, fake_crc, fake_captured, nullptr); +} + #else -#error "PERFAGENT_CUBIN_DEFER must be 0..5" +#error "PERFAGENT_CUBIN_DEFER must be 0..6" #endif diff --git a/shim/core/cubinqueue.cc b/shim/core/cubinqueue.cc index f47e232..277ede9 100644 --- a/shim/core/cubinqueue.cc +++ b/shim/core/cubinqueue.cc @@ -15,17 +15,19 @@ CubinQueue::CubinQueue(const CubinQueueLimits &limits) : limits_(limits) {} CubinQueue::~CubinQueue() { std::lock_guard g(mu_); + for (Entry &e : pending_) free(e.bytes); + pending_.clear(); + pending_bytes_ = 0; for (Entry &e : q_) free(e.bytes); q_.clear(); queued_bytes_ = 0; } -bool CubinQueue::capture(const CubinView &view, CubinCrcFn crc, CubinCapturedFn on_captured, - void *ctx) { +bool CubinQueue::capture(const CubinView &view, uint64_t module_id, uint64_t load_ns) { // Everything before the memcpy is a decision NOT to copy - a bound, or a // missing input. Nothing here defers the copy, and nothing here may: see // cubinqueue.h and core/cubin_defer_test.cc. - if (view.empty() || crc == nullptr) return false; + if (view.empty()) return false; const size_t len = view.size(); if (len > limits_.max_cubin_bytes) { // Refused whole, never truncated. A truncated cubin parses into a @@ -47,67 +49,110 @@ bool CubinQueue::capture(const CubinView &view, CubinCrcFn crc, CubinCapturedFn return false; } - // Over the copy, never over the view: a CRC and the bytes it names must - // come from one read of one buffer. - const uint64_t key = crc(copy, len); - if (key == 0) { - // Zero is the ABI's "no module". Sending it would produce a module - // record that joins to nothing, with every other counter green. - free(copy); - crc_failed_.fetch_add(1, std::memory_order_relaxed); - return false; - } - - // One lock from the intern check to the push, so two threads loading the - // same module cannot both decide they are the first and fire two records - // for one CRC. on_captured runs inside it and is documented as forbidden - // from re-entering the queue; it is a probe fire, which is a nop when no - // consumer is attached and a ~1-2us uprobe trap when one is. + // And that is all the callback does. The CRC is NOT taken here: on the + // vendor side it takes the CUPTI guard, and this thread may be inside + // cuLaunchKernel holding a driver lock the guard's holder is waiting + // for (issue #211). settle() computes it, over this same copy, on a + // thread that holds nothing the driver needs. std::unique_lock g(mu_); - - if (interned_.count(key) != 0) { + if (pending_.size() >= limits_.max_pending_entries || + pending_bytes_ + len > limits_.max_pending_bytes) { + // More modules inside one settle tick than the pending bounds hold. + // Unlike a full offer queue this loses the record as well as the + // bytes, because the CRC the record carries is never computed -- + // which is why it has a counter of its own. g.unlock(); free(copy); - reload_skipped_.fetch_add(1, std::memory_order_relaxed); + pending_full_.fetch_add(1, std::memory_order_relaxed); return false; } + Entry e; + e.crc = 0; + e.module_id = module_id; + e.load_ns = load_ns; + e.bytes = copy; + e.len = len; + pending_.push_back(e); + pending_bytes_ += len; + return true; +} - // Fired here, before the push, because this is the last moment the copy - // is owned by nobody else: after the push the drain thread may free it, - // and gpu_module_load_v1.bytes_ptr would then name freed memory. The - // record stays accurate exactly because it is emitted here. - if (on_captured != nullptr) on_captured(ctx, key, copy, len); - - if (q_.size() >= limits_.max_entries || queued_bytes_ + len > limits_.max_queued_bytes) { - // The OFFER is dropped; the copy has already served its purpose of - // being taken in time. The module goes unresolvable rather than the - // application going slow, and this counter is the size of that. - // - // Deliberately NOT interned: a later re-load of the same module gets - // another chance at a queue that may since have drained. - g.unlock(); - free(copy); - queue_full_.fetch_add(1, std::memory_order_relaxed); - return false; +size_t CubinQueue::settle(CubinCrcFn crc, CubinCapturedFn on_captured, void *ctx) { + if (crc == nullptr) return 0; + size_t budget; + { + std::lock_guard g(mu_); + budget = pending_.size(); } + size_t done = 0; + for (; done < budget; done++) { + Entry e; + { + std::lock_guard g(mu_); + if (pending_.empty()) break; + e = pending_.front(); + pending_.pop_front(); + pending_bytes_ -= e.len; + } - // Interning is bounded like everything else. Past the bound we stop - // remembering rather than stop capturing: the consequence of forgetting - // is a duplicate offer, which the consumer answers with a counted no-op - // (CubinsDuplicate), and the consequence of not capturing is a module - // nothing can resolve. The cheaper mistake is the one we make. - if (interned_.size() < limits_.max_interned) interned_.insert(key); + // Over the copy, never over the view: a CRC and the bytes it names + // must come from one read of one buffer. The mutex is NOT held: this + // is a vendor call behind the CUPTI guard, and a capture() blocked on + // mu_ behind it is the #211 cycle again with one more lock in it. + e.crc = crc(e.bytes, e.len); + if (e.crc == 0) { + // Zero is the ABI's "no module". Sending it would produce a + // module record that joins to nothing, with every other counter + // green. + free(e.bytes); + crc_failed_.fetch_add(1, std::memory_order_relaxed); + continue; + } - Entry e; - e.crc = key; - e.bytes = copy; - e.len = len; - q_.push_back(e); - queued_bytes_ += len; - g.unlock(); + // Only settle() touches interned_, and only one thread settles, so + // two loads of one module can no longer both decide they are first: + // they are settled one after the other on the same thread. + if (interned_.count(e.crc) != 0) { + free(e.bytes); + reload_skipped_.fetch_add(1, std::memory_order_relaxed); + continue; + } - captured_.fetch_add(1, std::memory_order_relaxed); - return true; + // Fired here, before the push, because this is the last moment the + // copy is owned by nobody else: after the push drain() may free it, + // and gpu_module_load_v1.bytes_ptr would then name freed memory. + if (on_captured != nullptr) + on_captured(ctx, e.module_id, e.load_ns, e.crc, e.bytes, e.len); + + std::unique_lock g(mu_); + if (q_.size() >= limits_.max_entries || queued_bytes_ + e.len > limits_.max_queued_bytes) { + // The OFFER is dropped; the copy has already served its purpose + // of being taken in time, and the record above has told the + // consumer the module exists. The module goes unresolvable + // rather than the application going slow, and this counter is + // the size of that. + // + // Deliberately NOT interned: a later re-load of the same module + // gets another chance at a queue that may since have drained. + g.unlock(); + free(e.bytes); + queue_full_.fetch_add(1, std::memory_order_relaxed); + continue; + } + q_.push_back(e); + queued_bytes_ += e.len; + g.unlock(); + + // Interning is bounded like everything else. Past the bound we stop + // remembering rather than stop capturing: the consequence of + // forgetting is a duplicate offer, which the consumer answers with a + // counted no-op (CubinsDuplicate), and the consequence of not + // capturing is a module nothing can resolve. The cheaper mistake is + // the one we make. + if (interned_.size() < limits_.max_interned) interned_.insert(e.crc); + captured_.fetch_add(1, std::memory_order_relaxed); + } + return done; } size_t CubinQueue::drain(CubinOfferFn offer, unsigned timeout_ms) { @@ -154,4 +199,9 @@ size_t CubinQueue::depth() const { return q_.size(); } +size_t CubinQueue::pending() const { + std::lock_guard g(mu_); + return pending_.size(); +} + } // namespace perfagent diff --git a/shim/core/cubinqueue.h b/shim/core/cubinqueue.h index b013646..044277c 100644 --- a/shim/core/cubinqueue.h +++ b/shim/core/cubinqueue.h @@ -47,6 +47,21 @@ // The CRC runs over the COPY, never over the view, for the same reason: the // crc function is handed the owned bytes, so a CRC and the bytes it names can // never come from two different reads of a buffer that changed in between. +// +// # Three steps, not two: the CRC is not computed in the callback +// +// capture() copies and returns. settle() -- on the adapter's own thread -- +// computes the CRC, de-duplicates, fires gpu_module_load_v1 and hands the +// entry to drain(). The CRC used to be taken inside capture(), and that +// deadlocked the profiled process (issue #211): on the vendor side the CRC is +// cuptiGetCubinCrc, which takes the adapter's one CUPTI guard, and the +// callback can fire from inside cuLaunchKernel (lazy module loading) on an +// application thread that holds a driver lock. The timer thread, holding the +// guard across cuptiActivityFlushAll, waited for a CUPTI worker that waited +// for that driver lock. An application thread inside a vendor callback must +// never wait for the guard, and the way to make that structural rather than +// a rule is for the callback half to take no CRC function at all: capture() +// has no parameter through which one could be passed. #ifndef PERFAGENT_CUBINQUEUE_H #define PERFAGENT_CUBINQUEUE_H @@ -99,11 +114,18 @@ class CubinView { // to nothing and would be an unresolvable module with every counter green. typedef uint64_t (*CubinCrcFn)(const void *bytes, size_t len); -// Called once per captured module, on the callback's thread, while the copy -// is owned by nobody else. This is where gpu_module_load_v1 is fired, and it -// is the only moment at which the record's bytes_ptr provably points at live -// adapter-owned bytes. It must not call back into the queue. -typedef void (*CubinCapturedFn)(void *ctx, uint64_t crc, const void *bytes, size_t len); +// Called once per settled module, on settle()'s thread, while the copy is +// owned by nobody else: it has left the pending list and has not yet been +// handed to drain(), so nothing can have offered and freed it. This is where +// gpu_module_load_v1 is fired, and it is the only moment at which the +// record's bytes_ptr provably points at live adapter-owned bytes. It must not +// call back into the queue. +// +// module_id and load_ns are what capture() was given, carried through +// unchanged: the record's load time is the callback's, not the settle tick's, +// so moving the CRC off the application's thread does not move the timestamp. +typedef void (*CubinCapturedFn)(void *ctx, uint64_t module_id, uint64_t load_ns, uint64_t crc, + const void *bytes, size_t len); // The send. Signature-compatible with cubin_offer_to_consumer by design, so // the production wiring is a name and the tests are a substitute. @@ -135,6 +157,17 @@ struct CubinQueueLimits { size_t max_queued_bytes = 64u * 1024 * 1024; // their total size_t max_cubin_bytes = 8u * 1024 * 1024; // one module, the memcpy bound size_t max_interned = 4096; // distinct CRCs remembered + + // The pending list: copies taken in the callback and not yet settled. + // It is emptied every settle tick, so it only fills when more than + // max_entries modules (or max_queued_bytes of them) load inside ONE + // tick. It is bounded separately from the offer queue on purpose: a + // retained queue that is full must still let a new module reach settle(), + // because settle() is where its gpu_module_load_v1 fires, and the + // consumer learning THAT a module loaded is worth more than its bytes. + // The price is that the worst case briefly holds both bounds' worth. + size_t max_pending_entries = 512; + size_t max_pending_bytes = 64u * 1024 * 1024; }; // Bounded, mutex-guarded, and drained by somebody else's thread. @@ -146,13 +179,25 @@ class CubinQueue { CubinQueue(const CubinQueue &) = delete; CubinQueue &operator=(const CubinQueue &) = delete; - // The callback half. Copies view NOW, computes the CRC over the copy, - // calls on_captured while the copy is still exclusively ours, and pushes. - // Returns true only when the copy was made AND enqueued for offer. + // The callback half. Copies view NOW and parks the copy as pending. + // Calls nothing it was not built with -- no CRC, no record, no vendor + // API -- so an application thread inside a vendor callback waits on at + // most this queue's own leaf mutex. Returns true when the copy was made + // and parked. // // Every early return has a counter; none is silent. - bool capture(const CubinView &view, CubinCrcFn crc, CubinCapturedFn on_captured, - void *ctx); + bool capture(const CubinView &view, uint64_t module_id, uint64_t load_ns); + + // The settle half, on the adapter's own thread and never on an + // application's: for each pending copy, computes the CRC over it, skips a + // CRC already seen, fires on_captured, and enqueues the copy for drain(). + // Settles at most the entries present on entry, and holds the mutex + // across neither crc nor on_captured -- crc is a vendor call that takes + // the CUPTI guard, and a capture() waiting on our mutex behind it would + // rebuild issue #211 one lock further down. Returns how many it settled. + // + // Must be called from one thread at a time: the intern set is its alone. + size_t settle(CubinCrcFn crc, CubinCapturedFn on_captured, void *ctx); // The drain-thread half. Offers at most the entries present on entry, so // a producer that keeps capturing cannot pin this thread here, and never @@ -170,16 +215,22 @@ class CubinQueue { uint64_t cubin_too_large() const { return too_large_.load(std::memory_order_relaxed); } uint64_t cubin_crc_failed() const { return crc_failed_.load(std::memory_order_relaxed); } uint64_t cubin_alloc_failed() const { return alloc_failed_.load(std::memory_order_relaxed); } + uint64_t cubin_pending_full() const { return pending_full_.load(std::memory_order_relaxed); } uint64_t cubins_sent() const { return sent_.load(std::memory_order_relaxed); } - // Queued-but-not-yet-offered. A producer waiting for its offers to land - // polls this; an operator reads it as the size of what a hard exit would - // take with it. + // Settled-but-not-yet-offered. A producer waiting for its offers to land + // polls this. size_t depth() const; + // Captured-but-not-yet-settled. depth() + pending() is the size of what + // a hard exit would take with it. + size_t pending() const; + private: struct Entry { - uint64_t crc; + uint64_t crc; // 0 while pending: the CRC is settle()'s to compute + uint64_t module_id; + uint64_t load_ns; void *bytes; size_t len; }; @@ -187,8 +238,11 @@ class CubinQueue { const CubinQueueLimits limits_; mutable std::mutex mu_; + std::deque pending_; + size_t pending_bytes_ = 0; std::deque q_; size_t queued_bytes_ = 0; + // settle()'s alone, so it is not under mu_. std::unordered_set interned_; std::atomic captured_{0}; @@ -198,6 +252,7 @@ class CubinQueue { std::atomic too_large_{0}; std::atomic crc_failed_{0}; std::atomic alloc_failed_{0}; + std::atomic pending_full_{0}; std::atomic sent_{0}; }; diff --git a/shim/core/cubinqueue_test.cc b/shim/core/cubinqueue_test.cc index d3d505a..3df8109 100644 --- a/shim/core/cubinqueue_test.cc +++ b/shim/core/cubinqueue_test.cc @@ -3,8 +3,9 @@ // // What this proves without a GPU: that the copy is taken inside capture() // and is a real copy (mutating the source afterwards does not change what is -// offered); that the CRC is computed over the copy; that gpu_module_load_v1's -// moment is inside capture() while the bytes are exclusively ours; that a +// offered); that capture() computes no CRC and fires no record, and settle() +// does both over the copy (issue #211); that gpu_module_load_v1's moment is +// inside settle() while the bytes are exclusively ours; that a // re-load of the same CRC is skipped and counted; that every bound drops the // OFFER and counts it rather than dropping the copy or blocking; and that // drain() does not hold the mutex across an offer, which is the property that @@ -78,8 +79,17 @@ struct Captured { Captured g_captured; -void record_capture(void *ctx, uint64_t crc, const void *bytes, size_t len) { +struct CapturedIds { + uint64_t module_id; + uint64_t load_ns; +}; +CapturedIds g_captured_ids; + +void record_capture(void *ctx, uint64_t module_id, uint64_t load_ns, uint64_t crc, + const void *bytes, size_t len) { (void)ctx; + g_captured_ids.module_id = module_id; + g_captured_ids.load_ns = load_ns; g_captured.crc = crc; g_captured.bytes.assign((const char *)bytes, len); g_captured.len = len; @@ -88,6 +98,18 @@ void record_capture(void *ctx, uint64_t crc, const void *bytes, size_t len) { void reset_captured() { g_captured = Captured{0, std::string(), 0, 0}; } +// The two halves as the adapter runs them on a quiet process: the callback +// captures, the next tick settles. Returns what the old one-step capture +// returned -- whether the module reached the offer queue -- so the tests that +// are about bounds and offers read exactly as they did. +bool capture_and_settle(CubinQueue &q, const CubinView &v, perfagent::CubinCrcFn crc, + perfagent::CubinCapturedFn on_captured) { + const uint64_t before = q.modules_captured(); + if (!q.capture(v, 0, 0)) return false; + q.settle(crc, on_captured, nullptr); + return q.modules_captured() > before; +} + std::string body(size_t n, char seed) { std::string s(n, '\0'); for (size_t i = 0; i < n; i++) s[i] = (char)(unsigned char)(i * 31 + (unsigned char)seed); @@ -110,7 +132,7 @@ void test_the_copy_is_taken_inside_the_callback() { { const CubinView view(vendor.data(), vendor.size()); - assert(q.capture(view, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, view, fnv1a, record_capture)); } // The vendor's buffer now holds something else entirely, exactly as it @@ -131,13 +153,13 @@ void test_the_copy_is_taken_inside_the_callback() { // gpu_module_load_v1 must be fired while the copy is exclusively ours, and // its bytes_ptr must name that copy - which is only checkable by asserting // the pointer handed to on_captured holds the right bytes at that instant. -void test_the_record_fires_inside_capture_over_live_bytes() { +void test_the_record_fires_inside_settle_over_live_bytes() { reset_offers(); reset_captured(); CubinQueue q; const std::string vendor = body(1024, 'B'); const CubinView view(vendor.data(), vendor.size()); - assert(q.capture(view, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, view, fnv1a, record_capture)); assert(g_captured.calls == 1); assert(g_captured.len == vendor.size()); @@ -146,7 +168,7 @@ void test_the_record_fires_inside_capture_over_live_bytes() { // ...and it fired BEFORE anything could have been offered, which is what // makes bytes_ptr valid at the moment the probe reads it. assert(g_offers.empty()); - printf(" gpu_module_load_v1 fires inside capture, over the owned copy\n"); + printf(" gpu_module_load_v1 fires inside settle, over the owned copy\n"); } // The CRC comes from the copy. A crc function that is handed the vendor's @@ -166,10 +188,17 @@ void test_the_crc_runs_over_the_copy() { g_crc_saw.clear(); CubinQueue q; std::string vendor = body(2048, 'C'); - const CubinView view(vendor.data(), vendor.size()); - assert(q.capture(view, recording_crc, record_capture, nullptr)); + { + const CubinView view(vendor.data(), vendor.size()); + assert(q.capture(view, 0, 0)); + } + // Between the callback and the settle tick the vendor reuses its buffer. + // The CRC is computed after this, so it is only right if it is computed + // over the copy. memset(&vendor[0], 0, vendor.size()); + assert(q.settle(recording_crc, record_capture, nullptr) == 1); assert(q.drain(record_offer, 1000) == 1); + assert(g_offers[0].crc == fnv1a(body(2048, 'C').data(), 2048)); assert(g_crc_saw_len == 2048); assert(g_crc_saw == g_offers[0].bytes); printf(" the CRC is computed over the copy, not over the borrowed buffer\n"); @@ -185,7 +214,7 @@ void test_a_crc_of_zero_is_counted_and_never_sent() { CubinQueue q; const std::string vendor = body(64, 'D'); const CubinView view(vendor.data(), vendor.size()); - assert(!q.capture(view, zero_crc, record_capture, nullptr)); + assert(!capture_and_settle(q, view, zero_crc, record_capture)); assert(q.cubin_crc_failed() == 1); assert(q.modules_captured() == 0); assert(g_captured.calls == 0); @@ -201,7 +230,7 @@ void test_a_reload_of_the_same_crc_is_skipped_and_counted() { const std::string vendor = body(512, 'E'); { const CubinView v(vendor.data(), vendor.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); } // The same module loaded again - CUDA's lazy loading does this - from a // different buffer with identical contents, which is what a content @@ -209,7 +238,7 @@ void test_a_reload_of_the_same_crc_is_skipped_and_counted() { const std::string again = body(512, 'E'); { const CubinView v(again.data(), again.size()); - assert(!q.capture(v, fnv1a, record_capture, nullptr)); + assert(!capture_and_settle(q, v, fnv1a, record_capture)); } assert(q.modules_captured() == 1); assert(q.module_reload_skipped() == 1); @@ -230,7 +259,7 @@ void test_a_full_queue_drops_the_offer_and_counts_it() { for (int i = 0; i < 5; i++) { const std::string vendor = body(128, (char)('a' + i)); const CubinView v(vendor.data(), vendor.size()); - q.capture(v, fnv1a, record_capture, nullptr); + capture_and_settle(q, v, fnv1a, record_capture); } assert(q.modules_captured() == 2); assert(q.cubin_queue_full() == 3); @@ -254,7 +283,7 @@ void test_the_byte_ceiling_is_a_second_full_condition() { for (int i = 0; i < 4; i++) { const std::string vendor = body(128, (char)('m' + i)); const CubinView v(vendor.data(), vendor.size()); - q.capture(v, fnv1a, record_capture, nullptr); + capture_and_settle(q, v, fnv1a, record_capture); } assert(q.modules_captured() == 2); // 128 + 128 = 256; a third would pass 300 assert(q.cubin_queue_full() == 2); @@ -271,7 +300,7 @@ void test_an_oversized_cubin_is_never_copied() { CubinQueue q(lim); const std::string vendor = body(1025, 'F'); const CubinView v(vendor.data(), vendor.size()); - assert(!q.capture(v, fnv1a, record_capture, nullptr)); + assert(!capture_and_settle(q, v, fnv1a, record_capture)); assert(q.cubin_too_large() == 1); assert(q.modules_captured() == 0); assert(g_captured.calls == 0); @@ -281,7 +310,7 @@ void test_an_oversized_cubin_is_never_copied() { // policy rather than an off-by-one. const std::string ok = body(1024, 'F'); const CubinView v2(ok.data(), ok.size()); - assert(q.capture(v2, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v2, fnv1a, record_capture)); assert(q.cubin_too_large() == 1); printf(" a cubin over the per-cubin ceiling is refused before the memcpy\n"); } @@ -292,7 +321,7 @@ void test_a_refused_offer_is_counted_as_a_send_failure() { CubinQueue q; const std::string vendor = body(256, 'G'); const CubinView v(vendor.data(), vendor.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); g_offer_result = kCubinOfferRefused; assert(q.drain(record_offer, 1000) == 1); assert(q.cubin_send_failed() == 1); @@ -308,7 +337,7 @@ void test_drain_bounds_itself_and_empties() { for (int i = 0; i < 8; i++) { const std::string vendor = body(64, (char)('0' + i)); const CubinView v(vendor.data(), vendor.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); } assert(q.depth() == 8); assert(q.drain(record_offer, 250) == 8); @@ -343,14 +372,14 @@ void test_a_slow_offer_does_not_block_a_capture() { const std::string first = body(64, 'H'); { const CubinView v(first.data(), first.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); } std::thread drainer([&] { q.drain(slow_offer, 1000); }); while (!g_slow_offer_entered.load()) std::this_thread::sleep_for(std::chrono::milliseconds(1)); const std::string second = body(64, 'I'); const CubinView v(second.data(), second.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); g_capture_finished.store(true); drainer.join(); assert(q.modules_captured() == 2); @@ -366,7 +395,7 @@ void test_a_healthy_run_reads_zero_on_every_drop_counter() { for (int i = 0; i < 4; i++) { const std::string vendor = body(1000 + i, (char)('p' + i)); const CubinView v(vendor.data(), vendor.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); } assert(q.drain(record_offer, 1000) == 4); assert(q.modules_captured() == 4); @@ -390,7 +419,7 @@ void test_destruction_releases_undrained_copies() { for (int i = 0; i < 3; i++) { const std::string vendor = body(256, (char)('x' + i)); const CubinView v(vendor.data(), vendor.size()); - assert(q.capture(v, fnv1a, record_capture, nullptr)); + assert(capture_and_settle(q, v, fnv1a, record_capture)); } assert(q.depth() == 3); } @@ -416,7 +445,7 @@ void test_an_undrained_queue_retains_everything_intact() { const char *payloads[] = {"module-alpha", "module-bravo", "module-charlie"}; for (int i = 0; i < 3; i++) { const CubinView v(payloads[i], strlen(payloads[i])); - assert(q.capture(v, fnv1a, nullptr, nullptr)); + assert(capture_and_settle(q, v, fnv1a, nullptr)); } // Many ticks' worth of doing nothing. Nothing may age out on its own: // there is no deadline in this policy, only the bounds. @@ -443,11 +472,200 @@ void test_the_retention_bounds_are_the_consumers_store_bounds() { printf(" the retention bounds are the consumer's module-store bounds\n"); } +// ---------------------------------------------------------- issue #211 +// +// The deadlock: the CRC used to be computed inside capture(), on the +// application's thread, and on the vendor side the CRC takes the CUPTI guard. +// The callback can fire from inside cuLaunchKernel holding a driver lock, and +// the guard's holder (the timer, inside cuptiActivityFlushAll) was waiting on +// a CUPTI worker that needed that lock. These tests pin the two properties +// that break the cycle. + +int g_crc_calls = 0; +uint64_t counting_crc(const void *bytes, size_t len) { + g_crc_calls++; + return fnv1a(bytes, len); +} + +// capture() calls no CRC and fires no record. It has no parameter through +// which it could, so this is as much a check that settle() is where both +// happen as that capture() does not do them. +void test_capture_computes_no_crc_and_fires_no_record() { + reset_offers(); + reset_captured(); + g_crc_calls = 0; + CubinQueue q; + const std::string vendor = body(512, 'J'); + { + const CubinView v(vendor.data(), vendor.size()); + assert(q.capture(v, 7, 1234)); + } + assert(g_crc_calls == 0); + assert(g_captured.calls == 0); + assert(q.pending() == 1); + assert(q.depth() == 0); + assert(q.modules_captured() == 0); + // Nothing is offerable before it is settled: a drain here offers nothing, + // because an offer without a CRC would join to nothing. + assert(q.drain(record_offer, 1000) == 0); + + assert(q.settle(counting_crc, record_capture, nullptr) == 1); + assert(g_crc_calls == 1); + assert(g_captured.calls == 1); + assert(q.pending() == 0); + assert(q.depth() == 1); + assert(q.modules_captured() == 1); + printf(" capture computes no CRC and fires no record; settle does both\n"); +} + +// The record's module_id and load_ns are the callback's, not the settle +// tick's: moving the CRC off the application's thread must not move the +// timestamp the consumer orders module loads by. +void test_the_record_carries_the_callbacks_module_id_and_load_ns() { + reset_offers(); + reset_captured(); + CubinQueue q; + const std::string vendor = body(256, 'K'); + { + const CubinView v(vendor.data(), vendor.size()); + assert(q.capture(v, 42, 987654321)); + } + assert(q.settle(fnv1a, record_capture, nullptr) == 1); + assert(g_captured_ids.module_id == 42); + assert(g_captured_ids.load_ns == 987654321); + printf(" the record carries the callback's module_id and load_ns\n"); +} + +// The cycle itself, at the queue's level. The CRC function blocks the way +// the vendor's does when the guard is held, and while it is blocked a +// capture() must still complete. If settle() held mu_ across the CRC, the +// capture would wait on mu_ behind it -- an application thread waiting, one +// lock further down, on the very call #211 moved off its thread. +std::atomic g_crc_entered{false}; +std::atomic g_blocked_capture_finished{false}; + +uint64_t blocking_crc(const void *bytes, size_t len) { + g_crc_entered.store(true); + for (int i = 0; i < 2000 && !g_blocked_capture_finished.load(); i++) { + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + return fnv1a(bytes, len); +} + +void test_a_capture_completes_while_settle_is_inside_the_crc() { + reset_offers(); + reset_captured(); + CubinQueue q; + const std::string first = body(64, 'L'); + { + const CubinView v(first.data(), first.size()); + assert(q.capture(v, 1, 1)); + } + std::thread settler([&] { q.settle(blocking_crc, nullptr, nullptr); }); + while (!g_crc_entered.load()) std::this_thread::sleep_for(std::chrono::milliseconds(1)); + + const auto t0 = std::chrono::steady_clock::now(); + const std::string second = body(64, 'M'); + { + const CubinView v(second.data(), second.size()); + assert(q.capture(v, 2, 2)); + } + const auto waited = std::chrono::steady_clock::now() - t0; + g_blocked_capture_finished.store(true); + settler.join(); + // The crc gives up after ~2 s on its own, so a capture that waited for it + // would still finish -- the assertion is that it did not wait. + assert(waited < std::chrono::milliseconds(500)); + // settle() took only what was present on entry; the second capture waits + // for the next tick rather than extending this one. + assert(q.modules_captured() == 1); + assert(q.pending() == 1); + printf(" a capture completes while settle is inside the CRC (no lock held across it)\n"); +} + +// Two loads of one module inside one tick are both pending when settle runs; +// the second is a reload of the first, not a second module. +void test_two_pending_loads_of_one_crc_settle_once() { + reset_offers(); + reset_captured(); + CubinQueue q; + const std::string a = body(300, 'N'); + const std::string b = body(300, 'N'); + { + const CubinView v(a.data(), a.size()); + assert(q.capture(v, 1, 1)); + } + { + const CubinView v(b.data(), b.size()); + assert(q.capture(v, 2, 2)); + } + assert(q.settle(fnv1a, record_capture, nullptr) == 2); + assert(q.modules_captured() == 1); + assert(q.module_reload_skipped() == 1); + assert(g_captured.calls == 1); + assert(q.depth() == 1); + printf(" two pending loads of one CRC settle as one module and one reload\n"); +} + +// The pending bounds, and their own counter: past them the module is lost +// whole -- no record, because its CRC is never computed. +void test_a_full_pending_list_is_counted_separately() { + reset_offers(); + reset_captured(); + CubinQueueLimits lim; + lim.max_pending_entries = 2; + CubinQueue q(lim); + for (int i = 0; i < 4; i++) { + const std::string vendor = body(128, (char)('q' + i)); + const CubinView v(vendor.data(), vendor.size()); + q.capture(v, (uint64_t)i, 0); + } + assert(q.pending() == 2); + assert(q.cubin_pending_full() == 2); + assert(q.cubin_queue_full() == 0); + assert(q.settle(fnv1a, record_capture, nullptr) == 2); + assert(g_captured.calls == 2); + printf(" a full pending list drops 2 and counts them as pending_full, not queue_full\n"); +} + +// The point of bounding pending separately: a RETAINED queue that is full +// must not stop a new module's record, because the record is what tells the +// consumer the module exists. +void test_a_full_offer_queue_does_not_stop_capture_or_the_record() { + reset_offers(); + reset_captured(); + CubinQueueLimits lim; + lim.max_entries = 1; + CubinQueue q(lim); + for (int i = 0; i < 3; i++) { + const std::string vendor = body(128, (char)('u' + i)); + const CubinView v(vendor.data(), vendor.size()); + assert(q.capture(v, (uint64_t)i, 0)); + q.settle(fnv1a, record_capture, nullptr); + } + assert(q.cubin_pending_full() == 0); + assert(q.cubin_queue_full() == 2); + assert(g_captured.calls == 3); + printf(" a full offer queue still lets every module reach settle and its record\n"); +} + +// Destruction frees pending copies as well as settled ones. +void test_destruction_releases_pending_copies() { + { + CubinQueue q; + const std::string vendor = body(256, 'Z'); + const CubinView v(vendor.data(), vendor.size()); + assert(q.capture(v, 1, 1)); + assert(q.pending() == 1); + } + printf(" destroying a queue with pending copies releases them\n"); +} + } // namespace int main() { test_the_copy_is_taken_inside_the_callback(); - test_the_record_fires_inside_capture_over_live_bytes(); + test_the_record_fires_inside_settle_over_live_bytes(); test_the_crc_runs_over_the_copy(); test_a_crc_of_zero_is_counted_and_never_sent(); test_a_reload_of_the_same_crc_is_skipped_and_counted(); @@ -461,6 +679,13 @@ int main() { test_destruction_releases_undrained_copies(); test_an_undrained_queue_retains_everything_intact(); test_the_retention_bounds_are_the_consumers_store_bounds(); + test_capture_computes_no_crc_and_fires_no_record(); + test_the_record_carries_the_callbacks_module_id_and_load_ns(); + test_a_capture_completes_while_settle_is_inside_the_crc(); + test_two_pending_loads_of_one_crc_settle_once(); + test_a_full_pending_list_is_counted_separately(); + test_a_full_offer_queue_does_not_stop_capture_or_the_record(); + test_destruction_releases_pending_copies(); printf("cubinqueue_test: OK\n"); return 0; } diff --git a/shim/nvidia/cupti_adapter.cc b/shim/nvidia/cupti_adapter.cc index 69a8fef..22358f6 100644 --- a/shim/nvidia/cupti_adapter.cc +++ b/shim/nvidia/cupti_adapter.cc @@ -314,30 +314,41 @@ uint64_t cupti_cubin_crc(const void *bytes, size_t len) { p.size = CUpti_GetCubinCrcParamsSize; p.cubinSize = len; p.cubin = bytes; - // Takes the CUPTI guard, like every other call this adapter initiates. - // It runs on the APPLICATION's cuModuleLoad path, so it can now wait - // behind one of the timer thread's CUPTI calls -- see issue-99's report - // for the size of that and why it is the right trade. + // Takes the CUPTI guard, like every other call this adapter initiates, + // and for that reason runs ONLY on the timer thread, from + // CubinQueue::settle -- never in the MODULE_LOADED callback. + // + // It used to run in the callback, on the application's thread, on the + // reasoning that waiting behind one of the timer's CUPTI calls was a + // bounded cost. It was not bounded (issue #211): under lazy loading the + // callback fires from inside cuLaunchKernel, with the application thread + // holding a driver lock, while the timer held the guard across + // cuptiActivityFlushAll waiting for a CUPTI worker that needed that same + // driver lock. Three threads, a permanent deadlock of the profiled + // process, 1 run in 8 on a PyTorch workload with nobody attached. if (perfagent::cupti::GetCubinCrc(&p) != CUPTI_SUCCESS) return 0; return p.cubinCrc; } -// Fired by CubinQueue::capture, on the application's thread, at the one -// instant the copy is owned by nobody else: the CRC is computed and the entry -// has not been pushed, so the drain thread cannot yet have offered and freed -// it. That is what keeps bytes_ptr accurate. +// Fired by CubinQueue::settle, on the timer thread, at the one instant the +// copy is owned by nobody else: the CRC is computed and the entry has not +// been handed to drain(), so nothing can yet have offered and freed it. That +// is what keeps bytes_ptr accurate. module_id and load_ns are the callback's, +// carried through the pending list, so the record still says when the module +// loaded rather than when the timer got to it. // // bytes_ptr stays in the ABI and stays true, and it is still NOT a transport: // it points into this process's address space and the consumer would need // CAP_SYS_PTRACE to read it. The bytes travel over the cubin channel // (core/cubin.h); this record announces THAT a module loaded, with its CRC // and size. -void on_cubin_captured(void *ctx, uint64_t crc, const void *bytes, size_t len) { +void on_cubin_captured(void *, uint64_t module_id, uint64_t load_ns, uint64_t crc, + const void *bytes, size_t len) { gpu_module_load_v1 r{}; r.cubin_crc = crc; - r.module_id = (uint64_t)(uintptr_t)ctx; + r.module_id = module_id; r.size_bytes = (uint64_t)len; - r.load_ns = mono_ns(); + r.load_ns = load_ns; r.bytes_ptr = (uint64_t)(uintptr_t)bytes; // Live, or logged for replay -- never both, and that is what keeps a // launched run from receiving every module twice. The log holds exactly @@ -398,10 +409,10 @@ void on_module_loaded(const CUpti_ResourceData *rd) { // asked to be able to profile. if (!consumer_attached()) g_module_retained.fetch_add(1, std::memory_order_relaxed); const perfagent::CubinView view(m->pCubin, m->cubinSize); - // moduleId travels as the context, not as a captured pointer: nothing - // about this call may outlive the callback except the owned copy. - g_cubins->capture(view, cupti_cubin_crc, on_cubin_captured, - (void *)(uintptr_t)m->moduleId); + // The copy and nothing else. No CRC here: it is a guarded CUPTI call and + // this thread may be inside cuLaunchKernel holding a driver lock (issue + // #211). capture() takes no CRC function, so it cannot be put back. + g_cubins->capture(view, m->moduleId, mono_ns()); } // Defined further down with the rest of the process setup; declared here @@ -1959,6 +1970,7 @@ void report(const char *why) { "modules_captured=%llu module_reload_skipped=%llu module_retained_unattached=%llu " "module_no_bytes=%llu cubin_too_large=%llu cubin_crc_failed=%llu " "cubin_alloc_failed=%llu cubin_queue_full=%llu cubin_queue_depth=%zu " + "cubin_pending_full=%llu cubin_pending_depth=%zu " "cubins_sent=%llu cubin_send_failed=%llu " "cubins_offered=%llu cubin_transport_send_failed=%llu " "clock_steps=%llu clock_offset_ns=%lld sem_at_exit=%u\n", @@ -1992,6 +2004,11 @@ void report(const char *why) { (unsigned long long)(g_cubins ? g_cubins->cubin_alloc_failed() : 0), (unsigned long long)(g_cubins ? g_cubins->cubin_queue_full() : 0), g_cubins ? g_cubins->depth() : (size_t)0, + // Copied in a callback, not yet settled on the timer thread (issue + // #211). A non-zero pending_full is modules lost whole: no record, + // because their CRC was never computed. + (unsigned long long)(g_cubins ? g_cubins->cubin_pending_full() : 0), + g_cubins ? g_cubins->pending() : (size_t)0, (unsigned long long)(g_cubins ? g_cubins->cubins_sent() : 0), (unsigned long long)(g_cubins ? g_cubins->cubin_send_failed() : 0), (unsigned long long)perfagent::cubins_offered(), @@ -2238,6 +2255,17 @@ void on_tick() { } g_names_was_attached = now_attached; + // The SETTLE half of a cubin capture: the CRC, the de-duplication and + // gpu_module_load_v1, for every module copied since the last tick. It + // runs here, on the timer thread, attached or not -- an unattached + // process still needs its modules' records in the replay log -- because + // the CRC is a guarded CUPTI call and the MODULE_LOADED callback is the + // one place it must never be made (issue #211). It runs before the + // replay so a module copied in the tick a consumer attaches is announced + // exactly once: live if the consumer is already there, from the log if + // not. + if (g_cubins) g_cubins->settle(cupti_cubin_crc, on_cubin_captured, nullptr); + // The SEND half of a cubin capture, and the whole reason this is here and // not in the MODULE_LOADED callback: a connect() plus an up-to-8 MiB // handover on the application's cuModuleLoad path would stall the @@ -2365,8 +2393,11 @@ void at_exit_handler() { if (g_lb) g_lb->flush(); if (g_eb) g_eb->flush(); if (g_pcb) g_pcb->flush(); - // One last offer pass, for modules captured inside the final drain - // interval. Bounded: drain() offers at most the entries present when it + // One last settle and offer pass, for modules captured inside the final + // drain interval. The timer is joined above, so this thread is the only + // settler -- which settle() requires. + if (g_cubins) g_cubins->settle(cupti_cubin_crc, on_cubin_captured, nullptr); + // The offer pass. Bounded: drain() offers at most the entries present when it // was entered, the queue holds at most CubinQueueLimits::max_entries, and // an offer to an absent listener refuses immediately -- so this cannot // turn process exit into a multi-second wait. diff --git a/shim/stub/stub.cc b/shim/stub/stub.cc index cec84aa..418e90e 100644 --- a/shim/stub/stub.cc +++ b/shim/stub/stub.cc @@ -95,8 +95,8 @@ static uint64_t mono_ns() { // // A GPU-free MODULE_LOADED. It reads a checked-in cubin from disk, runs it // through the SAME CubinQueue the CUPTI adapter uses -- same capture, same -// CubinView, same crc-over-the-copy, same drain-thread offer -- and fires -// gpu_module_load_v1 from inside capture(). That drives the cubin transport +// CubinView, same crc-over-the-copy, same drain-thread settle and offer -- +// and fires gpu_module_load_v1 from inside settle(). That drives the cubin transport // (Task 3), the module store (Task 4), the consumer's decode (Task 7) and the // projection's source labels (Task 9) end to end on a machine with no GPU, // which is what the phase gate needs. @@ -133,17 +133,18 @@ static uint64_t stub_cubin_crc(const void *bytes, size_t len) { static unsigned long g_stub_module_seq = 0; -// Fired by CubinQueue::capture, before the copy is visible to any drain, so +// Fired by CubinQueue::settle, before the copy is visible to any drain, so // bytes_ptr names live adapter-owned bytes at the moment the probe reads it. // It is still not a transport: the consumer cannot read another process's // address space without CAP_SYS_PTRACE. The bytes go over the cubin channel. -static void stub_on_cubin_captured(void *ctx, uint64_t crc, const void *bytes, size_t len) { +static void stub_on_cubin_captured(void *, uint64_t module_id, uint64_t load_ns, uint64_t crc, + const void *bytes, size_t len) { if (!gpu_module_load_v1_enabled()) return; gpu_module_load_v1 r{}; r.cubin_crc = crc; - r.module_id = (uint64_t)(uintptr_t)ctx; + r.module_id = module_id; r.size_bytes = (uint64_t)len; - r.load_ns = mono_ns(); + r.load_ns = load_ns; r.bytes_ptr = (uint64_t)(uintptr_t)bytes; gpu_module_load_v1_emit(&r, 1, g_stub_module_seq++); } @@ -186,8 +187,7 @@ static unsigned stub_capture_modules(perfagent::CubinQueue &q) { } const uint64_t crc = stub_cubin_crc(bytes.data(), bytes.size()); const perfagent::CubinView view(bytes.data(), bytes.size()); - const bool ok = q.capture(view, stub_cubin_crc, stub_on_cubin_captured, - (void *)(uintptr_t)(n + 1)); + const bool ok = q.capture(view, n + 1, mono_ns()); fprintf(stderr, "stub: module id=%u path=%s size=%zu crc=0x%016llx captured=%s\n", n + 1, path.c_str(), bytes.size(), (unsigned long long)crc, ok ? "yes" : "no"); @@ -281,7 +281,7 @@ perfagent_stub_run(unsigned launches, unsigned period_us, unsigned sample_period uint64_t first_exec_ns = 0, last_exec_ns = 0; // The same queue the CUPTI adapter runs, wired the same way: capture on - // the caller's thread, offer on the drain thread. + // the caller's thread, settle and offer on the drain thread. perfagent::CubinQueue cubins; const unsigned cubin_timeout_ms = perfagent::cubin_timeout_ms(2000); @@ -311,9 +311,11 @@ perfagent_stub_run(unsigned launches, unsigned period_us, unsigned sample_period drainer.on_tick([&] { lb.flush(); eb.flush(); - // The offer, on the drain thread and nowhere else. In an application - // this is what keeps a connect() and a multi-megabyte handover off - // the cuModuleLoad path; here it is what proves that wiring works. + // The CRC and the module record, then the offer, on the drain thread + // and nowhere else. In an application this is what keeps a guarded + // CUPTI call (issue #211), a connect() and a multi-megabyte handover + // off the cuModuleLoad path; here it is what proves that wiring works. + cubins.settle(stub_cubin_crc, stub_on_cubin_captured, nullptr); cubins.drain(perfagent::cubin_offer_to_consumer, cubin_timeout_ms); // Late-attach replay: only on the unattached -> attached transition, // so an already-attached consumer does not get names re-sent every @@ -334,13 +336,13 @@ perfagent_stub_run(unsigned launches, unsigned period_us, unsigned sample_period drainer.start(100); // Modules load before kernels run, as they do in a CUDA process. The - // capture is synchronous; the offers are not, so wait -- bounded -- for + // copy is synchronous; the settle and the offer are not, so wait -- bounded -- for // the drain thread to have emptied the queue before the launches start, // so a gate can assert on a module the consumer provably already holds. const unsigned modules = stub_capture_modules(cubins); if (modules) { const uint64_t wait_until = mono_ns() + 5000000000ULL; - while (cubins.depth() && mono_ns() < wait_until) { + while ((cubins.pending() || cubins.depth()) && mono_ns() < wait_until) { std::this_thread::sleep_for(std::chrono::milliseconds(5)); } }