Make ROCSHMEM_DISABLE_MIXED_IPC a synonym for ROCSHMEM_RO_DISABLE_IPC, ROCSHMEM_DISABLE_IPC (#273)

* Make ROCSHMEM_DISABLE_IPC a synonym for ROCSHMEM_RO_DISABLE_IPC

* Introduce ROCSHMEM_DISABLE_MIXED_IPC and deprecate old variants

[ROCm/rocshmem commit: db8e5f1086]

projects/rocshmem/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.

projects/rocshmem/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.

projects/rocshmem/src/envvar.cpp

@@ -47,6 +47,8 @@ namespace envvar {
     const var<std::string> requested_dev("USE_IB_HCA", "");
     const var<uint32_t> sq_size("SQ_SIZE", "", 1024);
     const var<std::string> backend("BACKEND", "");
+    const var<bool> disable_mixed_ipc("DISABLE_MIXED_IPC", "", false);
+    const var<bool> disable_ipc("DISABLE_IPC", "", false);
   }  // inline namespace _base
 
   namespace bootstrap {

projects/rocshmem/src/envvar.hpp

@@ -405,6 +405,8 @@ namespace envvar {
     extern const var<size_t> heap_size;
     extern const var<size_t> max_num_teams;
     extern const var<std::string> backend;
+    extern const var<bool> disable_mixed_ipc;
+    extern const var<bool> disable_ipc;
 
     /**
      * @brief Maximum number of contexts for the application

projects/rocshmem/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<void**>(&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<void**>(&pes_with_ipc_avail), shm_size * sizeof(int)));
     std::copy(shm_ranks.begin(), shm_ranks.end(), pes_with_ipc_avail);
   }