Orientation: antirez/ds4

Entry 13 goes to device 1, not device 0. d only ever advances; the while loop never rewinds it, so once the walker leaves a device it never comes back — even when a later, smaller layer would fit in the gap.

Each device therefore gets a contiguous run of layers. Inference walks layers in order, so contiguous placement means the activation crosses a device boundary once per device instead of ping-ponging. A tighter bin-packing would waste less VRAM and cost far more PCIe traffic.

The same one-way walk produces the CPU behaviour the comment calls out: once an entry falls through to DS4_LAYER_PACK_CPU, d is parked at n_gpus and every later entry lands on CPU too. The tail of the model runs on CPU, not a random scatter of leftovers.

ROCm is not a runtime backend at all. It is a build-time substitution, and the trail runs through the Makefile, not the header.

# Makefile — strix-halo target
strix-halo:
	$(MAKE) -B ds4 ds4-server ds4-bench ds4-eval ds4-agent \
		CORE_OBJS="ds4.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_rocm.o ds4_rocm_compat.o ds4_rocm_unavailable.o ds4_layer_pack.o" \
		CFLAGS="$(CFLAGS) -DDS4_ROCM_BUILD" \
		DS4_LINK="$(HIPCC) $(ROCM_CFLAGS)" \
		DS4_LINK_LIBS="$(ROCM_LDLIBS)"

rocm: strix-halo

make strix-halo swaps ds4_cuda.o out of CORE_OBJS for ds4_rocm.o, links with hipcc, and defines DS4_ROCM_BUILD. ds4_rocm.h then rewrites the CUDA API onto HIP with a wall of macros:

/* ds4_rocm.h */
#include <hip/hip_runtime.h>
#include <hipblas/hipblas.h>
#include <hip/hip_fp16.h>
#include <hipcub/hipcub.hpp>
#include <rocwmma/rocwmma-version.hpp>
#include <rocwmma/rocwmma.hpp>
...
#define cudaError_t hipError_t
#define cudaStream_t hipStream_t
#define cudaEvent_t hipEvent_t
#define cudaDeviceProp hipDeviceProp_t
#define cudaMemLocation hipMemLocation
#define cudaMalloc hipMalloc
#define cudaMallocHost hipHostMalloc
#define cudaMallocManaged hipMallocManaged
#define cudaFree hipFree
#define cudaFreeHost hipFreeHost
#define cudaMemset hipMemset

So an AMD binary still selects DS4_BACKEND_CUDA at runtime. The only place the difference surfaces to a user is here:

/* ds4.c:47258 */
const char *ds4_backend_name(ds4_backend backend) {
    switch (backend) {
    case DS4_BACKEND_METAL: return "metal";
    case DS4_BACKEND_CUDA:
#ifdef DS4_ROCM_BUILD
        return "rocm";
#else
        return "cuda";
#endif
    case DS4_BACKEND_CPU:   return "cpu";
    }
    return "unknown";
}

Three consequences worth carrying: the enum is a link-time choice, not a dispatch table, so one binary cannot switch backends; #ifdef DS4_ROCM_BUILD is the string to grep when AMD behaviour diverges (it appears in the CLI, the server and ds4_gpu.h); and the CPU build is a fourth path with its own object set, which is why make cpu compiles every .c a second time into *_cpu.o.

In C, the configuration another language would put in a plugin registry usually lives in the build system. When a header does not contain the concept you were promised, look there next.