I'm seeing a segfault/backtrace for NVSHMEM -> libfabric -> ucx control flow for a 2-node test run of GROMACS on one of our supercomputers with OpenMPI 5.0.2 on Cray Slingshot 11. I think what I'm really looking for is clear runtime error messages that tell me what is wrong (API, ABI, whatever version mismatches, etc.) before I ever get to a segfault. I've labelled this a bug on the sole basis that I shouldn't be able to segfault, but it could be that the error resides with i.e., the use of fi_getinfo() "upstream" of the segfault happening (i.e., that NVSHMEM should handle their runtime check differently?).

I've talked to NVIDIA engineers about this, and the problem really isn't clear to them. I did some experiments with runtime swapping of libfabric versions. NVSHMEM was built from source against libfabric 1.20.1 from spack.

Here is what happens if I use libfabric 1.18.1 at runtime instead: /lustre/scratch5/treddy/march_april_2024_testing/nvshmem_custom_dl/nvshmem_src_2.10.1-3/src/modules/transport/libfabric/libfabric.cpp:1524: non-zero status: -61 No providers matched fi_getinfo query: -61: No data available and then a segfault at ucx again.

The local C++ code they're using looks like this:

1521     status = fi_getinfo(FI_VERSION(NVSHMEMT_LIBFABRIC_MAJ_VER, NVSHMEMT_LIBFABRIC_MIN_VER), NULL,
1522                         NULL, 0, &info, &returned_fabrics);
1523                 
1524     NVSHMEMI_NZ_ERROR_JMP(status, NVSHMEMX_ERROR_INTERNAL, out,
1525                           "No providers matched fi_getinfo query: %d: %s\n", status,
1526                           fi_strerror(status * -1)); 

where those two version variables are set to 1 and 5, respectively. I know I've had some success in the past using libfabric 1.18.1 if I build NVSHMEM against that directly and use it at runtime. Is there a good reason I wouldn't be able to use 1.20.1, and if so how should the NVSHMEM folks guard against it?

Also, pretty sure that in this case I rebuilt NVSHMEM against 1.18.1 and still had the same problem, so it isn't quite that simple for me to work around (I did change some other things like newer OpenMPI, I was on a release candidate of 5.x series prior, and using newer NVSHMEM now).

Probably the most useful thing is this--do I have a shot at debugging this/fixing this without basically needing to rebuild my whole dependency chain? If I could just adjust LD_LIBRARY_PATH to somehow fix this that would be amazing, but I suspect I won't be so lucky.

Clean rebuild of dependency chain with libfabric at 1.18.1, OpenMPI at v5.0.3, and NVSHMEM at 2.10.1 produces a similar backtrace:

pmix is at 4.2.9.

Changing this from version 5 to 18 in the NVSHMEM source had no effect, same errors: src/modules/transport/libfabric/libfabric.h:#define NVSHMEMT_LIBFABRIC_MIN_VER 18 (which gets called by FI_VERSION).

I've simplified the reproducer to remove GROMACS entirely, using only the cuFFTMp example at: https://github.com/NVIDIA/CUDALibrarySamples/tree/master/cuFFTMp/samples/r2c_c2r_slabs_GROMACS.

Interactive run script for 2 nodes (4 A100 GPUs each)

Diff on above Makefile:

Output:

@tylerjereddy We occasionally see similar segfault at finalization phase inside ucp_worker_destroy(), the exact reason has not been identified with the most like guess being some race-condition inside ucx.

I don't see any libfabric related symbols from your trace. I would suggest run with debug build of libfabric and ucx to help to locate where the segfault happens.

@j-xiong I swapped in debug version of ucx and libfabric and added the log below the fold. Also added was FI_LOG_LEVEL=debug. This assumes that LD_LIBRARY_PATH swapping is sufficient and that I don't need to rebuild/re-link things.

I'm not sure I see anything more helpful on the backtraces? The added debug log info may help though.

So far, it looks like NVSHMEM tries to call fi_getinfo(), and this function fails to find any providers, but somehow the control flow continues to the point of a segfault.

From the extra fabric debug info, if I do grep -E "cannot open shared object file" out.txt | wc -l I see 200 such "missing" libs. For example, I don't see this shared lib in a standard ucx install:

libfabric:60833:1713994966::core:core:ofi_reg_dl_prov():610<debug> dlopen(libucx-fi.so): libucx-fi.so: cannot open shared object file: No such file or directory

