GPU Collision Handling | Elias Füricht

For my master's thesis at TU Wien I worked on a GPU-based collision detection and resolution pipeline for potentially non-convex elastic polyhedral meshes. The goal was to enable the handling of collisions inside a Virtual Element Method simulation framework, currently developed at the Research Unit of Computer Graphics at TU Wien with a major focus on performant GPU execution.

The biggest challenge came with the elements that we were working with: deformable polyhedrals. They can become non-convex during simulation, which limits the number of algorithms usable for CD by a lot. On top of that we are working with irregular topology (arbitrary number of faces and vertices, potentially differing per element) as well as deformation which are both limiting constraints when it comes to a fully GPU-based solution.

Method

The design I have decided on in the end consists of a BVH broad phase using cuBQL, followed by an image-based narrow phase adapting and extending Faure's rasterization method, which enables collision detection for non-convex meshes based on traversing depth-sorted fragment lists and recording overlaps. The collision resolution is handled using a position-based approach derived from position-based dynamics followed by the computation of collision and friction forces.

Implementation - Focus Narrow Phase

The majority of development went into designing a performant custom CUDA compute rasterizer for the narrow phase. For that I parallelized the work per collision as well as per pixel of the rasterization - so one block owns a colliding pair, and each thread of that block is responsible for a subset of the pixels of the rasterization process.

NarrowPhaseKernel - Thread-local pixel ownershipCUDA

constexpr unsigned int totalPixelsPerBlock = NP_COLL_TEX_RES * NP_COLL_TEX_RES;
constexpr unsigned int pixelsPerThread = (totalPixelsPerBlock + NP_THR_PER_BLOCK - 1u) / NP_THR_PER_BLOCK;
NPFragment stackFrags[pixelsPerThread][NP_COLL_TEX_FRAG_LIST_MAX];

// Each thread rasterizes its own subset of pixels.
for (unsigned int p = 0; p < pixelsPerThread; p++)
{
    const unsigned int pixelFlattenedIndex = threadIdx.x + p * blockDim.x;
    if (pixelFlattenedIndex >= totalPixelsPerBlock) break;
    // ... rasterize fragments into thread-local stackFrags[][] ...
}

This design enabled, for low enough collision resolutions (up to 32x32), to keep fragment lists in thread-local low-latency memory. Through that design achieved throughputs for the narrow phase upwards of 50% with almost no DRAM accesses, seen in the profiling screenshots.

Compute, memory and cache throughputs for the narrow phase kernel.

Narrow-phase occupancy at 32 by 32 raster resolution: 41.67 percent theoretical and 41.30 percent achieved occupancy. — Almost achieving occupancy and active warp limits (limited due to high register usage)

Another design necessary to facilitate the irregular topology of the elements as well as increase performance was to batch and preproject the triangles and store them in shared memory for all threads of the block to access.

NarrowPhaseKernel - Triangle batchingCUDA

constexpr unsigned int triangleBatchSize = 256;
__shared__ SharedMemTri shTriBatch[triangleBatchSize];

for (unsigned int batchStart = 0; batchStart < triTotal; batchStart += triangleBatchSize)
{
    const unsigned int batchEnd = min(batchStart + (unsigned int)triangleBatchSize, triTotal);
    const unsigned int batchSize = batchEnd - batchStart;

    // Threads cooperatively prepare the batch.
    for (unsigned int j = threadIdx.x; j < batchSize; j += blockDim.x)
    {
        const unsigned int i = batchStart + j;
        // ... construct and project triangle i ...
        SharedMemTri& st = shTriBatch[j];
        // ... cache pixel bounds, edge-function increments, and depths in st ...
    }
    __syncthreads();

    // Reuse the preprojected triangles for each thread's pixels.
    for (unsigned int t = 0; t < batchSize; t++)
    {
        const SharedMemTri& tri = shTriBatch[t];
        // ... rasterize against this thread's owned pixels ...
    }
}

There were lots of other interesting challenges and design decisions during my thesis, but explaining all of them here would be too much - for in-depth explanations please refer to my written thesis linked at the top of this page.

Method

Implementation - Focus Narrow Phase

More videos and images