plugin: set tunerPluginLoaded on the already-loaded path so refcount balances

[EylonKrause]

Description

ncclTunerPluginLoad has two success paths. The first-load path sets comm->tuner, increments tunerPluginRefCount, and sets comm->tunerPluginLoaded = 1. The "already loaded" fast path — taken by every subsequent communicator — sets comm->tuner and increments tunerPluginRefCount but does not set comm->tunerPluginLoaded. ncclTunerPluginUnload gates the decrement on that per-comm flag:

if (comm->tunerPluginLoaded && 0 == (--tunerPluginRefCount)) { ... ncclClosePluginLib(...); }

Because && short-circuits, fast-path comms (flag == 0) never execute --tunerPluginRefCount. With N communicators sharing one tuner, the refcount is incremented N times but decremented at most once, so it never reaches 0: the plugin is never dlclosed (it stays loaded for the life of the process) and the accounting is permanently skewed.

Related Issues

None.

Changes & Impact

• src/plugin/tuner.cc (ncclTunerPluginLoad): set comm->tunerPluginLoaded = 1 on the already-loaded fast path, matching the first-load path, so every comm that increments the refcount can also decrement it in ncclTunerPluginUnload. The single-comm path is unaffected (it already set the flag). Non-breaking.

Performance Impact

None.

Testing

• Builds clean with make src.build; all_reduce_perf regression: Out of bounds values : 0 OK.
• The asymmetry is evident between the two success paths and the unload predicate; with ≥2 communicators sharing one tuner, tunerPluginRefCount reaches 0 (and the plugin is closed) only after this fix.

[EylonKrause]

Disclosure: this contribution was authored with an AI coding assistant (Claude) and reviewed before submission.