Maybe the 200 missing *-fi.so shared libs is a hint? I'll see what I can do about experimenting with more explicit fabric options in my libfabric builds with spack, though I don't think I've needed that previously.

Edit: spack install libfabric@1.18.1+debug fabrics=mlx,sockets,tcp,udp,ucx ^ucx+debug and using that version via LD_LIBRARY_PATH seems to be of no use either.

@tylerjereddy Those warnings about missing *-fi.so are fine. Those files only exist for providers built as dl libraries, which most providers are by default not.

Based on the log, the libfabric installation doesn't have ucx provider at all. The top part of the stack trace indicates that the ucx error occurred in another thread (probably spawned by other part of NVSHMEM or OpenMPI).

 0  libucs.so.0(ucs_handle_error+0x294) [0x14e151015394]
 1  libucs.so.0(+0x30564) [0x14e151015564]
 2  libucs.so.0(+0x3082e) [0x14e15101582e]
 3  /lib64/libpthread.so.0(+0x168c0) [0x14e15303c8c0]
 4  nvshmem_transport_libfabric.so.1(+0xb9fd) [0x14e12c1b99fd]

@j-xiong interesting...

If I do fi_info -l I see the providers I built libfabric with listed near the top:

ucx:
    version: 118.10
udp:
    version: 118.10
tcp:
    version: 118.10
sockets:
    version: 118.10
<snip>

However, fi_info -p ucx returns the same -61 that NVSHMEM reports in its C-level query, while fi_info -p for udp and sockets both succeed, dumping a bunch of output.

So, what's the conclusion here? The -l command suggests consistency with how I built libfabric fabrics=mlx,sockets,tcp,ucx,udp, but the -p probes suggest that there's a runtime issue picking up ucx only? Is the sane debugging approach to then just try a bunch different builds of ucx and load them in via LD_LIBRARY_PATH?

I'll take a shot at turning on even more ucx options at compile time, I see they have a detailed backtrace options in addition to debug, and some threading options I suppose I could try turning on.

Turning on a ton more ucx compile options does change the backtrace as shown below. Actually, that's a bit strange, I don't even see ucx on the backtrace proper anymore (it is still absent from fi_info -p ucx though).

spack install ucx+cuda+gdrcopy+debug+backtrace_detail+assertions+dm+logging+parameter_checking+thread_multiple %gcc@12.2.0

The output of fi_info with the -l option is generated differently from without the -l option. With the -l option, no attempt is made to check if the providers are usable. It just returns the name and version of the available providers.

The fact that fi_info -p ucx returns -61 indicates that either the ucx provider doesn't exist or is not usable. The provider appears in the output of fi_info -l suggests the latter. What puzzled me was that the output with FI_LOG_LEVEL=debug doesn't have any line related to the ucx provider. For example, for the sockets provider we can see:

libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var pe_waittime
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var conn_timeout
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var max_conn_retry
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var def_conn_map_sz
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var def_av_sz
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var def_cq_sz
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var def_eq_sz
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var pe_affinity
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var keepalive_enable
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var keepalive_time
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var keepalive_intvl
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var keepalive_probes
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var iface
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var max_buf_sz
libfabric:60831:1713994966::sockets:core:fi_param_define_():251<debug> registered var dgram_drop_rate

The ucx provider has similar parameter definitions at the very beginning of the provider initialization code and we expect to see similar output for the ucx provider.

Could you try again with FI_LOG_LEVEL=debug fi_info -p ucx?

Sure, see the attached file below (it is too large to paste in its entirety). I do see libfabric:114048:1714069851::ucx:core:ucx_getinfo():256<info> no ucx device is found. before the -61 return value, maybe that's what you're looking for.

out.txt

Are there some common causes of ucx providers not being usable?

Yes, that's the info I want to see.

The ucx provider looks for devices under /sys/class/infiniband with vendor id 0x15b3 (NVIDIA). It appears that no such device exists on your system.

See the code here: https://github.com/ofiwg/libfabric/blob/v1.18.x/prov/ucx/src/ucx_init.c#L207

The /sys/class/infiniband folder is indeed empty, but is that surprising given that this is a Cray Slingshot 11 machine that doesn't use IFB? (the fact that we're using SS11 has been the source of many problems, and is indeed the reason we need super-new OpenMPI as well, for libfabric support).

