Skip to content

Handle profiler proxy initialization failures - #2326

Open
pawelchcki wants to merge 1 commit into
NVIDIA:masterfrom
pawelchcki:repro_for_nccp_shm_bug
Open

Handle profiler proxy initialization failures#2326
pawelchcki wants to merge 1 commit into
NVIDIA:masterfrom
pawelchcki:repro_for_nccp_shm_bug

Conversation

@pawelchcki

@pawelchcki pawelchcki commented Aug 6, 2026

Copy link
Copy Markdown

Summary

Do not enable profiler proxy events when proxy initialization fails. This prevents SaveProxyProfiler() from dereferencing a null proxy pool after /dev/shm allocation failure.

Validation

  • Profiler v6 requesting ncclProfileColl | ncclProfileKernelCh, two T4 ranks.
  • 64 MiB /dev/shm: previously segfaulted in ncclLocalOpAppend; now logs allocation warnings and exits 0.
  • 1 GiB /dev/shm: exits 0.
Minimal reproducer

Save as repro.cc:

#ifdef PLUGIN
#include "nccl.h"
#include "nccl_common.h"
#include "plugin/nccl_profiler.h"
static ncclResult_t init(void** c, uint64_t, int* m, const char*, int, int, int,
                         ncclDebugLogger_t) {
  static int context; *c = &context;
  *m = ncclProfileColl | ncclProfileKernelCh;
  return ncclSuccess;
}
static ncclResult_t start(void* c, void** e, ncclProfilerEventDescr_t*) {
  *e = c; return ncclSuccess;
}
static ncclResult_t stop(void*) { return ncclSuccess; }
static ncclResult_t record(void*, ncclProfilerEventState_t,
                           ncclProfilerEventStateArgs_t*) { return ncclSuccess; }
extern "C" __attribute__((visibility("default"))) ncclProfiler_t ncclProfiler_v6 =
    {"shm-repro", init, start, stop, record, stop};
#else
#include <cuda_runtime.h>
#include <nccl.h>
#define OK(x) do { if ((x) != 0) return 1; } while (0)
int main() {
  int devs[] = {0, 1}; ncclComm_t comms[2]; float* ptrs[2];
  OK(ncclCommInitAll(comms, 2, devs));
  for (int i = 0; i < 2; ++i) {
    OK(cudaSetDevice(i)); OK(cudaMalloc(&ptrs[i], sizeof(float)));
  }
  OK(ncclGroupStart());
  for (int i = 0; i < 2; ++i) {
    OK(cudaSetDevice(i));
    OK(ncclAllReduce(ptrs[i], ptrs[i], 1, ncclFloat, ncclSum, comms[i], 0));
  }
  OK(ncclGroupEnd());
  for (int i = 0; i < 2; ++i) {
    OK(cudaSetDevice(i)); OK(cudaDeviceSynchronize());
  }
}
#endif

Build from the NCCL source root and run in a two-GPU container with 64 MiB /dev/shm:

g++ -DPLUGIN -DNCCL_OS_LINUX=1 -Isrc/include -Ibuild/include -I/usr/local/cuda/include -shared -fPIC repro.cc -o profiler.so
g++ -Ibuild/include -I/usr/local/cuda/include repro.cc -Lbuild/lib -L/usr/local/cuda/lib64 -lnccl -lcudart -o repro
NCCL_PROFILER_PLUGIN=$PWD/profiler.so ./repro

@pawelchcki pawelchcki changed the title [codex] Handle profiler proxy initialization failures Handle profiler proxy initialization failures Aug 6, 2026
@pawelchcki
pawelchcki marked this pull request as ready for review August 6, 2026 20:31
@xiaofanl-nvidia

Copy link
Copy Markdown
Collaborator

@pawelchcki please help add DCO signoff as per automatic checks. Thanks!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants