PedramAlizadeh · Oct 15, 2026 · Oct 6, 2026 · Oct 7, 2026 · Oct 7, 2026 · Oct 13, 2026
@@ -25,6 +25,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * `RCCL_TUNER_EMBEDDED_CONFIG` (default `1`): the built-in CSV tuner configurations in `tuner/*.csv` are now compiled into `librccl.so` and used as the stock tuning defaults, so a cluster no longer has to stage the CSV on every node or on NFS. Set `0` to ignore the embedded configurations. `NCCL_TUNER_CONFIG_FILE` and the on-disk tuner directories still take precedence, and `NCCL_TUNER_PLUGIN=none` still disables the tuner entirely.
 
 ### Changed
+* `install.sh --all_unrolls` (`-DBUILD_ALL_UNROLLS=ON`) now adds the unroll factors the multi-arch default skips (8 and 16) on top of each architecture's existing set. It no longer replaces that set or clears the gfx1250 pin on unroll 32. The flag stays off by default, so a normal build does not compile unroll 8 or 16.
 * Host-side device API declarations (`ncclDevCommCreate`, `ncclDevCommDestroy`, `ncclCommQueryProperties`, `ncclGetPeerDevicePointer`, the `*CreateRequirement` helpers, the `ncclDevCommRequirements` and `ncclCommProperties` structs and their initializers) moved to `nccl_device/host.h`. `nccl_device.h` still includes it; code that includes individual `nccl_device/*.h` headers directly must also include `nccl_device/host.h`.
 * Profiler `ncclProfileKernelPhase` events are emitted only by symmetric kernels, matching NCCL 2.32. Regular and P2P kernels report only `ncclProfileKernelCh` start and stop.
 * Updated the RMA plugin interface to v16, which adds a per-communicator `getRmaProperties` query. Plugins built against v13, v14 or v15 still load.

@@ -80,7 +80,7 @@ include(cmake/NcclCompat.cmake)
 rccl_define_nccl_compat_options()
 
 option(BUILD_ADDRESS_SANITIZER                 "Enable address sanitizer"                      ${RCCL_DEFAULT_BUILD_ADDRESS_SANITIZER})
-option(BUILD_ALL_UNROLLS                       "Generate every unroll factor (1,2,4,8,16,32) for the targeted GPU arch(es)" OFF)
+option(BUILD_ALL_UNROLLS                       "Also generate unroll factors 8 and 16 on top of the per-arch default set" OFF)
 option(BUILD_BFD                               "Enable custom backtrace (if bfd.h exists)"     OFF)
 option(BUILD_LOCAL_GPU_TARGET_ONLY             "Build only for GPUs detected on this machine"  OFF)
 option(BUILD_SHARED_LIBS                       "Build as shared library"                       ON)

@@ -74,7 +74,7 @@ function display_help()
     echo "RCCL build & installation helper script"
     echo " Options:"
     echo "       --address-sanitizer     Build with address sanitizer enabled"
-    echo "       --all_unrolls           Build every unroll factor (1,2,4,8,16,32) for the targeted GPU arch(es) instead of the per-arch default set"
+    echo "       --all_unrolls           Also build unroll factors 8 and 16 on top of the per-arch default set"
     echo "       --amdgpu_targets        Only compile for specified GPU architecture(s). For multiple targets, separate by ';' (builds for all supported GPU architectures by default)"
     echo "       --cmake-options         Pass additional CMake options (e.g. --cmake-options \"-DFOO=BAR -DBAZ=ON\")"
     echo "       --debug                 Build debug library"
@@ -399,7 +399,7 @@ if [[ "${build_local_gpu_only}" == true ]]; then
     cmake_common_options="${cmake_common_options} -DBUILD_LOCAL_GPU_TARGET_ONLY=ON"
 fi
 
-# Build every unroll factor for the targeted arch(es)
+# Also build unroll 8 and 16 on top of the per-arch default
 if [[ "${build_all_unrolls}" == true ]]; then
     cmake_common_options="${cmake_common_options} -DBUILD_ALL_UNROLLS=ON"
 fi

@@ -18,7 +18,7 @@
 # Unroll factors whose device functions are compiled for one arch only, and the single
 # source of truth for that restriction: get_arch_guard(), the specialized_files.txt guard
 # the device linker filters on, the local-arch unroll set, and the ncclDevFuncUnrollArch[]
-# table the host checks before selecting an unroll. --all_unrolls clears it.
+# table the host checks before selecting an unroll. --all_unrolls does not clear it.
 unroll_arch_requirement = {
   "32": "gfx1250",
 }
@@ -132,7 +132,8 @@ def paste(sep, *args):
 is_ifc             = 1 if sys.argv[2] == "ON" else 0
 is_local_arch_only = 1 if sys.argv[4] == "ON" else 0
 is_rocshmem        = 1 if sys.argv[5] == "ON" else 0
-# BUILD_ALL_UNROLLS / install.sh --all_unrolls: generate every unroll for the targeted archs.
+# BUILD_ALL_UNROLLS / install.sh --all_unrolls: also generate unroll 8 and 16.
+# Off by default, so a normal build does not compile those factors.
 build_all_unrolls  = 1 if sys.argv[6] == "ON" else 0
 
 func_pattern = sys.argv[7:8]