Not surprising at all. That just confirms that libfabric not picking up the ucx provider is an expected behavior.

Back to the segfault, what is going on inside nvshmemt_libfabric_finalize()?

[nid001249:102664:0:102664] Caught signal 11 (Segmentation fault: address not mapped to object at address 0x10)
[nid001249:102665:0:102665] Caught signal 11 (Segmentation fault: address not mapped to object at address 0x10)
[nid001249:102666:0:102666] Caught signal 11 (Segmentation fault: address not mapped to object at address 0x10)
==== backtrace (tid: 102663) ====
 0 0x00000000000168c0 __funlockfile()  ???:0
 1 0x000000000000b9fd nvshmemt_libfabric_finalize()  libfabric.cpp:0
 2 0x000000000000d312 nvshmemt_init()  ???:0
 3 0x00000000000e6c89 nvshmemi_transport_init()  :0  
 4 0x00000000000ec01c nvshmemi_common_init()  :0  
 5 0x00000000000ede99 nvshmemi_try_common_init()  :0  
 6 0x00000000000ee30e nvshmemx_host_init()  ???:0

Using good old "printf peppering," it seems to be this segment of that function that is called last (control flow never exceeds checkpoint 7b):

Based on grepping the output log with grep -E "checkpoint" out.txt, where this is repeated for each rank it seems:

** nvshmemt_libfabric_finalize checkpoint 1
** nvshmemt_libfabric_finalize checkpoint 2
** nvshmemt_libfabric_finalize checkpoint 3
** nvshmemt_libfabric_finalize checkpoint 4
** nvshmemt_libfabric_finalize checkpoint 6
** nvshmemt_libfabric_finalize checkpoint 7
** nvshmemt_libfabric_finalize checkpoint 7b

So, crash at fi_close() call, or one of the structure member accesses therein?

So the failure happened when closing the endpoints.

Question: Since fi_getinfo() previously returned -61, is there another fi_getinfo() call that succeeded later? If so, which provider was actually being used? If not, how it comes to the point to close the endpoint which is not opened before?

Hi all,

The segfault is happening in an error path inside of nvshmemt_init. It seems we are missing a null check for endpoints in an error path. (Instancing a bug for this internally). We should fail gracefully here, but do not.

I see two issues before this one that lead to us entering this error path in the first place though:

1. Since you are using slingshot 11, you need to have the CXI provider built and enabled in your libfabric installation. We can't find the CXI provider which supports the features called out here. It seems that provider wasn't compiled into your libfabric build @tylerjereddy. Looking at the source, I don't see the CXI provider in this repository until somewhere between 1.20.2 and 1.21.0. I don't know anything beyond this about compiling or loading the CXI provider with versions of libfabric <1.21.0. I have only used pre-compiled libfabric + CXI.

2. I see errors with respect to not having cuda support. This is also a compile-time option.

   [ ](libfabric:114111:1714069859::core:core:ofi_hmem_init():414 Hmem iface FI_HMEM_CUDA not supported)

These errors will disable HMEM CUDA support required by NVSHMEM to register device memory. This is a configure-time option in libfabric. if you pass --with-cuda[=dir] to configure, it will fail if it can't find the directory.

CXI is closed source, so for now I switch to our HPC module libfabric/1.15.2.0, which does have CXI, but apparently not the CUDA requirement mentioned above (see output log below). spack doesn't even have a CUDA option for libfabric config, guess I'll need to check with local HPC to see if they can do it for me or send me their build instructions for CXI-enabled libfabric so I can also add CUDA.

out_with_cxi.txt

@tylerjereddy https://github.com/ofiwg/libfabric/issues/9835 the CXI provider is open source, but it cannot be built on any machine we have come across without significant hacks.

Ok, I'm referring to the build system we usually use which indicates: https://github.com/spack/spack/blob/develop/var/spack/repos/builtin/packages/libfabric/package.py#L87

CXI is a closed source package and only exists when an external.

maybe they haven't updated yet, but the effect is the same I guess

On latest libfabric main branch:

