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)); } }