@@ -306,12 +307,16 @@ def calc_unroll_and_pipeline_for_local_arch():
 # except for gfx950. For gfx950, we also disable pipelining.
 local_unroll, local_pipeline = calc_unroll_and_pipeline_for_local_arch()
 
-# --all_unrolls widens the unroll set and compiles every one for the targeted archs, so the
-# pin no longer holds. The arch's pipeline decision stands. Must follow the call above, the
-# only reader of unrolls_requiring_arch.
+# --all_unrolls appends the unrolls the multi-arch default skips (8 and 16). The
+# per-arch set and any arch pin stay as calc_unroll_and_pipeline_for_local_arch
+# decided them, and the arch's pipeline decision stands. A build without the flag
+# does not compile the skipped factors. Must follow the call above, the only
+# reader of unrolls_requiring_arch.
 if build_all_unrolls:
-  local_unroll = all_unrolls
-  unroll_arch_requirement.clear()
+  local_unroll = list(local_unroll)
+  for unroll in [u for u in all_unrolls if u not in default_unrolls]:
+    if unroll not in local_unroll:
+      local_unroll.append(unroll)
 
 # rocSHMEM/GDA-based collectives: only generated when ENABLE_ROCSHMEM build is requested
 gda_colls = {"AlltoAllGda", "AlltoAllvGda"}

@@ -245,26 +245,62 @@ def test_device_tables_are_index_aligned_across_unrolls(self):
                 % (unroll, base_unroll),
             )
 
-    def test_all_unrolls_opt_in_generates_every_unroll(self):
-        # BUILD_ALL_UNROLLS fills in the skipped unrolls and drops the gfx1250 restriction.
+    def test_default_build_omits_unroll_8_and_16(self):
+        # Multi-arch default is 1/2/4 plus gfx1250-pinned 32. Unroll 8 and 16 stay
+        # out, so a default build does not compile them.
+        self.assertEqual({"1", "2", "4", "32"}, set(_unroll_tables(self.header)))
+        generated = self._unroll_table("ncclDevFuncUnrollGenerated", r"true|false")
+        self.assertEqual(
+            {
+                "1": "true",
+                "2": "true",
+                "4": "true",
+                "8": "false",
+                "16": "false",
+                "32": "true",
+            },
+            generated,
+        )
+        pinned = self._unroll_table("ncclDevFuncUnrollArch", r'nullptr|"gfx\w+"')
+        self.assertEqual(
+            {
+                "1": "nullptr",
+                "2": "nullptr",
+                "4": "nullptr",
+                "8": "nullptr",
+                "16": "nullptr",
+                "32": '"gfx1250"',
+            },
+            pinned,
+        )
+
+    def test_all_unrolls_adds_8_and_16_without_unpinning(self):
+        # BUILD_ALL_UNROLLS appends 8 and 16 to the per-arch default. Unroll 32
+        # stays compiled for gfx1250 only, and the other defaults stay generated.
         with tempfile.TemporaryDirectory(prefix="rccl_devtable_all_") as tmpdir:
             header = _generate(tmpdir, all_unrolls="ON")
             host = _read_generated(tmpdir, "host_table.cpp")
         self.assertEqual({"1", "2", "4", "8", "16", "32"}, set(_unroll_tables(header)))
-        # Both host tables flip under the flag, and no other test reads them on this path.
         generated = self._unroll_table("ncclDevFuncUnrollGenerated", r"true|false", host)
         self.assertEqual(
             ["true"] * 6,
             list(generated.values()),
-            "BUILD_ALL_UNROLLS must mark every unroll generated",
+            "BUILD_ALL_UNROLLS must mark the default unrolls plus 8 and 16 generated",
         )
         pinned = self._unroll_table("ncclDevFuncUnrollArch", r'nullptr|"gfx\w+"', host)
         self.assertEqual(
-            ["nullptr"] * 6,
-            list(pinned.values()),
-            "BUILD_ALL_UNROLLS compiles every unroll for the target, so none stays pinned",
+            {
+                "1": "nullptr",
+                "2": "nullptr",
+                "4": "nullptr",
+                "8": "nullptr",
+                "16": "nullptr",
+                "32": '"gfx1250"',
+            },
+            pinned,
+            "BUILD_ALL_UNROLLS must leave the gfx1250 pin on unroll 32",
         )