• ./autogen.sh
• ./configure --prefix=/lustre/scratch5/treddy/march_april_2024_testing/libfabric_install_custom --enable-debug --enable-cxi=/usr --disable-opx --disable-psm2 --disable-efa --disable-rxm --disable-sockets --enable-psm3 --enable-tcp --enable-verbs --with-cuda=/lustre/scratch5/.mdt1/treddy/march_april_2024_testing/github_projects/spack/opt/spack/linux-sles15-zen2/gcc-12.2.0/nvhpc-23.11-tkgy4rjxilbay253jj65msbj2vrbq673/Linux_x86_64/23.11/cuda/12.3

I get:

configure: *** Configuring cxi provider
checking cxi_prov_hw.h usability... no
checking cxi_prov_hw.h presence... yes
configure: WARNING: cxi_prov_hw.h: present but cannot be compiled
configure: WARNING: cxi_prov_hw.h:     check for missing prerequisite headers?
configure: WARNING: cxi_prov_hw.h: see the Autoconf documentation
configure: WARNING: cxi_prov_hw.h:     section "Present But Cannot Be Compiled"
configure: WARNING: cxi_prov_hw.h: proceeding with the compiler's result
configure: WARNING:     ## ------------------------------------------ ##
configure: WARNING:     ## Report this to ofiwg@lists.openfabrics.org ##
configure: WARNING:     ## ------------------------------------------ ##

The libfabric install we have was done by HPE directly apparently, and was done on a version that preceeds direct support for CXI, so may be patched in some way as Seth had suggested to me at some point. --enable-cxi=/usr is based on HPC support basically just finding some of the relevant files for CXI in subdirs there.

@tylerjereddy the configure issues are fixed in #9793 if you want to try and cherry-pick those commits. Make sure to re-run autogen.sh.

@raffenet Cool, I ended up having to use https://github.com/thomasgillis/libfabric/tree/dev-cxi after following some more cross-linked issues/PRs that complain about LANL-specific problems, and this does seem to get me a new segfault at runtime at least, and I think the problem is still related to the CUDA provider based on the fi_info -p cuda output in the log below.

out_cxi_and_cuda_libfabric.txt

It looks like libfabric does correctly identify the number of CUDA devices based on the verbose output there, so some progress is being made. Just before the call to fi_info -p cuda "fails" with -61, I see libfabric:60885:1714254829::core:core:fi_getinfo_():1304<info> fi_getinfo: provider ofi_mrail returned -61 (No data available). Do I need that ofi_mrail provider as well here, or?

Edit: here's my current configure line for libfabric:

./configure --prefix=/lustre/scratch5/treddy/march_april_2024_testing/libfabric_install_custom --enable-cxi=/usr --disable-opx --disable-psm2 --disable-efa --disable-rxm --disable-sockets --enable-psm3 --enable-tcp --enable-verbs --with-cuda=/lustre/scratch5/.mdt1/treddy/march_april_2024_testing/github_projects/spack/opt/spack/linux-sles15-zen2/gcc-12.2.0/nvhpc-23.11-tkgy4rjxilbay253jj65msbj2vrbq673/Linux_x86_64/23.11/cuda/12.3

cc @hppritcha as well perhaps

I checked that addition of --enable-mrail to the config line above followed by clean rebuild doesn't help, fi_info -p cuda and finfo -p ofi_mrail both still fail the same way it seems.

out_cxi_and_cuda_and_mrail_libfabric.txt

I think I'm still misunderstanding something though, because I can get the same kinds of errors with fi_info -p junk, rather than getting an error about a provider name that can't possibly exist. That's pretty confusing! What's the idiomatic way to probe the CUDA support of libfabric then? Something must be different about the CUDA support..

I wonder if this is relevant:

libfabric:9451:1714321826::cxi:fabric:cxip_nic_get_ss_env_get_vni():23<info> nid001204: SLINGSHOT_VNIS not found

@tylerjereddy You don't need to enable the mrail provider. The line fi_getinfo: provider ofi_mrail returned -61 in the debug trace is normal which just says the mrail provider can't be used.

