Optimized Parallel GPU Merge Kernel

• Parallel GPU merge kernel is a specialized algorithm that efficiently merges k sorted sequences on GPUs by partitioning work into tiles and using warp-level operations.
• It exploits GPU features such as coalesced memory accesses, binary search partitioning, and register-based minBlockHeap merging for optimal throughput.
• The design overcomes traditional performance bottlenecks by eliminating shared memory bank conflicts and ensuring balanced load distribution across warps.

A parallel GPU merge kernel is a specialized algorithmic component designed to perform k-way merging of sorted sequences on Graphics Processing Units (GPUs), exploiting their massive parallelism and unique compute and memory hierarchy characteristics. At its core, this kernel facilitates the merging of $K$ input runs into a single sorted sequence, mapping the problem efficiently onto the GPU architecture to maximize memory bandwidth and parallel throughput while minimizing synchronization and avoiding shared memory bank conflicts. The GPU Multiway Mergesort (MMS) algorithm exemplifies the state of the art in this area, offering an asymptotically optimal implementation in terms of global memory accesses and achieving practical performance advantages over previous comparison-based GPU sorting algorithms (Casanova et al., 2017).

1. Problem Formulation: Parallel k-way Merge on GPU

Given $K$ sorted runs $A^1, \dots, A^K$ of total length $N$, the objective is to produce a single output run by merging all input runs. In the GPU setting, $P$ warps are launched, each responsible for a contiguous segment (termed a "tile") of the output, of size roughly $N/P$. Each warp operates in lockstep, issuing coalesced global-memory accesses of width $B$ (corresponding to the warp width). This design partitions the k-way merge into $P$ independent sub-merges, each of size $N/P$, promoting perfect load balance and maximizing the efficiency of memory coalescing.

2. Partitioning: Discovering Merge Tile Boundaries

Partitioning is accomplished by mapping each output tile to precise indices in each input run. For warp $t$ ($K$0) and per-run starting indices $K$1, the requirement is:

$K$2

where $K$3 denotes the global rank. This boundary discovery utilizes a generalization of the merge-path technique with binary search: for each run $K$4, a binary search on $K$5 finds $K$6, relying on K–1 binary searches or SIMD comparisons in other runs to determine global rank. Each warp’s total partitioning cost is $K$7 global memory transfers, which is asymptotically negligible provided $K$8.

3. The minBlockHeap Merge Kernel Architecture

Upon determining per-run boundaries, each warp merges its assigned tile using the minBlockHeap, a binary heap structure where each node contains $K$9 sorted elements and has $A^1, \dots, A^K$0 leaves (each representing a block from an input run). The construction proceeds as follows:

• Build Phase: Every leaf node reads a block of $A^1, \dots, A^K$1 elements from its run using coalesced accesses. The heap is built bottom-up: at every internal node $A^1, \dots, A^K$2 (with children $A^1, \dots, A^K$3, $A^1, \dots, A^K$4), the fillEmptyNode procedure merges the $A^1, \dots, A^K$5 child elements, propagating the $A^1, \dots, A^K$6 smallest into $A^1, \dots, A^K$7 and recursing to refill the emptied child until a leaf is reached.
• Steady-State Output: The root’s $A^1, \dots, A^K$8 elements—always the smallest available—are written back to global memory as a coalesced store. The root is then marked empty, and fillEmptyNode is invoked recursively to refill from the appropriate input, continuing until all $A^1, \dots, A^K$9 elements are emitted.

Critical to performance, all merge steps within fillEmptyNode are realized using a parallel bitonic network across registers, leveraging warp-shuffle instructions to avoid any dependence on shared memory. This design totally eliminates shared memory bank conflicts and synchronization primitives such as __syncthreads() within warps.

4. I/O and Work Complexity Analysis

For each merge round, every warp reads and writes $N$0 elements, and the total global I/O cost per merge round is therefore $N$1. The algorithm conducts $N$2 rounds, with $N$3 representing device memory size. The total global I/O complexity is thus:

$N$4

Choosing $N$5 yields:

$N$6

This matches the Parallel External Memory (PEM) model’s lower bound for sorting (Casanova et al., 2017). Per output block, the only additional computation is the $N$7 bitonic warp-shuffle steps in fillEmptyNode, giving total parallel work of

$N$8

with $N$9 (warp width).

5. Hardware Constraints and GPU-Specific Optimizations

The choice of $P$0 is governed by shared memory availability. For NVIDIA Kepler (48 KiB per SM), $P$1 is optimal; for Maxwell (96 KiB per SM), $P$2. The active warps $P$3 are selected to saturate all SMs, typically 512–1024 warps. By employing warp-shuffle instructions for all intra-tile communication, all shared-memory bank conflicts are averted. The base-case merge employs a warp-level shearsort for $P$4 elements, using a bank-conflict-free transpose. The base-case size is tuned so that $P$5 avoids problematic powers of $P$6, circumventing wasteful final merges.

6. Empirical Performance and Comparative Evaluation

MMS is benchmarked against Thrust and modernGPU comparison-based mergesorts, as well as CUB’s radix sort, on three platforms: GTX 770 Kepler, K40m Kepler, and M4000 Maxwell. Throughputs of 400–600 M keys/s (Kepler) and ~700 M keys/s (Maxwell) are attained. For small $P$7, performance matches that of MGPU/Thrust; for large $P$8, MMS achieves 5–15% higher performance due to lower I/O costs. While MMS, being comparison-based, is within 30–40% of CUB’s radix sort (1.2 G keys/s), it excels on workloads that induce worst-case bank conflicts in competitors: in such scenarios, MMS delivers a 30–45% speedup and is unaffected by the performance collapses that degrade MGPU/Thrust by 2–3x in shared-memory merges.

7. Theoretical and Practical Significance

The MMS parallel GPU merge kernel establishes an asymptotically optimal solution both in global memory accesses and in practical throughput for comparison-based sorting on GPUs. By eliminating shared memory bank conflicts and achieving near-perfect coalescing via partitioned tiles and register-only merges, it overcomes performance bottlenecks in existing algorithms. This positions the MMS merge kernel as a foundational primitive in high-throughput, memory-efficient sorting frameworks on GPU architectures (Casanova et al., 2017).