-        self.assertNotIn(
+        self.assertIn(
             "#if (defined(__gfx1250__) || defined(__gfx1250_strict__))\n", header
         )
 

@@ -970,7 +970,7 @@ extern bool const ncclDevFuncUnrollGenerated[NCCL_NUM_UNROLLS];
 // Arch each unroll factor's device functions were compiled for, or nullptr when
 // the unroll carries no arch restriction. Neither table implies the other, so both
 // have to be consulted: an entry names its arch even for an unroll this build left
-// out, and BUILD_ALL_UNROLLS clears every entry because it compiles them all.
+// out. BUILD_ALL_UNROLLS does not clear the pin.
 extern char const* const ncclDevFuncUnrollArch[NCCL_NUM_UNROLLS];
 
 // `ncclDevFuncId()` needs to be in sync with 'all_colls' in generate.py

@@ -2654,8 +2654,7 @@ ncclResult_t commSetUnrollFactor(struct ncclComm* comm) {
     case ncclUnrollWrongArch:
       WARN("RCCL_UNROLL_FACTOR %d (unroll %d) is compiled for %s only, so its device function table is empty on "
            "arch %s and dispatching to it would crash. "
-           "Rebuild with install.sh --all_unrolls (-DBUILD_ALL_UNROLLS=ON) to compile it for every target, or select a "
-           "different one.",
+           "Select a different unroll factor.",
            comm->unroll, (1 << comm->unroll), ncclDevFuncUnrollArch[comm->unroll], comm->archName);
       return ncclInvalidArgument;
     case ncclUnrollUsable:

@@ -2900,9 +2900,9 @@ TEST(SkipPresetTopoMatching, Gfx1250_SkipsRomeModelMatching)
 // test_runner/configs), and Rcclwrap.* is the only listed pattern this file
 // matches. A suite of their own would never be run.
 // ---------------------------------------------------------------------------
-// Which unroll factors are arch-pinned is a build property, not a constant:
-// BUILD_ALL_UNROLLS compiles every one for the targeted archs and pins none. Ask
-// the table for a pinned factor rather than hardcoding 32.
+// Which unroll factors are arch-pinned is a build property, not a constant.
+// BUILD_ALL_UNROLLS adds unroll 8 and 16 and leaves an existing pin in place.
+// Ask the table for a pinned factor rather than hardcoding 32.
 constexpr char kOtherArch[] = "gfx1200";
 
 TEST(Rcclwrap, UnrollFactor_RejectsArchRestrictedUnrollOnOtherArch)

@@ -345,16 +345,49 @@ def test_all_unrolls_opt_in_adds_the_skipped_unrolls(tmp_path_factory, generated
     )
 
     counts = _count_by(records, "unroll")
-    assert set(counts) == {"1", "2", "4", "8", "16", "32"}, (
-        "--all_unrolls did not generate every unroll: %s" % sorted(counts)
+    baseline = _count_by(generated["OFF"]["records"], "unroll")
+    assert set(counts) - set(baseline) == {"8", "16"}, (
+        "--all_unrolls should add only unroll 8 and 16: %s vs %s" % (sorted(counts), sorted(baseline))
+    )
+    assert set(baseline) <= set(counts), (
+        "--all_unrolls dropped a default unroll: %s vs %s" % (sorted(counts), sorted(baseline))
     )
     assert len(set(counts.values())) == 1, "--all_unrolls broke unroll lockstep: %s" % counts
-
-    baseline = _count_by(generated["OFF"]["records"], "unroll")
     assert set(counts.values()) == set(baseline.values()), (
         "--all_unrolls changed the per-unroll kernel count: %s vs %s" % (counts, baseline)
     )
 
+    # Counts alone cannot tell this apart from the old flag, which replaced the
+    # set with every unroll and cleared the gfx1250 pin: both emit {1,2,4,8,16,32}
+    # at the same per-unroll count. The pin catches the cleared-pin half only; at
+    # local_gpu_only=OFF append and replace give the same set. SendRecv LL128
+    # (reg=1) is a wider guard on purpose.
+    pin = re.compile(
+        r"^\(defined\(__gfx1250__\) \|\| defined\(__gfx1250_strict__\)\)"
+        r"(?: && defined\(ENABLE_LL128\))?$"
+    )
+    pinned_32 = sendrecv_ll128_32 = 0
+    for fp in sorted(glob.glob(os.path.join(str(d), "specialized", "*.cpp"))):
+        with open(fp) as f:
+            text = f.read()
+        match = _DEFINE_RE.search(text)
+        assert match is not None, "specialized file has no DEFINE_ncclDevFunc: %s" % fp
+        guards = re.findall(r"^#if (.*)$", text, re.M)
+        guard = guards[0] if guards else ""
+        unroll = match.group("unroll")
+        if unroll == "32" and match.group("coll") == "SendRecv" and match.group("reg") == "1":
+            assert not pin.match(guard), "SendRecv LL128 unroll 32 must keep its wider guard"
+            sendrecv_ll128_32 += 1
+        elif unroll == "32":
+            assert pin.match(guard), "unroll 32 lost the gfx1250 pin: %s guard %r" % (fp, guard)
+            pinned_32 += 1
+        else:
+            assert not pin.match(guard), "unroll %s picked up the gfx1250 pin: %s" % (unroll, fp)
+    assert pinned_32 > 0, "--all_unrolls emitted no gfx1250-pinned unroll 32 kernels"
+    assert sendrecv_ll128_32 == 1, (
+        "expected one SendRecv LL128 unroll-32 kernel, found %d" % sendrecv_ll128_32
+    )
+
 
 @pytest.mark.main_generator
 def test_rocshmem_on_is_off_plus_gda(generated):