There is no provider called cuda so passing -p cuda is equivalent to use -p junk. Usually, the provider would pick the GPU/device interface (or HMEM interface in libfabric's term) based on either application input (e.g. iface passed to fi_mr_regattr) or some autodetect logic. You can force to use cuda only by setting FI_HMEM=cuda. That setting will prevent other HMEM interfaces from being initialized.

I get the same backtraces with export FI_HMEM=cuda, darn:

out_fi_hmem_cuda.txt

@seth-howell I built NVSHMEM with NVSHMEM_DEBUG=1 (along with verbose libfabric settings) and tried to pull out some relevant portions of the 150 MB log file for the 2-node cuFFTMp reproducer. In this scenario, it looks like we get things erroring out instead of hitting a segfault. This output isn't in order, just grepped based on things that look suspicious.

Looks like a mixture of remote memory access (RMA), GDRCopy, CXI, and CUDA-related complaints, but not sure if it is clear what needs to be addressed.

libfabric:98917:1714598377::cxi:ep_data:cxip_rma_emit_dma():329<warn> nid001448: TXC (0x1b32:0): Failed to emit dma command: -11:Resource temporarily unavailable
libfabric:98917:1714598377::cxi:ep_data:cxip_rma_common():635<warn> nid001448: TXC (0x1b32:0): DMA RMA write failed: buf=0x150278b80600 len=80 rkey=0 roffset=0x18b806a0 nic=0x1b72 pid=0 pid_idx=117

nid001449:67129:67129 [2] NVSHMEM INFO [6] nvshmemi_get_cucontext->cuCtxGetDevice->0(CUDA_ERROR_INVALID_CONTEXT 201) cuStreamCreateWithPriority my_stream 0x4e147b0
nid001448:98917:98917 [1] NVSHMEM INFO [1] status 0 cudaErrorInvalidValue 1 cudaErrorInvalidSymbol 13 cudaErrorInvalidMemcpyDirection 21 cudaErrorNoKernelImageForDevice 209
libfabric:98918:1714598371::core:core:cuda_gdrcopy_dev_unregister():362<warn> gdr_unmap failed! error: Invalid argument
/lustre/scratch5/treddy/march_april_2024_testing/nvshmem_custom_dl/nvshmem_src_2.10.1-3/src/modules/transport/libfabric/libfabric.cpp:nvshmemt_libfabric_rma:517: Received an error when trying to post an RMA operation.
/lustre/scratch5/treddy/march_april_2024_testing/nvshmem_custom_dl/nvshmem_src_2.10.1-3/src/include/internal/common/nvshmem_internal.h:nvshmemi_process_multisend_rma:302: aborting due to error in process_channel_dma
libfabric:67128:1714598429::core:core:cuda_gdrcopy_dev_register():333<warn> gdr_map failed! error: Cannot allocate memory

/lustre/scratch5/treddy/march_april_2024_testing/nvshmem_custom_dl/nvshmem_src_2.10.1-3/src/modules/transport/libfabric/libfabric.cpp 1706 GDRCopy requested, but unused by transport. Disabling
libfabric:98717:1714598347::core:core:fi_param_get_():372<info> variable hmem_cuda_use_gdrcopy=<not set>
libfabric:67130:1714598364::core:core:fi_param_get_():399<info> read bool var hmem_cuda_use_gdrcopy=1

Here's a shorter ouput log with different debug/verbosity settings and also export GDRCOPY_ENABLE_LOGGING=1 turned on. It seems to have extra output with ERR: mh is not mapped yet before hitting the segfault via NVSHMEM->CXI->CUDA (gdrcopy).

It looks like an NVIDIA engineer responded to the ERR: mh is not mapped yet type error at https://github.com/NVIDIA/gdrcopy/issues/242#issuecomment-1396372455, but needed a custom build of gdrcopy to get to the bottom of it. I'm assuming that we're ok to use spack find -v gdrcopy

gdrcopy@2.3+cuda build_system=makefile cuda_arch=none patches=c5efec1

When I build gdrcopy master branch by hand, in the interest of pulling in https://github.com/NVIDIA/gdrcopy/pull/248, and then rebuild NVSHMEM 2.10.1 to point at that version of gdrcopy, I do see different behavior.

In particular, I see the cuFFTMp reproducer code hang on two nodes instead of segfaulting. Not sure if this might be diagnostically useful cc @pakmarkthub.

I guess the Slingshot 11 network was having problems (based on feedback from local HPC). When I re-run today with gdrcopy master branch I see this output/backtrace:

I believe using the correct CXI provider build of libfabric resolved the original segfaults here.

We're now seeing other crashes in gdrcopy, quite possibly from libfabric incantations, per the cross-linked issue above. However, I think it is best if I close this and open a separate issue if I manage to get a minimum viable reproducer for libfabric side.