diff --git a/README.md b/README.md index 4b8126c725..a5f8261a84 100644 --- a/README.md +++ b/README.md @@ -121,8 +121,8 @@ rocSHMEM has the following enviroment variables: Defines the size of the rocSHMEM symmetric heap Note the heap is on the GPU memory. - ROCSHMEM_RO_DISABLE_IPC (default : 0) - Disables IPC support for the reverse offload backend. + ROCSHMEM_RO_DISABLE_MIXED_IPC (default : 0) + Disables IPC support for the RO or GDA backends. ROCSHMEM_MAX_NUM_CONTEXTS (default : 1024) Maximum number of contexts used in library. diff --git a/docs/compile_and_run.rst b/docs/compile_and_run.rst index 3d78c3f556..767c923283 100644 --- a/docs/compile_and_run.rst +++ b/docs/compile_and_run.rst @@ -84,9 +84,9 @@ You can control the behavior of rocSHMEM by using the following environment vari * - ROCSHMEM_UNIQUEID_WITH_MPI - 0 - Defines whether rocSHMEM is expected to use MPI when using the uniqueId based initialization. - * - ROCSHMEM_RO_DISABLE_IPC + * - ROCSHMEM_DISABLE_MIXED_IPC - 0 - - Defines whether to force using the RO conduit even when IPC is available. + - Defines whether to force using the network conduit even when IPC is available. * - ROCSHMEM_GDA_ALTERNATE_QP_PORTS - 1 - Enables/Disables having QPs alternate their mappings across rocSHMEM contexts. This helps saturate bandwidth on multiport bonded interfaces. diff --git a/src/envvar.cpp b/src/envvar.cpp index 436ffa2d74..9b94e31d28 100644 --- a/src/envvar.cpp +++ b/src/envvar.cpp @@ -47,6 +47,8 @@ namespace envvar { const var requested_dev("USE_IB_HCA", ""); const var sq_size("SQ_SIZE", "", 1024); const var backend("BACKEND", ""); + const var disable_mixed_ipc("DISABLE_MIXED_IPC", "", false); + const var disable_ipc("DISABLE_IPC", "", false); } // inline namespace _base namespace bootstrap { diff --git a/src/envvar.hpp b/src/envvar.hpp index b80d84e7a9..583901ac08 100644 --- a/src/envvar.hpp +++ b/src/envvar.hpp @@ -405,6 +405,8 @@ namespace envvar { extern const var heap_size; extern const var max_num_teams; extern const var backend; + extern const var disable_mixed_ipc; + extern const var disable_ipc; /** * @brief Maximum number of contexts for the application diff --git a/src/ipc_policy.cpp b/src/ipc_policy.cpp index d970d17925..313607b1e1 100644 --- a/src/ipc_policy.cpp +++ b/src/ipc_policy.cpp @@ -108,9 +108,14 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, */ free(vec_ipc_handle); - if (!envvar::ro::disable_ipc) { - int thread_comm_rank {-1}; - + if (envvar::ro::disable_ipc || envvar::disable_ipc) { + if (0 == my_pe) { + printf("ROCSHMEM_RO_DISABLE_IPC and RO_DISABLE_IPC environment variables have been deprecated.\n" + " Please use ROCSHMEM_DISABLE_MIXED_IPC as a replacement.\n"); + } + } + auto disable_ipc = envvar::disable_mixed_ipc || envvar::ro::disable_ipc || envvar::disable_ipc; + if (!disable_ipc) { CHECK_HIP(hipMalloc(reinterpret_cast(&pes_with_ipc_avail), shm_size * sizeof(int))); MPI_Group thread_grp; @@ -187,9 +192,14 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, */ free(vec_ipc_handle); - if (!envvar::ro::disable_ipc) { - int thread_comm_rank {-1}; - + if (envvar::ro::disable_ipc || envvar::disable_ipc) { + if (0 == my_pe) { + printf("ROCSHMEM_RO_DISABLE_IPC and RO_DISABLE_IPC environment variables have been deprecated.\n" + " Please use ROCSHMEM_DISABLE_MIXED_IPC as a replacement.\n"); + } + } + auto disable_ipc = envvar::disable_mixed_ipc || envvar::ro::disable_ipc || envvar::disable_ipc; + if (!disable_ipc) { CHECK_HIP(hipMalloc(reinterpret_cast(&pes_with_ipc_avail), shm_size * sizeof(int))); std::copy(shm_ranks.begin(), shm_ranks.end(), pes_with_ipc_avail); }