Potential deadlock in CavityManager::migrate() without color-based scheduling

[will-zzy]

Hi,

I may have found a potential deadlock issue while integrating my own edge-split kernel with RXMesh dynamic topology.

In my setup, the kernel can hang inside cavity.prologue(), more specifically in CavityManager::migrate()->lock_neighbour_patches() during edge split.

From debugging, it seems that multiple adjacent patches may enter migrate()->lock_neighbour_patches() in the same round without an effective color-based gating step on the path I am using. Since each block has already acquired the lock for its own patch, and then tries to acquire multiple neighboring patch locks according to PatchStash slot order, the lock acquisition order is not globally consistent. This appears to create a circular-wait scenario, and in practice the kernel hangs before the split synchronization point.

My understanding is that the issue is not in the spin lock itself, but in the scheduling/rounding discipline:
only patches of the current color should be allowed to enter CavityManager in that round, and the host-side loop should drive the round with something like: num_inner_iter % rx.get_num_colors()

In cavity_manager.cuh:

struct CavityColorIteration
{
    uint32_t value;
};

Then restore the color-based filter here

uint32_t color = INVALID32;
if (s_patch_id != INVALID32 && iteration.value != INVALID32) {
    color = m_context.m_patches_info[s_patch_id].color;
    if (color != iteration.value) {
        push(s_patch_id);
        s_patch_id = INVALID32;
    }
}

In my local experiment, I introduced an explicit color-round argument to the CavityManager construction path used by my split kernel, so that the existing color filter could actually be applied on that path.

template <uint32_t blockThreads, CavityOp cop>
device forceinline CavityManager<blockThreads, cop>::CavityManager(
cooperative_groups::thread_block& block,
Context& context,
ShmemAllocator& shrd_alloc,
bool preserve_cavity,
CavityColorIteration iteration,
bool allow_touching_cavities,
uint32_t current_p)

and the process is like:

        while (!rx->is_queue_empty()) {
            if (++inner_iter > kMaxInnerIter) {
                break;
            }
            const uint32_t color_iteration = inner_iter % num_colors;
            printf("[pass %d] before split launch\n", inner_iter);
            if (use_provenance) {
                face_subdivide_split_with_provenance<float, blockThreads>
                    <<<lb_split.blocks,
                       lb_split.num_threads,
                       lb_split.smem_bytes_dyn>>>(
                        rx->get_context(),
                        *coords,
                        *v_boundary,
                        *edge_status,
                        *edge_flip,
                        *persistent_id,
                        *parents,
                        color_iteration,
                        debug_stage_limit,
                        debug_prologue_stage_limit,
                        debug_migrate_stage_limit);
            } else {
                face_subdivide_split<float, blockThreads>
                    <<<lb_split.blocks,
                       lb_split.num_threads,
                       lb_split.smem_bytes_dyn>>>(
                        rx->get_context(),
                        *coords,
                        *v_boundary,
                        *edge_status,
                        *edge_flip,
                        color_iteration,
                        debug_stage_limit,
                        debug_prologue_stage_limit,
                        debug_migrate_stage_limit);
            }
            CUDA_ERROR(cudaGetLastError());
            printf("[pass %d] after split launch\n", inner_iter);
            CUDA_ERROR(cudaDeviceSynchronize());
            printf("[pass %d] after split sync\n", inner_iter);

            printf("[pass %d] before cleanup1\n", inner_iter);
            rx->cleanup();

            printf("[pass %d] after cleanup1 call\n", inner_iter);

            CUDA_ERROR(cudaDeviceSynchronize());
            printf("[pass %d] after cleanup1 sync\n", inner_iter);

            printf("[pass %d] before slice_patches\n", inner_iter);
            if (use_provenance) {
                rx->slice_patches(
                    *coords,
                    *v_boundary,
                    *edge_status,
                    *edge_flip,
                    *persistent_id,
                    *parents);
            } else {
                rx->slice_patches(*coords, *v_boundary, *edge_status, *edge_flip);
            }
            CUDA_ERROR(cudaDeviceSynchronize());
            printf("[pass %d] after slice_patches\n", inner_iter);
            printf("[pass %d] before cleanup2\n", inner_iter);
            rx->cleanup();
            CUDA_ERROR(cudaDeviceSynchronize());
            printf("[pass %d] after cleanup2\n", inner_iter);
        }

Please let me know if my understanding is incorrect. Thank you for your time and for maintaining RXMesh.

[will-zzy]

