diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index bdfaedd0927..9f33a036695 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -151,6 +151,11 @@ rccl_assert_source_line("${CMAKE_CURRENT_LIST_DIR}/fakes/tuning_fakes.cc" "int64_t g_paramMaxNchannels = -2;" "the MaxNchannels fake default") rccl_assert_source_line("${CMAKE_CURRENT_LIST_DIR}/fakes/wrap_fakes.cc" "static int64_t DefaultParamForceCe() { return 1; }" "the ForceCe fake default") +rccl_assert_source_line("${_rccl_src_root}/init.cc" + "NCCL_PARAM(SetCpuStackSize, \"SET_CPU_STACK_SIZE\", 1)" "fakes/linux_fakes.cc: ncclParamSetCpuStackSize()") +rccl_assert_source_line("${CMAKE_CURRENT_LIST_DIR}/fakes/linux_fakes.cc" + "int64_t ncclParamSetCpuStackSize() { return g_loadParam(\"SET_CPU_STACK_SIZE\", 1); }" + "the SetCpuStackSize fake default") # RasDiagnosticsTests.cpp (rccl-UnitTests) classifies report lines by kRasFailureMarkers and kRasHealthMarkers; # each marker must stay both in those lists and in every source file that emits it (comma-separated). set(_rccl_ras_tests "${CMAKE_CURRENT_LIST_DIR}/../RasDiagnosticsTests.cpp") @@ -733,6 +738,26 @@ function(rccl_define_micro_source_lists) ../common/main_altrsmi.cpp PARENT_SCOPE) + # rccl-UnitTestsMicroLinux -- src/os/linux*.cc, misc/{socket,shmutils,signals}.cc. Own binary: + # fakes/os_fakes.cc, which most targets link, defines ncclOs* symbols that os/linux.cc defines. + set(TEST_MICRO_LINUX_SOURCE_FILES + linux-test.cc + linux-ipcsocket-test.cc + linux-socket-pair-test.cc + socket-test.cc + shmutils-test.cc + signals-test.cc + fakes/dl_fakes.cc + fakes/env_fakes.cc + fakes/hip_fakes.cc + fakes/libc_fakes.cc + fakes/linux_fakes.cc + fakes/nccl_fakes.cc + ${RCCL_HOST_HIPIFY_DIR}/src/misc/crypt.cc + ${RCCL_HOST_HIPIFY_DIR}/src/misc/utils.cc + ../common/main_altrsmi.cpp + PARENT_SCOPE) + # rccl-UnitTestsMicroRmaProxy -- REAL src/rma/rma_proxy.cc. set(TEST_MICRO_RMA_PROXY_SOURCE_FILES rma-proxy-test.cc @@ -2079,6 +2104,23 @@ foreach(_variant "" "-uncached" "-faultinj") endif() endforeach() +# --------------------------------------------------------------------------- +# rccl-UnitTestsMicroLinux +# --------------------------------------------------------------------------- +rccl_mark_oracle_tus_generated(${RCCL_HOST_HIPIFY_DIR} crypt) +rccl_add_micro_binary(rccl-UnitTestsMicroLinux + SOURCES ${TEST_MICRO_LINUX_SOURCE_FILES} + EXTRA_DEFS + ${_MICRO_MODE_DEFS} + RCCL_NCCL_FAKES_OMIT_IPC_FD_CLOSE + LINUX_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/os/linux.cc" + LINUX_IPCSOCKET_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/os/linux_ipcsocket.cc" + LINUX_SOCKET_PAIR_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/os/linux_socket_pair.cc" + SOCKET_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/socket.cc" + SHMUTILS_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/shmutils.cc" + SIGNALS_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/misc/signals.cc" + CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_linux.yaml") + # --------------------------------------------------------------------------- # rccl-UnitTestsMicroConnect -- ncclTopoPreset / ncclTopoPostset unstubbed. # --------------------------------------------------------------------------- diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index b9750ffc546..634837f8be5 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -297,6 +297,14 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: `argcheck-test.cc`); suite `Argcheck*`. Its own binary because the test seams HIP, `ncclDevrFindWindow` and `bootstrapAllGather` around a unit other binaries link whole. See `test_categories_micro_argcheck.yaml`. +- **`rccl-UnitTestsMicroLinux`**: `os/linux.cc` (`LINUX_CC_PATH`, from `linux-test.cc`), + plus `os/linux_ipcsocket.cc`, `os/linux_socket_pair.cc`, `misc/socket.cc`, + `misc/shmutils.cc` and `misc/signals.cc`, one test TU each; suites `Linux*`, + `Socket*`, `Shmutils*`, `Signals*`. Its own binary: `fakes/os_fakes.cc`, which most + targets link, defines `ncclOsInitialize`, `ncclOsTopoGetStrFromSys` and other `ncclOs*` + symbols `os/linux.cc` defines. The units run on real sockets and files through + `fakes/libc_seam.h` with `UseRealLibc()`. See `test_categories_micro_linux.yaml`. + - **`rccl-UnitTestsMicroCpuset`** — header-only `include/cpuset.h` (from `cpuset-test.cc`); suite `Cpuset*`. Its own binary because the test defines `os/linux.cc`'s `ncclOsCpuZero` / `ncclOsCpuSet` / `ncclOsCpuIsSet` one-liners, which would @@ -912,8 +920,9 @@ symbol. | `src/transport/p2p.cc` registration entry points, macro-renamed in one unit via `fakes/transport_p2p_reg_seam.h` (`rccl-UnitTestsMicro`, which compiles the real file) | `fakes/transport_p2p_reg_fakes.cc` | | selected entry points from `src/{transport,proxy}.cc`, `src/transport/{generic,coll_net,net,nvls}.cc`, `src/{graph/paths,tuning/tuning,plugin/net}.cc` (targets that do not compile the real files) | `fakes/transport_stubs.cc` | | libc (`gethostname`, `dladdr`), process-wide for units linking libc directly | `fakes/libc_interposers.cc` | -| libc macro-renamed in one unit via `fakes/libc_seam.h` (sockets/stdio/exit for `ras/client.cc` and `ras/client_support.cc`; `gethostname`/`access` for `diagnostics/ib_write_bw.cc`; heap for `graph/rccl_graph_gen.cc`; files, `ioctl`, `atexit`, clock, aligned allocation, thread start and sleep for `transport/net_telemetry.cc`; `opendir` for `misc/gdr_peermem.cc`; `open`/`read`/`close` for `transport/net_socket.cc`; `clock_gettime` for `transport/net_ib/wqe_lat_mon.cc` and `transport/net_ib_cast/scheduler.cc`) | `fakes/libc_fakes.cc` | -| libdl macro-renamed in one unit via `fakes/dl_seam.h` (the `misc/*symbols.cc` builders) | `fakes/dl_fakes.cc` | +| libc macro-renamed in one unit via `fakes/libc_seam.h` (sockets/stdio/exit for `ras/client.cc` and `ras/client_support.cc`; `gethostname`/`access` for `diagnostics/ib_write_bw.cc`; heap for `graph/rccl_graph_gen.cc`; files, `ioctl`, `atexit`, clock, aligned allocation, thread start and sleep for `transport/net_telemetry.cc`; `opendir` for `misc/gdr_peermem.cc`; `open`/`read`/`close` for `transport/net_socket.cc`; `clock_gettime` for `transport/net_ib/wqe_lat_mon.cc` and `transport/net_ib_cast/scheduler.cc`; the OS surface of `os/linux*.cc` and `misc/{socket,signals}.cc`) | `fakes/libc_fakes.cc` | +| libdl macro-renamed in one unit via `fakes/dl_seam.h` (the `misc/*symbols.cc` builders and `os/linux.cc`) | `fakes/dl_fakes.cc` | +| `src/init.cc`'s `SetCpuStackSize` param (`rccl-UnitTestsMicroLinux`, which cannot link `init_fakes.cc`) | `fakes/linux_fakes.cc` | | reusable RAS poll-entry, message, and diagnostics-context seams | `fakes/ras_fakes.cc` | | core/lifecycle floor + data symbols | `fakes/nccl_stubs.cc` | | reusable `nccl*` seams | `fakes/nccl_fakes.cc` | @@ -1159,7 +1168,8 @@ triage it into the right bucket: `fakes/utils_fakes.cc`, since `src/misc/utils.cc` owns them). Only when the symbol has no owning TU does it belong in the test itself. - **It's a logging or env-param helper** → already covered by - `ncclDebugLog` (prints to stderr) / `ncclLoadParam` (parses the value + `ncclDebugLog` (prints to stderr; a test that asserts on what was logged + hooks `g_ncclDebugLog`) / `ncclLoadParam` (parses the value `SetMicroEnv` scripted for that name, and the param's compile-time default when the name is unmapped). If a new logging primitive appears, follow the same pattern. @@ -1461,6 +1471,7 @@ cmake --build build -j"$(nproc)" ./build/rccl-UnitTestsMicroSymKernels # sym_kernels.cc tests ./build/rccl-UnitTestsMicroGinHostProxy # src/gin/gin_host_proxy.cc + proxy_gpucontext/ ./build/rccl-UnitTestsMicroTaskPrep # src/enqueue/task_prep/ + task_sched/ tests +./build/rccl-UnitTestsMicroLinux # src/os/linux*.cc, misc/{socket,shmutils,signals}.cc ./build/rccl-UnitTestsMicroConnect # src/graph/connect.cc Preset/Postset tests ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics.cc + src/diagnostics/ tests ./build/rccl-UnitTestsMicroGraph # src/graph/{xml,topo,paths,search,rome_models}.cc tests diff --git a/projects/rccl/test/host/fakes/hip_fakes.cc b/projects/rccl/test/host/fakes/hip_fakes.cc index 87df182f411..ae6df8c6c14 100644 --- a/projects/rccl/test/host/fakes/hip_fakes.cc +++ b/projects/rccl/test/host/fakes/hip_fakes.cc @@ -83,6 +83,10 @@ ASSERT_HOOK_MATCHES_PROD(g_hipFuncSetAttribute, hipFuncSetAttribute); RCCL_HIP_PROFILE_INTERCEPTORS_HOOKED(RCCL_ASSERT_INTERCEPTOR_HOOK) #undef RCCL_ASSERT_INTERCEPTOR_HOOK +ASSERT_HOOK_MATCHES_PROD(g_hipHostRegister, hipHostRegister); +ASSERT_HOOK_MATCHES_PROD(g_hipHostGetDevicePointer, hipHostGetDevicePointer); +ASSERT_HOOK_MATCHES_PROD(g_hipHostUnregister, hipHostUnregister); + #undef ASSERT_HOOK_MATCHES_PROD // =========================================================================== diff --git a/projects/rccl/test/host/fakes/libc_fakes.cc b/projects/rccl/test/host/fakes/libc_fakes.cc index 69f731fd09e..2fcd375d16a 100644 --- a/projects/rccl/test/host/fakes/libc_fakes.cc +++ b/projects/rccl/test/host/fakes/libc_fakes.cc @@ -15,6 +15,7 @@ #include #include #include +#include #include #include @@ -374,9 +375,53 @@ std::function g DefaultPthreadCreate; std::function g_nanosleep = DefaultNanosleep; +static int RealFcntl(int fd, int cmd, int arg) { return ::fcntl(fd, cmd, arg); } +static void RealExit(int status) { ::exit(status); } + +std::function g_poll = ::poll; +std::function g_getsockopt = ::getsockopt; +std::function g_getsockname = ::getsockname; +std::function g_shutdown = ::shutdown; +std::function g_recvmsg = ::recvmsg; +std::function g_sendmsg = ::sendmsg; +std::function g_pipe = ::pipe; +std::function g_getifaddrs = ::getifaddrs; +std::function g_freeifaddrs = ::freeifaddrs; +std::function g_getrlimit = ::getrlimit; +std::function g_setrlimit = ::setrlimit; +std::function g_schedGetaffinity = ::sched_getaffinity; +std::function g_realpath = ::realpath; +std::function g_mkstemp = ::mkstemp; +std::function g_fallocate = ::fallocate; +std::function g_mmap = ::mmap; +std::function g_munmap = ::munmap; +std::function g_unlink = ::unlink; +std::function g_signal = ::signal; + ASSERT_HOOK_MATCHES_PROD(g_accept, ::accept); ASSERT_HOOK_MATCHES_PROD(g_recv, ::recv); ASSERT_HOOK_MATCHES_PROD(g_send, ::send); +ASSERT_HOOK_MATCHES_PROD(g_poll, ::poll); +ASSERT_HOOK_MATCHES_PROD(g_getsockopt, ::getsockopt); +ASSERT_HOOK_MATCHES_PROD(g_getsockname, ::getsockname); +ASSERT_HOOK_MATCHES_PROD(g_shutdown, ::shutdown); +ASSERT_HOOK_MATCHES_PROD(g_recvmsg, ::recvmsg); +ASSERT_HOOK_MATCHES_PROD(g_sendmsg, ::sendmsg); +ASSERT_HOOK_MATCHES_PROD(g_pipe, ::pipe); +ASSERT_HOOK_MATCHES_PROD(g_getifaddrs, ::getifaddrs); +ASSERT_HOOK_MATCHES_PROD(g_freeifaddrs, ::freeifaddrs); +ASSERT_HOOK_MATCHES_PROD(g_getrlimit, ::getrlimit); +ASSERT_HOOK_MATCHES_PROD(g_setrlimit, ::setrlimit); +ASSERT_HOOK_MATCHES_PROD(g_schedGetaffinity, ::sched_getaffinity); +ASSERT_HOOK_MATCHES_PROD(g_realpath, ::realpath); +ASSERT_HOOK_MATCHES_PROD(g_fopen, ::fopen); +ASSERT_HOOK_MATCHES_PROD(g_opendir, ::opendir); +ASSERT_HOOK_MATCHES_PROD(g_mkstemp, ::mkstemp); +ASSERT_HOOK_MATCHES_PROD(g_fallocate, ::fallocate); +ASSERT_HOOK_MATCHES_PROD(g_mmap, ::mmap); +ASSERT_HOOK_MATCHES_PROD(g_munmap, ::munmap); +ASSERT_HOOK_MATCHES_PROD(g_unlink, ::unlink); +ASSERT_HOOK_MATCHES_PROD(g_signal, ::signal); #undef ASSERT_HOOK_MATCHES_PROD @@ -436,6 +481,25 @@ void ResetLibcFakes() { g_posixMemalign = DefaultPosixMemalign; g_pthreadCreate = DefaultPthreadCreate; g_nanosleep = DefaultNanosleep; + g_poll = ::poll; + g_getsockopt = ::getsockopt; + g_getsockname = ::getsockname; + g_shutdown = ::shutdown; + g_recvmsg = ::recvmsg; + g_sendmsg = ::sendmsg; + g_pipe = ::pipe; + g_getifaddrs = ::getifaddrs; + g_freeifaddrs = ::freeifaddrs; + g_getrlimit = ::getrlimit; + g_setrlimit = ::setrlimit; + g_schedGetaffinity = ::sched_getaffinity; + g_realpath = ::realpath; + g_mkstemp = ::mkstemp; + g_fallocate = ::fallocate; + g_mmap = ::mmap; + g_munmap = ::munmap; + g_unlink = ::unlink; + g_signal = ::signal; g_writtenData.clear(); g_stdoutData.clear(); @@ -483,6 +547,34 @@ void ResetLibcFakes() { errno = 0; } +void UseRealLibc() { + ResetLibcFakes(); + g_write = ::write; + g_read = ::read; + g_close = ::close; + g_socket = ::socket; + g_connect = ::connect; + g_bind = ::bind; + g_listen = ::listen; + g_accept = ::accept; + g_fcntl = RealFcntl; + g_recv = ::recv; + g_send = ::send; + g_setsockopt = ::setsockopt; + g_getaddrinfo = ::getaddrinfo; + g_freeaddrinfo = ::freeaddrinfo; + g_getnameinfo = ::getnameinfo; + g_fwrite = ::fwrite; + g_fflush = ::fflush; + g_perror = ::perror; + g_exit = RealExit; + g_free = ::free; + g_open = [](const char* path, int flags, mode_t mode) { return ::open(path, flags, mode); }; + g_fopen = ::fopen; + g_opendir = ::opendir; + g_readlink = ::readlink; +} + // --------------------------------------------------------------------------- // Trampolines. fakes/libc_seam.h macro-renames each libc call in the unit // under test to the matching micro_* name; these dispatch through the slots @@ -561,6 +653,29 @@ int micro_pthread_create(pthread_t* thread, const pthread_attr_t* attr, void* (* return g_pthreadCreate(thread, attr, fn, arg); } int micro_nanosleep(const struct timespec* req, struct timespec* rem) { return g_nanosleep(req, rem); } +int micro_poll(struct pollfd* fds, nfds_t nfds, int timeout) { return g_poll(fds, nfds, timeout); } +int micro_getsockopt(int fd, int level, int optname, void* optval, socklen_t* optlen) { + return g_getsockopt(fd, level, optname, optval, optlen); +} +int micro_getsockname(int fd, struct sockaddr* addr, socklen_t* len) { return g_getsockname(fd, addr, len); } +int micro_shutdown(int fd, int how) { return g_shutdown(fd, how); } +ssize_t micro_recvmsg(int fd, struct msghdr* msg, int flags) { return g_recvmsg(fd, msg, flags); } +ssize_t micro_sendmsg(int fd, const struct msghdr* msg, int flags) { return g_sendmsg(fd, msg, flags); } +int micro_pipe(int* fds) { return g_pipe(fds); } +int micro_getifaddrs(struct ifaddrs** ifap) { return g_getifaddrs(ifap); } +void micro_freeifaddrs(struct ifaddrs* ifa) { g_freeifaddrs(ifa); } +int micro_getrlimit(int resource, struct rlimit* lim) { return g_getrlimit(resource, lim); } +int micro_setrlimit(int resource, const struct rlimit* lim) { return g_setrlimit(resource, lim); } +int micro_sched_getaffinity(pid_t pid, size_t size, cpu_set_t* set) { return g_schedGetaffinity(pid, size, set); } +char* micro_realpath(const char* path, char* resolved) { return g_realpath(path, resolved); } +int micro_mkstemp(char* tmpl) { return g_mkstemp(tmpl); } +int micro_fallocate(int fd, int mode, off_t offset, off_t len) { return g_fallocate(fd, mode, offset, len); } +void* micro_mmap(void* addr, size_t len, int prot, int flags, int fd, off_t offset) { + return g_mmap(addr, len, prot, flags, fd, offset); +} +int micro_munmap(void* addr, size_t len) { return g_munmap(addr, len); } +int micro_unlink(const char* path) { return g_unlink(path); } +sighandler_t micro_signal(int sig, sighandler_t handler) { return g_signal(sig, handler); } // Always forwards to the real vfprintf so stderr output still happens; records only the FILE* and never the // formatted text, since no test needs the diagnostic's wording, only whether and where one was printed. diff --git a/projects/rccl/test/host/fakes/libc_fakes.h b/projects/rccl/test/host/fakes/libc_fakes.h index ffdd36e9b35..fe7d9e882f4 100644 --- a/projects/rccl/test/host/fakes/libc_fakes.h +++ b/projects/rccl/test/host/fakes/libc_fakes.h @@ -15,11 +15,17 @@ // restores every default and clears every record. // // Add a symbol here when a unit under test reaches it -- with a working -// default and a reset, never as a hardcoded always-succeed. poll remains -// deliberately absent until a caller needs a shared recording surface. +// default and a reset, never as a hardcoded always-succeed. poll and the other +// OS slots call straight through and record nothing; add a record when a +// caller needs one. #include +#include #include +#include +#include +#include +#include #include #include #include @@ -121,6 +127,28 @@ extern std::function g_posixMemalign; extern std::function g_pthreadCreate; extern std::function g_nanosleep; +// The OS surface os/linux*.cc and misc/{socket,signals}.cc reach. Each defaults to the real call +// and records nothing; a test makes one fail or vary with ScopedHook. +extern std::function g_poll; +extern std::function g_getsockopt; +extern std::function g_getsockname; +extern std::function g_shutdown; +extern std::function g_recvmsg; +extern std::function g_sendmsg; +extern std::function g_pipe; +extern std::function g_getifaddrs; +extern std::function g_freeifaddrs; +extern std::function g_getrlimit; +extern std::function g_setrlimit; +extern std::function g_schedGetaffinity; +extern std::function g_realpath; +extern std::function g_mkstemp; +extern std::function g_fallocate; +extern std::function g_mmap; +extern std::function g_munmap; +extern std::function g_unlink; +extern std::function g_signal; + // --------------------------------------------------------------------------- // Observation points fed by the default seams. A test that installs its own // hook over a seam stops feeding the corresponding record. @@ -192,4 +220,8 @@ void ScriptRecvData(std::string data); // Restores every seam to its default and clears every record above. void ResetLibcFakes(); +// ResetLibcFakes, then points every seam above at the real call, recording nothing: for units tested +// on real sockets and files, whose production threads would race on the records. +void UseRealLibc(); + #endif // RCCL_TEST_HOST_FAKES_LIBC_FAKES_H_ diff --git a/projects/rccl/test/host/fakes/libc_seam.h b/projects/rccl/test/host/fakes/libc_seam.h index b1103efb18a..35e64a4de61 100644 --- a/projects/rccl/test/host/fakes/libc_seam.h +++ b/projects/rccl/test/host/fakes/libc_seam.h @@ -42,6 +42,12 @@ #include #include #include +#include +#include +#include +#include +#include +#include #include #include #include @@ -88,6 +94,25 @@ int micro_clock_gettime(clockid_t, struct timespec*); int micro_posix_memalign(void**, size_t, size_t); int micro_pthread_create(pthread_t*, const pthread_attr_t*, void* (*)(void*), void*); int micro_nanosleep(const struct timespec*, struct timespec*); +int micro_poll(struct pollfd*, nfds_t, int); +int micro_getsockopt(int, int, int, void*, socklen_t*); +int micro_getsockname(int, struct sockaddr*, socklen_t*); +int micro_shutdown(int, int); +ssize_t micro_recvmsg(int, struct msghdr*, int); +ssize_t micro_sendmsg(int, const struct msghdr*, int); +int micro_pipe(int*); +int micro_getifaddrs(struct ifaddrs**); +void micro_freeifaddrs(struct ifaddrs*); +int micro_getrlimit(int, struct rlimit*); +int micro_setrlimit(int, const struct rlimit*); +int micro_sched_getaffinity(pid_t, size_t, cpu_set_t*); +char* micro_realpath(const char*, char*); +int micro_mkstemp(char*); +int micro_fallocate(int, int, off_t, off_t); +void* micro_mmap(void*, size_t, int, int, int, off_t); +int micro_munmap(void*, size_t); +int micro_unlink(const char*); +sighandler_t micro_signal(int, sighandler_t); } // extern "C" #define write micro_write @@ -128,3 +153,22 @@ int micro_nanosleep(const struct timespec*, struct timespec*); #define posix_memalign micro_posix_memalign #define pthread_create micro_pthread_create #define nanosleep micro_nanosleep +#define poll micro_poll +#define getsockopt micro_getsockopt +#define getsockname micro_getsockname +#define shutdown micro_shutdown +#define recvmsg micro_recvmsg +#define sendmsg micro_sendmsg +#define pipe micro_pipe +#define getifaddrs micro_getifaddrs +#define freeifaddrs micro_freeifaddrs +#define getrlimit micro_getrlimit +#define setrlimit micro_setrlimit +#define sched_getaffinity micro_sched_getaffinity +#define realpath micro_realpath +#define mkstemp micro_mkstemp +#define fallocate micro_fallocate +#define mmap micro_mmap +#define munmap micro_munmap +#define unlink micro_unlink +#define signal micro_signal diff --git a/projects/rccl/test/host/fakes/libc_seam_undef.h b/projects/rccl/test/host/fakes/libc_seam_undef.h index 56f5aee166e..bd5e0de0f9a 100644 --- a/projects/rccl/test/host/fakes/libc_seam_undef.h +++ b/projects/rccl/test/host/fakes/libc_seam_undef.h @@ -22,6 +22,25 @@ #undef opendir #undef fopen #undef open +#undef signal +#undef unlink +#undef munmap +#undef mmap +#undef fallocate +#undef mkstemp +#undef realpath +#undef sched_getaffinity +#undef setrlimit +#undef getrlimit +#undef freeifaddrs +#undef getifaddrs +#undef pipe +#undef sendmsg +#undef recvmsg +#undef shutdown +#undef getsockname +#undef getsockopt +#undef poll #undef access #undef gethostname #undef free diff --git a/projects/rccl/test/host/fakes/linux_fakes.cc b/projects/rccl/test/host/fakes/linux_fakes.cc new file mode 100644 index 00000000000..99c80c92110 --- /dev/null +++ b/projects/rccl/test/host/fakes/linux_fakes.cc @@ -0,0 +1,15 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// The one non-libc, non-HIP symbol rccl-UnitTestsMicroLinux needs that no shared fakes file on its +// link line provides: src/init.cc's SetCpuStackSize param, which os/linux.cc reads. init_fakes.cc +// cannot be linked there. Defined with init.cc's real int64_t type, not os/linux.cc's int declaration. + +#include + +#include "nccl_fakes.h" + +int64_t ncclParamSetCpuStackSize() { return g_loadParam("SET_CPU_STACK_SIZE", 1); } diff --git a/projects/rccl/test/host/fakes/nccl_fakes.cc b/projects/rccl/test/host/fakes/nccl_fakes.cc index 7cf78b2952f..ca05d3d2e1d 100644 --- a/projects/rccl/test/host/fakes/nccl_fakes.cc +++ b/projects/rccl/test/host/fakes/nccl_fakes.cc @@ -28,9 +28,11 @@ #include "comm.h" // ncclCommGraphRegister / Deregister #include "strongstream.h" // ncclStrongStream* #include "mem_manager.h" // ncclMemTrack / ncclMemUntrack / ncclDynMemMarkExportToPeer -#include "ipcsocket.h" // ncclIpcFd +#include "ipcsocket.h" +#include #include +#include #include "nccl_fakes.h" // controllable seam hooks @@ -78,19 +80,32 @@ uint32_t ncclDebugLevelMask = 0; // NCCL_LOG_NONE uint64_t ncclDebugMask = 0; thread_local int ncclDebugNoWarn = 0; -void ncclDebugLog(ncclDebugLogLevel /*level*/, - unsigned long /*flags*/, +static void DefaultNcclDebugLog(ncclDebugLogLevel, unsigned long, const char* filefunc, int line, const char* msg) +{ + std::fprintf(stderr, "[fake] %s:%d %s\n", filefunc ? filefunc : "?", line, msg); +} + +std::function g_ncclDebugLog = + DefaultNcclDebugLog; + +void ncclDebugLog(ncclDebugLogLevel level, + unsigned long flags, const char* filefunc, int line, const char* fmt, ...) { - std::fprintf(stderr, "[fake] %s:%d ", filefunc ? filefunc : "?", line); + // debug.cc demotes a WARN raised under NOWARN to INFO. + if (ncclDebugNoWarn != 0 && level == NCCL_LOG_WARN) level = NCCL_LOG_INFO; std::va_list ap; va_start(ap, fmt); - std::vfprintf(stderr, fmt, ap); + std::va_list again; + va_copy(again, ap); + std::string msg(std::max(0, std::vsnprintf(nullptr, 0, fmt, ap)), '\0'); + std::vsnprintf(&msg[0], msg.size() + 1, fmt, again); + va_end(again); va_end(ap); - std::fputc('\n', stderr); + g_ncclDebugLog(level, flags, filefunc, line, msg.c_str()); } // src/debug.cc, same owner as the logging symbols above. A no-op rather than a @@ -219,10 +234,9 @@ ncclResult_t ncclProxyClientQueryFdBlocking(struct ncclComm* comm, return g_proxyClientQueryFdBlocking(comm, proxyConn, localFd, rmtFd); } -// src/os/linux_ipcsocket.cc (NCCL 2.32): p2p.cc and dev_runtime.cc close -// imported fds through this instead of close(2). A real close, as production -// does, so suites that hand back a dup'd fd can still prove it was closed. +#ifndef RCCL_NCCL_FAKES_OMIT_IPC_FD_CLOSE int ncclIpcFdClose(ncclIpcFd fd) { return ::close(fd); } +#endif // --- Controllable seam: ncclRegLocalIsValid ----------------------------- // Default preserves the old stub: report the record as not locally valid @@ -626,6 +640,7 @@ int64_t ncclParamMultiSegmentRegister() { return g_loadParam("MULTI_SEGMENT_REGI void ResetNcclFakes() { + g_ncclDebugLog = DefaultNcclDebugLog; g_ncclProxyClientGetFdBlocking = DefaultNcclProxyClientGetFdBlocking; g_strongStreamAcquire = DefaultStrongStreamAcquire; g_strongStreamRelease = DefaultStrongStreamRelease; diff --git a/projects/rccl/test/host/fakes/nccl_fakes.h b/projects/rccl/test/host/fakes/nccl_fakes.h index fb920cdb77b..c4a3c66619b 100644 --- a/projects/rccl/test/host/fakes/nccl_fakes.h +++ b/projects/rccl/test/host/fakes/nccl_fakes.h @@ -108,6 +108,13 @@ extern std::function g_memTrackImportFromPeer; +// ncclDebugLog formats each line and hands it here, already demoted from WARN to INFO under NOWARN as +// debug.cc does. The default prints it to stderr; a test that asserts on what was logged installs a +// recorder. +extern std::function + g_ncclDebugLog; + // Defaults succeed without tracking; UntrackDynamic reports a live persistent entry. struct ncclMemUntrackInfo; extern std::function struct FnSigOf> { using type = R(A...); }; template using FnSigOf_t = typename FnSigOf::type; +// glibc declares most of libc noexcept; std::function cannot carry that, so compare without it. +template struct DropNoexcept { using type = F; }; +template +struct DropNoexcept { using type = R(A...); }; + template constexpr bool HookMatchesProd() { - return std::is_same_v, std::remove_pointer_t>; + return std::is_same_v, typename DropNoexcept>::type>; } } // namespace rccl_test_host diff --git a/projects/rccl/test/host/linux-ipcsocket-test.cc b/projects/rccl/test/host/linux-ipcsocket-test.cc new file mode 100644 index 00000000000..f1762dd1739 --- /dev/null +++ b/projects/rccl/test/host/linux-ipcsocket-test.cc @@ -0,0 +1,280 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/os/linux_ipcsocket.cc on real Unix datagram sockets, abstract and +// path-bound. Its libc calls go through fakes/libc_seam.h, hooked only to fail or to forge a reply. + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "ipcsocket.h" +#include "os.h" +#include "param.h" +#include "utils.h" + +#include "linux-test-common.h" + +#include "fakes/param_redirect.h" +#include "fakes/libc_seam.h" +#include LINUX_IPCSOCKET_CC_PATH +#include "fakes/libc_seam_undef.h" + +namespace { + +class LinuxIpcsocketMicrotest : public LinuxTestBase { + protected: + // Abstract names are shared by every process in the network namespace, and pids repeat across + // containers, so each process draws a random token. + void SetUp() override { + static const uint64_t token = ((uint64_t)std::random_device{}() << 32 | std::random_device{}()) & ~0xffffull; + hash_ = token | (++counter_ & 0xffff); + // Close skips descriptor 0, so with stdin closed a socket must not land on it. + if (!FdOpen(0)) devNull_ = open("/dev/null", O_RDONLY); + ASSERT_TRUE(FdOpen(0)); + } + void TearDown() override { + for (auto* s : {&a_, &b_}) ncclIpcSocketClose(s); + if (devNull_ >= 0) close(devNull_); + } + std::string Name(int rank) { + char name[NCCL_IPC_SOCKNAME_LEN]; + snprintf(name, sizeof(name), "/tmp/nccl-socket-%d-%lx", rank, (unsigned long)hash_); + return name; + } + // a_ is rank 0 and b_ rank 1, both on hash_. + void Open(volatile uint32_t* abortFlag = nullptr) { + ASSERT_EQ(ncclIpcSocketInit(&a_, 0, hash_, abortFlag), ncclSuccess); + ASSERT_EQ(ncclIpcSocketInit(&b_, 1, hash_, abortFlag), ncclSuccess); + } + // Passes a pipe's write end from a_ to b_ and checks the received copy writes into the pipe. + void PassDescriptor() { + int p[2]; + ASSERT_EQ(pipe(p), 0); + ASSERT_EQ(ncclIpcSocketSendFd(&a_, p[1], 1, hash_), ncclSuccess); + ncclIpcFd got = -1; + ASSERT_EQ(ncclIpcSocketRecvFd(&b_, &got), ncclSuccess); + EXPECT_NE(got, p[1]); + EXPECT_EQ(write(got, "z", 1), 1); + char c = 0; + EXPECT_EQ(read(p[0], &c, 1), 1); + EXPECT_EQ(c, 'z'); + EXPECT_EQ(ncclIpcFdClose(got), 0); + EXPECT_FALSE(FdOpen(got)); + close(p[0]); + close(p[1]); + } + static uint64_t counter_; + uint64_t hash_ = 0; + int devNull_ = -1; + ncclIpcSocket a_ = {-1, "", nullptr}, b_ = {-1, "", nullptr}; +}; +uint64_t LinuxIpcsocketMicrotest::counter_ = 0; + +TEST_F(LinuxIpcsocketMicrotest, Init_NullHandleIsAnInternalError) { + EXPECT_EQ(ncclIpcSocketInit(nullptr, 0, hash_, nullptr), ncclInternalError); +} + +TEST_F(LinuxIpcsocketMicrotest, Init_AbstractSocketLeavesNoFile) { + ASSERT_EQ(ncclIpcSocketInit(&a_, 3, hash_, nullptr), ncclSuccess); + EXPECT_TRUE(FdOpen(a_.fd)); + EXPECT_EQ(a_.socketName, Name(3)); + EXPECT_EQ(a_.abortFlag, nullptr); + EXPECT_NE(access(a_.socketName, F_OK), 0); + EXPECT_EQ(fcntl(a_.fd, F_GETFL) & O_NONBLOCK, 0); +} + +TEST_F(LinuxIpcsocketMicrotest, Init_PathSocketReplacesAStaleFileAndCloseRemovesIt) { + params_.Set("IPC_USE_ABSTRACT_SOCKET", 0); + std::string name = Name(0); + FILE* stale = fopen(name.c_str(), "w"); + ASSERT_NE(stale, nullptr); + fclose(stale); + ASSERT_EQ(ncclIpcSocketInit(&a_, 0, hash_, nullptr), ncclSuccess); + struct stat st; + ASSERT_EQ(stat(name.c_str(), &st), 0); + EXPECT_TRUE(S_ISSOCK(st.st_mode)); + int fd = a_.fd; + ASSERT_EQ(ncclIpcSocketClose(&a_), ncclSuccess); + EXPECT_FALSE(FdOpen(fd)); + EXPECT_NE(access(name.c_str(), F_OK), 0); + a_.fd = -1; +} + +TEST_F(LinuxIpcsocketMicrotest, Init_AbortFlagMakesTheSocketNonBlocking) { + volatile uint32_t abortFlag = 0; + ASSERT_EQ(ncclIpcSocketInit(&a_, 0, hash_, &abortFlag), ncclSuccess); + EXPECT_EQ(a_.abortFlag, &abortFlag); + EXPECT_NE(fcntl(a_.fd, F_GETFL) & O_NONBLOCK, 0); +} + +TEST_F(LinuxIpcsocketMicrotest, Init_NameInUseIsASystemError) { + ASSERT_EQ(ncclIpcSocketInit(&a_, 0, hash_, nullptr), ncclSuccess); + EXPECT_EQ(ncclIpcSocketInit(&b_, 0, hash_, nullptr), ncclSystemError); + EXPECT_EQ(b_.fd, -1); + EXPECT_STREQ(b_.socketName, ""); +} + +TEST_F(LinuxIpcsocketMicrotest, Init_SocketFailureIsASystemError) { + ScopedHook socketHook(g_socket, [](int family, int type, int) { + EXPECT_EQ(family, AF_UNIX); + EXPECT_EQ(type, SOCK_DGRAM); + errno = EMFILE; + return -1; + }); + EXPECT_EQ(ncclIpcSocketInit(&a_, 0, hash_, nullptr), ncclSystemError); + EXPECT_EQ(a_.fd, -1); +} + +TEST_F(LinuxIpcsocketMicrotest, GetFd) { + ncclIpcFd fd = -1; + EXPECT_EQ(ncclIpcSocketGetFd(nullptr, &fd), ncclInvalidArgument); + Open(); + EXPECT_EQ(ncclIpcSocketGetFd(&a_, nullptr), ncclSuccess); + ASSERT_EQ(ncclIpcSocketGetFd(&a_, &fd), ncclSuccess); + EXPECT_EQ(fd, a_.fd); +} + +TEST_F(LinuxIpcsocketMicrotest, Close_RejectsANullSocket) { + EXPECT_EQ(ncclIpcSocketClose(nullptr), ncclInternalError); +} + +// Descriptor 0 counts as "not open", like -1, so Close leaves it. Pins AICOMRCCL-2976; a fix flips this pin. +TEST_F(LinuxIpcsocketMicrotest, Close_KnownBug_TreatsDescriptorZeroAsUnopened) { + ncclIpcSocket zero = {0, "", nullptr}; + EXPECT_EQ(ncclIpcSocketClose(&zero), ncclSuccess); + EXPECT_TRUE(FdOpen(0)); +} + +TEST_F(LinuxIpcsocketMicrotest, SendFd_PassesADescriptorOverAnAbstractSocket) { + Open(); + PassDescriptor(); +} + +TEST_F(LinuxIpcsocketMicrotest, SendFd_PassesADescriptorOverAPathSocket) { + params_.Set("IPC_USE_ABSTRACT_SOCKET", 0); + Open(); + EXPECT_EQ(access(b_.socketName, F_OK), 0); + PassDescriptor(); + for (auto* s : {&a_, &b_}) { + ASSERT_EQ(ncclIpcSocketClose(s), ncclSuccess); // while the param still says path sockets: unlinks + EXPECT_NE(access(s->socketName, F_OK), 0); + s->fd = -1; + } +} + +TEST_F(LinuxIpcsocketMicrotest, SendMsg_CarriesAHeaderWithOrWithoutADescriptor) { + Open(); + char hdr[6] = "hello"; + ASSERT_EQ(ncclIpcSocketSendMsg(&a_, hdr, sizeof(hdr), -1, 1, hash_), ncclSuccess); + char got[6] = {}; + ASSERT_EQ(ncclIpcSocketRecvMsg(&b_, got, sizeof(got), nullptr), ncclSuccess); + EXPECT_STREQ(got, "hello"); + + ASSERT_EQ(ncclIpcSocketSendMsg(&a_, hdr, sizeof(hdr), STDOUT_FILENO, 1, hash_), ncclSuccess); + std::memset(got, 0, sizeof(got)); + ncclIpcFd fd = -1; + ASSERT_EQ(ncclIpcSocketRecvMsg(&b_, got, sizeof(got), &fd), ncclSuccess); + EXPECT_STREQ(got, "hello"); + EXPECT_TRUE(FdOpen(fd)); + close(fd); +} + +TEST_F(LinuxIpcsocketMicrotest, RecvFd_MessageWithoutADescriptorIsASystemError) { + Open(); + ASSERT_EQ(ncclIpcSocketSendMsg(&b_, nullptr, 0, -1, 0, hash_), ncclSuccess); + ncclIpcFd fd = -1; + EXPECT_EQ(ncclIpcSocketRecvFd(&a_, &fd), ncclSystemError); +} + +TEST_F(LinuxIpcsocketMicrotest, RecvFd_ForeignControlMessageIsASystemError) { + Open(); + ScopedHook recvmsgHook(g_recvmsg, [](int, msghdr* msg, int) -> ssize_t { + cmsghdr* c = CMSG_FIRSTHDR(msg); + c->cmsg_len = CMSG_LEN(sizeof(int)); + c->cmsg_level = SOL_SOCKET; + c->cmsg_type = SCM_CREDENTIALS; + return 1; + }); + ncclIpcFd fd = -1; + EXPECT_EQ(ncclIpcSocketRecvFd(&a_, &fd), ncclSystemError); + ScopedHook wrongLevelHook(g_recvmsg, [](int, msghdr* msg, int) -> ssize_t { + cmsghdr* c = CMSG_FIRSTHDR(msg); + c->cmsg_len = CMSG_LEN(sizeof(int)); + c->cmsg_level = SOL_IP; + c->cmsg_type = SCM_RIGHTS; + return 1; + }); + EXPECT_EQ(ncclIpcSocketRecvFd(&a_, &fd), ncclSystemError); +} + +TEST_F(LinuxIpcsocketMicrotest, RecvMsg_RetriesTransientErrorsAndFailsOnOthers) { + Open(); + const int transient[] = {EINTR, EAGAIN, EWOULDBLOCK}; + int failures = 0; + ScopedHook recvmsgHook(g_recvmsg, [&](int fd, msghdr* msg, int flags) -> ssize_t { + if (failures < 3) { + errno = transient[failures++]; + return -1; + } + return ::recvmsg(fd, msg, flags); + }); + char hdr[2] = "x", got[2] = {}; + ASSERT_EQ(ncclIpcSocketSendMsg(&b_, hdr, sizeof(hdr), -1, 0, hash_), ncclSuccess); + ASSERT_EQ(ncclIpcSocketRecvMsg(&a_, got, sizeof(got), nullptr), ncclSuccess); + EXPECT_EQ(failures, 3); + EXPECT_STREQ(got, "x"); + + ScopedHook failingHook(g_recvmsg, FailWith(ENOTSOCK)); + EXPECT_EQ(ncclIpcSocketRecvMsg(&a_, got, sizeof(got), nullptr), ncclSystemError); +} + +TEST_F(LinuxIpcsocketMicrotest, RecvMsg_AbortStopsWaiting) { + volatile uint32_t abortFlag = 1; + Open(&abortFlag); + ncclIpcFd fd = -1; + EXPECT_EQ(ncclIpcSocketRecvFd(&a_, &fd), ncclInternalError); // non-blocking: EAGAIN, then the flag + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(LinuxIpcsocketMicrotest, SendMsg_NoPeerIsASystemError) { + Open(); + EXPECT_EQ(ncclIpcSocketSendFd(&a_, STDOUT_FILENO, 7, hash_), ncclSystemError); +} + +TEST_F(LinuxIpcsocketMicrotest, SendMsg_RetriesTransientErrorsUntilAborted) { + volatile uint32_t abortFlag = 0; + Open(&abortFlag); + int attempts = 0; + ScopedHook sendmsgHook(g_sendmsg, [&](int fd, const msghdr* msg, int flags) -> ssize_t { + if (++attempts < 3) { + errno = attempts == 1 ? EINTR : EAGAIN; + return -1; + } + return ::sendmsg(fd, msg, flags); + }); + char hdr[2] = "y", got[2] = {}; + ASSERT_EQ(ncclIpcSocketSendMsg(&a_, hdr, sizeof(hdr), -1, 1, hash_), ncclSuccess); + EXPECT_EQ(attempts, 3); + ASSERT_EQ(ncclIpcSocketRecvMsg(&b_, got, sizeof(got), nullptr), ncclSuccess); + EXPECT_STREQ(got, "y"); + + abortFlag = 1; + ScopedHook wouldBlockHook(g_sendmsg, FailWith(EWOULDBLOCK)); + EXPECT_EQ(ncclIpcSocketSendMsg(&a_, hdr, sizeof(hdr), -1, 1, hash_), ncclInternalError); +} + +} // namespace diff --git a/projects/rccl/test/host/linux-socket-pair-test.cc b/projects/rccl/test/host/linux-socket-pair-test.cc new file mode 100644 index 00000000000..5fc0044b35b --- /dev/null +++ b/projects/rccl/test/host/linux-socket-pair-test.cc @@ -0,0 +1,106 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/os/linux_socket_pair.cc on real pipes; only a pipe() failure is hooked. + +#include + +#include +#include +#include +#include +#include + +#include "checks.h" +#include "os_socket_pair.h" + +#include "linux-test-common.h" + +#include "fakes/libc_seam.h" +#include LINUX_SOCKET_PAIR_CC_PATH +#include "fakes/libc_seam_undef.h" + +namespace { + +class LinuxSocketPairMicrotest : public LinuxTestBase { + protected: + void TearDown() override { ncclOsSocketPairClose(pair_); } + ncclSocketPairDescriptor pair_[2] = {NCCL_SOCKET_PAIR_INVALID, NCCL_SOCKET_PAIR_INVALID}; +}; + +TEST_F(LinuxSocketPairMicrotest, Create_WriteEndFeedsReadEnd) { + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + size_t n = 0; + ASSERT_EQ(ncclOsSocketPairWrite(pair_[1], "hello", 5, &n), ncclSuccess); + EXPECT_EQ(n, 5u); + char buf[8] = {}; + ASSERT_EQ(ncclOsSocketPairRead(pair_[0], buf, sizeof(buf), &n), ncclSuccess); + EXPECT_EQ(n, 5u); + EXPECT_STREQ(buf, "hello"); +} + +TEST_F(LinuxSocketPairMicrotest, Read_ZeroBytesAfterTheWriterCloses) { + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + ::close(pair_[1]); + pair_[1] = NCCL_SOCKET_PAIR_INVALID; + char buf[4]; + size_t n = 99; + ASSERT_EQ(ncclOsSocketPairRead(pair_[0], buf, sizeof(buf), &n), ncclSuccess); + EXPECT_EQ(n, 0u); +} + +TEST_F(LinuxSocketPairMicrotest, Close_ClosesBothEndsAndSkipsInvalidOnes) { + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + int fds[2] = {pair_[0], pair_[1]}; + ASSERT_EQ(ncclOsSocketPairClose(pair_), ncclSuccess); + EXPECT_FALSE(FdOpen(fds[0])); + EXPECT_FALSE(FdOpen(fds[1])); + EXPECT_EQ(pair_[0], NCCL_SOCKET_PAIR_INVALID); + EXPECT_EQ(pair_[1], NCCL_SOCKET_PAIR_INVALID); + EXPECT_EQ(ncclOsSocketPairClose(pair_), ncclSuccess); + + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + int readEnd = pair_[0]; + ::close(pair_[1]); + pair_[1] = NCCL_SOCKET_PAIR_INVALID; + ASSERT_EQ(ncclOsSocketPairClose(pair_), ncclSuccess); + EXPECT_FALSE(FdOpen(readEnd)); + EXPECT_TRUE(log_.Text().empty()); +} + +// Both ends are still reset, and only the first failure is reported. +TEST_F(LinuxSocketPairMicrotest, Close_ReportsTheFirstFailure) { + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + ::close(pair_[0]); + ::close(pair_[1]); + EXPECT_EQ(ncclOsSocketPairClose(pair_), ncclSystemError); + EXPECT_EQ(pair_[0], NCCL_SOCKET_PAIR_INVALID); + EXPECT_EQ(pair_[1], NCCL_SOCKET_PAIR_INVALID); + EXPECT_EQ(log_.Warns(), 1); + + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + int readEnd = pair_[0]; + ::close(pair_[1]); + EXPECT_EQ(ncclOsSocketPairClose(pair_), ncclSystemError); // the second end fails + EXPECT_FALSE(FdOpen(readEnd)); +} + +TEST_F(LinuxSocketPairMicrotest, Create_PipeFailureIsASystemError) { + ScopedHook pipeHook(g_pipe, FailWith(EMFILE)); + EXPECT_EQ(ncclOsSocketPairCreate(pair_), ncclSystemError); + EXPECT_EQ(pair_[0], NCCL_SOCKET_PAIR_INVALID); +} + +TEST_F(LinuxSocketPairMicrotest, ReadAndWrite_FailuresAreSystemErrors) { + ASSERT_EQ(ncclOsSocketPairCreate(pair_), ncclSuccess); + size_t n = 99; + char buf[4]; + EXPECT_EQ(ncclOsSocketPairWrite(pair_[0], "x", 1, &n), ncclSystemError); // the read end + EXPECT_EQ(ncclOsSocketPairRead(pair_[1], buf, sizeof(buf), &n), ncclSystemError); + EXPECT_EQ(n, 99u); +} + +} // namespace diff --git a/projects/rccl/test/host/linux-test-common.h b/projects/rccl/test/host/linux-test-common.h new file mode 100644 index 00000000000..4f32889ac5b --- /dev/null +++ b/projects/rccl/test/host/linux-test-common.h @@ -0,0 +1,165 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Shared by the rccl-UnitTestsMicroLinux test TUs: descriptor helpers, the base fixture, a log recorder +// and a scripted getifaddrs() list, all over the shared libc, nccl, env and HIP fakes. +#pragma once + +#include + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "../common/LogCapture.hpp" +#include "ScopedHook.h" +#include "fakes/env_fakes.h" +#include "fakes/dl_fakes.h" +#include "fakes/libc_fakes.h" +#include "fakes/nccl_fakes.h" + +inline bool FdOpen(int fd) { return ::fcntl(fd, F_GETFD) != -1; } + +inline int IntOpt(int fd, int level, int name) { + int v = -1; + socklen_t len = sizeof(v); + ::getsockopt(fd, level, name, &v, &len); + return v; +} + +// A hook that fails like a syscall: sets errno and returns -1. +inline auto FailWith(int err) { + return [err](auto...) { + errno = err; + return -1; + }; +} + +// What production logged, from any thread, as "WARN|INFO|ATTN \n" lines. +class LogRecord { + public: + int Warns() const { + std::lock_guard lock(mutex_); + return warns_; + } + bool Has(const std::string& needle) const { + std::lock_guard lock(mutex_); + return text_.find(needle) != std::string::npos; + } + std::string Text() const { + std::lock_guard lock(mutex_); + return text_; + } + void Clear() { + std::lock_guard lock(mutex_); + text_.clear(); + warns_ = 0; + } + + private: + void Append(ncclDebugLogLevel level, const char* msg) { + const char* tag = level == NCCL_LOG_WARN ? "WARN" : level == NCCL_LOG_ATTN ? "ATTN" : "INFO"; + std::lock_guard lock(mutex_); + text_ += std::string(tag) + " " + msg + "\n"; + if (level == NCCL_LOG_WARN) ++warns_; + } + mutable std::mutex mutex_; + std::string text_; + int warns_ = 0; + RcclUnitTesting::ScopedDebugLogging everyLevel_{-1, ~uint64_t(0)}; + ScopedHook hook_{ + g_ncclDebugLog, + [this](ncclDebugLogLevel level, unsigned long, const char*, int, const char* msg) { Append(level, msg); }}; +}; + +// Params reach the units through param_redirect.h's g_loadParam; one not Set() keeps its default. +class ParamScript { + public: + void Set(const char* name, int64_t value) { values_[name] = value; } + + private: + std::map values_; + ScopedHook hook_{g_loadParam, [this](const char* name, int64_t deftVal) { + auto it = values_.find(name); + return it == values_.end() ? deftVal : it->second; + }}; +}; + +// Real libc behind every seam, an empty scripted environment, and socket retries that do not sleep. +class LinuxTestBase : public ::testing::Test { + protected: + LinuxTestBase() { + UseRealLibc(); + ResetDlFakes(); + ResetEnvFakes(); + params_.Set("SOCKET_RETRY_SLEEP_MSEC", 0); + } + ~LinuxTestBase() override { + ResetEnvFakes(); + UseRealLibc(); + ResetDlFakes(); + } + LogRecord log_; + ParamScript params_; +}; + +// A scripted getifaddrs() list, installed for the object's lifetime. +class LinuxFakeIfaddrs { + public: + void Add(const char* name, int family, const char* addr = nullptr, const char* mask = nullptr, + unsigned flags = IFF_UP | IFF_RUNNING, uint32_t scope = 0) { + nodes_.emplace_back(); + Node& n = nodes_.back(); + n.name = name; + n.ifa.ifa_name = const_cast(n.name.c_str()); + n.ifa.ifa_flags = flags; + if (family == AF_UNSPEC) return; + Fill(&n.addr, family, addr, scope); + n.ifa.ifa_addr = (sockaddr*)&n.addr; + if (mask) { + Fill(&n.mask, family, mask, 0); + n.ifa.ifa_netmask = (sockaddr*)&n.mask; + } + } + static void Fill(sockaddr_storage* ss, int family, const char* text, uint32_t scope) { + std::memset(ss, 0, sizeof(*ss)); + ss->ss_family = family; + if (family == AF_INET) inet_pton(AF_INET, text, &((sockaddr_in*)ss)->sin_addr); + if (family == AF_INET6) { + inet_pton(AF_INET6, text, &((sockaddr_in6*)ss)->sin6_addr); + ((sockaddr_in6*)ss)->sin6_scope_id = scope; + } + } + int Listed() const { return list_.calls; } + int Freed() const { return free_.calls; } + + private: + struct Node { + ifaddrs ifa{}; + std::string name; + sockaddr_storage addr{}, mask{}; + }; + std::deque nodes_; + ScopedHook list_{g_getifaddrs, [this](ifaddrs** out) { + ifaddrs* next = nullptr; + for (auto it = nodes_.rbegin(); it != nodes_.rend(); ++it) { + it->ifa.ifa_next = next; + next = &it->ifa; + } + *out = next; + return 0; + }}; + ScopedHook free_{g_freeifaddrs, [](ifaddrs*) {}}; +}; diff --git a/projects/rccl/test/host/linux-test.cc b/projects/rccl/test/host/linux-test.cc new file mode 100644 index 00000000000..198bffa2e3b --- /dev/null +++ b/projects/rccl/test/host/linux-test.cc @@ -0,0 +1,1708 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/os/linux.cc on real sockets, files, shared memory and dlopen. Its libc +// calls go through fakes/libc_seam.h, real unless a test hooks one to fail or vary. + +#include + +#include "hip/hip_runtime.h" +#include "os.h" +#include "checks.h" +#include "comm.h" +#include "crypt.h" +#include "param.h" +#include "socket.h" +#include "utils.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "../common/ScopedEnv.hpp" +#include "fakes/hip_fakes.h" +#include "linux-test-common.h" + +#include "fakes/libc_seam.h" +#include "fakes/dl_seam.h" +#include LINUX_CC_PATH +#include "fakes/dl_seam_undef.h" +#include "fakes/libc_seam_undef.h" + +namespace { + +bool Exists(const std::string& path) { return access(path.c_str(), F_OK) == 0; } + +// Whether the thread whose /proc/.../stat this is sleeps ('S'), the state field after the name. +bool ThreadSleeps(const std::string& statPath) { + char buf[512] = ""; + FILE* f = ::fopen(statPath.c_str(), "r"); + if (!f) return false; + size_t n = fread(buf, 1, sizeof(buf) - 1, f); + fclose(f); + buf[n] = '\0'; + const char* end = std::strrchr(buf, ')'); + return end && end[1] == ' ' && end[2] == 'S'; +} + +using LinuxMicrotest = LinuxTestBase; + +// Dynamic loading + +using LinuxDlMicrotest = LinuxMicrotest; + +TEST_F(LinuxDlMicrotest, Dlopen_LoadsLocallyAndResolvesASymbol) { + int seenFlags = -1; + ScopedHook dlopenHook(g_dlopen, [&](const char* f, int flags) { + seenFlags = flags; + return ::dlopen(f, flags); + }); + ncclOsLibraryHandle h = ncclOsDlopen("libc.so.6"); + ASSERT_NE(h, nullptr); + EXPECT_EQ(seenFlags, RTLD_NOW | RTLD_LOCAL); + auto pid = reinterpret_cast(ncclOsDlsym(h, "getpid")); + ASSERT_NE(pid, nullptr); + EXPECT_EQ(pid(), getpid()); + ncclOsDlclose(h); + ncclOsDlclose(nullptr); +} + +TEST_F(LinuxDlMicrotest, Dlopen_ModeSelectsNowOrLazy) { + int seenFlags = -1; + ScopedHook dlopenHook(g_dlopen, [&](const char* f, int flags) { + seenFlags = flags; + return ::dlopen(f, flags); + }); + ncclOsLibraryHandle h = ncclOsDlopen("libc.so.6", NCCL_OS_DL_NOW); + EXPECT_EQ(seenFlags, RTLD_NOW); + ncclOsDlclose(h); + h = ncclOsDlopen("libc.so.6", NCCL_OS_DL_LAZY); + EXPECT_EQ(seenFlags, RTLD_LAZY); + ncclOsDlclose(h); +} + +TEST_F(LinuxDlMicrotest, Failures_KeepTheLoaderMessageForDlerror) { + EXPECT_EQ(ncclOsDlopen("libnccl-no-such-lib.so"), nullptr); + EXPECT_NE(std::string(ncclOsDlerror()).find("libnccl-no-such-lib.so"), std::string::npos); + + EXPECT_EQ(ncclOsDlopen("libnccl-no-such-lib-2.so", NCCL_OS_DL_NOW), nullptr); + EXPECT_NE(std::string(ncclOsDlerror()).find("libnccl-no-such-lib-2.so"), std::string::npos); + + ncclOsLibraryHandle h = ncclOsDlopen("libc.so.6"); + EXPECT_NE(std::string(ncclOsDlerror()).find("libnccl-no-such-lib-2.so"), std::string::npos); // success keeps it + EXPECT_EQ(ncclOsDlsym(h, "ncclNoSuchSymbol"), nullptr); + EXPECT_NE(std::string(ncclOsDlerror()).find("ncclNoSuchSymbol"), std::string::npos); + ncclOsDlclose(h); + + ScopedHook dlerrorHook(g_dlerror, []() -> char* { return nullptr; }); + EXPECT_EQ(ncclOsDlopen("libnccl-no-such-lib.so"), nullptr); + EXPECT_STREQ(ncclOsDlerror(), ""); +} + +TEST_F(LinuxDlMicrotest, Dlerror_IsPerThread) { + EXPECT_EQ(ncclOsDlopen("libnccl-no-such-lib.so"), nullptr); + std::string other = "unset"; + std::thread([&] { other = ncclOsDlerror(); }).join(); + EXPECT_EQ(other, ""); + EXPECT_STRNE(ncclOsDlerror(), ""); +} + +TEST_F(LinuxDlMicrotest, NvmlOpen_LoadsLibnvidiaMl) { + std::string seen; + ScopedHook dlopenHook(g_dlopen, [&](const char* f, int) { + seen = f; + return reinterpret_cast(0x1); + }); + ncclOsLibraryHandle h = nullptr; + ASSERT_EQ(ncclOsNvmlOpen(&h), ncclSuccess); + EXPECT_EQ(seen, "libnvidia-ml.so.1"); + EXPECT_EQ(h, reinterpret_cast(0x1)); +} + +TEST_F(LinuxDlMicrotest, NvmlOpen_MissingLibraryIsASystemError) { + ScopedHook dlopenHook(g_dlopen, [](const char*, int) -> void* { return nullptr; }); + ScopedHook dlerrorHook(g_dlerror, []() { return const_cast("not here"); }); + ncclOsLibraryHandle h = reinterpret_cast(0x1); + EXPECT_EQ(ncclOsNvmlOpen(&h), ncclSystemError); + EXPECT_EQ(h, nullptr); +} + +// Process and memory helpers + +using LinuxProcessMicrotest = LinuxMicrotest; + +TEST_F(LinuxProcessMicrotest, PidTidAndPageSize) { + EXPECT_EQ(ncclOsGetPid(), (uint64_t)getpid()); + EXPECT_EQ(ncclOsGetTid(), (uint64_t)syscall(SYS_gettid)); + uint64_t otherTid = 0; + std::thread([&] { otherTid = ncclOsGetTid(); }).join(); + EXPECT_NE(otherTid, ncclOsGetTid()); + EXPECT_EQ(ncclOsGetPageSize(), (size_t)sysconf(_SC_PAGESIZE)); + EXPECT_EQ(ncclOsGetCommMempoolMaxSize(), 0u); // no cap +} + +TEST_F(LinuxProcessMicrotest, LocaltimeFillsTheCallersBuffer) { + time_t t = 86400 * 365; + std::tm buf{}, expected{}; + EXPECT_EQ(ncclOsLocaltime(&t, &buf), &buf); + localtime_r(&t, &expected); + EXPECT_EQ(buf.tm_year, expected.tm_year); + EXPECT_EQ(buf.tm_yday, expected.tm_yday); + EXPECT_EQ(buf.tm_hour, expected.tm_hour); +} + +TEST_F(LinuxProcessMicrotest, AlignedAllocHonoursTheAlignment) { + void* p = ncclOsAlignedAlloc(4096, 8192); + ASSERT_NE(p, nullptr); + EXPECT_EQ(reinterpret_cast(p) % 4096, 0u); + std::memset(p, 0xab, 8192); + ncclOsAlignedFree(p); +} + +TEST_F(LinuxProcessMicrotest, SetEnvNeverOverwrites) { + RcclUnitTesting::ScopedEnv env("NCCL_LINUX_TEST_VAR", nullptr); + ncclOsSetEnv("NCCL_LINUX_TEST_VAR", "first"); + EXPECT_STREQ(getenv("NCCL_LINUX_TEST_VAR"), "first"); + ncclOsSetEnv("NCCL_LINUX_TEST_VAR", "second"); + EXPECT_STREQ(getenv("NCCL_LINUX_TEST_VAR"), "first"); +} + +TEST_F(LinuxProcessMicrotest, StrSepAndRealpath) { + char text[] = "a:b"; + char* cursor = text; + EXPECT_STREQ(ncclOsStrSep(&cursor, ":"), "a"); + EXPECT_STREQ(ncclOsStrSep(&cursor, ":"), "b"); + EXPECT_EQ(cursor, nullptr); + + char resolved[PATH_MAX], expected[PATH_MAX]; + ASSERT_NE(::realpath("/tmp", expected), nullptr); + EXPECT_EQ(ncclOsRealpath("/tmp/../tmp/.", resolved), resolved); + EXPECT_STREQ(resolved, expected); + EXPECT_EQ(ncclOsRealpath("/nccl/no/such/dir", resolved), nullptr); +} + +// Stack size and file limits + +class LinuxInitMicrotest : public LinuxMicrotest { + protected: + void SetUp() override { + LinuxMicrotest::SetUp(); + ASSERT_EQ(pthread_getattr_default_np(&saved_), 0); + } + void TearDown() override { + pthread_setattr_default_np(&saved_); + pthread_attr_destroy(&saved_); + LinuxMicrotest::TearDown(); + } + static void SetDefaultStack(size_t bytes) { + pthread_attr_t attr; + pthread_attr_init(&attr); + pthread_attr_setstacksize(&attr, bytes); + pthread_setattr_default_np(&attr); + pthread_attr_destroy(&attr); + } + static size_t DefaultStack() { + pthread_attr_t attr; + size_t size = 0; + pthread_attr_init(&attr); + pthread_attr_getstacksize(&attr, &size); + pthread_attr_destroy(&attr); + return size; + } + pthread_attr_t saved_; +}; + +TEST_F(LinuxInitMicrotest, Initialize_ParamOffLeavesTheStackAlone) { + params_.Set("SET_CPU_STACK_SIZE", 0); + SetDefaultStack(1 << 20); + ASSERT_EQ(ncclOsInitialize(), ncclSuccess); + EXPECT_EQ(DefaultStack(), 1u << 20); + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(LinuxInitMicrotest, Initialize_RaisesASmallDefaultStackTo8MB) { + SetDefaultStack(1 << 20); + ScopedHook getrlimitHook(g_getrlimit, [](int resource, rlimit* lim) { + int rc = ::getrlimit((__rlimit_resource_t)resource, lim); + if (resource == RLIMIT_STACK) lim->rlim_cur = 2048 * 1024; + return rc; + }); + ASSERT_EQ(ncclOsInitialize(), ncclSuccess); + EXPECT_EQ(DefaultStack(), 8192u * 1024); +} + +TEST_F(LinuxInitMicrotest, Initialize_ReportsAnUnlimitedStackLimit) { + SetDefaultStack(2 << 20); // glibc's fallback when RLIMIT_STACK is unlimited + ScopedHook getrlimitHook(g_getrlimit, [](int resource, rlimit* lim) { + int rc = ::getrlimit((__rlimit_resource_t)resource, lim); + if (resource == RLIMIT_STACK) lim->rlim_cur = RLIM_INFINITY; + return rc; + }); + ASSERT_EQ(ncclOsInitialize(), ncclSuccess); + EXPECT_EQ(DefaultStack(), 8192u * 1024); + EXPECT_TRUE(log_.Has("(unlimited)")); // the only trace of the RLIM_INFINITY branch +} + +TEST_F(LinuxInitMicrotest, Initialize_KeepsAStackThatIsAlready8MB) { + ScopedHook getrlimitHook(g_getrlimit, [](int resource, rlimit* lim) { + return ::getrlimit((__rlimit_resource_t)resource, lim); + }); + SetDefaultStack(8192 * 1024); + ASSERT_EQ(ncclOsInitialize(), ncclSuccess); + SetDefaultStack(16 << 20); + ASSERT_EQ(ncclOsInitialize(), ncclSuccess); + EXPECT_EQ(DefaultStack(), 16u << 20); + EXPECT_EQ(getrlimitHook.calls, 0); // only a small stack looks up the limit +} + +TEST_F(LinuxInitMicrotest, Initialize_GetrlimitFailureIsASystemError) { + SetDefaultStack(1 << 20); + ScopedHook getrlimitHook(g_getrlimit, FailWith(EPERM)); + EXPECT_EQ(ncclOsInitialize(), ncclSystemError); +} + +TEST_F(LinuxInitMicrotest, SetFilesLimit_RaisesTheSoftLimitToTheHardLimit) { + rlimit saved; + ASSERT_EQ(::getrlimit(RLIMIT_NOFILE, &saved), 0); + rlimit lowered = saved; + lowered.rlim_cur = std::min(saved.rlim_max, 256); + ASSERT_EQ(::setrlimit(RLIMIT_NOFILE, &lowered), 0); + rlimit seen{}; + ScopedHook setrlimitHook(g_setrlimit, [&](int resource, const rlimit* lim) { + seen = *lim; + return ::setrlimit((__rlimit_resource_t)resource, lim); + }); + EXPECT_EQ(ncclOsSetFilesLimit(), ncclSuccess); + EXPECT_EQ(seen.rlim_cur, saved.rlim_max); + EXPECT_EQ(seen.rlim_max, saved.rlim_max); + ::setrlimit(RLIMIT_NOFILE, &saved); +} + +TEST_F(LinuxInitMicrotest, SetFilesLimit_SyscallFailuresAreSystemErrors) { + ScopedHook setrlimitHook(g_setrlimit, FailWith(EPERM)); + EXPECT_EQ(ncclOsSetFilesLimit(), ncclSystemError); + ScopedHook getrlimitHook(g_getrlimit, FailWith(EPERM)); + EXPECT_EQ(ncclOsSetFilesLimit(), ncclSystemError); + EXPECT_EQ(setrlimitHook.calls, 1); // stopped at getrlimit +} + +// Sockets + +class LinuxSocketMicrotest : public LinuxMicrotest { + protected: + void TearDown() override { + for (int fd : fds_) ::close(fd); + LinuxMicrotest::TearDown(); + } + int Track(int fd) { + if (fd >= 0) fds_.push_back(fd); + return fd; + } + // A listening loopback TCP socket; *addr gets its port. + int Listen(sockaddr_in* addr, bool nonblocking = false) { + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + std::memset(addr, 0, sizeof(*addr)); + addr->sin_family = AF_INET; + addr->sin_addr.s_addr = htonl(INADDR_LOOPBACK); + EXPECT_EQ(::bind(fd, (sockaddr*)addr, sizeof(*addr)), 0); + socklen_t len = sizeof(*addr); + EXPECT_EQ(::getsockname(fd, (sockaddr*)addr, &len), 0); + EXPECT_EQ(::listen(fd, 16), 0); + if (nonblocking) fcntl(fd, F_SETFL, fcntl(fd, F_GETFL) | O_NONBLOCK); + return fd; + } + // A loopback port nothing listens on. + sockaddr_in ClosedPort() { + sockaddr_in addr; + int fd = Listen(&addr); + ::close(fd); + fds_.pop_back(); + return addr; + } + static ncclSocket Sock(int fd, const sockaddr_in* addr = nullptr) { + ncclSocket sock; + std::memset(&sock, 0, sizeof(sock)); + sock.socketDescriptor = fd; + sock.acceptSocketDescriptor = NCCL_INVALID_SOCKET; + if (addr) std::memcpy(&sock.addr.sin, addr, sizeof(*addr)); + else sock.addr.sa.sa_family = AF_INET; + sock.salen = sizeof(sockaddr_in); + return sock; + } + void Pair(int sv[2]) { + ASSERT_EQ(::socketpair(AF_UNIX, SOCK_STREAM, 0, sv), 0); + Track(sv[0]); + Track(sv[1]); + } + std::vector fds_; +}; + +TEST_F(LinuxSocketMicrotest, DescriptorValidity) { + EXPECT_FALSE(ncclOsSocketDescriptorIsValid(NCCL_INVALID_SOCKET)); + EXPECT_TRUE(ncclOsSocketDescriptorIsValid(0)); + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + EXPECT_FALSE(ncclOsSocketIsValid(&sock)); + sock.socketDescriptor = 3; + EXPECT_TRUE(ncclOsSocketIsValid(&sock)); +} + +TEST_F(LinuxSocketMicrotest, PollSocket_WaitsForTheOpsDirectionWithThePollTimeout) { + params_.Set("SOCKET_POLL_TIMEOUT_MSEC", 7); + std::vector> calls; + ScopedHook pollHook(g_poll, [&](pollfd* p, nfds_t n, int timeout) { + EXPECT_EQ(n, 1u); + EXPECT_EQ(p->fd, 5); + calls.push_back({p->events, timeout}); + return 1; + }); + ncclOsPollSocket(5, NCCL_SOCKET_RECV); + ncclOsPollSocket(5, NCCL_SOCKET_SEND); + ASSERT_EQ(calls.size(), 2u); + EXPECT_EQ(calls[0], std::make_pair((short)POLLIN, 7)); + EXPECT_EQ(calls[1], std::make_pair((short)POLLOUT, 7)); +} + +TEST_F(LinuxSocketMicrotest, TryAccept_AcceptsAPendingConnection) { + sockaddr_in addr; + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.acceptSocketDescriptor = Listen(&addr); + sock.state = ncclSocketStateAccepting; + int client = Track(::socket(AF_INET, SOCK_STREAM, 0)); + ASSERT_EQ(::connect(client, (sockaddr*)&addr, sizeof(addr)), 0); + ASSERT_EQ(ncclOsSocketTryAccept(&sock), ncclSuccess); + Track(sock.socketDescriptor); + EXPECT_TRUE(FdOpen(sock.socketDescriptor)); + EXPECT_EQ(sock.state, ncclSocketStateAccepted); + sockaddr_in local; + socklen_t len = sizeof(local); + ::getsockname(client, (sockaddr*)&local, &len); + EXPECT_EQ(sock.addr.sin.sin_port, local.sin_port); // the peer's address +} + +TEST_F(LinuxSocketMicrotest, TryAccept_NothingPendingIsNotAnError) { + sockaddr_in addr; + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.acceptSocketDescriptor = Listen(&addr, /*nonblocking=*/true); + sock.state = ncclSocketStateAccepting; + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateAccepting); + EXPECT_EQ(sock.errorRetries, 0); + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(LinuxSocketMicrotest, TryAccept_OtherErrorsAreSystemErrors) { + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.acceptSocketDescriptor = Track(::socket(AF_INET, SOCK_STREAM, 0)); // not listening: EINVAL + sock.state = ncclSocketStateAccepting; + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSystemError); +} + +// Network errors that accept() may report for an already-failed peer are retried up to NCCL_SOCKET_RETRY_CNT. +TEST_F(LinuxSocketMicrotest, TryAccept_RetriesPendingNetworkErrorsUpToTheRetryCount) { + params_.Set("SOCKET_RETRY_CNT", 3); + for (int err : {ENETDOWN, EPROTO, ENOPROTOOPT, EHOSTDOWN, ENONET, EHOSTUNREACH, EOPNOTSUPP, ENETUNREACH, EINTR}) { + SCOPED_TRACE(err); + ScopedHook acceptHook(g_accept, FailWith(err)); + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.state = ncclSocketStateAccepting; + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSuccess); + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSuccess); + EXPECT_EQ(sock.errorRetries, 2); + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSystemError); + EXPECT_EQ(sock.state, ncclSocketStateAccepting); + } +} + +TEST_F(LinuxSocketMicrotest, TryAccept_WouldBlockIsNotRetried) { + for (int err : {EAGAIN, EWOULDBLOCK}) { + ScopedHook acceptHook(g_accept, FailWith(err)); + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + EXPECT_EQ(ncclOsSocketTryAccept(&sock), ncclSuccess); + EXPECT_EQ(sock.errorRetries, 0); + } + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(LinuxSocketMicrotest, SetFlags_RejectsAnInvalidSocket) { + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + EXPECT_EQ(ncclOsSocketSetFlags(&sock), ncclInvalidArgument); +} + +TEST_F(LinuxSocketMicrotest, SetFlags_BlockingSocketGetsNoDelayAndDefaultBuffers) { + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + int sndDefault = IntOpt(fd, SOL_SOCKET, SO_SNDBUF), rcvDefault = IntOpt(fd, SOL_SOCKET, SO_RCVBUF); + ncclSocket sock = Sock(fd); + ASSERT_EQ(ncclOsSocketSetFlags(&sock), ncclSuccess); + EXPECT_EQ(fcntl(fd, F_GETFL) & O_NONBLOCK, 0); + EXPECT_NE(IntOpt(fd, IPPROTO_TCP, TCP_NODELAY), 0); + EXPECT_EQ(IntOpt(fd, SOL_SOCKET, SO_SNDBUF), sndDefault); + EXPECT_EQ(IntOpt(fd, SOL_SOCKET, SO_RCVBUF), rcvDefault); +} + +TEST_F(LinuxSocketMicrotest, SetFlags_AsyncOrAbortableSocketsAreNonBlocking) { + volatile uint32_t abortFlag = 0; + for (int variant = 0; variant < 2; variant++) { + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + ncclSocket sock = Sock(fd); + if (variant == 0) sock.asyncFlag = 1; + else sock.abortFlag = &abortFlag; + ASSERT_EQ(ncclOsSocketSetFlags(&sock), ncclSuccess); + EXPECT_NE(fcntl(fd, F_GETFL) & O_NONBLOCK, 0) << variant; + } +} + +TEST_F(LinuxSocketMicrotest, SetFlags_AppliesTheBufferSizeParams) { + params_.Set("SOCKET_SNDBUF", 65536); + params_.Set("SOCKET_RCVBUF", 32768); + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + ncclSocket sock = Sock(fd); + ASSERT_EQ(ncclOsSocketSetFlags(&sock), ncclSuccess); + EXPECT_EQ(IntOpt(fd, SOL_SOCKET, SO_SNDBUF), 2 * 65536); // Linux doubles the request + EXPECT_EQ(IntOpt(fd, SOL_SOCKET, SO_RCVBUF), 2 * 32768); +} + +TEST_F(LinuxSocketMicrotest, SetFlags_SyscallFailuresAreSystemErrors) { + ncclSocket sock = Sock(1 << 20); // valid-looking, but not open: fcntl fails + sock.asyncFlag = 1; + EXPECT_EQ(ncclOsSocketSetFlags(&sock), ncclSystemError); + + int p[2]; + ASSERT_EQ(pipe(p), 0); + Track(p[0]); + Track(p[1]); + sock = Sock(p[0]); // not a socket: setsockopt fails + EXPECT_EQ(ncclOsSocketSetFlags(&sock), ncclSystemError); +} + +TEST_F(LinuxSocketMicrotest, SetFlags_BufferSizeFailuresAreSystemErrors) { + params_.Set("SOCKET_SNDBUF", 65536); + params_.Set("SOCKET_RCVBUF", 32768); + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + ncclSocket sock = Sock(fd); + for (int failing : {SO_SNDBUF, SO_RCVBUF}) { + ScopedHook setsockoptHook(g_setsockopt, [failing](int s, int level, int name, const void* v, socklen_t len) { + if (level == SOL_SOCKET && name == failing) { + errno = ENOBUFS; + return -1; + } + return ::setsockopt(s, level, name, v, len); + }); + EXPECT_EQ(ncclOsSocketSetFlags(&sock), ncclSystemError); + } +} + +TEST_F(LinuxSocketMicrotest, ResetAccept_ClosesThePeerAndMarksABadHandshake) { + int sv[2]; + Pair(sv); + ncclSocket sock = Sock(sv[0]); + sock.state = ncclSocketStateAccepted; + sock.finalizeCounter = 5; + ncclOsSocketResetAccept(&sock); + EXPECT_FALSE(FdOpen(sv[0])); + EXPECT_EQ(sock.socketDescriptor, NCCL_INVALID_SOCKET); + EXPECT_EQ(sock.state, ncclSocketStateBadHandshake); + EXPECT_EQ(sock.finalizeCounter, 0); +} + +TEST_F(LinuxSocketMicrotest, ResetFd_CreatesAConfiguredSocket) { + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.asyncFlag = 1; + ASSERT_EQ(ncclOsSocketResetFd(&sock), ncclSuccess); + Track(sock.socketDescriptor); + ASSERT_TRUE(ncclOsSocketIsValid(&sock)); + EXPECT_EQ(IntOpt(sock.socketDescriptor, SOL_SOCKET, SO_TYPE), SOCK_STREAM); + EXPECT_NE(fcntl(sock.socketDescriptor, F_GETFL) & O_NONBLOCK, 0); + EXPECT_NE(IntOpt(sock.socketDescriptor, IPPROTO_TCP, TCP_NODELAY), 0); +} + +// An existing descriptor keeps its number but becomes a fresh socket. +TEST_F(LinuxSocketMicrotest, ResetFd_ReusesTheExistingDescriptorNumber) { + sockaddr_in addr; + int fd = Listen(&addr); + int lowestFree = ::socket(AF_INET, SOCK_STREAM, 0); + ::close(lowestFree); + ncclSocket sock = Sock(fd); + ASSERT_EQ(ncclOsSocketResetFd(&sock), ncclSuccess); + EXPECT_EQ(sock.socketDescriptor, fd); + EXPECT_EQ(IntOpt(fd, SOL_SOCKET, SO_ACCEPTCONN), 0); + int probe = ::socket(AF_INET, SOCK_STREAM, 0); // the temporary descriptor was closed + EXPECT_EQ(probe, lowestFree); + ::close(probe); +} + +TEST_F(LinuxSocketMicrotest, ResetFd_SocketFailureLeavesTheDescriptorAlone) { + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.addr.sa.sa_family = AF_UNSPEC; + EXPECT_EQ(ncclOsSocketResetFd(&sock), ncclSystemError); + EXPECT_EQ(sock.socketDescriptor, NCCL_INVALID_SOCKET); +} + +TEST_F(LinuxSocketMicrotest, ResetFd_Dup2FailureClosesTheNewSocket) { + int probe = ::socket(AF_INET, SOCK_STREAM, 0); // the number the new socket will get + ::close(probe); + // Lower the soft RLIMIT_NOFILE so the dup2 target is past it on any host. + rlimit saved; + ASSERT_EQ(getrlimit(RLIMIT_NOFILE, &saved), 0); + rlimit tight = saved; + tight.rlim_cur = probe + 8; + ASSERT_EQ(setrlimit(RLIMIT_NOFILE, &tight), 0); + const int target = probe + 8; + ncclSocket sock = Sock(target); + ncclResult_t ret = ncclOsSocketResetFd(&sock); + setrlimit(RLIMIT_NOFILE, &saved); + EXPECT_EQ(ret, ncclSystemError); + EXPECT_EQ(sock.socketDescriptor, target); + EXPECT_FALSE(FdOpen(probe)); +} + +TEST_F(LinuxSocketMicrotest, ResetFd_PropagatesASetFlagsFailure) { + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.addr.sa.sa_family = AF_UNIX; // a UNIX stream socket rejects TCP_NODELAY + EXPECT_EQ(ncclOsSocketResetFd(&sock), ncclSystemError); + Track(sock.socketDescriptor); + EXPECT_TRUE(FdOpen(sock.socketDescriptor)); // created, then SetFlags failed +} + +TEST_F(LinuxSocketMicrotest, StartConnect_ConnectsToAListener) { + sockaddr_in addr; + Listen(&addr); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnected); +} + +TEST_F(LinuxSocketMicrotest, StartConnect_InProgressPollsUntilConnected) { + sockaddr_in addr; + Listen(&addr); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + ScopedHook connectHook(g_connect, [](int fd, const sockaddr* a, socklen_t l) { + fcntl(fd, F_SETFL, O_NONBLOCK); + int rc = ::connect(fd, a, l); + if (rc == 0) errno = EINPROGRESS; // loopback may finish at once + return -1; + }); + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnectPolling); + for (int i = 0; i < 1000 && sock.state == ncclSocketStateConnectPolling; i++) + ASSERT_EQ(ncclOsSocketPollConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnected); +} + +TEST_F(LinuxSocketMicrotest, StartConnect_RefusedIsRetriedWithAGrowingSleep) { + params_.Set("SOCKET_RETRY_CNT", 2); + sockaddr_in addr = ClosedPort(); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + int fd = sock.socketDescriptor; + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnecting); + EXPECT_EQ(sock.errorRetries, 1); + EXPECT_EQ(sock.socketDescriptor, fd); // reset in place + + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.errorRetries, 2); + EXPECT_EQ(ncclOsSocketStartConnect(&sock), ncclRemoteError); + EXPECT_EQ(sock.state, ncclSocketStateError); +} + +TEST_F(LinuxSocketMicrotest, StartConnect_SleepGrowsWithTheRetryCount) { + params_.Set("SOCKET_RETRY_CNT", 5); + params_.Set("SOCKET_RETRY_SLEEP_MSEC", 10); + sockaddr_in addr = ClosedPort(); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + sock.errorRetries = 2; + auto start = std::chrono::steady_clock::now(); + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.errorRetries, 3); + EXPECT_GE(std::chrono::steady_clock::now() - start, std::chrono::milliseconds(30)); // the third retry sleeps 3x +} + +TEST_F(LinuxSocketMicrotest, StartConnect_TransientErrnosAreRetried) { + sockaddr_in addr = ClosedPort(); + for (int err : {EINTR, EWOULDBLOCK, EAGAIN, ETIMEDOUT, EHOSTUNREACH, ECONNREFUSED}) { + SCOPED_TRACE(err); + ScopedHook connectHook(g_connect, FailWith(err)); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnecting); + EXPECT_EQ(sock.errorRetries, 1); + } +} + +// A caller with its own retry policy gets the reset but no count and no sleep. +TEST_F(LinuxSocketMicrotest, StartConnect_CustomRetryOnlyResets) { + params_.Set("SOCKET_RETRY_CNT", 0); + sockaddr_in addr = ClosedPort(); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + sock.customRetry = 1; + ASSERT_EQ(ncclOsSocketStartConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnecting); + EXPECT_EQ(sock.errorRetries, 0); +} + +TEST_F(LinuxSocketMicrotest, StartConnect_OtherErrorsFail) { + sockaddr_in addr; + Listen(&addr); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + sock.salen = 1; // EINVAL + EXPECT_EQ(ncclOsSocketStartConnect(&sock), ncclSystemError); + EXPECT_EQ(sock.state, ncclSocketStateError); +} + +TEST_F(LinuxSocketMicrotest, StartConnect_PropagatesAResetFailure) { + ScopedHook connectHook(g_connect, FailWith(ECONNREFUSED)); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0))); + sock.addr.sa.sa_family = AF_UNSPEC; // the replacement socket cannot be created + EXPECT_EQ(ncclOsSocketStartConnect(&sock), ncclSystemError); + EXPECT_EQ(sock.errorRetries, 1); // failed in the retry's reset, not as a connect error +} + +TEST_F(LinuxSocketMicrotest, PollConnect_NotReadyYetIsSuccess) { + ncclSocket sock = Sock(3); + sock.state = ncclSocketStateConnectPolling; + int pollRc = 0; + ScopedHook pollHook(g_poll, [&](pollfd* p, nfds_t, int timeout) { + EXPECT_EQ(p->events, POLLOUT); + EXPECT_EQ(timeout, 1); + errno = EINTR; + return pollRc; + }); + EXPECT_EQ(ncclOsSocketPollConnect(&sock), ncclSuccess); + pollRc = -1; // interrupted + EXPECT_EQ(ncclOsSocketPollConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnectPolling); +} + +TEST_F(LinuxSocketMicrotest, PollConnect_PollFailureIsASystemError) { + ncclSocket sock = Sock(3); + ScopedHook pollHook(g_poll, FailWith(ENOMEM)); + EXPECT_EQ(ncclOsSocketPollConnect(&sock), ncclSystemError); +} + +TEST_F(LinuxSocketMicrotest, PollConnect_ChecksThePendingSocketError) { + sockaddr_in addr = ClosedPort(); + ncclSocket sock = Sock(Track(::socket(AF_INET, SOCK_STREAM, 0)), &addr); + ScopedHook pollHook(g_poll, [](pollfd*, nfds_t, int) { return 1; }); + ScopedHook getsockoptHook(g_getsockopt, [](int, int level, int name, void* v, socklen_t*) { + EXPECT_EQ(level, SOL_SOCKET); + EXPECT_EQ(name, SO_ERROR); + *(int*)v = ECONNREFUSED; + return 0; + }); + ASSERT_EQ(ncclOsSocketPollConnect(&sock), ncclSuccess); + EXPECT_EQ(sock.state, ncclSocketStateConnecting); + + ScopedHook failingHook(g_getsockopt, FailWith(EBADF)); + EXPECT_EQ(ncclOsSocketPollConnect(&sock), ncclSystemError); +} + +TEST_F(LinuxSocketMicrotest, ProgressOpt_BlockingSendAndRecvMoveEverything) { + int sv[2]; + Pair(sv); + ncclSocket a = Sock(sv[0]), b = Sock(sv[1]); + char out[64], in[64] = {}; + for (int i = 0; i < 64; i++) out[i] = (char)i; + int offset = 0, closed = -1; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_SEND, &a, out, sizeof(out), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(offset, 64); + EXPECT_EQ(closed, 0); + offset = 16; // resumes at the offset + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(offset, 64); + EXPECT_EQ(std::memcmp(in + 16, out, 48), 0); +} + +// A blocking recv keeps reading until the buffer is full. +TEST_F(LinuxSocketMicrotest, ProgressOpt_BlockingRecvWaitsForTheRest) { + int sv[2]; + Pair(sv); + ncclSocket b = Sock(sv[1]); + std::vector flags; + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + flags.push_back(f); + ssize_t got = ::recv(fd, buf, n, f); + int err = errno; + if (flags.size() == 1) ::send(sv[0], "efgh", 4, 0); // the rest arrives only after the first read + errno = err; + return got; + }); + ::send(sv[0], "abcd", 4, 0); + char in[8]; + int offset = 0, closed = 0; + EXPECT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(flags, (std::vector{0, 0})); // two blocking reads + EXPECT_EQ(offset, 8); + EXPECT_EQ(std::string(in, 8), "abcdefgh"); +} + +TEST_F(LinuxSocketMicrotest, ProgressOpt_NonBlockingReturnsWhatIsThere) { + int sv[2]; + Pair(sv); + ncclSocket b = Sock(sv[1]); + char in[8]; + int offset = 0, closed = 0; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 0, &closed), ncclSuccess); + EXPECT_EQ(offset, 0); + ::send(sv[0], "abc", 3, 0); + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 0, &closed), ncclSuccess); + EXPECT_EQ(offset, 3); + EXPECT_EQ(closed, 0); +} + +// Async and abortable sockets never block, whatever the caller asks. +TEST_F(LinuxSocketMicrotest, ProgressOpt_AsyncOrAbortableNeverBlocks) { + int sv[2]; + Pair(sv); + volatile uint32_t abortFlag = 0; + char in[8]; + int offset = 0, closed = 0; + ncclSocket b = Sock(sv[1]); + b.asyncFlag = 1; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(offset, 0); + b = Sock(sv[1]); + b.abortFlag = &abortFlag; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(offset, 0); +} + +// A synchronous socket keeps reading while data comes in; an async one makes a single attempt. +TEST_F(LinuxSocketMicrotest, ProgressOpt_AsyncMakesOneAttempt) { + for (int async = 0; async < 2; async++) { + int sv[2]; + ASSERT_EQ(::socketpair(AF_UNIX, SOCK_SEQPACKET, 0, sv), 0); // one record per recv() + Track(sv[0]); + Track(sv[1]); + ::send(sv[1], "abcd", 4, 0); + ::send(sv[1], "efgh", 4, 0); + ncclSocket b = Sock(sv[0]); + b.asyncFlag = async; + char in[8]; + int offset = 0, closed = 0; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &b, in, sizeof(in), &offset, 0, &closed), ncclSuccess); + EXPECT_EQ(offset, async ? 4 : 8); + } +} + +TEST_F(LinuxSocketMicrotest, ProgressOpt_PeerCloseIsReportedNotFailed) { + int sv[2]; + Pair(sv); + ncclSocket a = Sock(sv[0]); + ::shutdown(sv[1], SHUT_RDWR); + char buf[4] = {}; + int offset = 0, closed = 0; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &a, buf, sizeof(buf), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(closed, 1); + closed = 0; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_SEND, &a, buf, sizeof(buf), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(closed, 1); // EPIPE, and no SIGPIPE + EXPECT_EQ(offset, 0); +} + +// Closing a UNIX stream socket with unread data resets the peer. +TEST_F(LinuxSocketMicrotest, ProgressOpt_ConnectionResetOnRecvIsAClose) { + int sv[2]; + ASSERT_EQ(::socketpair(AF_UNIX, SOCK_STREAM, 0, sv), 0); + Track(sv[0]); + ::send(sv[0], "x", 1, 0); + ::close(sv[1]); + ncclSocket a = Sock(sv[0]); + char buf[4]; + int offset = 0, closed = 0; + ASSERT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &a, buf, sizeof(buf), &offset, 1, &closed), ncclSuccess); + EXPECT_EQ(closed, 1); + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(LinuxSocketMicrotest, ProgressOpt_OtherErrorsAreRemoteErrors) { + int p[2]; + ASSERT_EQ(pipe(p), 0); + Track(p[0]); + Track(p[1]); + char buf[4] = {}; + int offset = 0, closed = 0; + ncclSocket sock = Sock(p[0]); + EXPECT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_RECV, &sock, buf, sizeof(buf), &offset, 1, &closed), ncclRemoteError); + sock = Sock(p[1]); + EXPECT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_SEND, &sock, buf, sizeof(buf), &offset, 1, &closed), ncclRemoteError); + EXPECT_EQ(closed, 0); +} + +TEST_F(LinuxSocketMicrotest, ProgressOpt_AbortStopsAfterTheTransfer) { + int sv[2]; + Pair(sv); + volatile uint32_t abortFlag = 1; + ncclSocket a = Sock(sv[0]); + a.abortFlag = &abortFlag; + int offset = 0, closed = 0; + EXPECT_EQ(ncclOsSocketProgressOpt(NCCL_SOCKET_SEND, &a, (void*)"abcd", 4, &offset, 1, &closed), ncclInternalError); + EXPECT_EQ(offset, 4); +} + +// Interface enumeration over a scripted getifaddrs() list. +class LinuxIfMicrotest : public LinuxMicrotest { + protected: + void TearDown() override { EXPECT_EQ(ifs_.Freed(), ifs_.Listed()); } + template + void Add(A... args) { + ifs_.Add(args...); + } + std::vector Find(const char* prefixList, int family = -1, int maxIfs = MAX_IFS) { + int found = -1; + EXPECT_EQ(ncclOsFindInterfaces(prefixList, names_, addrs_, family, MAX_IF_NAME_SIZE, maxIfs, &found), ncclSuccess); + std::vector out; + for (int i = 0; i < found; i++) out.push_back(names_ + i * MAX_IF_NAME_SIZE); + return out; + } + std::string Addr(int i) { + char host[INET6_ADDRSTRLEN] = ""; + if (addrs_[i].sa.sa_family == AF_INET) inet_ntop(AF_INET, &addrs_[i].sin.sin_addr, host, sizeof(host)); + else inet_ntop(AF_INET6, &addrs_[i].sin6.sin6_addr, host, sizeof(host)); + return host; + } + // The local interface in the remote address's subnet, or "". + std::string MatchSubnet(int family, const char* remote, uint32_t scope = 0) { + sockaddr_storage ss; + LinuxFakeIfaddrs::Fill(&ss, family, remote, scope); + union ncclSocketAddress remoteAddr, local; + std::memset(&local, 0, sizeof(local)); + std::memcpy(&remoteAddr, &ss, sizeof(remoteAddr)); + char ifName[MAX_IF_NAME_SIZE] = ""; + int found = -1; + EXPECT_EQ(ncclFindInterfaceMatchSubnet(ifName, &local, &remoteAddr, sizeof(ifName), &found), ncclSuccess); + if (!found) return ""; + addrs_[0] = local; + return ifName; + } + LinuxFakeIfaddrs ifs_; + char names_[MAX_IFS * MAX_IF_NAME_SIZE]; + union ncclSocketAddress addrs_[MAX_IFS]; +}; + +TEST_F(LinuxIfMicrotest, FindInterfaces_KeepsRunningIpInterfacesOncePerName) { + Add("noaddr", AF_UNSPEC); + Add("eth0", AF_INET, "10.0.0.1"); + Add("eth0", AF_INET6, "2001:db8::1"); // same name: only the first address is kept + Add("pkt0", AF_PACKET); + Add("down0", AF_INET, "10.0.0.2", nullptr, IFF_UP); // not RUNNING + Add("lo", AF_INET6, "::1"); // IPv6 loopback is skipped + Add("eth1", AF_INET6, "2001:db8::2"); + Add("lo", AF_INET, "127.0.0.1"); + EXPECT_EQ(Find(nullptr), (std::vector{"eth0", "eth1", "lo"})); + EXPECT_EQ(Addr(0), "10.0.0.1"); + EXPECT_EQ(addrs_[0].sin.sin_family, AF_INET); + EXPECT_EQ(Addr(1), "2001:db8::2"); + EXPECT_EQ(addrs_[1].sin6.sin6_family, AF_INET6); +} + +TEST_F(LinuxIfMicrotest, FindInterfaces_PrefixExclusionAndExactLists) { + Add("eth0", AF_INET, "10.0.0.1"); + Add("eth10", AF_INET, "10.0.0.2"); + Add("ib0", AF_INET, "10.0.1.1"); + Add("lo", AF_INET, "127.0.0.1"); + EXPECT_EQ(Find("eth"), (std::vector{"eth0", "eth10"})); + EXPECT_EQ(Find("ib,lo"), (std::vector{"ib0", "lo"})); + EXPECT_EQ(Find("^eth"), (std::vector{"ib0", "lo"})); + EXPECT_EQ(Find("=eth0"), (std::vector{"eth0"})); + EXPECT_EQ(Find("=eth"), (std::vector{})); + EXPECT_EQ(Find("^=eth0"), (std::vector{"eth10", "ib0", "lo"})); +} + +TEST_F(LinuxIfMicrotest, FindInterfaces_FamilyFilterAndMaxIfs) { + Add("eth0", AF_INET, "10.0.0.1"); + Add("eth1", AF_INET6, "2001:db8::2"); + Add("eth2", AF_INET, "10.0.0.3"); + EXPECT_EQ(Find(nullptr, AF_INET6), (std::vector{"eth1"})); + EXPECT_EQ(Find(nullptr, AF_INET), (std::vector{"eth0", "eth2"})); + EXPECT_EQ(Find(nullptr, -1, 2), (std::vector{"eth0", "eth1"})); +} + +// The address union is cleared past the stored sockaddr_in. +TEST_F(LinuxIfMicrotest, FindInterfaces_ClearsTheStoredAddress) { + Add("eth0", AF_INET, "10.0.0.1"); + std::memset(addrs_, 0xff, sizeof(addrs_)); + ASSERT_EQ(Find(nullptr).size(), 1u); + const char* tail = (const char*)&addrs_[0] + sizeof(sockaddr_in); + for (size_t i = 0; i < sizeof(addrs_[0]) - sizeof(sockaddr_in); i++) EXPECT_EQ(tail[i], 0); +} + +TEST_F(LinuxIfMicrotest, FindInterfaces_GetifaddrsFailureIsASystemError) { + ScopedHook getifaddrsHook(g_getifaddrs, FailWith(ENOMEM)); + int found = -1; + EXPECT_EQ(ncclOsFindInterfaces(nullptr, names_, addrs_, -1, MAX_IF_NAME_SIZE, MAX_IFS, &found), ncclSystemError); + EXPECT_EQ(found, 0); + EXPECT_EQ(ncclFindInterfaceMatchSubnet(names_, addrs_, addrs_, MAX_IF_NAME_SIZE, &found), ncclSystemError); +} + +TEST_F(LinuxIfMicrotest, MatchSubnet_IPv4UsesTheNetmask) { + Add("noaddr", AF_UNSPEC); + Add("pkt0", AF_PACKET); + Add("eth1", AF_INET6, "2001:db8::1", "ffff:ffff:ffff:ffff::"); + Add("eth0", AF_INET, "10.0.0.1", "255.255.0.0"); + Add("eth2", AF_INET, "10.0.9.9", "255.0.0.0"); + EXPECT_EQ(MatchSubnet(AF_INET, "10.0.200.3"), "eth0"); // first match wins + EXPECT_EQ(Addr(0), "10.0.0.1"); + EXPECT_EQ(MatchSubnet(AF_INET, "10.7.0.1"), "eth2"); + EXPECT_EQ(MatchSubnet(AF_INET, "11.0.0.1"), ""); +} + +TEST_F(LinuxIfMicrotest, MatchSubnet_IPv6ComparesThePrefixAndTheScope) { + Add("eth0", AF_INET, "10.0.0.1", "255.0.0.0"); + Add("eth1", AF_INET6, "2001:db8:0:1::1", "ffff:ffff:ffff:ffff::"); + Add("eth2", AF_INET6, "fe80::1", "ffff:ffff:ffff:ffff::", IFF_UP | IFF_RUNNING, 2); + EXPECT_EQ(MatchSubnet(AF_INET6, "2001:db8:0:1::99"), "eth1"); + EXPECT_EQ(Addr(0), "2001:db8:0:1::1"); + EXPECT_EQ(MatchSubnet(AF_INET6, "2001:db8:0:2::1"), ""); // differs in the last prefix byte + EXPECT_EQ(MatchSubnet(AF_INET6, "fe80::5", 2), "eth2"); + EXPECT_EQ(MatchSubnet(AF_INET6, "fe80::5", 3), ""); // same prefix, other link +} + +TEST_F(LinuxIfMicrotest, RealGetifaddrsFindsLoopback) { + ScopedHook realList(g_getifaddrs, ::getifaddrs); + ScopedHook realFree(g_freeifaddrs, ::freeifaddrs); + int found = 0; + ASSERT_EQ(ncclOsFindInterfaces("=lo", names_, addrs_, AF_INET, MAX_IF_NAME_SIZE, MAX_IFS, &found), ncclSuccess); + ASSERT_EQ(found, 1); + EXPECT_STREQ(names_, "lo"); + EXPECT_EQ(Addr(0), "127.0.0.1"); +} + +// Socket close + +using LinuxSocketCloseMicrotest = LinuxSocketMicrotest; + +TEST_F(LinuxSocketCloseMicrotest, ClosesAnOpenSocket) { + EXPECT_EQ(ncclSocketClose(nullptr), ncclSuccess); + int sv[2]; + Pair(sv); + ncclSocket sock = Sock(sv[0]); + sock.state = ncclSocketStateReady; + ASSERT_EQ(ncclSocketClose(&sock), ncclSuccess); + EXPECT_FALSE(FdOpen(sv[0])); + EXPECT_EQ(sock.state, ncclSocketStateClosed); + EXPECT_EQ(sock.socketDescriptor, NCCL_INVALID_SOCKET); + char c; + EXPECT_EQ(::recv(sv[1], &c, 1, 0), 0); // the peer saw the FIN +} + +// shutdown() reaches the peer even if another descriptor still references the socket. +TEST_F(LinuxSocketCloseMicrotest, ShutsDownADuplicatedSocket) { + int sv[2]; + Pair(sv); + int dupFd = Track(dup(sv[0])); + ncclSocket sock = Sock(sv[0]); + sock.state = ncclSocketStateReady; + ASSERT_EQ(ncclSocketClose(&sock), ncclSuccess); + char c; + EXPECT_EQ(::recv(sv[1], &c, 1, 0), 0); + EXPECT_TRUE(FdOpen(dupFd)); +} + +TEST_F(LinuxSocketCloseMicrotest, LeavesTheDescriptorOfANoneOrOutOfRangeState) { + for (ncclSocketState state : {ncclSocketStateNone, ncclSocketStateNum}) { + int sv[2]; + Pair(sv); + ncclSocket sock = Sock(sv[0]); + sock.state = state; + ASSERT_EQ(ncclSocketClose(&sock), ncclSuccess); + EXPECT_TRUE(FdOpen(sv[0])); + EXPECT_EQ(sock.state, ncclSocketStateClosed); + EXPECT_EQ(sock.socketDescriptor, NCCL_INVALID_SOCKET); + } + ncclSocket sock = Sock(NCCL_INVALID_SOCKET); + sock.state = ncclSocketStateReady; + EXPECT_EQ(ncclSocketClose(&sock), ncclSuccess); +} + +TEST_F(LinuxSocketCloseMicrotest, WaitDrainsUntilThePeerCloses) { + int sv[2]; + Pair(sv); + ::send(sv[1], "leftover", 8, 0); + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + int avail = 0; + ioctl(fd, FIONREAD, &avail); + if (avail == 0) ::shutdown(sv[1], SHUT_WR); // the peer closes once the leftover is drained + return ::recv(fd, buf, n, f); + }); + ncclSocket sock = Sock(sv[0]); + sock.state = ncclSocketStateReady; + EXPECT_EQ(ncclSocketClose(&sock, /*wait=*/true), ncclSuccess); + EXPECT_EQ(recvHook.calls, 9); // eight bytes, then the close + EXPECT_FALSE(FdOpen(sv[0])); +} + +#ifndef NCCL_TLS_BACKEND_OPENSSL3 // the fake TLS state below is only safe without a backend +TEST_F(LinuxSocketCloseMicrotest, WaitStopsOnAReadErrorAndFreesTlsState) { + int sv[2]; + Pair(sv); + ncclSocket sock = Sock(sv[0]); + sock.state = ncclSocketStateReady; + sock.crypto = (ncclSocketCrypto*)malloc(16); // reads now fail: no TLS backend + ScopedHook recvHook(g_recv, [](int, void*, size_t, int) -> ssize_t { return 0; }); // a plain read sees a close + ASSERT_EQ(ncclSocketClose(&sock, /*wait=*/true), ncclSuccess); + EXPECT_EQ(sock.crypto, nullptr); + EXPECT_FALSE(FdOpen(sv[0])); + EXPECT_EQ(recvHook.calls, 0); // the read went through crypt.cc +} +#endif + +// Shared mutex/condition variable and CPU sets + +using LinuxSyncMicrotest = LinuxMicrotest; + +TEST_F(LinuxSyncMicrotest, MutexCondShared_HonoursTheInitializedFlag) { + std::mutex m; + std::condition_variable cv; + int initialized = 1; + ncclOsSetMutexCondShared(m, cv, &initialized); // already set: untouched + EXPECT_EQ(initialized, 1); + initialized = 0; + ncclOsUnsetMutexCondShared(m, cv, &initialized); // not set: untouched + EXPECT_EQ(initialized, 0); + ncclOsSetMutexCondShared(m, cv, &initialized); + EXPECT_EQ(initialized, 1); + { std::lock_guard lock(m); } + ncclOsUnsetMutexCondShared(m, cv, &initialized); + EXPECT_EQ(initialized, 0); + ncclOsSetMutexCondShared(m, cv, nullptr); + ncclOsUnsetMutexCondShared(m, cv, nullptr); +} + +// A process-private condition variable would not wake a waiter in another process. +TEST_F(LinuxSyncMicrotest, MutexCondShared_WakesAWaiterInAnotherProcess) { + struct Shared { + std::mutex m; + std::condition_variable cv; + int ready = 0; + }; + void* mem = ::mmap(nullptr, sizeof(Shared), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, -1, 0); + ASSERT_NE(mem, MAP_FAILED); + Shared* s = new (mem) Shared; + ncclOsSetMutexCondShared(s->m, s->cv, nullptr); + const std::string waiterStat = "/proc/" + std::to_string(getpid()) + "/task/" + std::to_string(syscall(SYS_gettid)) + + "/stat"; + auto start = std::chrono::steady_clock::now(); + std::unique_lock lock(s->m); // held across fork: the child gets it only once we wait + pid_t child = fork(); + if (child == 0) { + { + std::lock_guard childLock(s->m); + // Notify only once the waiter sleeps in the kernel, where only a shared condition variable reaches it. + for (int i = 0; i < 1000000 && !ThreadSleeps(waiterStat); i++) sched_yield(); + s->ready = 1; + } + s->cv.notify_all(); + _exit(0); + } + // Likewise release the mutex only once the child sleeps on it, where only a shared mutex reaches it. + const std::string childStat = "/proc/" + std::to_string(child) + "/stat"; + for (int i = 0; i < 1000000 && !ThreadSleeps(childStat); i++) sched_yield(); + while (!s->ready && std::chrono::steady_clock::now() - start < std::chrono::seconds(10)) + s->cv.wait_for(lock, std::chrono::seconds(5)); + lock.unlock(); + auto waited = std::chrono::steady_clock::now() - start; + if (!s->ready) kill(child, SIGKILL); + int status = 0; + waitpid(child, &status, 0); + EXPECT_EQ(s->ready, 1); + EXPECT_LT(waited, std::chrono::seconds(4)); + ncclOsUnsetMutexCondShared(s->m, s->cv, nullptr); + ::munmap(mem, sizeof(Shared)); +} + +TEST_F(LinuxSyncMicrotest, CpuSetOperations) { + ncclAffinity a, b; + std::memset(&a, 0xff, sizeof(a)); + ncclOsCpuZero(a); + EXPECT_EQ(ncclOsCpuCount(a), 0); + ncclOsCpuSet(a, 1); + ncclOsCpuSet(a, 3); + ncclOsCpuSet(a, 200); + EXPECT_EQ(ncclOsCpuCount(a), 3); + EXPECT_TRUE(ncclOsCpuIsSet(a, 200)); + EXPECT_FALSE(ncclOsCpuIsSet(a, 2)); + ncclOsCpuZero(b); + ncclOsCpuSet(b, 3); + ncclOsCpuSet(b, 4); + ncclAffinity both = ncclOsCpuAnd(a, b); + EXPECT_EQ(ncclOsCpuCount(both), 1); + EXPECT_TRUE(ncclOsCpuIsSet(both, 3)); +} + +TEST_F(LinuxSyncMicrotest, Affinity_GetAndSetTheCallingThread) { + ncclAffinity mine, expected; + ASSERT_EQ(ncclOsGetAffinity(&mine), ncclSuccess); + ASSERT_EQ(sched_getaffinity(0, sizeof(expected), &expected), 0); + EXPECT_TRUE(CPU_EQUAL(&mine, &expected)); + EXPECT_EQ(ncclOsSetAffinity(mine), ncclSuccess); + int cpu = ncclOsGetCpu(); + EXPECT_GE(cpu, 0); + EXPECT_TRUE(ncclOsCpuIsSet(mine, cpu)); +} + +TEST_F(LinuxSyncMicrotest, Affinity_FailuresAreSystemErrors) { + ncclAffinity none; + CPU_ZERO(&none); + EXPECT_EQ(ncclOsSetAffinity(none), ncclSystemError); // EINVAL + ScopedHook schedGetaffinityHook(g_schedGetaffinity, FailWith(EFAULT)); + EXPECT_EQ(ncclOsGetAffinity(&none), ncclSystemError); +} + +// sysfs readers, over a scratch tree that the /sys hooks redirect to. + +class LinuxSysfsMicrotest : public LinuxMicrotest { + protected: + void SetUp() override { + char tmpl[] = "/tmp/nccl-linux-sysfs-XXXXXX"; + ASSERT_NE(mkdtemp(tmpl), nullptr); + root_ = tmpl; + } + void TearDown() override { + std::error_code ec; + if (!root_.empty()) std::filesystem::remove_all(root_, ec); + } + std::string Redirect(const char* path) const { + std::string p = path; + return p.rfind("/sys/", 0) == 0 ? root_ + p : p; + } + void Mkdirs(const std::string& rel) { + std::string path = root_; + size_t pos = 0; + while (pos != std::string::npos) { + pos = rel.find('/', pos + 1); + path = root_ + rel.substr(0, pos); + mkdir(path.c_str(), 0755); + } + } + void Write(const std::string& rel, const std::string& content) { + FILE* f = ::fopen((root_ + rel).c_str(), "w"); + ASSERT_NE(f, nullptr); + fwrite(content.data(), 1, content.size(), f); + fclose(f); + } + // A PCI device the way sysfs lays it out: /sys/class/pci_bus/ links into /sys/devices. + std::string AddPci(const std::string& busId) { + std::string bus = busId.substr(0, 7), dev = "/sys/devices/pci" + bus; + Mkdirs(dev + "/pci_bus/" + bus); + Mkdirs(dev + "/" + busId); + Mkdirs("/sys/class/pci_bus"); + symlink(("../../devices/pci" + bus + "/pci_bus/" + bus).c_str(), (root_ + "/sys/class/pci_bus/" + bus).c_str()); + return root_ + dev + "/" + busId; + } + std::string root_; + std::vector seenPaths_; + // The /sys hooks: production's sysfs reads land in root_. + ScopedHook realpath_{g_realpath, [this](const char* p, char* out) { + seenPaths_.push_back(p); + return ::realpath(Redirect(p).c_str(), out); + }}; + ScopedHook fopen_{ + g_fopen, [this](const char* p, const char* mode) { return ::fopen(Redirect(p).c_str(), mode); }}; + ScopedHook opendir_{g_opendir, [this](const char* p) { return ::opendir(Redirect(p).c_str()); }}; +}; + +TEST_F(LinuxSysfsMicrotest, GetPciPath_LowercasesTheBusIdAndResolves) { + std::string dev = AddPci("0000:1a:00.0"); + char* path = nullptr; + ASSERT_EQ(ncclOsGetPciPath("0000:1A:00.0", &path), ncclSuccess); + EXPECT_EQ(std::string(path), dev); + free(path); + ASSERT_EQ(seenPaths_.size(), 1u); + EXPECT_EQ(seenPaths_[0], "/sys/class/pci_bus/0000:1a/../../0000:1a:00.0"); +} + +TEST_F(LinuxSysfsMicrotest, GetPciPath_MissingDeviceIsASystemError) { + char* path = (char*)0x1; + EXPECT_EQ(ncclOsGetPciPath("0000:1b:00.0", &path), ncclSystemError); + EXPECT_EQ(path, nullptr); +} + +TEST_F(LinuxSysfsMicrotest, TopoGetStrFromSys_StripsTheNewlineAndTruncates) { + Mkdirs("/d"); + Write("/d/nl", "0x030000\n"); + Write("/d/empty", ""); + char buf[16]; + std::string dir = root_ + "/d"; + ASSERT_EQ(ncclOsTopoGetStrFromSys(dir.c_str(), "nl", buf, sizeof(buf)), ncclSuccess); + EXPECT_STREQ(buf, "0x030000"); + ASSERT_EQ(ncclOsTopoGetStrFromSys(dir.c_str(), "nl", buf, 4), ncclSuccess); // truncated to maxLen + EXPECT_STREQ(buf, "0x0"); + ASSERT_EQ(ncclOsTopoGetStrFromSys(dir.c_str(), "empty", buf, sizeof(buf)), ncclSuccess); + EXPECT_STREQ(buf, ""); + ASSERT_EQ(ncclOsTopoGetStrFromSys(dir.c_str(), "missing", buf, sizeof(buf)), ncclSuccess); + EXPECT_STREQ(buf, ""); +} + +// A value without a trailing newline loses its last character. +TEST_F(LinuxSysfsMicrotest, TopoGetStrFromSys_DropsTheLastCharacter) { + Mkdirs("/d"); + Write("/d/raw", "abc"); + char buf[16]; + std::string dir = root_ + "/d"; + ASSERT_EQ(ncclOsTopoGetStrFromSys(dir.c_str(), "raw", buf, sizeof(buf)), ncclSuccess); + EXPECT_STREQ(buf, "ab"); +} + +TEST_F(LinuxSysfsMicrotest, DeviceClass_ReadsTheClassAttribute) { + std::string dev = AddPci("0000:1a:00.0"); + FILE* f = ::fopen((dev + "/class").c_str(), "w"); + fputs("0x038000\n", f); + fclose(f); + char cls[32] = "junk"; + ASSERT_EQ(ncclOsGetPciDeviceClassByBusId("0000:1a:00.0", cls, sizeof(cls)), ncclSuccess); + EXPECT_STREQ(cls, "0x038000"); +} + +TEST_F(LinuxSysfsMicrotest, DeviceClass_MissingDeviceIsAQuietSystemError) { + char cls[32] = "junk"; + EXPECT_EQ(ncclOsGetPciDeviceClassByBusId("0000:1b:00.0", cls, sizeof(cls)), ncclSystemError); + EXPECT_STREQ(cls, ""); + EXPECT_EQ(log_.Warns(), 0); +} + +TEST_F(LinuxSysfsMicrotest, ComputePartition_ReadsTheAttributeOrEmpty) { + std::string dev = AddPci("0000:1a:00.0"); + char part[16] = "junk"; + ASSERT_EQ(ncclOsGetPciDeviceComputePartitionByBusId("0000:1a:00.0", part, sizeof(part)), ncclSuccess); + EXPECT_STREQ(part, ""); // attribute absent + FILE* f = ::fopen((dev + "/current_compute_partition").c_str(), "w"); + fputs("CPX\n", f); + fclose(f); + ASSERT_EQ(ncclOsGetPciDeviceComputePartitionByBusId("0000:1a:00.0", part, sizeof(part)), ncclSuccess); + EXPECT_STREQ(part, "CPX"); +} + +TEST_F(LinuxSysfsMicrotest, ComputePartition_BadArgumentsAndMissingDevice) { + char part[16] = "junk"; + EXPECT_EQ(ncclOsGetPciDeviceComputePartitionByBusId("0000:1a:00.0", nullptr, 16), ncclInvalidArgument); + EXPECT_EQ(ncclOsGetPciDeviceComputePartitionByBusId("0000:1a:00.0", part, 0), ncclInvalidArgument); + EXPECT_STREQ(part, "junk"); + EXPECT_EQ(ncclOsGetPciDeviceComputePartitionByBusId("0000:1b:00.0", part, sizeof(part)), ncclSystemError); + EXPECT_STREQ(part, ""); + EXPECT_EQ(log_.Warns(), 0); +} + +TEST_F(LinuxSysfsMicrotest, BcmLinks_ListsPeersThatArePciDevices) { + AddPci("0000:2b:00.0"); + AddPci("0000:2c:00.0"); + std::string dir = "/sys/kernel/pci_switch_link/virtual_switch_links/0000:1a:00.0"; + Mkdirs(dir + "/0000:2b:00.0"); + Mkdirs(dir + "/0000:2c:00.0"); + Mkdirs(dir + "/0000:3d:00.0"); // no such PCI device + Mkdirs(dir + "/short"); + int n = -1; + char* peers = nullptr; + ASSERT_EQ(ncclOsGetBcmLinks("0000:1A:00.0", &n, &peers), ncclSuccess); + ASSERT_EQ(n, 2); + std::set got{peers, peers + sizeof("0000:00:00.0")}; + EXPECT_EQ(got, (std::set{"0000:2b:00.0", "0000:2c:00.0"})); + free(peers); +} + +TEST_F(LinuxSysfsMicrotest, BcmLinks_NoSwitchDirectoryMeansNoLinks) { + int n = -1; + char* peers = (char*)0x1; + ASSERT_EQ(ncclOsGetBcmLinks("0000:1a:00.0", &n, &peers), ncclSuccess); + EXPECT_EQ(n, 0); + EXPECT_EQ(peers, nullptr); +} + +TEST_F(LinuxSysfsMicrotest, NumaNodeAffinity_StripsLineEndings) { + Mkdirs("/sys/devices/system/node/node3"); + Write("/sys/devices/system/node/node3/cpumap", "00000000,0000ffff\r\n"); + char str[64]; + int offset = -1; + ASSERT_EQ(ncclOsGetNumaNodeAffinity(3, str, sizeof(str), &offset), ncclSuccess); + EXPECT_STREQ(str, "00000000,0000ffff"); + EXPECT_EQ(offset, 0); + ASSERT_EQ(ncclOsGetNumaNodeAffinity(3, str, 6, &offset), ncclSuccess); // maxLen counts the NUL + EXPECT_STREQ(str, "00000"); +} + +TEST_F(LinuxSysfsMicrotest, NumaNodeAffinity_MissingNodeIsEmpty) { + char str[64] = "junk"; + int offset = -1; + ASSERT_EQ(ncclOsGetNumaNodeAffinity(7, str, sizeof(str), &offset), ncclSuccess); + EXPECT_STREQ(str, ""); +} + +// Shared memory, on real files. + +class LinuxShmMicrotest : public LinuxMicrotest { + protected: + void SetUp() override { + char tmpl[] = "/tmp/nccl-linux-shm-XXXXXX"; + ASSERT_NE(mkdtemp(tmpl), nullptr); + dir_ = tmpl; + } + void TearDown() override { + for (auto* h : handles_) ncclOsShmClose(h); + for (auto& p : created_) ::unlink(p.c_str()); + std::error_code ec; + if (!dir_.empty()) std::filesystem::remove_all(dir_, ec); + } + ncclResult_t Open(char* path, size_t size, int refcount, void** ptr, ncclShmHandleInternal** h, + void** dev = nullptr) { + ncclResult_t ret = ncclOsShmOpen(path, PATH_MAX, size, ptr, dev, refcount, h); + if (*h) handles_.push_back(*h); + if (refcount > 0 && path[0]) created_.push_back(path); + return ret; + } + void Forget(ncclShmHandleInternal* h) { handles_.erase(std::find(handles_.begin(), handles_.end(), h)); } + std::string dir_; + std::vector handles_; + std::vector created_; + // HIP host registration succeeds unless a test says otherwise. + ScopedHook capture_{g_hipThreadExchangeStreamCaptureMode, + [](hipStreamCaptureMode*) { return hipSuccess; }}; + ScopedHook register_{g_hipHostRegister, + [](void*, size_t, unsigned int) { return hipSuccess; }}; + ScopedHook devPtr_{g_hipHostGetDevicePointer, + [](void** d, void* h, unsigned int) { + *d = h; + return hipSuccess; + }}; + ScopedHook unregister_{g_hipHostUnregister, [](void*) { return hipSuccess; }}; +}; + +TEST_F(LinuxShmMicrotest, CreateInDevShmThenAttachUntilTheLastPeerUnlinks) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 4096, /*refcount=*/2, &ptr, &h), ncclSuccess); + EXPECT_EQ(std::string(path).rfind("/dev/shm/nccl-", 0), 0u); + EXPECT_TRUE(Exists(path)); + ASSERT_NE(ptr, nullptr); + EXPECT_EQ(h->shmPtr, ptr); + EXPECT_EQ(h->shmSize, 4096u); + EXPECT_EQ(h->realShmSize, 4096u + sizeof(int)); + EXPECT_EQ(h->refcount, (int*)((char*)ptr + 4096)); + EXPECT_EQ(*h->refcount, 2); + EXPECT_STREQ(h->shmPath, path); + EXPECT_NE(h->shmPath, path); // a copy + EXPECT_EQ(((char*)ptr)[0], 0); + std::strcpy((char*)ptr, "shared"); + + void* peer1 = nullptr; + ncclShmHandleInternal* h1 = nullptr; + ASSERT_EQ(Open(path, 4096, 0, &peer1, &h1), ncclSuccess); + EXPECT_STREQ((char*)peer1, "shared"); + EXPECT_EQ(h1->shmPath, nullptr); + EXPECT_EQ(*h->refcount, 1); + EXPECT_TRUE(Exists(path)); + + void* peer2 = nullptr; + ncclShmHandleInternal* h2 = nullptr; + ASSERT_EQ(Open(path, 4096, 0, &peer2, &h2), ncclSuccess); + EXPECT_EQ(*h->refcount, 0); + EXPECT_FALSE(Exists(path)); // every peer attached + + EXPECT_EQ(ncclOsShmClose(h), ncclSuccess); // refcount 0: no second unlink + Forget(h); + EXPECT_EQ(log_.Warns(), 0); +} + +TEST_F(LinuxShmMicrotest, CreateAtAGivenPath) { + std::string p = dir_ + "/seg"; + char path[PATH_MAX]; + std::strcpy(path, p.c_str()); + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 100, 1, &ptr, &h), ncclSuccess); + EXPECT_STREQ(path, p.c_str()); + struct stat st; + ASSERT_EQ(stat(path, &st), 0); + EXPECT_EQ((size_t)st.st_size, 100 + sizeof(int)); + EXPECT_EQ(st.st_mode & 0777, (mode_t)(S_IRUSR | S_IWUSR)); +} + +// open() does not truncate, so a creator reusing a path must clear what the file held. +TEST_F(LinuxShmMicrotest, CreateClearsAReusedFile) { + std::string p = dir_ + "/seg"; + FILE* f = ::fopen(p.c_str(), "w"); + fputs("stale contents", f); + fclose(f); + char path[PATH_MAX]; + std::strcpy(path, p.c_str()); + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + EXPECT_EQ(std::string((char*)ptr, 14), std::string(14, '\0')); +} + +// Closing a creator handle while peers have not attached yet removes the file (abort cleanup). +TEST_F(LinuxShmMicrotest, CloseUnlinksWhilePeersAreStillExpected) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + Forget(h); + EXPECT_EQ(ncclOsShmClose(h), ncclSuccess); + EXPECT_FALSE(Exists(path)); + EXPECT_EQ(ncclOsShmClose(nullptr), ncclSuccess); +} + +TEST_F(LinuxShmMicrotest, CloseReportsUnlinkAndMunmapFailures) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + ::unlink(path); + ScopedHook munmapHook(g_munmap, [](void* p, size_t n) { + ::munmap(p, n); + errno = EINVAL; + return -1; + }); + Forget(h); + EXPECT_EQ(ncclOsShmClose(h), ncclSystemError); + EXPECT_EQ(log_.Warns(), 2); // both failures, not just the first +} + +TEST_F(LinuxShmMicrotest, UnlinkRemovesTheFileOnce) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + ASSERT_EQ(ncclOsShmUnlink(h), ncclSuccess); + EXPECT_FALSE(Exists(path)); + EXPECT_EQ(h->shmPath, nullptr); + EXPECT_EQ(ncclOsShmUnlink(h), ncclSuccess); // nothing left to unlink + EXPECT_EQ(ncclOsShmUnlink(nullptr), ncclSuccess); + EXPECT_EQ(log_.Warns(), 0); +} + +TEST_F(LinuxShmMicrotest, UnlinkSkipsAFullyAttachedSegmentAndReportsFailures) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + *h->refcount = 0; + EXPECT_EQ(ncclOsShmUnlink(h), ncclSuccess); + EXPECT_TRUE(Exists(path)); + *h->refcount = 1; + ::unlink(path); + EXPECT_EQ(ncclOsShmUnlink(h), ncclSystemError); +} + +TEST_F(LinuxShmMicrotest, OpenFailuresReturnNoHandle) { + std::string missing = dir_ + "/no/such/dir/seg"; + char path[PATH_MAX]; + std::strcpy(path, missing.c_str()); + void* ptr = (void*)0x1; + ncclShmHandleInternal* h = (ncclShmHandleInternal*)0x1; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); + EXPECT_EQ(ptr, nullptr); + EXPECT_EQ(Open(path, 64, 0, &ptr, &h), ncclSystemError); +} + +TEST_F(LinuxShmMicrotest, InterruptedMkstempAndFallocateAreRetried) { + int mkstemps = 0, fallocates = 0; + ScopedHook mkstempHook(g_mkstemp, [&](char* tmpl) { + if (mkstemps++ == 0) { + errno = EINTR; + return -1; + } + return ::mkstemp(tmpl); + }); + ScopedHook fallocateHook(g_fallocate, [&](int fd, int mode, off_t off, off_t len) { + if (fallocates++ == 0) { + errno = EINTR; + return -1; + } + return ::fallocate(fd, mode, off, len); + }); + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + EXPECT_EQ(mkstemps, 2); + EXPECT_EQ(fallocates, 2); +} + +TEST_F(LinuxShmMicrotest, MkstempFailure) { + ScopedHook mkstempHook(g_mkstemp, FailWith(EACCES)); + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); +} + +// Close only unlinks once a refcount is mapped, so a failure after creation leaves the file. +// Pins AICOMRCCL-2974; a fix flips this pin. +TEST_F(LinuxShmMicrotest, ShmOpen_KnownBug_FallocateFailureLeavesTheCreatedFile) { + ScopedHook fallocateHook(g_fallocate, FailWith(ENOSPC)); + std::string p = dir_ + "/seg"; + char path[PATH_MAX]; + std::strcpy(path, p.c_str()); + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); + EXPECT_TRUE(Exists(p)); +} + +TEST_F(LinuxShmMicrotest, MmapAndCallocFailures) { + ScopedHook mmapHook(g_mmap, [](void*, size_t, int, int, int, off_t) { + errno = ENOMEM; + return MAP_FAILED; + }); + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); + EXPECT_EQ(ptr, nullptr); + ::unlink(path); + + ScopedHook callocHook(g_calloc, [](size_t, size_t) -> void* { return nullptr; }); + path[0] = '\0'; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); + EXPECT_STREQ(path, ""); // failed before creating anything +} + +TEST_F(LinuxShmMicrotest, LastAttachUnlinkFailureIsOnlyLogged) { + char path[PATH_MAX] = ""; + void* ptr = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h), ncclSuccess); + ScopedHook unlinkHook(g_unlink, FailWith(EBUSY)); + void* peer = nullptr; + ncclShmHandleInternal* hp = nullptr; + EXPECT_EQ(Open(path, 64, 0, &peer, &hp), ncclSuccess); + EXPECT_EQ(unlinkHook.calls, 1); + EXPECT_EQ(log_.Warns(), 0); +} + +TEST_F(LinuxShmMicrotest, DevicePointerRegistersTheMappingWithHip) { + std::vector calls; + unsigned int regFlags = 0; + ScopedHook hipExchangeCaptureModeHook(g_hipThreadExchangeStreamCaptureMode, [&](hipStreamCaptureMode* m) { + calls.push_back("exchange"); + EXPECT_EQ(*m, hipStreamCaptureModeRelaxed); + return hipSuccess; + }); + ScopedHook hipHostRegisterHook(g_hipHostRegister, [&](void*, size_t size, unsigned int flags) { + calls.push_back("register"); + regFlags = flags; + EXPECT_EQ(size, 64 + sizeof(int)); + return hipSuccess; + }); + ScopedHook hipHostGetDevicePointerHook(g_hipHostGetDevicePointer, [&](void** d, void*, unsigned int) { + calls.push_back("devptr"); + *d = (void*)0xd00d; + return hipSuccess; + }); + ScopedHook hipHostUnregisterHook(g_hipHostUnregister, [&](void*) { + calls.push_back("unregister"); + return hipSuccess; + }); + char path[PATH_MAX] = ""; + void *ptr = nullptr, *dev = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h, &dev), ncclSuccess); + EXPECT_EQ(dev, (void*)0xd00d); + EXPECT_EQ(h->devShmPtr, (void*)0xd00d); + EXPECT_EQ(regFlags & (hipHostRegisterPortable | hipHostRegisterMapped), + (unsigned)(hipHostRegisterPortable | hipHostRegisterMapped)); + Forget(h); + EXPECT_EQ(ncclOsShmClose(h), ncclSuccess); + EXPECT_EQ(calls, (std::vector{"exchange", "register", "devptr", "exchange", "unregister"})); +} + +TEST_F(LinuxShmMicrotest, HipFailuresUnwindTheSegment) { + auto openFails = [&] { + char path[PATH_MAX] = ""; + void *ptr = nullptr, *dev = (void*)0x1; + ncclShmHandleInternal* h = nullptr; + EXPECT_EQ(Open(path, 64, 1, &ptr, &h, &dev), ncclUnhandledCudaError); + EXPECT_EQ(h, nullptr); + EXPECT_EQ(dev, nullptr); + EXPECT_FALSE(Exists(path)); + }; + { + SCOPED_TRACE("capture mode"); + ScopedHook fail(g_hipThreadExchangeStreamCaptureMode, [](hipStreamCaptureMode*) { return hipErrorInvalidValue; }); + openFails(); + } + { + SCOPED_TRACE("register"); + ScopedHook fail(g_hipHostRegister, [](void*, size_t, unsigned int) { return hipErrorInvalidValue; }); + openFails(); + } + { + SCOPED_TRACE("device pointer"); + ScopedHook fail(g_hipHostGetDevicePointer, [](void**, void*, unsigned int) { return hipErrorInvalidValue; }); + openFails(); + } + { + SCOPED_TRACE("restoring the capture mode"); + int exchanges = 0; + ScopedHook fail(g_hipThreadExchangeStreamCaptureMode, + [&](hipStreamCaptureMode*) { return exchanges++ ? hipErrorInvalidValue : hipSuccess; }); + openFails(); + } +} + +TEST_F(LinuxShmMicrotest, CloseReportsAHipUnregisterFailure) { + ScopedHook hipHostUnregisterHook(g_hipHostUnregister, [](void*) { return hipErrorInvalidValue; }); + char path[PATH_MAX] = ""; + void *ptr = nullptr, *dev = nullptr; + ncclShmHandleInternal* h = nullptr; + ASSERT_EQ(Open(path, 64, 1, &ptr, &h, &dev), ncclSuccess); + Forget(h); + void* mapping = h->shmPtr; + size_t size = h->realShmSize; + EXPECT_EQ(ncclOsShmClose(h), ncclUnhandledCudaError); // returns before munmap and free + ::munmap(mapping, size); + free(h); +} + +TEST_F(LinuxShmMicrotest, HandleInitWithoutAMapping) { + ncclShmHandleInternal h; + char path[] = "/x"; + ncclOsShmHandleInit(9, path, 10, 14, nullptr, nullptr, /*create=*/false, &h); + EXPECT_EQ(h.shmDesc, 9); + EXPECT_EQ(h.refcount, nullptr); + EXPECT_EQ(h.shmPath, nullptr); + ncclOsShmHandleInit(9, path, 10, 14, nullptr, nullptr, /*create=*/true, &h); + EXPECT_STREQ(h.shmPath, "/x"); + free(h.shmPath); +} + +} // namespace diff --git a/projects/rccl/test/host/shmutils-test.cc b/projects/rccl/test/host/shmutils-test.cc new file mode 100644 index 00000000000..65635716fa9 --- /dev/null +++ b/projects/rccl/test/host/shmutils-test.cc @@ -0,0 +1,131 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/shmutils.cc over the real os/linux.cc shared memory, with local +// ranks played by threads for ncclShmemAllgather. + +#include + +#include +#include +#include +#include +#include +#include + +#include "comm.h" +#include "shmutils.h" + +#include "linux-test-common.h" + +#include SHMUTILS_CC_PATH + +namespace { + +class ShmutilsMicrotest : public LinuxTestBase { + protected: + void TearDown() override { + for (int k = 0; k < 2; k++) { + free(coll_.cnt[k]); + free(coll_.ptr[k]); + } + } + // Shared counters and slots for nRanks ranks. + void InitColl(int nRanks, size_t maxTypeSize) { + coll_.maxTypeSize = maxTypeSize; + coll_.round = 0; + for (int k = 0; k < 2; k++) { + coll_.cnt[k] = (size_t*)calloc(nRanks, CACHE_LINE_SIZE); + coll_.ptr[k] = calloc(nRanks, maxTypeSize); + } + } + static std::unique_ptr Comm(int localRank, int localRanks) { + std::unique_ptr comm(new ncclComm()); + comm->localRank = localRank; + comm->localRanks = localRanks; + return comm; + } + ncclShmemCollBuff coll_ = {}; +}; + +TEST_F(ShmutilsMicrotest, OpenAttachUnlinkAndClose) { + char path[64] = ""; + void* ptr = nullptr; + ncclShmHandle_t creator = nullptr, peer = nullptr; + ASSERT_EQ(ncclShmOpen(path, sizeof(path), 128, &ptr, nullptr, 2, &creator), ncclSuccess); + ASSERT_NE(creator, nullptr); + EXPECT_EQ(((ncclShmHandleInternal*)creator)->shmPtr, ptr); + std::strcpy((char*)ptr, "hi"); + void* peerPtr = nullptr; + ASSERT_EQ(ncclShmOpen(path, sizeof(path), 128, &peerPtr, nullptr, 0, &peer), ncclSuccess); + EXPECT_STREQ((char*)peerPtr, "hi"); + EXPECT_EQ(access(path, F_OK), 0); // one peer still to attach + ASSERT_EQ(ncclShmUnlink(creator), ncclSuccess); + EXPECT_NE(access(path, F_OK), 0); + EXPECT_EQ(ncclShmClose(peer), ncclSuccess); + EXPECT_EQ(ncclShmClose(creator), ncclSuccess); +} + +TEST_F(ShmutilsMicrotest, OpenFailureReturnsANullHandle) { + char path[64] = "/nccl/no/such/dir/seg"; + void* ptr = nullptr; + ncclShmHandle_t h = (ncclShmHandle_t)0x1; + EXPECT_EQ(ncclShmOpen(path, sizeof(path), 128, &ptr, nullptr, 1, &h), ncclSystemError); + EXPECT_EQ(h, nullptr); +} + +TEST_F(ShmutilsMicrotest, Allgather_RejectsMissingArgumentsAndAnOversizedType) { + InitColl(1, 8); + auto comm = Comm(0, 1); + char send[16] = {}, recv[16]; + EXPECT_EQ(ncclShmemAllgather(nullptr, &coll_, send, recv, 4), ncclInvalidArgument); + EXPECT_EQ(ncclShmemAllgather(comm.get(), &coll_, nullptr, recv, 4), ncclInvalidArgument); + EXPECT_EQ(ncclShmemAllgather(comm.get(), &coll_, send, nullptr, 4), ncclInvalidArgument); + EXPECT_EQ(ncclShmemAllgather(comm.get(), &coll_, send, recv, 9), ncclInvalidArgument); + EXPECT_EQ(coll_.round, 0); + EXPECT_EQ(ncclShmemAllgather(comm.get(), &coll_, send, recv, 8), ncclSuccess); +} + +TEST_F(ShmutilsMicrotest, Allgather_SingleRankAlternatesBuffersEachRound) { + InitColl(1, 8); + auto comm = Comm(0, 1); + for (int round = 1; round <= 3; round++) { + int send = 100 + round, recv = 0; + ASSERT_EQ(ncclShmemAllgather(comm.get(), &coll_, &send, &recv, sizeof(int)), ncclSuccess); + EXPECT_EQ(recv, send); + EXPECT_EQ(coll_.round, round); + int slot = (round - 1) % 2; + EXPECT_EQ(*(int*)coll_.cnt[slot], round); // this rank's arrival + EXPECT_EQ(*(int*)coll_.ptr[slot], send); + } +} + +// Each thread is a local rank with its own round counter over the shared slots. +TEST_F(ShmutilsMicrotest, Allgather_EveryRankGetsEveryRanksValueInRankOrder) { + constexpr int kRanks = 4, kRounds = 5; + InitColl(kRanks, 16); // slots wider than the type: packed into recvbuff + std::vector> got(kRanks); + std::vector threads; + for (int r = 0; r < kRanks; r++) { + threads.emplace_back([&, r] { + auto comm = Comm(r, kRanks); + ncclShmemCollBuff mine = coll_; + for (int round = 0; round < kRounds; round++) { + int send = 10 * round + r, recv[kRanks]; + EXPECT_EQ(ncclShmemAllgather(comm.get(), &mine, &send, recv, sizeof(int)), ncclSuccess); + got[r].insert(got[r].end(), recv, recv + kRanks); + } + EXPECT_EQ(mine.round, kRounds); + }); + } + for (auto& t : threads) t.join(); + std::vector expected; + for (int round = 0; round < kRounds; round++) + for (int r = 0; r < kRanks; r++) expected.push_back(10 * round + r); + for (int r = 0; r < kRanks; r++) EXPECT_EQ(got[r], expected) << "rank " << r; +} + +} // namespace diff --git a/projects/rccl/test/host/signals-test.cc b/projects/rccl/test/host/signals-test.cc new file mode 100644 index 00000000000..506cd51ef39 --- /dev/null +++ b/projects/rccl/test/host/signals-test.cc @@ -0,0 +1,112 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/signals.cc: the opt-in handler registration (signal() is hooked) +// and the handler itself, whose exit path runs in a death test. + +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "../common/LogCapture.hpp" +#include "debug.h" +#include "param.h" + +#include "linux-test-common.h" + +#undef HAVE_BFD // the libbfd backtrace needs libbfd, which this binary does not link +#include "fakes/param_redirect.h" +#include "fakes/libc_seam.h" +#include SIGNALS_CC_PATH +#include "fakes/libc_seam_undef.h" + +namespace { + +using RcclUnitTesting::CaptureStdout; + +const std::vector kCaught = {SIGILL, SIGBUS, SIGFPE, SIGSEGV, SIGUSR2}; + +class SignalsMicrotest : public LinuxTestBase { + protected: + void SetUp() override { + for (int sig : kCaught) sigaction(sig, nullptr, &saved_[sig]); + } + void TearDown() override { + for (int sig : kCaught) sigaction(sig, &saved_[sig], nullptr); + } + static sighandler_t Installed(int sig) { + struct sigaction sa; + sigaction(sig, nullptr, &sa); + return sa.sa_handler; + } + struct sigaction saved_[NSIG]; +}; + +TEST_F(SignalsMicrotest, Register_OffByDefault) { + ScopedHook signalHook(g_signal, [](int, sighandler_t) { return SIG_DFL; }); + RegisterSignalHandlers(); + EXPECT_EQ(signalHook.calls, 0); + EXPECT_TRUE(log_.Text().empty()); +} + +TEST_F(SignalsMicrotest, Register_InstallsTheHandlerForTheFaultSignalsAndUsr2) { + params_.Set("RCCL_ENABLE_SIGNALHANDLER", 1); + std::vector seen; + ScopedHook signalHook(g_signal, [&](int sig, sighandler_t h) { + seen.push_back(sig); + EXPECT_EQ(h, &sig_handler); + return ::signal(sig, h); + }); + RegisterSignalHandlers(); + EXPECT_EQ(seen, kCaught); + for (int sig : kCaught) EXPECT_EQ(Installed(sig), &sig_handler) << sig; +} + +TEST_F(SignalsMicrotest, Register_FailureIsOnlyLogged) { + params_.Set("RCCL_ENABLE_SIGNALHANDLER", 1); + ScopedHook signalHook(g_signal, [](int, sighandler_t) { return SIG_ERR; }); + RegisterSignalHandlers(); + EXPECT_EQ(signalHook.calls, (int)kCaught.size()); // a failure does not stop the rest +} + +TEST_F(SignalsMicrotest, Handler_Usr2PrintsABacktraceAndReturns) { + std::string out = CaptureStdout([] { + sig_handler(SIGUSR2); + fflush(stdout); + }); + std::string head = "\n[Process: " + std::to_string(getpid()) + "] Inside handler function signal: " + + strsignal(SIGUSR2) + " (" + std::to_string(SIGUSR2) + ")\n"; + ASSERT_EQ(out.rfind(head, 0), 0u) << out; + EXPECT_GT(out.size(), head.size()); // the backtrace frames follow +} + +TEST_F(SignalsMicrotest, Handler_RaisedUsr2ReturnsToTheProgram) { + params_.Set("RCCL_ENABLE_SIGNALHANDLER", 1); + RegisterSignalHandlers(); + std::string out = CaptureStdout([] { + raise(SIGUSR2); + fflush(stdout); + }); + EXPECT_NE(out.find("Inside handler function signal"), std::string::npos); +} + +TEST_F(SignalsMicrotest, Handler_FaultSignalsExitWithMinusOne) { + // A fresh process, so LeakSanitizer's exit check finds no leaks from earlier tests and keeps the code. + const std::string style = ::testing::GTEST_FLAG(death_test_style); + ::testing::GTEST_FLAG(death_test_style) = "threadsafe"; + for (int sig : {SIGILL, SIGBUS, SIGFPE, SIGSEGV}) { + EXPECT_EXIT(sig_handler(sig), ::testing::ExitedWithCode(255), "") << sig; + } + ::testing::GTEST_FLAG(death_test_style) = style; +} + +} // namespace diff --git a/projects/rccl/test/host/socket-test.cc b/projects/rccl/test/host/socket-test.cc new file mode 100644 index 00000000000..a369602b5fd --- /dev/null +++ b/projects/rccl/test/host/socket-test.cc @@ -0,0 +1,1089 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only microtests for src/misc/socket.cc over real loopback TCP, with the real os/linux.cc and +// the plaintext crypt.cc handshake underneath. Its libc calls go through fakes/libc_seam.h, real +// unless a test hooks one. + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "crypt.h" +#include "os.h" +#include "param.h" +#include "socket.h" +#include "utils.h" + +#include "linux-test-common.h" + +#include "fakes/param_redirect.h" +#include "fakes/libc_seam.h" +#include SOCKET_CC_PATH +#include "fakes/libc_seam_undef.h" + +namespace { + +constexpr uint64_t kMagic = 0x5eed5eed12345678ull; + +ncclSocketAddress V4(const char* ip, uint16_t port) { + ncclSocketAddress a; + std::memset(&a, 0, sizeof(a)); + a.sin.sin_family = AF_INET; + a.sin.sin_port = htons(port); + inet_pton(AF_INET, ip, &a.sin.sin_addr); + return a; +} + +ncclSocketAddress V6(const char* ip, uint16_t port) { + ncclSocketAddress a; + std::memset(&a, 0, sizeof(a)); + a.sin6.sin6_family = AF_INET6; + a.sin6.sin6_port = htons(port); + inet_pton(AF_INET6, ip, &a.sin6.sin6_addr); + return a; +} + +std::string Str(const ncclSocketAddress& a) { + char buf[SOCKET_NAME_MAXLEN + 1]; + return ncclSocketToString(&a, buf, sizeof(buf)); +} + +class SocketMicrotest : public LinuxTestBase { + protected: + void SetUp() override { + for (ncclSocket* s : {&listen_, &client_, &server_}) { + std::memset(s, 0, sizeof(*s)); + s->socketDescriptor = s->acceptSocketDescriptor = NCCL_INVALID_SOCKET; + } + } + void TearDown() override { + for (ncclSocket* s : {&client_, &server_, &listen_}) ncclSocketClose(s); + for (int fd : fds_) ::close(fd); + ncclSetEncryption(nullptr); + } + int Track(int fd) { + fds_.push_back(fd); + return fd; + } + void Listen(int asyncFlag = 0, volatile uint32_t* abortFlag = nullptr) { + ncclSocketAddress any = V4("127.0.0.1", 0); + ASSERT_EQ(ncclSocketInit(&listen_, &any, kMagic, ncclSocketTypeBootstrap, abortFlag, asyncFlag), ncclSuccess); + ASSERT_EQ(ncclSocketListen(&listen_), ncclSuccess); + } + ncclSocketAddress ListenAddr() { + ncclSocketAddress a; + EXPECT_EQ(ncclSocketGetAddr(&listen_, &a), ncclSuccess); + return a; + } + // client_ connects to listen_ with the given magic and type; blocking unless asyncFlag. + ncclResult_t Connect(uint64_t magic = kMagic, ncclSocketType type = ncclSocketTypeBootstrap, int asyncFlag = 0) { + ncclSocketAddress a = ListenAddr(); + EXPECT_EQ(ncclSocketInit(&client_, &a, magic, type, nullptr, asyncFlag), ncclSuccess); + return ncclSocketConnect(&client_); + } + // client_ and server_ end Ready. + void Pair() { + Listen(); + ASSERT_EQ(Connect(), ncclSuccess); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + ASSERT_EQ(client_.state, ncclSocketStateReady); + ASSERT_EQ(server_.state, ncclSocketStateReady); + } + // A raw TCP client of listen_ that sends the given hello bytes, or none. + int RawClient(const std::string& hello) { + ncclSocketAddress a = ListenAddr(); + int fd = Track(::socket(AF_INET, SOCK_STREAM, 0)); + EXPECT_EQ(::connect(fd, &a.sa, sizeof(a.sin)), 0); + if (!hello.empty()) EXPECT_EQ(::send(fd, hello.data(), hello.size(), 0), (ssize_t)hello.size()); + return fd; + } + static std::string Hello(uint64_t magic, ncclSocketType type) { + std::string h(NCCL_SOCKET_PLAIN_HELLO_BYTES, '\0'); + std::memcpy(&h[0], &magic, sizeof(magic)); + std::memcpy(&h[sizeof(magic)], &type, sizeof(type)); + return h; + } + ncclSocket listen_, client_, server_; + std::vector fds_; +}; + +// Basics + +TEST_F(SocketMicrotest, DefaultMagic_EnvOverridesOnceElseBuiltIn) { + // Each case needs a fresh process, as the value is cached on first use. + const std::string style = ::testing::GTEST_FLAG(death_test_style); + ::testing::GTEST_FLAG(death_test_style) = "threadsafe"; + auto magicIs = [](const char* env, uint64_t expected) { + if (env) SetMicroEnv("NCCL_SOCKET_MAGIC", env); + bool ok = ncclSocketDefaultMagic() == expected; + SetMicroEnv("NCCL_SOCKET_MAGIC", "0x99"); + ok = ok && ncclSocketDefaultMagic() == expected; // cached + exit(ok ? 0 : 1); + }; + EXPECT_EXIT(magicIs(nullptr, NCCL_SOCKET_MAGIC), ::testing::ExitedWithCode(0), ""); + EXPECT_EXIT(magicIs("", NCCL_SOCKET_MAGIC), ::testing::ExitedWithCode(0), ""); + EXPECT_EXIT(magicIs("0x1234", 0x1234), ::testing::ExitedWithCode(0), ""); + EXPECT_EXIT(magicIs("12abc", NCCL_SOCKET_MAGIC), ::testing::ExitedWithCode(0), ""); + ::testing::GTEST_FLAG(death_test_style) = style; +} + +TEST_F(SocketMicrotest, Move_TransfersOwnershipAndInvalidatesTheSource) { + Pair(); + ncclSocket moved; + int fd = client_.socketDescriptor; +#ifndef NCCL_TLS_BACKEND_OPENSSL3 // a TLS backend would rebind this fake state + client_.crypto = (ncclSocketCrypto*)0x1; +#endif + ncclSocketMove(&moved, &client_); + EXPECT_EQ(moved.socketDescriptor, fd); + EXPECT_EQ(moved.state, ncclSocketStateReady); + EXPECT_EQ(moved.magic, kMagic); +#ifndef NCCL_TLS_BACKEND_OPENSSL3 + EXPECT_EQ(moved.crypto, (ncclSocketCrypto*)0x1); +#endif + EXPECT_EQ(client_.socketDescriptor, NCCL_INVALID_SOCKET); + EXPECT_EQ(client_.acceptSocketDescriptor, NCCL_INVALID_SOCKET); + EXPECT_EQ(client_.state, ncclSocketStateNone); + EXPECT_EQ(client_.crypto, nullptr); + moved.crypto = nullptr; + ncclSocketClose(&moved); +} + +TEST_F(SocketMicrotest, ToPort_ReadsEitherFamily) { + ncclSocketAddress a = V4("10.0.0.1", 5000); + EXPECT_EQ(ncclSocketToPort(&a), 5000); + a = V6("::1", 6001); + EXPECT_EQ(ncclSocketToPort(&a), 6001); +} + +TEST_F(SocketMicrotest, GetAndSetFd) { + int fd = 0; + EXPECT_EQ(ncclSocketGetFd(nullptr, &fd), ncclInvalidArgument); + EXPECT_EQ(ncclSocketSetFd(3, nullptr), ncclInvalidArgument); + ASSERT_EQ(ncclSocketSetFd(7, &client_), ncclSuccess); + EXPECT_EQ(ncclSocketGetFd(&client_, nullptr), ncclSuccess); + ASSERT_EQ(ncclSocketGetFd(&client_, &fd), ncclSuccess); + EXPECT_EQ(fd, 7); + client_.socketDescriptor = NCCL_INVALID_SOCKET; +} + +TEST_F(SocketMicrotest, Shutdown_SignalsThePeerAndMarksTerminating) { + EXPECT_EQ(ncclSocketShutdown(nullptr, SHUT_RDWR), ncclSuccess); + ASSERT_EQ(ncclSocketShutdown(&client_, SHUT_RDWR), ncclSuccess); // no descriptor: state only + EXPECT_EQ(client_.state, ncclSocketStateTerminating); + Pair(); + ASSERT_EQ(ncclSocketShutdown(&client_, SHUT_WR), ncclSuccess); + EXPECT_EQ(client_.state, ncclSocketStateTerminating); + char c; + EXPECT_EQ(::recv(server_.socketDescriptor, &c, 1, 0), 0); +} + +TEST_F(SocketMicrotest, Shutdown_FailureIsASystemError) { + int p[2]; + ASSERT_EQ(pipe(p), 0); + Track(p[0]); + Track(p[1]); + ncclSocket sock = client_; + sock.socketDescriptor = p[0]; + EXPECT_EQ(ncclSocketShutdown(&sock, SHUT_RDWR), ncclSystemError); +} + +// Listen + +TEST_F(SocketMicrotest, Listen_RejectsNullOrUnopenedSockets) { + EXPECT_EQ(ncclSocketListen(nullptr), ncclInvalidArgument); + EXPECT_EQ(ncclSocketListen(&listen_), ncclInvalidArgument); +} + +TEST_F(SocketMicrotest, Listen_AnyPortGetsAnEphemeralPort) { + Listen(); + EXPECT_EQ(listen_.state, ncclSocketStateReady); + EXPECT_EQ(listen_.acceptSocketDescriptor, listen_.socketDescriptor); + EXPECT_NE(ncclSocketToPort(&listen_.addr), 0); + EXPECT_EQ(IntOpt(listen_.socketDescriptor, SOL_SOCKET, SO_ACCEPTCONN), 1); + EXPECT_EQ(IntOpt(listen_.socketDescriptor, SOL_SOCKET, SO_REUSEADDR), 0); +} + +TEST_F(SocketMicrotest, Listen_AFixedPortIsReusable) { + Listen(); + ncclSocketAddress fixed = ListenAddr(); + ncclSocketClose(&listen_); + ASSERT_EQ(ncclSocketInit(&listen_, &fixed), ncclSuccess); + ASSERT_EQ(ncclSocketListen(&listen_), ncclSuccess); + EXPECT_EQ(ncclSocketToPort(&listen_.addr), ncclSocketToPort(&fixed)); + EXPECT_EQ(IntOpt(listen_.socketDescriptor, SOL_SOCKET, SO_REUSEADDR), 1); + EXPECT_EQ(IntOpt(listen_.socketDescriptor, SOL_SOCKET, SO_REUSEPORT), 1); +} + +TEST_F(SocketMicrotest, Listen_PortInUseIsASystemError) { + Listen(); // without SO_REUSEPORT, so nobody can share the port + ncclSocketAddress taken = ListenAddr(); + ncclSocket other; + ASSERT_EQ(ncclSocketInit(&other, &taken), ncclSuccess); + EXPECT_EQ(ncclSocketListen(&other), ncclSystemError); + ncclSocketClose(&other); +} + +TEST_F(SocketMicrotest, Listen_SyscallFailuresAreSystemErrors) { + auto fail = FailWith(EACCES); + // A fixed port takes the setsockopt path. Init sets socket options too, so each hook goes in after it. + ncclSocket sock[4]; + auto init = [&](int i, uint16_t port) { + ncclSocketAddress a = V4("127.0.0.1", port); + EXPECT_EQ(ncclSocketInit(&sock[i], &a), ncclSuccess); + }; + init(0, 1); + { + ScopedHook hook(g_setsockopt, fail); + EXPECT_EQ(ncclSocketListen(&sock[0]), ncclSystemError); + EXPECT_EQ(hook.calls, 1); + } + init(1, 1); + { + ScopedHook hook(g_setsockopt, [](int fd, int level, int name, const void* v, socklen_t len) { + if (name != SO_REUSEPORT) return ::setsockopt(fd, level, name, v, len); + errno = EACCES; + return -1; + }); + EXPECT_EQ(ncclSocketListen(&sock[1]), ncclSystemError); + EXPECT_EQ(hook.calls, 2); // SO_REUSEADDR went through + } + init(2, 0); + { + ScopedHook hook(g_getsockname, fail); + EXPECT_EQ(ncclSocketListen(&sock[2]), ncclSystemError); + EXPECT_EQ(hook.calls, 1); + } + init(3, 0); + { + ScopedHook hook(g_listen, fail); + EXPECT_EQ(ncclSocketListen(&sock[3]), ncclSystemError); + EXPECT_EQ(hook.calls, 1); + } + for (ncclSocket& s : sock) ncclSocketClose(&s); +} + +// Addresses + +TEST_F(SocketMicrotest, ToString_NumericHostAndPort) { + EXPECT_EQ(Str(V4("127.0.0.1", 5000)), "127.0.0.1<5000>"); + EXPECT_EQ(Str(V6("::1", 80)), "::1<80>"); + char buf[6] = "junk"; + ncclSocketAddress a = V4("127.0.0.1", 5000); + EXPECT_STREQ(ncclSocketToString(&a, buf, sizeof(buf)), "127.0"); +} + +TEST_F(SocketMicrotest, ToString_UnprintableAddressesAreEmpty) { + char buf[64] = "junk"; + EXPECT_EQ(ncclSocketToString(nullptr, buf, sizeof(buf)), buf); + EXPECT_STREQ(buf, ""); + ncclSocketAddress a = V4("127.0.0.1", 1); + EXPECT_EQ(ncclSocketToString(&a, nullptr, 0), nullptr); + a.sa.sa_family = AF_UNIX; + std::strcpy(buf, "junk"); + EXPECT_STREQ(ncclSocketToString(&a, buf, sizeof(buf)), ""); + a = V4("127.0.0.1", 1); + ScopedHook getnameinfoHook(g_getnameinfo, [](const sockaddr*, socklen_t, char*, socklen_t, char*, socklen_t, int) { + return EAI_FAIL; + }); + std::strcpy(buf, "junk"); + EXPECT_STREQ(ncclSocketToString(&a, buf, sizeof(buf)), ""); +} + +TEST_F(SocketMicrotest, ToString_NumericHostFormIsOptional) { + std::vector flags; + ScopedHook getnameinfoHook(g_getnameinfo, [&](const sockaddr* sa, socklen_t len, char* host, socklen_t hl, + char* serv, socklen_t sl, int f) { + flags.push_back(f); + return ::getnameinfo(sa, len, host, hl, serv, sl, f | NI_NUMERICHOST); + }); + ncclSocketAddress a = V4("127.0.0.1", 1); + char buf[64]; + ncclSocketToString(&a, buf, sizeof(buf)); + ncclSocketToString(&a, buf, sizeof(buf), /*numericHostForm=*/0); + EXPECT_EQ(flags, (std::vector{NI_NUMERICSERV | NI_NUMERICHOST, NI_NUMERICSERV})); +} + +TEST_F(SocketMicrotest, EnvSocketFamily) { + EXPECT_EQ(ncclEnvSocketFamily(), -1); + EXPECT_TRUE(log_.Text().empty()); + const std::pair cases[] = {{"AF_INET", AF_INET}, {"AF_INET6", AF_INET6}, {"AF_UNIX", -1}}; + for (auto& c : cases) { + SetMicroEnv("NCCL_SOCKET_FAMILY", c.first); + EXPECT_EQ(ncclEnvSocketFamily(), c.second) << c.first; + } +} + +TEST_F(SocketMicrotest, GetAddrFromString_IPv4WithPort) { + ncclSocketAddress a; + ASSERT_EQ(ncclSocketGetAddrFromString(&a, "10.1.2.3:5000"), ncclSuccess); + EXPECT_EQ(Str(a), "10.1.2.3<5000>"); +} + +// Without a port, the missing -1 becomes port 65535. Pins AICOMRCCL-2977; a fix flips this pin. +TEST_F(SocketMicrotest, GetAddrFromString_KnownBug_MissingPortBecomes65535) { + ncclSocketAddress a; + ASSERT_EQ(ncclSocketGetAddrFromString(&a, "10.1.2.3"), ncclSuccess); + EXPECT_EQ(ncclSocketToPort(&a), 65535); +} + +TEST_F(SocketMicrotest, GetAddrFromString_ResolvedIPv6IsGlobalScope) { + ScopedHook getaddrinfoHook(g_getaddrinfo, [](const char*, const char* svc, const addrinfo* hints, addrinfo** res) { + EXPECT_EQ(hints->ai_family, AF_UNSPEC); + EXPECT_EQ(hints->ai_socktype, SOCK_STREAM); + int rc = ::getaddrinfo("fe80::1", svc, hints, res); + ((sockaddr_in6*)(*res)->ai_addr)->sin6_scope_id = 5; + ((sockaddr_in6*)(*res)->ai_addr)->sin6_flowinfo = 7; + return rc; + }); + ncclSocketAddress a; + ASSERT_EQ(ncclSocketGetAddrFromString(&a, "somehost:4000"), ncclSuccess); + EXPECT_EQ(a.sa.sa_family, AF_INET6); + EXPECT_EQ(ncclSocketToPort(&a), 4000); + EXPECT_EQ(a.sin6.sin6_scope_id, 0u); + EXPECT_EQ(a.sin6.sin6_flowinfo, 0u); +} + +TEST_F(SocketMicrotest, GetAddrFromString_IPv4Failures) { + ncclSocketAddress a; + ScopedHook getaddrinfoHook(g_getaddrinfo, + [](const char*, const char*, const addrinfo*, addrinfo**) { return EAI_NONAME; }); + EXPECT_EQ(ncclSocketGetAddrFromString(&a, nullptr), ncclInvalidArgument); + EXPECT_EQ(ncclSocketGetAddrFromString(&a, "x"), ncclInvalidArgument); + EXPECT_EQ(ncclSocketGetAddrFromString(&a, ":5000"), ncclInvalidArgument); + EXPECT_EQ(getaddrinfoHook.calls, 0); // rejected before resolving + EXPECT_EQ(ncclSocketGetAddrFromString(&a, "nohost:1"), ncclInvalidArgument); + EXPECT_EQ(getaddrinfoHook.calls, 1); + ScopedHook unixFamilyHook(g_getaddrinfo, [](const char*, const char* svc, const addrinfo* hints, addrinfo** res) { + int rc = ::getaddrinfo("127.0.0.1", svc, hints, res); + (*res)->ai_family = AF_UNIX; + return rc; + }); + EXPECT_EQ(ncclSocketGetAddrFromString(&a, "odd:1"), ncclInvalidArgument); + EXPECT_EQ(unixFamilyHook.calls, 1); +} + +TEST_F(SocketMicrotest, GetAddrFromString_BracketedIPv6) { + ncclSocketAddress a; + ASSERT_EQ(ncclSocketGetAddrFromString(&a, "[2001:db8::7]:8080"), ncclSuccess); + EXPECT_EQ(Str(a), "2001:db8::7<8080>"); + EXPECT_EQ(a.sin6.sin6_scope_id, 0u); + ASSERT_EQ(ncclSocketGetAddrFromString(&a, "[fe80::1%lo]:80"), ncclSuccess); + EXPECT_EQ(ncclSocketToPort(&a), 80); + EXPECT_EQ(a.sin6.sin6_scope_id, if_nametoindex("lo")); + char host[INET6_ADDRSTRLEN]; + EXPECT_STREQ(inet_ntop(AF_INET6, &a.sin6.sin6_addr, host, sizeof(host)), "fe80::1"); + EXPECT_EQ(ncclSocketGetAddrFromString(&a, "[::1:80"), ncclInvalidArgument); +} + +TEST_F(SocketMicrotest, GetAddr_OnlyFromAReadySocket) { + ncclSocketAddress a; + EXPECT_EQ(ncclSocketGetAddr(nullptr, &a), ncclInvalidArgument); + EXPECT_EQ(ncclSocketGetAddr(&listen_, &a), ncclInternalError); + Listen(); + ASSERT_EQ(ncclSocketGetAddr(&listen_, &a), ncclSuccess); + EXPECT_EQ(std::memcmp(&a, &listen_.addr, sizeof(a)), 0); +} + +// Interface selection, over a scripted getifaddrs() list + +class SocketIfMicrotest : public SocketMicrotest { + protected: + std::vector Find() { + char names[MAX_IFS * MAX_IF_NAME_SIZE]; + ncclSocketAddress addrs[MAX_IFS]; + int n = -1; + EXPECT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclSuccess); + std::vector out; + for (int i = 0; i < n; i++) out.push_back(names + i * MAX_IF_NAME_SIZE); + return out; + } + LinuxFakeIfaddrs ifs_; +}; + +TEST_F(SocketIfMicrotest, Auto_PrefersIbThenAnythingButDockerLoopbackAndBridges) { + ifs_.Add("lo", AF_INET, "127.0.0.1"); + ifs_.Add("docker0", AF_INET, "172.17.0.1"); + ifs_.Add("virbr0", AF_INET, "192.168.122.1"); + ifs_.Add("eth0", AF_INET, "10.0.0.1"); + ifs_.Add("ib0", AF_INET, "10.1.0.1"); + EXPECT_EQ(Find(), (std::vector{"ib0"})); +} + +TEST_F(SocketIfMicrotest, Auto_FallsBackToDockerThenLoopbackThenBridges) { + ifs_.Add("lo", AF_INET, "127.0.0.1"); + ifs_.Add("virbr0", AF_INET, "192.168.122.1"); + ifs_.Add("docker0", AF_INET, "172.17.0.1"); + ifs_.Add("eth0", AF_INET, "10.0.0.1"); + EXPECT_EQ(Find(), (std::vector{"eth0"})); + { + LinuxFakeIfaddrs noEth; + noEth.Add("lo", AF_INET, "127.0.0.1"); + noEth.Add("virbr0", AF_INET, "192.168.122.1"); + noEth.Add("docker0", AF_INET, "172.17.0.1"); + EXPECT_EQ(Find(), (std::vector{"docker0"})); + } + { + LinuxFakeIfaddrs noDocker; + noDocker.Add("virbr0", AF_INET, "192.168.122.1"); + noDocker.Add("lo", AF_INET, "127.0.0.1"); + EXPECT_EQ(Find(), (std::vector{"lo"})); + } + { + LinuxFakeIfaddrs bridgeOnly; + bridgeOnly.Add("virbr0", AF_INET, "192.168.122.1"); + EXPECT_EQ(Find(), (std::vector{"virbr0"})); + } + LinuxFakeIfaddrs none; + EXPECT_EQ(Find(), (std::vector{})); +} + +TEST_F(SocketIfMicrotest, Auto_CommIdPicksTheInterfaceInItsSubnet) { + ifs_.Add("eth0", AF_INET, "192.168.0.1", "255.255.255.0"); + ifs_.Add("eth1", AF_INET, "10.0.0.1", "255.0.0.0"); + SetMicroEnv("NCCL_COMM_ID", "10.9.9.9:5000"); + EXPECT_EQ(Find(), (std::vector{"eth1"})); + SetMicroEnv("NCCL_COMM_ID", "x"); // too short: ignored + EXPECT_EQ(Find(), (std::vector{"eth0", "eth1"})); + SetMicroEnv("NCCL_COMM_ID", "11.0.0.1:5000"); // no subnet match: fall through + EXPECT_EQ(Find(), (std::vector{"eth0", "eth1"})); +} + +TEST_F(SocketIfMicrotest, Auto_BadCommIdIsAnError) { + SetMicroEnv("NCCL_COMM_ID", "[::1"); + ifs_.Add("eth0", AF_INET, "10.0.0.1"); + char names[MAX_IFS * MAX_IF_NAME_SIZE]; + ncclSocketAddress addrs[MAX_IFS]; + int n = -1; + EXPECT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclInvalidArgument); +} + +TEST_F(SocketIfMicrotest, IfnameAndFamilyComeFromTheEnvironment) { + ifs_.Add("eth0", AF_INET, "10.0.0.1"); + ifs_.Add("eth0", AF_INET6, "2001:db8::1"); + ifs_.Add("ib0", AF_INET, "10.1.0.1"); + SetMicroEnv("NCCL_SOCKET_IFNAME", "eth"); + EXPECT_EQ(Find(), (std::vector{"eth0"})); + SetMicroEnv("NCCL_SOCKET_IFNAME", "e"); // one character: ignored + EXPECT_EQ(Find(), (std::vector{"ib0"})); + SetMicroEnv("NCCL_SOCKET_IFNAME", "eth0"); + SetMicroEnv("NCCL_SOCKET_FAMILY", "AF_INET6"); + char names[MAX_IFS * MAX_IF_NAME_SIZE]; + ncclSocketAddress addrs[MAX_IFS]; + int n = -1; + ASSERT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclSuccess); + ASSERT_EQ(n, 1); + EXPECT_EQ(addrs[0].sa.sa_family, AF_INET6); +} + +TEST_F(SocketIfMicrotest, EnumerationFailuresPropagate) { + ScopedHook getifaddrsHook(g_getifaddrs, FailWith(ENOMEM)); + char names[MAX_IFS * MAX_IF_NAME_SIZE]; + ncclSocketAddress addrs[MAX_IFS]; + int n = -1; + EXPECT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclSystemError); + SetMicroEnv("NCCL_SOCKET_IFNAME", "eth0"); + EXPECT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclSystemError); +} + +// Fails the n-th getifaddrs() call, to reach each fallback's error check. +TEST_F(SocketIfMicrotest, EachFallbackPropagatesAnEnumerationFailure) { + SetMicroEnv("NCCL_COMM_ID", "10.9.9.9:5000"); + for (int failAt = 1; failAt <= 6; failAt++) { + SCOPED_TRACE(failAt); + int calls = 0; + ScopedHook getifaddrsHook(g_getifaddrs, [&](ifaddrs** out) { + if (++calls == failAt) { + errno = ENOMEM; + return -1; + } + *out = nullptr; + return 0; + }); + ScopedHook freeifaddrsHook(g_freeifaddrs, [](ifaddrs*) {}); + char names[MAX_IFS * MAX_IF_NAME_SIZE]; + ncclSocketAddress addrs[MAX_IFS]; + int n = -1; + EXPECT_EQ(ncclFindInterfaces(names, addrs, MAX_IF_NAME_SIZE, MAX_IFS, &n), ncclSystemError); + } +} + +// Init + +TEST_F(SocketMicrotest, Init_WithoutAnAddressOnlyResetsState) { + EXPECT_EQ(ncclSocketInit(nullptr), ncclSuccess); + volatile uint32_t abortFlag = 0; + std::memset(&client_, 0x5a, sizeof(client_)); + ASSERT_EQ(ncclSocketInit(&client_, nullptr, kMagic, ncclSocketTypeProxy, &abortFlag, 1, 1), ncclSuccess); + EXPECT_EQ(client_.state, ncclSocketStateInitialized); + EXPECT_EQ(client_.socketDescriptor, NCCL_INVALID_SOCKET); + EXPECT_EQ(client_.acceptSocketDescriptor, NCCL_INVALID_SOCKET); + EXPECT_EQ(client_.errorRetries, 0); + EXPECT_EQ(client_.abortFlag, &abortFlag); + EXPECT_EQ(client_.asyncFlag, 1); + EXPECT_EQ(client_.customRetry, 1); + EXPECT_EQ(client_.magic, kMagic); + EXPECT_EQ(client_.type, ncclSocketTypeProxy); + EXPECT_EQ(client_.finalizeCounter, 0); + EXPECT_EQ(client_.crypto, nullptr); + ncclSocketAddress zero; + std::memset(&zero, 0, sizeof(zero)); + EXPECT_EQ(std::memcmp(&client_.addr, &zero, sizeof(zero)), 0); +} + +TEST_F(SocketMicrotest, Init_OpensASocketOfTheAddressFamily) { + ncclSocketAddress a = V4("127.0.0.1", 1); + ASSERT_EQ(ncclSocketInit(&client_, &a), ncclSuccess); + EXPECT_TRUE(FdOpen(client_.socketDescriptor)); + EXPECT_EQ(client_.salen, (int)sizeof(sockaddr_in)); + EXPECT_EQ(client_.magic, ncclSocketDefaultMagic()); + EXPECT_EQ(client_.type, ncclSocketTypeUnknown); + EXPECT_EQ(IntOpt(client_.socketDescriptor, SOL_SOCKET, SO_REUSEADDR), 0); + linger l{}; + socklen_t len = sizeof(l); + ::getsockopt(client_.socketDescriptor, SOL_SOCKET, SO_LINGER, &l, &len); + EXPECT_EQ(l.l_onoff, 0); + ncclSocketAddress b = V6("::1", 1); + ASSERT_EQ(ncclSocketInit(&server_, &b), ncclSuccess); + EXPECT_EQ(server_.salen, (int)sizeof(sockaddr_in6)); + EXPECT_EQ(IntOpt(server_.socketDescriptor, SOL_SOCKET, SO_DOMAIN), AF_INET6); +} + +TEST_F(SocketMicrotest, Init_AppliesTheReuseAddrAndLingerParams) { + params_.Set("RCCL_SOCKET_REUSEADDR", 1); + params_.Set("RCCL_SOCKET_LINGER", 0); + ncclSocketAddress a = V4("127.0.0.1", 1); + ASSERT_EQ(ncclSocketInit(&client_, &a), ncclSuccess); + EXPECT_EQ(IntOpt(client_.socketDescriptor, SOL_SOCKET, SO_REUSEADDR), 1); + linger l{}; + socklen_t len = sizeof(l); + ::getsockopt(client_.socketDescriptor, SOL_SOCKET, SO_LINGER, &l, &len); + EXPECT_EQ(l.l_onoff, 1); + EXPECT_EQ(l.l_linger, 0); +} + +TEST_F(SocketMicrotest, Init_RejectsANonIpAddress) { + ncclSocketAddress a = V4("127.0.0.1", 1); + a.sa.sa_family = AF_UNIX; + EXPECT_EQ(ncclSocketInit(&client_, &a), ncclInternalError); + EXPECT_EQ(client_.socketDescriptor, NCCL_INVALID_SOCKET); +} + +// Failures after the descriptor exists close it again. +TEST_F(SocketMicrotest, Init_FailuresCloseTheSocket) { + params_.Set("RCCL_SOCKET_REUSEADDR", 1); + params_.Set("RCCL_SOCKET_LINGER", 3); + ncclSocketAddress a = V4("127.0.0.1", 1); + for (int failing : {SO_REUSEADDR, SO_LINGER}) { + int created = -1; + ScopedHook setsockoptHook(g_setsockopt, [&](int fd, int level, int name, const void* v, socklen_t len) { + created = fd; + if (name == failing) { + errno = EINVAL; + return -1; + } + return ::setsockopt(fd, level, name, v, len); + }); + EXPECT_EQ(ncclSocketInit(&client_, &a), ncclSystemError); + EXPECT_FALSE(FdOpen(created)); + EXPECT_EQ(client_.state, ncclSocketStateClosed); + EXPECT_EQ(client_.socketDescriptor, NCCL_INVALID_SOCKET); + } +} + +TEST_F(SocketMicrotest, Init_SocketCreationFailureClosesNothing) { + rlimit saved; + ASSERT_EQ(getrlimit(RLIMIT_NOFILE, &saved), 0); + int lowest = open("/dev/null", O_RDONLY); // not dup(0): stdin may be closed + ASSERT_GE(lowest, 0); + ::close(lowest); + rlimit tight = saved; + tight.rlim_cur = lowest; // the next descriptor would exceed the limit + ASSERT_EQ(setrlimit(RLIMIT_NOFILE, &tight), 0); + ncclSocketAddress a = V4("127.0.0.1", 1); + ncclResult_t ret = ncclSocketInit(&client_, &a); + setrlimit(RLIMIT_NOFILE, &saved); + EXPECT_EQ(ret, ncclSystemError); + EXPECT_EQ(client_.state, ncclSocketStateClosed); +} + +// Connect, accept and the plaintext hello + +TEST_F(SocketMicrotest, ConnectAndAccept_HandshakeThenExchangeData) { + Pair(); + EXPECT_EQ(client_.finalizeCounter, (int)NCCL_SOCKET_PLAIN_HELLO_BYTES); + EXPECT_EQ(server_.magic, kMagic); + EXPECT_NE(server_.socketDescriptor, listen_.socketDescriptor); + EXPECT_EQ(server_.acceptSocketDescriptor, listen_.socketDescriptor); + int running = 0; + ASSERT_EQ(ncclSocketReady(&client_, &running), ncclSuccess); + EXPECT_EQ(running, 1); + char out[] = "ping", in[5] = {}; + ASSERT_EQ(ncclSocketSend(&client_, out, sizeof(out)), ncclSuccess); + ASSERT_EQ(ncclSocketRecv(&server_, in, sizeof(in)), ncclSuccess); + EXPECT_STREQ(in, "ping"); +} + +TEST_F(SocketMicrotest, Connect_RejectsBadSocketsAndStates) { + EXPECT_EQ(ncclSocketConnect(nullptr), ncclInvalidArgument); + EXPECT_EQ(ncclSocketConnect(&client_), ncclInvalidArgument); + ncclSocketAddress a = V4("127.0.0.1", 1); + ASSERT_EQ(ncclSocketInit(&client_, &a), ncclSuccess); + client_.state = ncclSocketStateReady; + EXPECT_EQ(ncclSocketConnect(&client_), ncclInternalError); + client_.state = ncclSocketStateError; + EXPECT_EQ(ncclSocketConnect(&client_), ncclRemoteError); +} + +#ifndef NCCL_TLS_BACKEND_OPENSSL3 +TEST_F(SocketMicrotest, Connect_EncryptionWithoutTlsFailsBeforeConnecting) { + Listen(); + ncclEncryptionConfig_t config = NCCL_ENCRYPTION_CONFIG_INITIALIZER; + config.mode = NCCL_ENCRYPTION_MODE_PSK; + config.psk = "0123456789abcdef0123456789abcdef"; + ASSERT_EQ(ncclSetEncryption(&config), ncclSuccess); + EXPECT_EQ(Connect(), ncclInvalidUsage); + EXPECT_EQ(client_.state, ncclSocketStateInitialized); +} +#endif + +TEST_F(SocketMicrotest, Connect_RefusedPastTheRetryCountIsARemoteError) { + params_.Set("SOCKET_RETRY_CNT", 1); + Listen(); + ncclSocketAddress a = ListenAddr(); + ncclSocketClose(&listen_); + ASSERT_EQ(ncclSocketInit(&client_, &a), ncclSuccess); + EXPECT_EQ(ncclSocketConnect(&client_), ncclRemoteError); + EXPECT_EQ(client_.state, ncclSocketStateError); + EXPECT_EQ(client_.errorRetries, 2); +} + +TEST_F(SocketMicrotest, Connect_AbortReturnsAnInternalError) { + Listen(); + ncclSocketAddress a = ListenAddr(); + volatile uint32_t abortFlag = 1; + ASSERT_EQ(ncclSocketInit(&client_, &a, kMagic, ncclSocketTypeBootstrap, &abortFlag), ncclSuccess); + EXPECT_EQ(ncclSocketConnect(&client_), ncclInternalError); +} + +TEST_F(SocketMicrotest, Connect_AsyncReturnsEarlyAndReadyFinishes) { + Listen(); + ASSERT_EQ(Connect(kMagic, ncclSocketTypeBootstrap, /*asyncFlag=*/1), ncclSuccess); + EXPECT_TRUE(client_.state == ncclSocketStateConnecting || client_.state == ncclSocketStateConnectPolling || + client_.state == ncclSocketStateConnected || client_.state == ncclSocketStateReady); + int running = 0; + for (int i = 0; i < 1000 && !running; i++) ASSERT_EQ(ncclSocketReady(&client_, &running), ncclSuccess); + EXPECT_EQ(running, 1); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateReady); +} + +TEST_F(SocketMicrotest, Connect_AsyncStopsInWhicheverConnectStateItReached) { + params_.Set("SOCKET_RETRY_CNT", 5); + Listen(); + for (int err : {EINPROGRESS, ECONNREFUSED}) { + ScopedHook connectHook(g_connect, FailWith(err)); + ScopedHook pollHook(g_poll, [](pollfd*, nfds_t, int) { return 0; }); // still connecting + ncclSocketClose(&client_); + ASSERT_EQ(Connect(kMagic, ncclSocketTypeBootstrap, /*asyncFlag=*/1), ncclSuccess); + EXPECT_EQ(client_.state, err == EINPROGRESS ? ncclSocketStateConnectPolling : ncclSocketStateConnecting); + } +} + +TEST_F(SocketMicrotest, PollConnect_CompletesAPendingConnect) { + EXPECT_EQ(ncclSocketPollConnect(nullptr), ncclInvalidArgument); + Listen(); + ncclSocketAddress a = ListenAddr(); + ASSERT_EQ(ncclSocketInit(&client_, &a), ncclSuccess); + ASSERT_EQ(::connect(client_.socketDescriptor, &a.sa, sizeof(a.sin)), 0); + client_.state = ncclSocketStateConnectPolling; + ASSERT_EQ(ncclSocketPollConnect(&client_), ncclSuccess); + EXPECT_EQ(client_.state, ncclSocketStateConnected); + ncclSocket bad = client_; + bad.socketDescriptor = 1 << 20; // poll reports POLLNVAL; getsockopt then fails + EXPECT_EQ(ncclSocketPollConnect(&bad), ncclSystemError); +} + +// An async connect whose hello cannot be sent yet stays Connected. +TEST_F(SocketMicrotest, Ready_HelloThatCannotBeSentYetStaysConnected) { + Pair(); + client_.asyncFlag = 1; + fcntl(client_.socketDescriptor, F_SETFL, O_NONBLOCK); + std::vector junk(1 << 16); + while (::send(client_.socketDescriptor, junk.data(), junk.size(), MSG_DONTWAIT) > 0) { + } + client_.state = ncclSocketStateConnected; + client_.finalizeCounter = 0; + int running = 1; + ASSERT_EQ(ncclSocketReady(&client_, &running), ncclSuccess); + EXPECT_EQ(running, 0); + EXPECT_EQ(client_.state, ncclSocketStateConnected); +} + +TEST_F(SocketMicrotest, Ready_NullStoppedOrFailedSockets) { + int running = 1; + ASSERT_EQ(ncclSocketReady(nullptr, &running), ncclSuccess); + EXPECT_EQ(running, 0); + for (ncclSocketState state : {ncclSocketStateError, ncclSocketStateClosed}) { + client_.state = state; + EXPECT_EQ(ncclSocketReady(&client_, &running), ncclRemoteError); + } + client_.state = ncclSocketStateNone; +} + +TEST_F(SocketMicrotest, Accept_RejectsBadArgumentsAndListenerStates) { + EXPECT_EQ(ncclSocketAccept(nullptr, &listen_), ncclInvalidArgument); + EXPECT_EQ(ncclSocketAccept(&server_, nullptr), ncclInvalidArgument); + listen_.state = ncclSocketStateInitialized; + EXPECT_EQ(ncclSocketAccept(&server_, &listen_), ncclInternalError); + listen_.state = ncclSocketStateError; + EXPECT_EQ(ncclSocketAccept(&server_, &listen_), ncclSystemError); + listen_.state = ncclSocketStateNone; +} + +TEST_F(SocketMicrotest, Accept_AbortReturnsAnInternalError) { + volatile uint32_t abortFlag = 1; + Listen(0, &abortFlag); // non-blocking: nothing pending, so the loop sees the flag + EXPECT_EQ(ncclSocketAccept(&server_, &listen_), ncclInternalError); +} + +TEST_F(SocketMicrotest, Accept_WrongMagicIsDroppedThenTheNextPeerAccepted) { + Listen(); + int bad = RawClient(Hello(kMagic ^ 1, ncclSocketTypeBootstrap)); + ASSERT_EQ(Connect(), ncclSuccess); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateReady); + char c; + EXPECT_EQ(::recv(bad, &c, 1, 0), 0); // the bad peer was closed +} + +TEST_F(SocketMicrotest, Accept_WithoutRetryReturnsAfterABadPeer) { + Listen(); + RawClient(Hello(kMagic, ncclSocketTypeProxy)); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_, /*retry=*/false), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateAccepting); + EXPECT_EQ(server_.socketDescriptor, NCCL_INVALID_SOCKET); +} + +// A blocking accept keeps reading a hello that arrives in pieces. +TEST_F(SocketMicrotest, Accept_BlockingWaitsForTheWholeHello) { + Listen(); + std::string hello = Hello(kMagic, ncclSocketTypeBootstrap); + int raw = RawClient(hello.substr(0, 3)); + bool sentRest = false; + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + ssize_t got = ::recv(fd, buf, n, f); + int err = errno; + if (got > 0 && !sentRest) { // the rest only arrives after a partial read + sentRest = true; + ::send(raw, hello.data() + 3, hello.size() - 3, 0); + } + errno = err; + return got; + }); + EXPECT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + EXPECT_TRUE(sentRest); + EXPECT_GE(recvHook.calls, 2); + EXPECT_EQ(server_.state, ncclSocketStateReady); +} + +TEST_F(SocketMicrotest, Accept_AsyncProgressesThroughReady) { + Listen(/*asyncFlag=*/1); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); // nobody yet + EXPECT_EQ(server_.state, ncclSocketStateAccepting); + std::string hello = Hello(kMagic, ncclSocketTypeBootstrap); + int raw = RawClient(""); + int running = 1; + for (int i = 0; i < 1000 && server_.state == ncclSocketStateAccepting; i++) + ASSERT_EQ(ncclSocketReady(&server_, &running), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateAccepted); // waiting for the hello + EXPECT_EQ(running, 0); + ::send(raw, hello.data(), hello.size(), 0); + for (int i = 0; i < 1000 && !running; i++) ASSERT_EQ(ncclSocketReady(&server_, &running), ncclSuccess); + EXPECT_EQ(running, 1); + // Re-entering accept on an accepted socket keeps it. + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateReady); +} + +TEST_F(SocketMicrotest, Ready_BadHandshakeGoesBackToAccepting) { + Listen(/*asyncFlag=*/1); + ASSERT_EQ(ncclSocketAccept(&server_, &listen_), ncclSuccess); + int bad = RawClient(Hello(kMagic ^ 1, ncclSocketTypeBootstrap)); + int running = 1; + char c; + for (int i = 0; i < 1000 && ::recv(bad, &c, 1, MSG_DONTWAIT) != 0; i++) // until the bad peer is closed + ASSERT_EQ(ncclSocketReady(&server_, &running), ncclSuccess); + EXPECT_EQ(server_.state, ncclSocketStateAccepting); + EXPECT_EQ(running, 0); + // A caller re-entering after a reset goes straight back to accepting the next peer. + RawClient(Hello(kMagic, ncclSocketTypeBootstrap)); + server_.state = ncclSocketStateBadHandshake; + ASSERT_EQ(ncclSocketReady(&server_, &running), ncclSuccess); + EXPECT_NE(server_.state, ncclSocketStateAccepting); + for (int i = 0; i < 1000 && !running; i++) ASSERT_EQ(ncclSocketReady(&server_, &running), ncclSuccess); + EXPECT_EQ(running, 1); +} + +TEST_F(SocketMicrotest, Accept_SetFlagsFailurePropagates) { + Listen(); + RawClient(Hello(kMagic, ncclSocketTypeBootstrap)); + ScopedHook setsockoptHook(g_setsockopt, FailWith(ENOMEM)); + EXPECT_EQ(ncclSocketAccept(&server_, &listen_), ncclSystemError); +} + +// Data transfer + +TEST_F(SocketMicrotest, Transfers_RejectNullSockets) { + int offset = 0; + char c; + bool closed; + EXPECT_EQ(ncclSocketProgress(NCCL_SOCKET_SEND, nullptr, &c, 1, &offset), ncclInvalidArgument); + EXPECT_EQ(ncclSocketWait(NCCL_SOCKET_SEND, nullptr, &c, 1, &offset), ncclInvalidArgument); + EXPECT_EQ(ncclSocketSend(nullptr, &c, 1), ncclInvalidArgument); + EXPECT_EQ(ncclSocketRecv(nullptr, &c, 1), ncclInvalidArgument); + EXPECT_EQ(ncclSocketTryRecv(nullptr, &c, 1, &closed, true), ncclInvalidArgument); + EXPECT_EQ(ncclSocketSendRecv(nullptr, &c, 1, &server_, &c, 1), ncclInternalError); + EXPECT_EQ(ncclSocketSendRecv(&client_, &c, 1, nullptr, &c, 1), ncclInternalError); +} + +TEST_F(SocketMicrotest, Transfers_NeedAReadySocket) { + char c = 0; + client_.state = ncclSocketStateTerminating; + EXPECT_EQ(ncclSocketSend(&client_, &c, 1), ncclInternalError); + client_.state = ncclSocketStateConnected; + EXPECT_EQ(ncclSocketRecv(&client_, &c, 1), ncclInternalError); + server_.state = ncclSocketStateReady; + EXPECT_EQ(ncclSocketSendRecv(&server_, &c, 1, &client_, &c, 1), ncclInternalError); + EXPECT_EQ(ncclSocketSendRecv(&client_, &c, 1, &server_, &c, 1), ncclInternalError); + client_.state = server_.state = ncclSocketStateNone; +} + +// A terminating socket can still drain what the peer sent. +TEST_F(SocketMicrotest, Recv_WorksWhileTerminating) { + Pair(); + char out[] = "tail", in[5] = {}; + ASSERT_EQ(ncclSocketSend(&client_, out, sizeof(out)), ncclSuccess); + ASSERT_EQ(ncclSocketShutdown(&server_, SHUT_WR), ncclSuccess); + ASSERT_EQ(ncclSocketRecv(&server_, in, sizeof(in)), ncclSuccess); + EXPECT_STREQ(in, "tail"); + ASSERT_EQ(ncclSocketSendRecv(&client_, out, sizeof(out), &server_, in, sizeof(in)), ncclSuccess); +} + +TEST_F(SocketMicrotest, Recv_PeerCloseIsARemoteError) { + Pair(); + ncclSocketClose(&client_); + char in[4]; + EXPECT_EQ(ncclSocketRecv(&server_, in, sizeof(in)), ncclRemoteError); +} + +TEST_F(SocketMicrotest, Wait_PollsUntilTheRestArrives) { + params_.Set("SOCKET_POLL_TIMEOUT_MSEC", 1000); + Pair(); + server_.asyncFlag = 1; // single non-blocking attempts, so the wait goes through poll + ScopedHook pollHook(g_poll, [&](pollfd* p, nfds_t n, int timeout) { + EXPECT_EQ(p->events, POLLIN); + EXPECT_EQ(timeout, 1000); + ::send(client_.socketDescriptor, "abcd", 4, 0); // the data only arrives once the wait polls + return ::poll(p, n, timeout); + }); + char in[4]; + int offset = 0; + EXPECT_EQ(ncclSocketWait(NCCL_SOCKET_RECV, &server_, in, sizeof(in), &offset), ncclSuccess); + EXPECT_EQ(pollHook.calls, 1); + EXPECT_EQ(offset, 4); + EXPECT_EQ(std::string(in, 4), "abcd"); +} + +TEST_F(SocketMicrotest, SendRecv_ExchangesBothWaysAtOnce) { + Pair(); + std::vector big(1 << 20, 'c'), fromServer(1 << 20); + std::thread peer([&] { + std::vector in(1 << 20); + std::vector out(1 << 20, 's'); + EXPECT_EQ(ncclSocketSendRecv(&server_, out.data(), out.size(), &server_, in.data(), in.size()), ncclSuccess); + EXPECT_EQ(in, big); + }); + EXPECT_EQ(ncclSocketSendRecv(&client_, big.data(), big.size(), &client_, fromServer.data(), fromServer.size()), + ncclSuccess); + peer.join(); + EXPECT_EQ(fromServer, std::vector(1 << 20, 's')); +} + +TEST_F(SocketMicrotest, SendRecv_PropagatesErrors) { + Pair(); + ncclSocketClose(&server_); + char c = 'x', in; + EXPECT_EQ(ncclSocketSendRecv(&client_, &c, 0, &client_, &in, 1), ncclRemoteError); + int p[2]; + ASSERT_EQ(pipe(p), 0); + Track(p[0]); + Track(p[1]); + ncclSocket notSock = client_; + notSock.socketDescriptor = p[0]; + EXPECT_EQ(ncclSocketSendRecv(¬Sock, &c, 1, &client_, &in, 0), ncclRemoteError); +} + +TEST_F(SocketMicrotest, MultiOp_ValidatesEveryOp) { + char buf[4]; + ncclSocketOp ops[2] = {{NCCL_SOCKET_SEND, &client_, buf, 4, 0}, {NCCL_SOCKET_RECV, &client_, buf, 4, 0}}; + EXPECT_EQ(ncclSocketMultiOp(nullptr, 1), ncclInvalidArgument); + EXPECT_EQ(ncclSocketMultiOp(ops, 0), ncclInvalidArgument); + ops[1].sock = nullptr; + EXPECT_EQ(ncclSocketMultiOp(ops, 2), ncclInvalidArgument); + ops[1].sock = &client_; + ops[1].op = 7; + EXPECT_EQ(ncclSocketMultiOp(ops, 2), ncclInvalidArgument); + ops[1].op = NCCL_SOCKET_RECV; + ops[1].size = -1; + EXPECT_EQ(ncclSocketMultiOp(ops, 2), ncclInvalidArgument); + ops[1].size = 4; + ops[1].ptr = nullptr; + EXPECT_EQ(ncclSocketMultiOp(ops, 2), ncclInvalidArgument); +} + +TEST_F(SocketMicrotest, MultiOp_RunsSendsAndReceivesToCompletion) { + Pair(); + std::vector a(1 << 20, 'a'), b(1 << 20, 'b'), gotA(1 << 20), gotB(1 << 20); + ncclSocketOp clientOps[3] = {{NCCL_SOCKET_SEND, &client_, a.data(), (int)a.size(), 9}, + {NCCL_SOCKET_RECV, &client_, gotB.data(), (int)gotB.size(), 9}, + {NCCL_SOCKET_SEND, &client_, nullptr, 0, 9}}; + ncclSocketOp serverOps[2] = {{NCCL_SOCKET_RECV, &server_, gotA.data(), (int)gotA.size(), 0}, + {NCCL_SOCKET_SEND, &server_, b.data(), (int)b.size(), 0}}; + std::thread peer([&] { EXPECT_EQ(ncclSocketMultiOp(serverOps, 2), ncclSuccess); }); + EXPECT_EQ(ncclSocketMultiOp(clientOps, 3), ncclSuccess); + peer.join(); + EXPECT_EQ(clientOps[0].offset, (int)a.size()); + EXPECT_EQ(clientOps[2].offset, 0); + EXPECT_EQ(gotA, a); + EXPECT_EQ(gotB, b); + ncclSocketOp onlyEmpty = {NCCL_SOCKET_RECV, &client_, nullptr, 0, 0}; + EXPECT_EQ(ncclSocketMultiOp(&onlyEmpty, 1), ncclSuccess); + ncclSocketClose(&server_); + char c; + ncclSocketOp closedPeer = {NCCL_SOCKET_RECV, &client_, &c, 1, 0}; + EXPECT_EQ(ncclSocketMultiOp(&closedPeer, 1), ncclRemoteError); +} + +TEST_F(SocketMicrotest, TryRecv_NonBlockingReportsInProgress) { + Pair(); + char in[4]; + bool closed = true; + EXPECT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, /*blocking=*/false), ncclInProgress); + EXPECT_FALSE(closed); + ::send(client_.socketDescriptor, "ab", 2, 0); + for (int i = 0; i < 1000; i++) { + int avail = 0; + ::ioctl(server_.socketDescriptor, FIONREAD, &avail); + if (avail) break; + } + std::string rest = "cd"; + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + ssize_t got = ::recv(fd, buf, n, f); + int err = errno; + if (got < 0 && !rest.empty()) { // the rest trickles in, a byte per empty read + ::send(client_.socketDescriptor, rest.data(), 1, 0); + rest.erase(0, 1); + } + errno = err; + return got; + }); + EXPECT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, /*blocking=*/false), ncclSuccess); + EXPECT_TRUE(rest.empty()); + EXPECT_EQ(std::string(in, 4), "abcd"); // once data started, it finishes the message +} + +TEST_F(SocketMicrotest, TryRecv_BlockingWaitsForTheFirstBytes) { + Pair(); + int empty = 0; + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + ssize_t got = ::recv(fd, buf, n, f); + int err = errno; + if (got < 0 && empty++ == 0) ::send(client_.socketDescriptor, "wxyz", 4, 0); // only after a miss + errno = err; + return got; + }); + char in[4]; + bool closed = true; + EXPECT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, /*blocking=*/true), ncclSuccess); + EXPECT_GE(empty, 1); + EXPECT_FALSE(closed); + EXPECT_EQ(std::string(in, 4), "wxyz"); +} + +TEST_F(SocketMicrotest, TryRecv_ReportsAClosedPeer) { + Pair(); + bool closed = false; + char in[4]; + ::send(client_.socketDescriptor, "ab", 2, 0); + ncclSocketClose(&client_); + ASSERT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, /*blocking=*/true), ncclSuccess); + EXPECT_TRUE(closed); // closed mid-message + closed = false; + ASSERT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, /*blocking=*/true), ncclSuccess); + EXPECT_TRUE(closed); // closed before any data +} + +TEST_F(SocketMicrotest, TryRecv_PropagatesErrors) { + int p[2]; + ASSERT_EQ(pipe(p), 0); + Track(p[0]); + Track(p[1]); + ncclSocket notSock = client_; + notSock.socketDescriptor = p[0]; + bool closed; + char in[4]; + EXPECT_EQ(ncclSocketTryRecv(¬Sock, in, sizeof(in), &closed, true), ncclRemoteError); + Pair(); + volatile uint32_t abortFlag = 0; + server_.abortFlag = &abortFlag; + ::send(client_.socketDescriptor, "ab", 2, 0); + ScopedHook recvHook(g_recv, [&](int fd, void* buf, size_t n, int f) { + ssize_t got = ::recv(fd, buf, n, f); + if (got > 0) abortFlag = 1; // aborted once "ab" is in + return got; + }); + EXPECT_EQ(ncclSocketTryRecv(&server_, in, sizeof(in), &closed, true), ncclInternalError); + EXPECT_EQ(std::string(in, 2), "ab"); +} + +#ifndef NCCL_TLS_BACKEND_OPENSSL3 +// With TLS state attached, I/O goes through crypt.cc, which this build has no backend for. +TEST_F(SocketMicrotest, Progress_TlsStateRoutesThroughCrypt) { + Pair(); + client_.crypto = (ncclSocketCrypto*)malloc(16); + ScopedHook sendHook(g_send, [](int fd, const void* buf, size_t n, int f) { return ::send(fd, buf, n, f); }); + ScopedHook recvHook(g_recv, [](int fd, void* buf, size_t n, int f) { return ::recv(fd, buf, n, f); }); + char buf[4] = {}; + int offset = 0; + EXPECT_EQ(ncclSocketProgress(NCCL_SOCKET_SEND, &client_, buf, sizeof(buf), &offset), ncclInternalError); + EXPECT_EQ(ncclSocketProgress(NCCL_SOCKET_RECV, &client_, buf, sizeof(buf), &offset), ncclInternalError); + EXPECT_EQ(sendHook.calls + recvHook.calls, 0); // neither reached the plain socket + offset = sizeof(buf); // nothing left to move: no call + log_.Clear(); + EXPECT_EQ(ncclSocketProgress(NCCL_SOCKET_SEND, &client_, buf, sizeof(buf), &offset), ncclSuccess); + EXPECT_TRUE(log_.Text().empty()); +} +#endif + +} // namespace diff --git a/projects/rccl/test/test_categories_micro_linux.yaml b/projects/rccl/test/test_categories_micro_linux.yaml new file mode 100644 index 00000000000..ddc1042441c --- /dev/null +++ b/projects/rccl/test/test_categories_micro_linux.yaml @@ -0,0 +1,49 @@ +# rccl-UnitTestsMicroLinux — shared/ctest schema. +# Host-only micro-tests for src/os/linux.cc and the units on it (no GPU, no librccl). +# RCCL does not support Windows: exclude_windows matches all tests so generated suites are empty on Windows. + +test_categories: + quick: µ_linux_quick + description: "Smoke — host-only os/linux.cc micro-tests, sub-second." + test_patterns: + - "Linux*" + - "Socket*" + - "Shmutils*" + - "Signals*" + exclude_windows: + - "*" + labels: + - "rccl" + - "quick" + - "smoke" + - "precheck" + - "microtest" + standard: + <<: *micro_linux_quick + labels: + - "rccl" + - "standard" + - "regression" + - "microtest" + comprehensive: + <<: *micro_linux_quick + labels: + - "rccl" + - "comprehensive" + - "comprehensive_regression" + - "microtest" + full: + <<: *micro_linux_quick + labels: + - "rccl" + - "full" + - "full_regression" + - "microtest" + +execution_settings: + timeout_multiplier: 1 + category_timeouts: + quick: 600 + standard: 600 + comprehensive: 600 + full: 600