I also observed that, during the query stage, newly created cavities should preferably be owned by the current patch. Otherwise, after cleanup and hash table calibration, a validation failure may occur: a non-owned entry may still reference an owner edge that has already been deleted.
So it may be better, at least in my case, to add an ownership check for the queried element before creating the cavity:

    query.dispatch<Op::EVDiamond>(
        block,
        shrd_alloc,
        [&](const EdgeHandle& eh, const VertexIterator& iter) {
            assert(iter.size() == 4);
            if (edge_status(eh) == TO_SPLIT) {
                if (!cavity.patch_info().is_owned(rxmesh::LocalEdgeT(eh.local_id()))) {
                    edge_status(eh) = SKIP;
                    return;
                }

                const VertexHandle va = iter[0];
                const VertexHandle vb = iter[2];
                const VertexHandle vc = iter[1];
                const VertexHandle vd = iter[3];

                if (!va.is_valid() || !vb.is_valid() || !vc.is_valid() ||
                    !vd.is_valid()) {
                    edge_status(eh) = SKIP;
                    return;
                }

                if (v_boundary(va) || v_boundary(vb) || v_boundary(vc) ||
                    v_boundary(vd)) {
                    edge_status(eh) = SKIP;
                    return;
                }

                if (va == vb || vb == vc || vc == va || va == vd ||
                    vb == vd || vc == vd) {
                    edge_status(eh) = SKIP;
                    return;
                }

                cavity.create(eh);
            }
            
        });

[Ahdhn]

Thanks @will-zzy for raising this issue. We have discussed the different considerations for the scheduling of patches in this paper (Section 4.2 for the pros and cons of different designs, Section 5.3 for the implementation details, and Appendix B for the locking algorithm). In a nutshell, coloring does not guarantee that patches of a certain color can be processed concurrently. For example, patches A and B of the same color are supposed to not be connected, but changes in patch A may lead to creating mesh elements that now touch patch B.

This is why we decided to use spin locking with a simple backoff strategy. The main idea is to make the common case faster. The common case here is that we usually have way more patches than we have SMs on the GPU. Thus, the probability of having two neighboring patches scheduled at the same time is low. Thus, the spin lock (for locking all neighboring patches) should--with high probability--be able to lock all neighboring patches without spinning for too long.

Now, in cases of the long tail of the execution, i.e., when most patches are processed and we are left with a few patches, the spin lock might reach a deadlock, as you pointed out. This is where the backoff strategy is most useful. It basically allows a patch to retire without spinning endlessly. However, this might lead to livelock. For example, patches A and B are neighbors, they are scheduled at the same time, both enter the spin lock, both back off at the same time since they are conflicting with one another, both retire, and then both try again. This is a livelock. We prevent this by having a priority-based backoff strategy where a patch with a higher index will retire if there is another patch with a lower index also spinning on the same resource (i.e., the same patch). This allows the lower-index patch to proceed. The worst case scenario here is that we will have only one patch to proceed but this at least guarantee forward progress. Of course, we don't have a formal proof that this will work all the time, but I have not seen evidence otherwise.

Now, there might be a bug in the code that requires more investigation before considering graph coloring (it might be something related to the CUDA memory consistency model). I think the coloring trick just hides the bugs and does not fix them.

Can you please provide code that reproduces the bug? I will be happy to investigate it further.

[Ahdhn]

Regarding the ownership check, it is already part of the query dispatch. However, there might be edge cases if you are doing query and updates at the same time. Can you provide a code that reproduces the issue?

[will-zzy]

Thanks @will-zzy for raising this issue. We have discussed the different considerations for the scheduling of patches in this paper (Section 4.2 for the pros and cons of different designs, Section 5.3 for the implementation details, and Appendix B for the locking algorithm). In a nutshell, coloring does not guarantee that patches of a certain color can be processed concurrently. For example, patches A and B of the same color are supposed to not be connected, but changes in patch A may lead to creating mesh elements that now touch patch B.

This is why we decided to use spin locking with a simple backoff strategy. The main idea is to make the common case faster. The common case here is that we usually have way more patches than we have SMs on the GPU. Thus, the probability of having two neighboring patches scheduled at the same time is low. Thus, the spin lock (for locking all neighboring patches) should--with high probability--be able to lock all neighboring patches without spinning for too long.

Now, in cases of the long tail of the execution, i.e., when most patches are processed and we are left with a few patches, the spin lock might reach a deadlock, as you pointed out. This is where the backoff strategy is most useful. It basically allows a patch to retire without spinning endlessly. However, this might lead to livelock. For example, patches A and B are neighbors, they are scheduled at the same time, both enter the spin lock, both back off at the same time since they are conflicting with one another, both retire, and then both try again. This is a livelock. We prevent this by having a priority-based backoff strategy where a patch with a higher index will retire if there is another patch with a lower index also spinning on the same resource (i.e., the same patch). This allows the lower-index patch to proceed. The worst case scenario here is that we will have only one patch to proceed but this at least guarantee forward progress. Of course, we don't have a formal proof that this will work all the time, but I have not seen evidence otherwise.

Now, there might be a bug in the code that requires more investigation before considering graph coloring (it might be something related to the CUDA memory consistency model). I think the coloring trick just hides the bugs and does not fix them.

Can you please provide code that reproduces the bug? I will be happy to investigate it further.

@Ahdhn
Thank you very much for your further response.

I have prepared a minimal reproducible example, which includes:

The RXMesh repository (commit e27a88e)
The RXMesh Python wrapper
A Python test script simulating my subdivide workflow
A mesh with one million faces for testing
A README with installation and reproduction instructions

If you encounter any issues while reproducing the problem, or if any part of the setup or observed behavior appears inconsistent on your side, please let me know. I would be happy to provide any additional information or clarification that may be helpful.