arXiv CSopen access

Eiger: An Efficient Library for GPU-based Data Analytics

arXiv:2607.04489v1 [cs.DB] 5 Jul 2026

Eiger: An Efficient Library for GPU-based Data Analytics Bowen Wu

Marko Kabić

Sven Hepkema

Systems Group, ETH Zurich Switzerland [email protected]

Systems Group, ETH Zurich Switzerland [email protected]

Systems Group, ETH Zurich Switzerland [email protected]

Vasilis Mageirakos

Christos Kozyrakis

Gustavo Alonso

Systems Group, ETH Zurich Switzerland [email protected]

NVIDIA & Stanford University United States [email protected]

Systems Group, ETH Zurich Switzerland [email protected]

ABSTRACT GPUs have become an increasingly attractive platform for accelerating analytical workloads due to their massive parallelism and high memory bandwidth. Recent studies show that in systems with fast CPU-GPU interconnects and fast networks, query processing within the GPU, rather than data movement, is the dominant bottleneck. This highlights the need for a more efficient implementation of relational operators on GPUs than the widely used library, cuDF. While offering rich functionality, cuDF commits to a single, statically chosen implementation for most operators and makes little use of runtime information about the data, limiting performance across diverse workloads and different GPUs. In this paper, we present Eiger, a high-performance library for GPU-based data analytics that improves single-GPU query processing through runtime workload adaptivity. Adaptivity in Eiger rests on two design principles. First, Eiger provides multiple implementation variants and tunable knobs for most operators, covering not only joins and group-bys but also expensive yet often overlooked operations, such as expression evaluation, string processing, and multi-key sorting, for which it contributes an array of new optimization techniques. Second, Eiger profiles intermediate data during query execution using lightweight statistics, such as value ranges and HyperLogLog++ sketches, and uses them to select implementations, tune configuration knobs, and compress data on the fly, thereby overcoming the limitations of traditional static query optimization. The breadth of operators and variants also enables a more comprehensive performance analysis, encompassing a wider range of operations and workloads than previous work. We evaluate Eiger with operator microbenchmarks on two GPU architectures and with the complete TPC-H benchmark (up to scale factor 100). Across the 22 queries, Eiger reduces the total runtime by up to 1.8× compared to the state-of-the-art cuDF library; for individual queries, Eiger achieves up to 6.1× better performance.

1

INTRODUCTION

Recent years have seen rapid growth in the use of graphics processing units (GPUs) to accelerate database applications [12, 54, 57]. GPUs, with their massive parallelism and high-bandwidth memory, are well-suited for many, if not all, operations in data analytics. Recently, driven by large language model training and inference, GPUs have been deployed at large scales and are equipped with increasingly higher memory bandwidth, larger memory capacity, faster CPU-GPU interconnects, and high-bandwidth RDMA networks.

All of these enable GPU-based databases to process data at one to two orders of magnitude faster than CPU-based databases [54, 57]. Recent performance studies by Kabić et al. [19] and Wu et al. [54] evaluate GPU-based database systems on two different deployments: (1) relations initially stored in CPU DRAM requiring data transfer to the GPU, and (2) relations pre-partitioned and loaded into high-bandwidth memory (HBM) of multiple GPUs in a cluster. Across both settings using the TPC-H benchmark, they find that query processing time on the GPU dominates overall execution time when fast interconnects (e.g., NVLink-C2C) and RDMA networks (e.g., InfiniBand) are available. This challenges the common assumption that data movement is the main bottleneck and suggests that the cost of query execution on GPUs will become even more critical as faster connectivity becomes mainstream [4, 8]. These observations motivate research focused on improving single-GPU query processing performance. Existing systems are typically built on cuDF [40], an open-source GPU DataFrame library, or on tensor operators from machine learning libraries [12, 60]. These libraries commit to a single, statically chosen implementation for most operators. cuDF, for instance, executes both joins and group-bys with non-partitioned hash-based implementations by default. This one-size-fits-all design is at odds with decades of CPU database experience, which shows that no single implementation performs best for all workloads [39]: hashbased joins and group-bys perform well when the hash table is cache-resident, but degrade otherwise, where sort- or partitionbased methods are faster. We identify the lack of runtime workload adaptivity as the key factor that prevents existing libraries from achieving higher performance. This lack of adaptivity appears in two ways. First, existing work provides only a limited set of implementation variants for each database operator, restricting the choices available for different workloads. Second, existing libraries use little runtime information about intermediate data, such as value ranges, distinct counts, and string lengths, even though exactly this information determines which variant, configuration, or data representation is the most efficient. Knowing the number of distinct values (NDV) in a grouping column, for example, would let the engine pick the best group-by algorithm and size its hash table accordingly. The two aspects are complementary and both necessary for adaptivity: the first provides the set of available choices, while the second enables informed runtime decisions. In this paper, we present Eiger, a GPU-based library designed to provide high-performance single-GPU query processing through

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso

runtime workload adaptivity. To this end, Eiger is designed around two principles. (1) Multiple operator implementations. Eiger provides multiple implementation variants for each operator, often with tunable knobs. For joins, Eiger implements non-partitioned hash, partitioned hash, and sort-merge algorithms, combinable with two payload materialization strategies. For group-bys, it provides hash-, sort-, and partition-based algorithms. For expression evaluation, it offers a per-tuple interpretation backend that keeps intermediate results in registers or shared memory, and a batch-based backend composed of type-specialized and vectorized kernels. For string matching, both the number of threads per string and the width of packed memory accesses are tunable, and substring matching can switch between a quadratic-time algorithm and the linear-time Knuth-Morris-Pratt algorithm. For sorting, Eiger chooses between radix sort and merge sort, the latter optionally accelerated by prefix extraction for strings. Table 1 summarizes all variants together with the workload properties that determine the best choice. As this catalog shows, Eiger practices the principle not only in traditionally expensive operators, such as joins and group-bys, but also in expensive operations that are often overlooked by the literature on GPU-accelerated databases. For the latter, Eiger contributes new techniques: packed (multi-byte) string accesses that handle arbitrary alignment, expression linearization based on the Sethi–Ullman ordering to minimize intermediate results, and an order-preserving dictionary encoding algorithm for GPUs. (2) Runtime adaptive execution. Instead of relying solely on static query optimization, Eiger defers the choice of implementations and configuration parameters to query runtime whenever possible. Unlike traditional CPU databases, where statistics are collected only for base tables beforehand, Eiger profiles intermediate data during query execution using a curated set of lightweight statistics: minimum, maximum, and mean values, and a HyperLogLog++ sketch estimating the NDV. These statistics drive three kinds of runtime decisions: algorithm selection (e.g., the estimated group cardinality selects the group-by algorithm and sizes its hash table), knob tuning (e.g., string length statistics set the parallelism and packed access width for string matching), and on-the-fly data compression (e.g., a mechanism we call smart key fusion compresses multiple sort or group-by keys via frame-of-reference+bitpacking or order-preserving dictionary encoding and fuses them into a single 4- or 8-byte key, making the much faster radix sort applicable in place of merge sort). Our key insight is that GPUs compute such statistics at close to memory bandwidth, so profiling is cheap enough to easily pay off: computing an HLL++ sketch over 228 keys takes about one millisecond, while choosing the wrong sort implementation in the same setting costs hundreds of milliseconds (Section 5.6). We evaluate Eiger with operator-level microbenchmarks and end-to-end queries on two different platforms (NVIDIA GH200 and A100); the results show that the best implementation depends not only on the workload but also on the GPU architecture. At the operator level, Eiger evaluates expressions up to 5× faster than cuDF, and smart key fusion accelerates multi-key sorting by up to 13× over merge sort. On the complete set of 22 TPC-H queries with scale factors 10, 30, and 100, Eiger outperforms a cuDF-based

execution engine by up to 1.8× in total runtime and up to 6.1× on individual queries. We summarize our key contributions as follows. • We argue that runtime workload adaptivity, i.e., the combination of multiple implementation variants per operator and runtime statistics to choose among them, is the key ingredient to push single-GPU data analytics performance further, and we present Eiger, a high-performance library for GPU-based data analytics that demonstrates this idea. • We propose optimization techniques for expensive but often neglected operations: packed accesses and cooperative-group-based parallelism for string matching, Sethi–Ullman linearization and register-resident intermediate results for expression evaluation, and smart key fusion with order-preserving dictionary encoding for multi-key sorting and group-bys. • We introduce a runtime profiling and adaptive execution mechanism for GPU-based query execution, and show in three concrete scenarios (algorithm selection, knob tuning, and on-the-fly compression) how cheap GPU-side statistics translate into large performance gains. • We conduct a comprehensive performance study with operatorlevel microbenchmarks on two GPU architectures and all 22 TPC-H queries at three scale factors, characterizing when each implementation variant wins and thereby providing a basis for cost models in future GPU query optimizers.

2 BACKGROUND 2.1 GPU Architecture and Programming The primary execution unit of a GPU is the Streaming Multiprocessor (SM), which comprises control units, a large register file, CUDA cores, and on-chip fast memory, including the co-locating L1 cache/shared memory, and specialized caches such as the constant cache. Multiple SMs are connected via a unified L2 cache to off-chip global memory, which has higher latency and lower throughput. GPU programs are executed by a large number of lightweight threads, organized hierarchically into thread blocks and grids. Threads within a block can cooperate using shared memory and synchronization primitives, while blocks are scheduled independently across SMs. At the hardware level, threads are executed in fixed-size groups called warps, following a single-instruction, multiple-thread (SIMT) execution model. This organization enables efficient latency hiding by rapidly switching between ready warps. From a programming perspective, GPU computation is expressed in terms of kernels, which define functions executed in parallel. A more recent programming abstraction, called cooperative group, allows the programmer to program collaboration among threads more easily, no longer bound to the warp level or block level.

2.2

GPU-based Query Processing

Query processing on the GPU either assumes the data (input, output, and any intermediate results) are completely resident in the GPU memory, or the data could move between the CPU memory, GPU memory, and sometimes even through a network. In the latter case, a buffer manager and/or a network manager is usually needed to administer the data movement. In this study, we focus on the

Eiger: An Efficient Library for GPU-based Data Analytics

former case because single-GPU in-memory processing serves as a foundation for all different scenarios. Unlike in CPU-based systems, which often implement a pipelined volcano or vector model, the most popular execution model for GPU-based query processing is full materialization. In other words, each operator waits for the last operator to completely finish before starting. This is also the model assumed in this work and that of the most popular GPU-based operator library, cuDF. The relations are usually stored in columns and almost never converted to a row format during execution. When an input relation has multiple columns, the operator often initializes a column of tuple IDs (TIDs) as the delegate of the columns that are not directly involved in the operation. Later, those columns will be materialized with the TIDs, which is also known as late materialization.

3

EIGER OPERATORS AND OPTIMIZATIONS

In this section, we explain operator implementations in Eiger. Table 1 summarizes them and lists the factors influencing when to use each. The factors are grouped into two categories: static ones, which can be known before the query execution, and runtime ones, which are only available during the query execution.

3.1

Data Layout and Data Types

Eiger adopts columnar storage, where each column of a relation is stored in one contiguous memory region, with the exception of strings (see Section 3.4). Programming-wise, Eiger provides the typeless column_view (immutable) and mutable_column_view (mutable) classes that can be passed directly as arguments to the device kernel without extra memory copy. A (mutable_)table_view is a collection of (mutable_)column_view, but cannot be passed to the device kernel. For GPU kernels that work on relations with arbitrary numbers of columns, we provide the table_device_view, which can be converted from the table_view through very cheap memory copies. The table_device_view ensures that all meta information about the table is stored in the device global memory instead of in the thread’s local memory. The latter may lead to substantial register pressure and consequently harm performance. Compared to cuDF, which uses the 32-bit integer as the “size type” to store the number of rows, Eiger adopts the 64-bit integer as the size type to be future-proof. Our rationale is that, with the memory capacity of GPUs steadily increasing, the table size that can be handled can easily outgrow (or may have already outgrown) the 32-bit integer range. One example is that cuDF has already encountered the issue that the 32-bit integer is not enough to index all the bytes of a string column [11]. However, moving from 32-bit to 64-bit size type is not free because handling 64-bit tuple IDs (TIDs) can be more expensive in some scenarios. For example, reading a 64-bit TID column results in 2 times more data being loaded from memory compared to a 32-bit TID column. This performance penalty is amplified when we need to load or store multiple passes of TIDs in an algorithm (e.g., sorting). Another example is that some implementations utilize shared memory for caching intermediate results, and using 64-bit TIDs may result in less data being cached, affecting the efficiency of the algorithm.

3.2

Joins

Eiger implements multiple join algorithms, including non-partitioned hash join, partitioned hash join, and sort-merge join. For sortmerge join and partitioned hash join, Eiger adopts the implementations detailed in the work of Wu et al. [55]. For the hash join, Eiger chooses the same implementation as cuDF, which uses the static_multiset provided by cuCollection [33] as the hash table implementation. The static_multiset only stores the tuple IDs and uses the user-provided hash function and equality comparator to hash and compare the actual data. This has two advantages over using the static_multimap in cuCollection, which stores keyvalue pairs. First, using static_multiset allows the hash join to work with arbitrary data types, such as strings, and avoids being restricted to only 4 or 8-byte data types, currently a limitation of static_multimap. Second, static_multimap suffers from bad performance when the key-value pair to insert cannot be packed into a compare-and-swap-supported data type. For example, an 8-byte key and 4-byte value cannot be packed into a 128-bit CAS-supported data type without extra padding. Using a static_multiset avoids this issue by only inserting the TIDs, which is a single type. Eiger includes two strategies proposed in previous work [55] to materialize the payload columns of a join, that is, gather-from -untransformed-relations (GFUR) and gather-from-transformedrelations (GFTR). While GFTR only applies to the partitioned hash join and sort-merge join, GFUR is applicable to all three algorithms. The main difference between GFTR and GFUR is whether we let the payload columns go through the same transformation (sort/partition) as the join keys. For example, in the sort-merge join, GFTR sorts all the payload columns according to the key column so that the materialization can avoid expensive random accesses later.

3.3

Group-by

For group-by, Eiger provides three algorithms (sort-based, hashbased, and partition-hash-based) and the two ways of handling payload columns (GFUR and GFTR) [55]. Compared to previous work [55], Eiger introduces a few important improvements. First, Eiger uses the static_set to implement the partitioned-hashbased group-by instead of the static_map used in previous work, following a similar reasoning as for the join implementation. Second, Eiger extends the hash-based group-by to support multiple numeric and non-numeric grouping keys, often seen in TPC-H queries. Third, multiple grouping keys may be combined to improve efficiency (see Section 4.3). Fourth, the partition-based group-by applies a hash function to the keys before partitioning to ensure more even partition sizes.

3.4

Strings

String processing often shows up in queries, but has not been discussed in-depth in the context of GPU-accelerated databases. Existing work, which claims to support full TPC-H queries, often gets around string processing by converting strings into numerics using dictionary encoding [31, 54]. We argue that although dictionary encoding works well for low-cardinality string columns, processing uncompressed strings is still important because not all operations (e.g., substring matching) can work with dictionary-encoded numeric values.

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso

Table 1: Summary of implementations in Eiger and factors to be considered for algorithm selection. Operation

Implementations/Techniques Sort-merge join Partitioned hash join Non-partitioned hash join Gather from untransformed relations (GFUR) Gather from transformed relations (GFTR) Sort-based (optional dictionary encoding opt) Partitioned hash-based Non-partitioned hash-based Packed access width Varied parallelism KMP algorithm for substring matching Per-tuple interpretation (reg/shared memory) Batch-based processing Radix sort Smart key fusion (§4.3) Merge sort (w/ prefix extraction for strings (§3.4.5))

Join (§3.2) Materialize (§3.2, §3.3) Group-by (§3.3)

String matching (§3.4) Expression eval (§3.5)

Sort (§3.6)

† Readily accessible at the start of the operation;

h t t p

t t p s

t p s :

p s : /

s : / /

: / / e

/ / e t / e t e t t

‡

Static workload properties Data types Sorted-ness of join keys Data types

Runtime workload properties Cardinality of the relations† Distribution of the join keys‡ Output cardinality (match ratio)‡ Cardinality of gather map† Distribution of gather map values‡

Data types Aggregation functions

Cardinality of the relation† Group cardinality (§4.1)‡

Length of the pattern string

Distribution of the string lengths‡ (§4.2)

Register usage Intermediate results size

Memory usage†

Data types of sort keys

Distribution of string prefixes (§4.1)‡ Key compressability (§4.3)‡

Needs additional statistics about the data, such as min, max, ndv.

pattern table

h t t p s : / / e t h z . c h target string 4-byte boundary

Figure 1: String matching with packed accesses. 3.4.1 Layout. In-memory string columns typically adopt one of two layouts, the old Arrow format [53] and the German string format [32]. Like cuDF, Eiger chooses the Arrow format mainly for two reasons. First, the German string format needs more space to store the string column due to the explicit length field if the majority of strings cannot be inlined into the header. Second, depending on the length of the string, reading a German string may need three steps, which means the program often contains more branches. Our preliminary experimental study shows that the CUDA program for processing the German string format is often compiled to use more instructions than the Arrow format. Our study also shows that the German string format is only beneficial for prefix matching when strings are long and the pattern is fewer than 4 bytes. On the other hand, the Arrow format only requires an offset array in addition to the string data, which results in only one additional lookup of the offset array. Because all the strings in a column are stored contiguously, it is also easy to exploit more optimization techniques (e.g., packed access), as we will show next. 3.4.2 Packed access. Accessing and processing one character (1 byte) at a time leads to high instruction counts and less efficient usage of memory bandwidth. To avoid this, whenever possible, Eiger uses a technique called packed access, which leverages wider load instructions (e.g., 4-byte) to load multiple characters at the

same time and perform operations altogether (e.g., comparison). This reduces the number of instructions in the binary code and improves the efficiency of memory access. The challenge of enabling wider access is dealing with data alignment. Since strings are concatenated one-by-one in memory, the starting address of a string may not be aligned to be loaded as a wider data type. Even if we compare the unaligned part byte-bybyte, when the string finally becomes aligned, the pattern string may become unaligned at that byte. To solve this, if we want to use 𝑘byte accesses, then we preprocess the pattern string by duplicating it 𝑘 times so that each character in the pattern string is aligned at 𝑘 bytes in one of the duplicates. All duplicates are stored in the shared memory of the thread block and can be created on the fly when reading the data from memory. During string comparison, when the beginning unaligned part of the target string is consumed, it then compares the rest with a duplicate of the pattern whose next byte to be compared is aligned. In this way, the rest of the comparison can be done with packed accesses except for a few trailing bytes. Figure 1 illustrates how Eiger handles the alignment issue. The first two characters “ht” are compared byte-by-byte. Then, starting from the third character “t”, we can compare 4 bytes at a time with the third duplicate of the pattern, whose third character is also aligned. 3.4.3 Cooperative groups. As mentioned in Section 2.1, the cooperative group abstraction allows the programmer to flexibly assign a set of threads to a task. Eiger uses this to accelerate string matching operations. Some pattern-matching operations, such as prefix and suffix matching, can benefit from more parallelism per string when the pattern string or target string is long. Moreover, when the target strings are also long, using more threads improves the cache locality. Therefore, we utilize cooperative groups to vary the level of parallelism used to process each string. In contrast to the warp programming paradigm, where the whole warp (32 threads) is allocated for a task, using cooperative groups gives finer control over the parallelism (2, 4, 8 threads, and so on), depending on the

Eiger: An Efficient Library for GPU-based Data Analytics

size of the task. This technique can be combined with packed access to further improve performance. 3.4.4 Substring matching. Substring matching is commonly seen in queries. In addition to the vanilla quadratic-time string matching algorithm used in cuDF, Eiger also includes the Knuth-Morris-Pratt (KMP) algorithm [23] that works in linear time. The KMP algorithm has been shown to work well for GPUs [5]. It requires a preprocessing step on the pattern string to construct the prefix function. We let each thread block perform the preprocessing independently and store the prefix function in shared memory. For the vanilla string matching algorithm, we can apply the packed access optimization; for the KMP implementation, we always use the byte-by-byte access. 3.4.5 String sorting. Since strings have variable lengths, lineartime radix sorting does not work. Instead, Eiger uses merge sorting with an important optimization for long strings. Before starting the merge sort, as an optional additional step, we extract the 4-byte prefix from all strings and store them contiguously in a prefix array. During the merge sort, the custom comparison functor will first compare the prefixes from the prefix array, and only when they have a tie, the functor will continue to check the actual strings. This method greatly accelerates the comparison by providing more efficient memory access. Note that extracting the prefixes will not always benefit the performance, especially when the prefixes are not discriminating enough. In Section 4, we discuss how we decide if it is worth extracting prefixes during the query execution time.

3.5

Expression evaluation

Expressions are an integral part of many database operators, most importantly, selection and projection. Recent work has pointed out that for a substantial number of queries on the TPC-H benchmark, the filtering operation in selections and projections is the most expensive part of the query when running on a GPU [19]. Eiger has two expression evaluation backends: a per-tuple interpretation (PTI) backend and a batch-based (BB) backend. Both backends take as input a source table and an expression tree, which dictates how the expression should be evaluated to an output column. Before evaluation, the expression tree is traversed and linearized into a sequence of steps, each consisting of an operation, a column reference, or a constant. Unlike cuDF, which uses the post-order traversal to linearize the expression tree, Eiger adopts the Sethi-Ullman ordering [44], which minimizes the peak number of intermediate results needed to evaluate the tree. Minimizing intermediate results benefits both backends, as we will explain soon. The PTI backend launches a single kernel that works as an expression tree interpreter to evaluate the expression for each tuple in the source table. Therefore, before launching the kernel, the linearized expression tree needs to be transferred to the device memory. The advantage of this backend is that it can prevent intermediate results from being written back to the device memory, which causes extra memory transactions. Despite the advantage, the challenge lies in storing and managing the intermediate results. A careless implementation may result in the intermediate results being stored in the thread-local memory, which is backed by the device memory under high register pressure. cuDF uses the shared

memory for this purpose, where each intermediate result is written into a designated slot. The slot assignment can already be done during the expression tree linearization on the CPU. Eiger adopts a similar design, but additionally implements a specialization where the intermediate results are stored in the registers. When linearizing the expression tree, if the peak size of intermediate results is smaller than a threshold, we can store them in registers, which have an order of magnitude lower access latency than shared memory. Thanks to the Sethi-Ullman ordering, we use fewer registers (equivalently, intermediate results) to evaluate an expression compared to the post-order traversal. The disadvantage of the PTI backend is the high interpretation overhead. Due to the generality of the interpretation kernel, the decision of which operator to dispatch is made at runtime and for each tuple, leading to high instruction counts and pressure on the ALU pipeline in CUDA cores. Instead of relying on one generic interpreter kernel, the BB backend evaluates each step in the linearized expression tree using a separate kernel (except for column reference and constants). The benefit of this is that there is almost no interpretation overhead. Each kernel is specialized for the type of operation and its operands through C++ templates. All operations can be categorized into unary operations, binary operations, and ternary condition operations (e.g., pred ? a : b). Since the kernel always deals with one to three input columns and produces one output column, we can leverage the vectorized load and store to further improve memory bandwidth utilization and reduce overheads. The disadvantage of this backend is that it consumes extra device memory to store the intermediate results, and multiple kernels lead to more reads and writes into the device memory. Choosing between PTI and BB depends heavily on the properties of the expression (e.g., intermediate result size), input relations (e.g., tuple size), and the GPU architecture (e.g., memory bandwidth vs. compute throughput).

3.6

Sorting

Sorting is another often-neglected operator when studying the GPU-based databases, especially for the case of multi-key multidata-type sorting. Many existing works also do not explain in detail how they handle sorting when it comes to multiple keys [31, 45, 54]. Eiger implements different ways to sort a table depending on the column data types. Eiger prefers to use the radix sort whenever possible, due to its substantially better performance than the merge sort. Eiger uses merge sort when lexicographical comparisons are not possible, or the sorting is based on multiple keys.1 We categorize the sort implementations into the following cases. Case 1 A single numerical sort key. We use the radix sort for the reason explained above. Case 2 Multiple numerical sort keys. We will start with trying to combine the keys into one single key using the smart key fusion (SKF) algorithm we propose in Section 4.3. If combining them is possible, then we can process the rest like Case 1; otherwise, we will handle the sorting as Case 3 (see below). Case 3 At least one non-numerical sort key. We use the merge sort with a custom comparator. The custom comparator maintains a table_device_view (see Section 3.1) and sort orders (ascending or 1 Currently, the cub::DeviceRadixSort function in CUB only works for a single key.

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso

descending) for each sort key. Given two tuple IDs, the comparator retrieves the rows from the table_device_view and compares them based on the specified sort orders.

4

RUNTIME ADAPTIVE EXECUTION

In the previous section, we explained the implementations of Eiger operators in detail. A recurring pattern is that each implementation has its own favorable input characteristics for which it works most efficiently. If the query runtime can better understand input data, it can select a more efficient implementation or choose better configuration parameters. CPU-based databases often rely on statically maintained table or column statistics for this purpose [34, 38] and lack runtime profiling of the intermediate relations. For GPU-based databases, we argue that runtime profiling is beneficial in certain scenarios based on the following observations. First, many statistics, such as min, max, and mean, can be calculated extremely fast on the GPU (close to the memory bandwidth, which is currently a few TB/s). Although these statistics are basic, they offer valuable information about the data being handled. Second, choosing suboptimal implementations or configuration parameters is very costly, as evidenced in the experimental evaluation (Section 5). Combining these two observations, as we will show in the experiments, the overhead of statistics computation for intermediate results is easily offset by the huge gain in choosing the optimal algorithm. Based on this, we propose that GPU-based query execution should embrace runtime profiling. In the rest of this section, we illustrate the idea with a few scenarios and show how calculating a curated set of statistics can boost the end-to-end operator performance.

4.1

Scenario I: Number of Distinct Values (NDV)

The performance of group-by is closely related to the group cardinality. We propose using the lightweight and GPU-friendly HyperLogLog++ [14] (HLL++) algorithm to estimate this group cardinality before executing the group-by operator for large input sizes. The construction of the HLL++ sketch consumes almost negligible time compared to the actual group-by operator, but is helpful for identifying the optimal algorithm. In the case of hash-based group-by, if we estimate the group cardinality to be small, we can allocate a much smaller hash table to save memory space. Similar ideas can be used to check whether the prefixes of a column of strings are diverse enough. This helps us decide if we should extract the prefixes to accelerate the sorting (see Section 3.4.5). For strings that have common prefixes, such as HTTP websites, we avoid constructing the prefix array.

4.2

Scenario II: String Matching

The string matching algorithm in Eiger has two degrees of freedom, namely parallelism per string (except for substring matching) and packed access width. To determine which configuration works better, we can calculate the min, max, and mean of the string lengths, which are stored as a separate array from the actual strings. With them, we can apply the following heuristic to pick the parameters: If both the average string length and the pattern length are long, we can increase both the parallelism and the packed access width. If only the average string length is long but the pattern is short, we should not use more parallelism, but can increase the packed

Algorithm 1: Order-preserving Dictionary Encoding Input: Column 𝐶𝑖𝑛 to be encoded Output: Encoded column 𝐶𝑜𝑢𝑡 , number of distinct values 𝑛𝑑𝑣 Function Main(𝐶𝑖𝑛 ): ⊲ (1) Assign group IDs and elect the group leader 2 𝐻𝑇 ← HashTable( |𝐶𝑖𝑛 | ); 3 𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐 ← alloc_on_device(𝑠𝑖𝑧𝑒_𝑡, |𝐶𝑖𝑛 | ); 4 𝑛𝑑𝑣 ← 0; 5 compute_group_id(𝐶𝑖𝑛 , 𝐻𝑇 , 𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐, 𝑛𝑑𝑣); ⊲ (2) Retrieve key-ID pairs 6 [𝑢𝑛𝑖𝑞𝑢𝑒𝐾𝑒𝑦𝑠, 𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝑠]← 𝐻𝑇 .retrieveAll(); ⊲ (3) Argsort the unique keys to get order-preserving dictionary 7 𝑎𝑟𝑔𝑠𝑜𝑟𝑡𝐼𝑑𝑠 ← argsort(𝑢𝑛𝑖𝑞𝑢𝑒𝐾𝑒𝑦𝑠 ); ⊲ (4) Scatter the final encoded values 8 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 ← alloc_on_device(𝑠𝑖𝑧𝑒_𝑡, |𝐶𝑖𝑛 | ); 9 scatter(𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝑠, 𝑎𝑟𝑔𝑠𝑜𝑟𝑡𝐼𝑑𝑠, 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 ); ⊲ (5) Encode all the keys 10 gather(𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐, 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠, 𝐶𝑜𝑢𝑡 ); 11 return 𝐶𝑜𝑢𝑡 , 𝑛𝑑𝑣 1

Function compute_group_id(𝑘𝑒𝑦, 𝐻𝑇 , 𝑔𝐼𝑑𝐿𝑜𝑐, 𝑛𝑑𝑣): for each thread 𝑖 do ⊲ Insert or return the value if key has existed 14 𝑠𝑡𝑎𝑡𝑢𝑠, 𝑔𝐼𝑑𝐿𝑜𝑐 [𝑖 ] ← 𝐻𝑇 .insert_and_find(𝑘𝑒𝑦 [𝑖 ], 𝑖 ); 15 if 𝑠𝑡𝑎𝑡𝑢𝑠 is Success then 16 atomicAdd(&𝑛𝑑𝑣, 1);

12

13

Function scatter(𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝑠, 𝑎𝑟𝑔𝑠𝑜𝑟𝑡𝐼𝑑𝑠, 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠): for each thread 𝑖 do 19 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 [𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝑠 [𝑎𝑟𝑔𝑠𝑜𝑟𝑡𝐼𝑑𝑠 [𝑖 ] ] ] ← 𝑖;

17

18

Function gather(𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐, 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠, 𝐶𝑜𝑢𝑡 ): for each thread 𝑖 do 22 𝐶𝑜𝑢𝑡 [𝑖 ] ← 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 [𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐 [𝑖 ] ];

20

21

access width. For the other cases, fixing the parallelism and access width to 1 is the most efficient.

4.3

Scenario III: Smart Keys Fusion

Databases often need to handle multiple keys in operators like join, group-by, and sorting. Managing multiple keys efficiently on the GPU is particularly challenging. The first reason is that each comparison (== or <) requires reading multiple keys, which are not stored together in the memory due to the columnar layout. The second reason is that it is harder to implement efficient algorithms for multiple keys due to the limited shared memory capacity. To solve this, we propose a smart key fusion mechanism (SKF) that can assess the compressibility of each key and combine them into a single key after compressing them. It works like the following. At query execution time, for each key, SKF will evaluate the range of values by finding the minimum and maximum. If the range (max − min) is narrower than a preset threshold, we will compress the key by frame-of-reference encoding and bitpacking. Otherwise, we will try to calculate the HLL++ sketch on the key and find the number of distinct values (NDV). If NDV is below a certain threshold, we will use dictionary encoding to compress the keys.

SMJ PHJ

HASH

Thpt (Billion Tuples/s)

25 20 15 10 5

212

216

220 |R|

224

228

SMJ/UR PHJ/UR HJ

212

216

SMJ/TR PHJ/TR

220 |R|

Thpt (Billion Tuples/s)

11 10 9 8 7 6 5 4 3

224

5.1

Experimental Setup

Hardware and system software. We use two GPU platforms: • A100. NVIDIA A100-SXM4 (40 GB HBM2e memory). CUDA 12.4 and NVIDIA driver 550.90.07. The host system is x86_64 with an Intel Xeon CPU (12 cores; 2 sockets; SMT enabled). This is provided as the a2-highgpu-1g machine type from Google Cloud. • GH200. NVIDIA GH200 (96 GB HBM3 memory). CUDA 13.1 and NVIDIA driver 550.54.15. The host system is aarch64 with an Arm Neoverse-V2 CPU (288 cores; 4 sockets). This is provided by the Swiss National Supercomputing Center. If the results from both platforms show similar conclusions, we will only include the results from GH200. Baselines and methodology. Our baseline is cuDF v25.12. Unless otherwise stated, all benchmarks operate on GPU-resident data and report operator throughput as measured by our benchmark harness. We use the nvbench benchmark framework for our methods, which repeats the runs until the measurement has a noise level below 0.5%. For cuDF measurements, we use their pylibcudf for better programmability without introducing extra Python-incurred overhead. In addition, we use a pre-allocated memory pool to reduce allocator-induced instability during the experiments, and we report the minimum time over repeated runs to be conservative. We additionally force deallocation using del and gc.collect() in the benchmark scripts to avoid OOM errors during sweeps. For microbenchmarks that involve numeric values, we measure both 4-byte data types (I32) and 8-byte data types (I64).

6 5

SMJ/UR PHJ/UR HJ

SMJ/TR PHJ/TR

4 3 2 212

216

220 |R|

224

8

216

220 |R|

228

(c) 4-byte join key (A100).

SMJ/UR PHJ/UR HJ

7

224

228

SMJ/TR PHJ/TR

6 5 4 3

228

212

216

220 |R|

224

228

(b) 8-byte join key (GH200). Thpt (Billion Tuples/s)

This section evaluates Eiger using microbenchmarks for joins, groupby, expression evaluation, string processing, and sorting, as well as end-to-end with the queries of the TPC-H benchmark.

212

HASH

Figure 2: Narrow join with fixed |𝑆 | = 228 (GH200).

Thpt (Billion Tuples/s)

EXPERIMENTAL EVALUATION

SMJ PHJ

(b) 8-byte join key.

(a) 4-byte join key (GH200).

5

20 18 16 14 12 10 8 6 4

(a) 4-byte join key.

Thpt (Billion Tuples/s)

After the compression, we are able to represent each key with fewer bits. Then, we will invoke a single kernel to fuse the compressed keys by bit shifting into 4-byte integers or 8-byte integers. During fusion, we will alter the keys based on the data type (e.g., signed vs. unsigned) and sorting order requirements to ensure correctness. Out of the above process, when used for sorting, the dictionary encoding is particularly tricky because it must preserve the order of the keys in the encoded values for the obvious correctness reason. To achieve this, we propose an order-preserving dictionary encoding algorithm for the GPU (Algorithm 1). In step 1, for each unique key, a leader tuple is picked, and the followers record the leader tuple ID in their own entry in the 𝑔𝑟𝑜𝑢𝑝𝐼𝑑𝐿𝑜𝑐 array (function compute_group_id). The function also returns the number of distinct values 𝑛𝑑𝑣. In step 2, we retrieve the pairs (𝑢𝑛𝑖𝑞𝑢𝑒𝐾𝑒𝑦, 𝑔𝑟𝑜𝑢𝑝𝐼𝑑) from the hash map (line 6). There are in total 𝑛𝑑𝑣 pairs. The 𝑔𝑟𝑜𝑢𝑝𝐼𝑑s record the tuple IDs of leader tuples. In step 3, we argsort the unique keys retrieved from the last step (line 7). The 𝑎𝑟𝑔𝑠𝑜𝑟𝑡𝐼𝑑𝑠 tells which position a unique key stays in before the argsort. In step 4, we initialize the 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 array with |𝐶𝑖𝑛 | elements and then fill in the entry of each group leader with their order-preserving encoded value (lines 8-9). Lastly, in step 5, all items 𝐶𝑖𝑛 fetch the encoded value stored in their leader’s entry in 𝑠𝑝𝑎𝑟𝑠𝑒𝐺𝑟𝑜𝑢𝑝𝐼𝑑𝑠 (line 10).

Thpt (Billion Tuples/s)

Eiger: An Efficient Library for GPU-based Data Analytics

4.0

SMJ/UR PHJ/UR HJ

3.5

SMJ/TR PHJ/TR

3.0 2.5 2.0 1.5 212

216

220 |R|

224

(d) 8-byte join key (A100).

Figure 3: Wide join with fixed |𝑆 | = 228 (3a-3c) or |𝑆 | = 227 (3d).

5.2

Joins

We evaluate the hash join (HJ), partitioned hash join (PHJ), and sortmerge join (SMJ). Our hash join uses the same static_multiset implementation from cuCollection with cuDF; therefore, we skip comparing to cuDF in this experiment. For PHJ and SMJ, we consider two materialization strategies: GFUR (gather-from-untransformedrelations) and GFTR (gather-from-transformed-relations). HJ is only compatible with GFUR. When both relations have at most two columns, GFUR and GFTR are equivalent. For brevity, we use the notation “SMJ/UR” to indicate the “join algorithm/materialization strategy” in the text. We denote the smaller side of the join as relation 𝑅 and the larger side as relation 𝑆. The size of the relation is indicated by | · |. The throughput of the join is defined as (|𝑅| + |𝑆 |)/Time. 5.2.1 Varying the Left Table Size. In this experiment, we fix |𝑆 | = 228 and vary |𝑅|, the size of the smaller relation 𝑅. This is to simulate a typical type of join, where at least one side of the join is very large in number of rows. Depending on whether we need to materialize

Wide join. In this case, each relation has four columns, and materialization of three non-key columns is necessary. Figure 3 repeats the experiment with four columns per table (two additional payload columns per relation compared to the narrow join). Here, A100 and GH200 show different behaviors. For 4-byte keys, PHJ/TR achieves the best performance across all configurations for both GH200 and A100 GPUs. HJ is secondbest only when the left table is sufficiently small; this threshold differs between platforms: HJ remains competitive up to |𝑅| ≤ 222 on A100, but only up to |𝑅| ≤ 219 on GH200, after which it quickly falls behind. Compared to the narrow join, HJ no longer has the worst performance for small-to-medium |𝑅| due to its more efficient materialization. When the build relation is small, materializing the payloads from it is relatively efficient, and materializing the payloads from the probe side always involves sequential accesses. Across all sizes, GFUR-based approaches are consistently slower than their GFTR counterparts. For 8-byte keys, on GH200, HJ performs best for small left tables and is then overtaken by PHJ/TR as the left table grows; on A100, HJ remains the best until |𝑅| = 220 . For both platforms, PHJ remains faster than SMJ regardless of the materialization strategy, and SMJ/TR is the slowest algorithm. 5.2.2 Varying the Right Table Size. In the following set of experiments, we do the opposite of the previous experiment by fixing |𝑅| = 216 and varying the size of |𝑆 | while keeping |𝑆 | ≥ |𝑅|. This allows us to understand medium-sized joins better, where neither of the two relations is too large. Figure 4 shows the results for GH200. A100 has similar results. In the narrow join scenario, for both 4-byte and 8-byte join keys, HJ performs the best for small |𝑆 | (up to 219 for 4-byte keys, and up to 222 for 8-byte keys). This agrees with the results of the previous experiment that HJ has a greater advantage for I64 keys. Interestingly, the results show that SMJ outperforms PHJ for medium 𝑆

SMJ PHJ

218

HASH

222 |S|

Thpt (Billion Tuples/s)

17.5 15.0 12.5 10.0 7.5 5.0 2.5 0.0

(a) Narrow join: 4-byte join key. 8 6

SMJ/UR PHJ/UR HJ

SMJ/TR PHJ/TR

4 2 0

218

222 |S|

14 12

226

(c) Wide join: 4-byte join key.

SMJ PHJ

HASH

218

222 |S|

10 8 6 4 2 0

226

226

(b) Narrow join: 8-byte join key. Thpt (Billion Tuples/s)

Narrow join. In this case, each relation has two columns and materialization of non-key columns can be “inlined” into the join, so no extra materialization is needed. Figure 2 shows that with 4-byte keys and payloads, both SMJ and PHJ significantly outperform HJ across all left table sizes. This comes as a surprise because HJ is commonly considered very efficient for joins with a small build relation. The reason for the inferior performance is that HJ, in this case, inserts and probes with 8-byte tuple IDs (see Section 3.1 and Section 3.2), while SMJ and PHJ process the 4-byte keys and payloads without using TIDs. Compared to SMJ, PHJ is consistently faster. In general, the throughput of PHJ and SMJ increases with larger |𝑅|, while the throughput of HJ decreases. For 8-byte keys and payloads, SMJ performance decreases substantially compared to the 4-byte case and becomes close to HJ for small-to-medium |𝑅|; however, SMJ still outperforms HJ at large left table sizes. PHJ performance remains the best across all |𝑅|s. This is surprising because even though both HJ and PHJ deal with 8-byte data types, HJ still does not have an advantage over PHJ for small-to-medium |𝑅|s.

Thpt (Billion Tuples/s)

the payloads, we study the narrow join (no materialization) and the wide join (needs materialization).

Thpt (Billion Tuples/s)

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso

6 5

SMJ/UR PHJ/UR HJ

SMJ/TR PHJ/TR

218

222 |S|

4 3 2 1 0

226

(d) Wide join: 8-byte join key.

Figure 4: Vary |𝑆 | with fixed |𝑅| = 216 (GH200).

(≤ 226 for I32 and ≤ 224 for I64). This is in contrast with the results from the last section, where PHJ is always better than SMJ for a very large relation 𝑆. In the wide join scenario, the HJ has the best or second-best performance across |𝑆 | for I32 and consistently outperforms other methods for I64. The advantage of SMJ over PHJ remains in the wide join for small-to-medium-sized relation 𝑆. For 4-byte keys, SMJ/UR is more efficient than SMJ/TR for |𝑆 | ≤ 224 but is less efficient for larger |𝑆 |s. For 8-byte keys, SMJ/UR is always better than SMJ/TR. PHJ/TR and PHJ/UR, for both I32 and I64, start with similar performance, but PHJ/TR performs better as |𝑆 | grows. This experiment reveals how different join algorithms perform under medium problem sizes (roughly measured by |𝑅|+|𝑆 |), whereas previous work [55] often focuses on large problem sizes. We demonstrate that in this case, the relative performance of different algorithms is significantly different from the large-sized joins. 5.2.3 Varying the Match Ratio. In this set of experiments, we study the effect of the join match ratio, which is defined as the percentage of tuples from relation 𝑆 that have a match in relation 𝑅. The match ratio can influence the efficiency of materialization. Due to page limitations, we only show the subset of results where |𝑆 | = 228 . The only exception is Figure 5d, where we set |𝑆 | = 227 to avoid running out of memory. Previous work [55] has studied in-depth the case where both 𝑅 and 𝑆 are large; therefore, here we study a medium-sized 𝑅, with |𝑅| = 216 . Each table has four columns. Figure 5 shows the results of this set of experiments. For 4-byte join keys, when the match ratio is lower than a certain threshold (30% for GH200 and 20% for A100), PHJ/UR is the most efficient because the random gathering cost is low. Beyond this threshold, PHJ/TR is the most efficient, followed by HJ and SMJ/TR. For 8-byte keys, PHJ/UR remains the best algorithm for higher match ratios on both platforms before being overtaken by PHJ/TR and HJ.

5 0

0.2

0.4 0.6 0.8 Match ratio

1.0

Thpt (Billion Tuples/s)

(a) 4-byte join key (GH200). 12

SMJ/UR PHJ/UR HJ

10

SMJ/TR PHJ/TR

8 6 4 2 0

0.2

0.4 0.6 0.8 Match ratio

1.0

(c) 4-byte join key (A100).

0.2

SMJ/TR PHJ/TR

0.4 0.6 0.8 Match ratio

1.0

(b) 8-byte join key (GH200). SMJ/UR PHJ/UR HJ

8

SMJ/TR PHJ/TR

6 4 2 0

0.2

0.4 0.6 0.8 Match ratio

50

30 20 10

Group-By

5.3.1 The Effect of Group Cardinality. For 4-byte keys and payloads, hash-based group-by is the most efficient up to 𝑔 = 214 on both A100 and GH200, and in that regime it substantially outperforms cuDF, which also implements a hash-based group-by. Across group cardinalities, sort-based approaches are comparatively stable. GFTR outperforms GFUR for most group cardinalities (with the noted exception at 𝑔 = 4). SORT-OPT provides an additional advantage over vanilla SORT/TR only up to a platform-dependent threshold: up to 𝑔 = 214 on A100 and up to 𝑔 = 210 on GH200. PARTITION/TR performs the best when 𝑔 > 220 on both platforms. For 8-byte keys and payloads, the hash-based implementation remains the most efficient up to 𝑔 = 218 on A100 and up to 𝑔 = 214 on GH200; within this range, it again substantially outperforms cuDF. Sort-based approaches remain stable; however, unlike the 4-byte

2

6

10

14

18

2 2 2 2 Number of groups

SORT/TR SORT-OPT SORT/UR Hash

30 25

22

2

PARTITION/TR PARTITION/UR Auto cuDF

15 10 5 0

22

26

210 214 218 222 Number of groups

50

PARTITION/TR PARTITION/UR Auto cuDF

40 30 20 10 0

26

20

22

26

210 214 218 222 Number of groups

226

(b) 8-byte group key (GH200). SORT/TR SORT-OPT SORT/UR Hash

30 25

PARTITION/TR PARTITION/UR Auto cuDF

20 15 10 5 0

226

(c) 4-byte group key (A100).

Figure 5: Effect of match ratio on join performance.

5.3

2

2

SORT/TR SORT-OPT SORT/UR Hash

60

(a) 4-byte group key (GH200).

(d) 8-byte join key (A100).

We evaluate hash-based, sort-based, and partition-based group-by implementations, including GFTR/GFUR variants and an optimized sort-based variant (SORT-OPT) that applies dictionary-encoding optimization on top of SORT/TR. To demonstrate the effectiveness of the runtime adaptive execution for group-by (Section 4.1), we also include a series in which the algorithm is selected by calculating the HLL++ sketch. The time to compute the sketch is also included. Unless otherwise stated, the number of rows is fixed to 228 , and the number of aggregations is fixed to 2. We denote the number of groups (i.e., group cardinality) as 𝑔. The results are shown in Figure 6.

PARTITION/TR PARTITION/UR Auto cuDF

40

0

1.0

When comparing the GH200 and A100 GPUs, we find that on GH200, a higher match ratio (hence a higher materialization cost) is needed to justify the usage of the GFTR technique. This could imply that the random memory access performance is improved from A100 to GH200.

SORT/TR SORT-OPT SORT/UR Hash

60

Thpt (Billion Tuples/s)

10

SMJ/UR PHJ/UR HJ

Thpt (Billion Tuples/s)

15

20.0 17.5 15.0 12.5 10.0 7.5 5.0 2.5 0.0

Thpt (Billion Tuples/s)

20

SMJ/TR PHJ/TR

Thpt (Billion Tuples/s)

SMJ/UR PHJ/UR HJ

Thpt (Billion Tuples/s)

25

Thpt (Billion Tuples/s)

Thpt (Billion Tuples/s)

Eiger: An Efficient Library for GPU-based Data Analytics

22

26

210 214 218 222 Number of groups

226

(d) 8-byte group key (A100).

Figure 6: Group-by microbenchmarks. Table 2: Expressions for microbenchmark Id

Expression

E1 E2 E3 E4 E5 E6

2 * col0 col0 * (1 - col1) col0 * (1 - col1) * (1 + col2) col0 * (1 - col1) - col2 * col3 col0 >134217728 col0 >161061273 and col1 <375809638 (col0 == 0 and col1 == 1) or (col0 == 1 and col1 == 0)

E7

#Cols

Description

1 2 3 4 1 2

Simple arithmetic TPC-H Q1-like arithmetic expression TPC-H Q1-like arithmetic expression TPC-H Q9-like arithmetic expression Simple predicate Simple logical and

2

Simple chained predicates

case, GFTR no longer consistently outperforms GFUR. SORT-OPT achieves the best overall performance and retains its advantage over vanilla SORT/UR up to 𝑔 = 222 . PARTITION/TR again performs the best when 𝑔 > 220 on both platforms. 5.3.2 Runtime-Adaptive Algorithm Selection. The series labeled as “Auto” in Figure 6 refers to the selection of the group algorithm guided by HLL++, which is part of our runtime adaptive execution mechanism introduced in Section 4. It includes the time to compute the HLL++ sketch to estimate the number of groups and the time to execute the selected algorithm based on the estimation. The results show that our proposed runtime adaptive mechanism can almost always find the optimal algorithm. Moreover, it shows that HLL++ adds negligible overhead, and the estimation is accurate.

5.4

Expression Evaluation

We compare our two expression evaluation backends (Section 3.5) against cuDF’s AST-based expression evaluation. We do not choose the cuDF just-in-time (JIT) based evaluator because JIT incurs a high compilation overhead and is usually only good for very complex or recurring expressions. The expressions we use to evaluate are listed in Table 2, which contains a rich mix of arithmetic (E1-E4)

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso

Time (ms) 25 20

BB PTI cuDF

Time (ms) 17.5 15.0

BB PTI cuDF

Thpt (109 strings/s) 100 [1,2] 80

12.5

60

15

10.0

40

10

7.5

20

5.0

0

5 0

2.5 E1 E2 E3 E4 E5 E6 E7

(a) 4-byte columns (A100) Time (ms) 35 30

BB PTI cuDF

25 20 15 10 5 0

E1 E2 E3 E4 E5 E6 E7

(c) 8-byte columns (A100)

0.0

E1 E2 E3 E4 E5 E6 E7

(b) 4-byte columns (GH200) Time (ms) 20.0 17.5 15.0 12.5 10.0 7.5 5.0 2.5 0.0

BB PTI cuDF

Thpt (109 strings/s) Ours cuDF

[1,4]

[1,1]

[1,4] [1,2]

[1,4] [1,2]

[1,2]

30 45 60 75 90 105120135 String Length (bytes)

(a) 5-byte (aligned)

prefix

matching

Thpt (109 strings/s) 50 40

[1,2] [2,8]

20

[2,4]

[2,8]

[2,8]

10 0

E1 E2 E3 E4 E5 E6 E7

(d) 8-byte columns (GH200)

Figure 7: Expression evaluation microbenchmarks.

60

75 90 105 120 135 String Length (bytes)

(c) 50-byte prefix matching (aligned)

40 35 30 [1,1] 25 20 15 10 5 0 51

Ours cuDF [1,2]

[2,4] [1,2]

[2,8] [2,8]

68 85 102 119 136 String Length (bytes)

(d) 50-byte prefix matching (unaligned)

Figure 8: String prefix matching microbenchmarks (GH200). Thpt (109 strings/s)

Thpt (109 strings/s) Ours cuDF

80 [1,2]

and logical (E5-E7) operations. From E1 to E4 and from E5 to E7, the number of intermediate results needed to evaluate the expression increases. The results are shown in Figure 7. Figures 7a and Figure 7b show the performance when all the columns involved are 4-byte wide. As the number of columns and the complexity of the expression increase, all implementations experience a longer execution time. Both of our implementations outperform cuDF. Of seven expressions, PTI and BB outperform cuDF by up to 1.4× and 3× on A100 and 1.6× and 5× on GH200. BB has a bigger performance advantage over PTI and cuDF on GH200. For columns of 8-byte width (Figure 7c and 7d), PTI still consistently outperforms cuDF for E1-E6, but loses to cuDF for E7. On A100, BB only outperforms cuDF for E1 and logical expressions (E5-E7). As the size of intermediate results doubles for 8-byte columns when processing E1-E4, the execution time of BB also increases by around a factor of 2, while the PTI and cuDF experience a much less dramatic performance degradation. For E5-E7, BB is less affected because the intermediate results are boolean values, which remain the same size for both the 4- and 8-byte cases. On GH200, due to the much higher memory bandwidth, materializing the intermediate results becomes less costly, which makes BB consistently outperform cuDF. The relative performance between BB and PTI is very different depending on the GPUs. BB is more sensitive to the size of the intermediate result in A100. This implies that when choosing a suitable implementation, the hardware specification must also be taken into account. In general, BB is preferred for expressions with predominantly logical operations, where the size of intermediate results is small. For GPUs with a lower memory bandwidth, like A100, PTI is preferred for arithmetic-heavy expressions; for GPUs with a

(b) 5-byte prefix matching (unaligned) Thpt (109 strings/s)

Ours cuDF

[1,4]

30

Ours

140 [1,1] cuDF 120 100 80 [1,2] [1,1] 60 [1,1] [1,2] [1,4] 40 [1,4] [1,4] 20 0 17 34 51 68 85 102119136 String Length (bytes)

Ours cuDF

[1,1]

80

60

60

[1,4]

40 [1,1]

20 0

100

[2,8] [2,4]

40

[4,8] [2,8]

[4,8]

30 45 60 75 90 105120135 String Length (bytes)

(a) Full matching (aligned)

20 0

[1,2]

[1,1]

[1,4]

[2,4]

[2,8] [2,8]

[2,8]

17 34 51 68 85 102119136 String Length (bytes)

(b) Full matching (unaligned)

Figure 9: String exact matching microbenchmark (GH200). higher memory bandwidth, BB is preferred for low intermediate result sizes (for example, 4-byte columns).

5.5

String Processing

We evaluate prefix/suffix, exact, and substring matching. If not specifically stated, we measure with 227 strings on both platforms for a fair comparison. Without loss of generality, when generating the column of strings, we always generate fixed-length strings for easier analysis. In order not to be misled by the benefit of accidental alignment, we also study the case with unfavorable alignments. 5.5.1 Prefix/Suffix/Exact Matching. In this set of experiments, we study the performance of prefix/suffix and exact string matching. For the Arrow string format, prefix and suffix matching are equivalent; therefore, we only include the results for prefix matching. Figures 8a–8d show the results of prefix matching over two pattern lengths, 5 and 50. Figures 8a and Figure 8c increase the length of the string from 30/60 bytes with a step size of 15, which

Eiger: An Efficient Library for GPU-based Data Analytics CG Size 1 1.00

CG Size 1.35

1.11

0.78

1.25

1 1.00

Thpt (109 strings/s) 0.76

0.77

0.65

1.0 0.9

1.00

2 0.51

0.53

0.44

2 0.70

0.40

0.66

0.63

0.59

0.75

4 0.41

0.45

0.37

0.34

0.50

0.7

4 0.56

0.53

0.55

0.51

1 2 4 8 Access Width (bytes)

1 2 4 8 Access Width (bytes)

(a) String length = 30

(b) String length = 45

CG Size 1 1.00

1.48

1.24

1.32

1.6

1 1.00

1.42

1.57

1.28

1.35

1.40

1.0

20

[1,1]

[1,2] [1,2]

1.94

2.14

2.13

3.19

3.99

2.77

3.72

[1,1]

1 2 4 8 Access Width (bytes)

1 2 4 8 Access Width (bytes)

(c) String length = 90

(d) String length = 135

[1,2]

30 45 60 75 90 105120135 String Length (bytes)

20

1

Figure 10: Effect of parallelism and packed access width on full string matching (GH200).

often creates favorable alignments. On the other hand, Figure 8b and Figure 8d increment the string length from 17/51 bytes with a step size of 17, creating less favorable alignments. Our implementation exposes two tuning parameters for prefix/suffix/full matching: the number of threads per string (parallelism) and the packed access width; the data points of our method in the figures are annotated with the best-performing configuration in the form [𝑥, 𝑦] where 𝑥 is parallelism and 𝑦 is access width. This assists in demonstrating how the best configuration changes with the change of string lengths. On GH200, Eiger is consistently faster than cuDF, and the advantage becomes more pronounced as the length of the string increases: Eiger’s throughput degrades more slowly with the length of the string. On A100, performance (not shown) is generally similar to cuDF across the same sweep, without a pronounced degradation trend relative to cuDF. For short patterns in prefix/suffix matching, using one thread per string is consistently best across both platforms. The effect of access-width follows alignment: on GH200, when the string length is a multiple of 4, 4-byte access is most efficient, otherwise 2-byte access. For the long pattern case, as the string length increases, it is beneficial to increase the parallelism as well as the access width for both favorable and unfavorable alignments. Whether the alignment is favorable or not does not have an impact on the performance. For exact matching (Figure 9), we make observations very similar to those of prefix/suffix matching. To better understand the effect of parallelism (cooperative group size (CG size)) and packed access width, we created the speedup heatmaps in Figure 10. For both GH200 and A100, increasing parallelism and access width tends to

10 5 0

0

[1,4]

[1,1] [1,4]

60

[1,4]

75 90 105 120 135 String Length (bytes)

(b) 50-byte pattern (GH200) Thpt (109 strings/s)

Ours Ours (KMP) cuDF

15

4.01

4.09

10

[1,4]

[1,1]

30 [1,1] 4

Ours Ours (KMP) cuDF

[1,4]

20

[1,1]

10 0

40 30

30

(a) 5-byte pattern (GH200)

2

4 2.21

40 [1,1]

Thpt (109 strings/s)

3

2 2.50

1.65

50

25

1.2

4 0.98

0.6

CG Size

1.4

2 1.00

0.8

Thpt (109 strings/s) Ours Ours (KMP) cuDF

[1,1] [1,1] [1,2] [1,1]

[1,2] [1,2]

[1,2]

30 45 60 75 90 105120135 String Length (bytes)

(c) 5-byte pattern (A100)

20.0 [1,4] 17.5 15.0 12.5 10.0 7.5 5.0 2.5 0.0 60

Ours Ours (KMP) cuDF

[1,4] [1,1] [1,8] [1,4]

[1,4]

75 90 105 120 135 String Length (bytes)

(d) 50-byte pattern (A100)

Figure 11: Substring matching microbenchmarks. help when the string length increases. For shorter patterns (e.g., pattern length=30), additional threads provide little benefit, while 2-byte access can help. At pattern length 45, one thread and 1-byte access is most efficient, possibly because string lengths are odd. For larger patterns, the optimal parallelism differs by GPUs: at pattern length 90, GH200 prefers 2 threads with larger access width, whereas A100 already prefers 4 threads and 8-byte access; at pattern length 135, performance correlates positively with both parallelism and access width, and this trend appears earlier on A100. 5.5.2 Substring Matching. We evaluate substring matching with a “vanilla” implementation, where the packed access width is configurable, and the parallelism is fixed to 1, and a KMP-based implementation without tuning knobs. Figures 11a–11b and Figures 11c–11d show the results for GH200 and A100, respectively. Our vanilla approach (“Ours” in the figures) is slightly more efficient than cuDF for all string lengths and pattern lengths on both GPUs. The speedup is more prominent for larger pattern lengths. The KMP variant does not provide consistent benefits: on A100, it does not significantly help; on GH200, it slightly improves performance for string length below 120 when the pattern length is short, but when the pattern length becomes 50, it does not help and can even hurt for short strings. 5.5.3 String Sorting. We evaluate string sorting by varying string length, using 226 strings. We generate each string by randomly picking each character from “a” to “z”. Figure 12 shows that the optimized sorting using prefix extraction is more efficient than the naive merge sort on both platforms. The time of optimized sort also includes the prefix extraction time. As the string length increases, the optimized sort almost maintains the same performance, whereas the naive merge sort experiences a decrease in performance. The performance benefit comes from extracting the prefix before the actual sorting, which allows most

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso Time (ms) 1000 800

Time (ms)

Naive sort Opt. sort

300 200

400

150 50 0

16 32 64 128 256 String Length (bytes)

16

32 64 128 256 String Length (bytes)

(a) A100

(b) GH200

Figure 12: String sorting microbenchmarks. Table 3: Breakdown of sorting 256-byte strings (GH200). Algorithm

Kernels

Time (ms)

Naive merge sort

Merge Block-level sort

261.579 69.541

Optimized merge sort

Merge Block-level sort Prefix extraction

104.062 2.911 1.867

of the comparisons to be made based on the 4-byte prefixes only. This avoids reading bytes from random locations in the memory during comparison, which suffers from a worse memory access pattern. This is confirmed by the breakdown shown in Table 3. The prefix extraction has a very minimal overhead but greatly reduces the execution time of the merge sort.

Sorting

In this experiment, we focus on the performance of multi-key sorting, commonly seen in queries, and demonstrate the effectiveness of Eiger’s smart key fusion (SKF) mechanism. We fix the number of rows to be 228 and vary the number of key columns and the type of keys (I32 and I64). Half of the key columns can be compressed because of the narrow value distribution (i.e., max − min is small), while the other half has a sparse distribution but a low number of distinct values, where HyperLogLog++ can identify this case, and dictionary encoding will be used to compress the column. Figure 13 shows the results of this experiment. The time reported for SKF also includes the time spent collecting statistics and fusing Time (ms) 2500 w/o SKF 2000

w/ SKF

1500 1000 500 0

Algorithm

Kernels

Time (ms)

Merge sort

Merge Block-level sort

612.311 35.791

SKF + Radix sort

Radix sort Dictionary encoding Fusing keys Computing HLL++

12.224 6.936 3.393 1.061

100

200

5.6

Table 4: Breakdown of sorting tables (4xI32) (GH200).

250

600

0

Naive sort Opt. sort

2xI32 2xI64 4xI32 4xI64 #Keys and key type

(a) A100

Time (ms) 800 700 600 500 400 300 200 100 0

w/o SKF w/ SKF

2xI32 2xI64 4xI32 4xI64 #Keys and key type

(b) GH200

Figure 13: Sort microbenchmark. SKF = smart key fusion. 2xI32 means two columns of 4-byte keys.

the keys. The implementation without SKF is the merge sort. The results show that SKF significantly improves performance by up to 13× on A100 and 12× on GH200 compared to the merge sort approach, even with the overhead of statistics computation and key fusion. The performance gain comes from the fact that fusing keys together enables us to use the more efficient radix sort, which has a lower time complexity than the merge sort. Fusing the keys together also results in fewer bytes being read, in contrast to the merge sort, where each comparison needs to read multiple keys. We detail the time breakdown in Table 4. It is obvious that the radix sort is more than an order of magnitude more efficient than the merge sort in this case, and the overhead of calculating statistics and compressing keys is well paid off.

5.7

TPC-H Benchmark

In this section, we evaluate Eiger using the standard TPC-H query benchmark [52]. Table 5 shows the execution time per-query at three different scale factors, as well as the comparison with cuDF. For the cuDF baseline, we use Maximus [18], an open-source query execution engine that integrates cuDF for GPU execution. Data movement times between the CPU and GPU are excluded, since the evaluation aims at comparing the efficiency of GPU execution. As discussed in Section 3, Eiger provides different implementations for almost all operators to adapt to a wide range of workloads. Therefore, we include two different measurements from Eiger, baseline and best. Eiger (baseline) uses the same algorithm and configuration for all occurrences of the same operation. Specifically, hash-based algorithms are used for join and grouped aggregation, PTI-based expression evaluation is used for selection and projection, and smart-key-fusion (SKF) optimization is disabled for sorting and grouped aggregation. The selection of baseline implementations mirrors the cuDF algorithms. On the other hand, Eiger (best) combines the best-performing implementation of each operator. For brevity, we only show Eiger (best) vs. cuDF in Table 5, and Eiger (best) vs. Eiger (baseline) is shown in Figure 14. For the total runtime of 22 queries, Eiger (best) is 1.7× (SF=10), 1.8× (SF=30) and 1.8× (SF=100) faster than cuDF. The most significant speedup is observed for Q16, Q13, and Q17. For Q16, cuDF uses the expensive merge sort for the “count(distinct)” aggregation while Eiger uses a more efficient hash-based approach. For Q13 and Q17, Eiger wins by using the the partition hash join. Compared to Eiger (baseline), Eiger (best) is 1.2× (SF=10), 1.4× (SF=30), and 1.5× (SF=100) faster. This shows the importance of having multiple implementations per operator, especially for larger datasets. To further demonstrate that the choice of operator implementations can greatly influence performance, we plot the operator time breakdown for SF=100 in Figure 14. The results reveal three main

Eiger: An Efficient Library for GPU-based Data Analytics

Table 5: TPC-H evaluation on GH200. Time is in ms. SF=10 Q1 Q2 Q3 Q4 Q5 Q6 Q7 Q8 Q9 Q10 Q11 Q12 Q13 Q14 Q15 Q16 Q17 Q18 Q19 Q20 Q21 Q22

SF=30

SF=100

Eiger (best)

cuDF

Speedup

Eiger (best)

cuDF

Speedup

Eiger (best)

cuDF

Speedup

14.36 5.66 6.11 3.91 8.28 1.6 8.57 7.15 13.94 7.19 1.81 4.73 6.52 1.86 2.11 5.23 4.18 7.59 4.56 4.52 35.36 1.95

24.01 6.73 11.11 4.67 13.45 4.42 14.52 12.78 17.95 9.61 3.36 8.81 23.02 4.4 4 17.46 9.49 8.37 7.76 6.73 52.29 5.25

1.7× 1.2× 1.8× 1.2× 1.6× 2.8× 1.7× 1.8× 1.3× 1.3× 1.9× 1.9× 3.5× 2.4× 1.9× 3.3× 2.3× 1.1× 1.7× 1.5× 1.5× 2.7×

40.54 7.96 13.14 10.55 19.77 4.41 16.69 15.94 34.15 16.13 3.52 12.16 16.82 4.85 4.65 9.53 8.66 19.68 12.39 9.27 95.07 4.43

67.59 8.86 23.66 10.81 34.49 11.29 24.13 31.27 50.08 20.63 5.91 21.36 68.32 8.68 7.22 45.88 23.66 20.77 19.27 11.73 146.75 8.84

1.7× 1.1× 1.8× 1× 1.7× 2.6× 1.4× 2× 1.5× 1.3× 1.7× 1.8× 4.1× 1.8× 1.6× 4.8× 2.7× 1.1× 1.6× 1.3× 1.5× 2×

132.22 11.4 34.2 33.38 58.94 14.08 43.31 85.69 94.09 53.18 7.58 40.41 49.59 14.58 13.2 21.85 23.75 64.24 38.58 21.45 307.4 12.66

219.51 15.89 74.75 31.83 110.44 34.59 73.45 98.62 165.71 60.69 14.99 67.72 224.66 24.3 17.64 133.31 73.7 64.46 59.86 26.63 483.89 22.26

1.7× 1.4× 2.2× 1× 1.9× 2.5× 1.7× 1.2× 1.8× 1.1× 2× 1.7× 4.5× 1.7× 1.3× 6.1× 3.1× 1× 1.6× 1.2× 1.6× 1.8×

influential factors that contribute to the performance gain. (1) Some queries (e.g., Q4, Q6, Q7, Q12, Q14, Q15, Q20) benefit from a better filter performance using the batch-based expression evaluation. The reason is that the predicates in TPC-H are commonly simple, and the intermediate results written back to the memory are only 1-byte booleans. This makes BB more efficient than PTI in evaluating expressions by reducing the interpretation overhead while not significantly increasing the memory load and store. (2) Another group of queries (e.g., Q13, Q21) benefit from a more efficient group-by and distinct implementation. Q13 benefits from using the partition-based group-by with GFTR materialization. The group-by in the subquery has a high group cardinality, equal to the total number of rows in the customer table, making partition-based group-by the best implementation (see Section 5.3). Q21’s distinct operators get more efficient from the smart-key-fusion (SKF) that fuses the l_orderkey and l_suppkey and using a sort-based implementation. (3) A third group of queries benefit substantially from more efficient joins, including Q2, Q3, Q5, Q7, Q8, Q9, Q11, Q13, Q16, Q17, and Q20. Joins in these queries gain more efficiency by using the sort-merge join or partition hash join with GFUR or GFTR materialization strategies. This justifies Eiger’s design principles as well as proposed optimization techniques. Together with the microbenchmark results, we demonstrate that Eiger offers better performance due to its more efficient operator implementation, richer choice of algorithms, and the runtime adaptive execution mechanism.

6

RELATED WORK

GPU-accelerated query engines. A multitude of research prototypes and industrial systems [2, 7, 12, 15, 16, 18, 22, 25–27, 30, 31, 35, 41, 45, 47, 54, 56–60] have explored query processing on GPUs from many different angles. Eiger distinguishes itself from this body of work by focusing on the operators themselves: it offers multiple implementation variants per operator, adapts the choice and configuration to the workload at runtime, and can be integrated into these engines to further improve their performance. Among

these systems, Themis [15] is also adaptive, but it targets load imbalance across threads and warps within an operator, whereas Eiger adapts the choice of algorithms and configurations to the data being processed. Studies of GPU operators. Many studies have focused on performance analysis [6] and optimizing various database operations on a single GPU, including join [10, 15, 20, 28, 29, 36, 43, 48, 50, 55], group-by [21, 24, 42, 51, 55], string processing [49], encoding and decoding [3, 13, 17, 46], etc. Wu et al. [55] and Sioulas et al. [48] propose two variants of partitioned hash join. The former is implemented in Eiger because of its compatibility with the GFTR technique. Wu et al. in the same work also propose three families of group-by implementations, including hash-based, sort-based, and partition-based. Eiger incorporates and further improves them and introduces the use of the HyperLogLog++ sketch to guide the algorithm selection. Shanbhag et al. [45] propose a block-based processing paradigm and implement basic filtering, projection, join, and group-by operations. However, their implementation is not generic enough to handle arbitrary expressions or input. Sitaridi et al. [49] propose multiple techniques for substring matching and alternative string formats. Eiger improves string processing performance by leveraging capabilities of modern GPUs and the CUDA programming model, such as packed accesses and cooperative groups. Adaptive query processing. Runtime adaptivity has a long tradition in CPU databases, ranging from Eddies that reroute individual tuples among operators during execution [1] to mid-query re-optimization and a broad spectrum of other adaptive query processing techniques [9]. This line of work adapts at the level of the query plan, reordering operators, switching plans, or deferring plan choices, and is therefore naturally situated inside a database system. Eiger, as a library, adapts within operators instead: it chooses the implementation, configuration, and data representation of each individual operator. The signals also differ: while prior CPU techniques mostly react to information observed as a byproduct of execution, such as cardinalities and selectivities, Eiger proactively computes dedicated statistics over intermediate data (min/max/mean and HyperLogLog++ sketches) and even compresses the data on the fly, both of which are affordable because GPUs perform these computations at close to memory bandwidth. Compared to existing work, Eiger offers multiple implementation variants for the same operator and features a runtime adaptive execution mechanism that selects among them based on lightweight data statistics. Instead of focusing on the few commonly discussed operators, such as joins, this work also studies in-depth the implementation of often overlooked but expensive operations, such as expression evaluation, string processing, and multi-key sorting. Furthermore, we present a more comprehensive performance analysis than previous work, covering a wider range of operators and workloads and characterizing when each implementation variant wins, which helps future query optimizers develop accurate cost models.

7

CONCLUSION

In this work, we present Eiger, a high-performance library for GPUbased data analytics built around runtime workload adaptivity. Eiger realizes this idea through two complementary design principles: it

Bowen Wu, Marko Kabić, Sven Hepkema, Vasilis Mageirakos, Christos Kozyrakis, and Gustavo Alonso Q1

Q2

GPU time (ms)

40

Q4

Q5

60

30

40

20

50

20

10

25 0

baseline

best

0

0

0

baseline

Q12

best

baseline

Q13

best

25

50

100

40

80

30

60

15

20

40

10

10

20

5

5

0

0

0

0

baseline

best

baseline

best

Q15

10

baseline

best

150

best

50 25

0

0

0

best

60

20

Filter Project

baseline

baseline

best Reorder Sort

0

baseline

best

20

20

0

0

baseline

TopK GroupBy

best

ScalarAgg EqJoin

best

best

0

baseline

Distinct Concat

5

0

baseline

best

best

Q22 10.0 7.5

200

10

100

0

0

baseline

baseline

12.5

300

best

0

Q21 400

10

SemiJoin AntiJoin

10

30 20

baseline

15

Q20

20 10

40

10

30

10

best

best

Q19 40

40

40

baseline

Q18

20

20 50

baseline

Q11

50

30

100

20

Q17 60

15

baseline

baseline

20

20

best

0

Q16

Q10

75

40

0

30

Q9 200

60

25

baseline

Q14

Q8 100

20

75 50

5

Q7 80

30

100

10

75

Q6

125

100

GPU time (ms)

Q3

80

15

125

best

baseline

5.0 2.5

best

0.0

baseline

best

CrossJoin

Figure 14: TPC-H SF=100 breakdown (GH200). provides multiple implementation variants and tunable knobs for each operator, including expensive but often overlooked operations such as expression evaluation, string processing, and multi-key sorting, and it profiles intermediate data during query execution with lightweight statistics to select implementations, tune knobs, and compress data on the fly. Our evaluation shows that this design pays off: Eiger outperforms the state-of-the-art cuDF library by up to 1.8× on the complete TPC-H benchmark and up to 6.1× on individual queries. Beyond raw performance, our analysis characterizes how each variant and configuration behaves across workloads and GPU architectures, providing a foundation for cost models in future GPU query optimizers. We hope that Eiger helps build an understanding of what an adaptive GPU-based library for data analytics should look like in the era of composable database systems [18, 37].

ACKNOWLEDGMENTS This work was supported by a grant from the Swiss AI initiative and the Swiss National Supercomputing Centre (CSCS) under project ID sm94 and a donation from NVIDIA Corporation.

REFERENCES [1] Ron Avnur and Joseph M. Hellerstein. 2000. Eddies: continuously adaptive query processing. SIGMOD Rec. 29, 2 (May 2000), 261–272. https://doi.org/10.1145/ 335191.335420 [2] Daniel Bauer, Luis Garces-Erice, Deepak Majeti, Zoltan Arnold Nagy, Sean Rooney, Greg Kimball, Devavret Makkar, Todd Mostak, and Karthikeyan Natarajan. 2026. Accelerating Presto with GPUs. arXiv:2606.24647 [cs.DB] https: //arxiv.org/abs/2606.24647 [3] Nils Boeschen, Tobias Ziegler, and Carsten Binnig. 2024. GOLAP: A GPU-inData-Path Architecture for High-Speed OLAP. Proc. ACM Manag. Data 2, 6, Article 237 (Dec. 2024), 26 pages. https://doi.org/10.1145/3698812 [4] Rani Borkar and Nidhi Chappell. (last accessed) 2026. Microsoft Azure delivers the first large scale cluster with NVIDIA GB300 NVL72 for OpenAI workloads. [Online] Available from: https://azure.microsoft.com/en-us/blog/microsoftazure-delivers-the-first-large-scale-cluster-with-nvidia-gb300-nvl72-foropenai-workloads/. [5] Beatrice Branchini, Pierluigi Negro, Ian Di Dio Lavore, and Marco D. Santambrogio. 2025. Harnessing GPU Acceleration for Exact DNA Sequence Matching via the KMP Algorithm. In 2025 IEEE International Symposium on Circuits and Systems (ISCAS). 1–5. https://doi.org/10.1109/ISCAS56072.2025.11043676 [6] Jiashen Cao, Rathijit Sen, Matteo Interlandi, Joy Arulraj, and Hyesoon Kim. 2023. GPU Database Systems Characterization and Optimization. Proc. VLDB Endow. 17, 3 (nov 2023), 441–454. https://doi.org/10.14778/3632093.3632107

[7] Periklis Chrysogelos, Manos Karpathiotakis, Raja Appuswamy, and Anastasia Ailamaki. 2019. HetExchange: encapsulating heterogeneous CPU-GPU parallelism in JIT compiled engines. Proc. VLDB Endow. 12, 5 (Jan. 2019), 544–556. https://doi.org/10.14778/3303753.3303760 [8] NVIDIA Corporation. (last accessed) 2026. NVIDIA Kicks Off the Next Generation of AI With Rubin — Six New Chips, One Incredible AI Supercomputer. [Online] Available from: https://nvidianews.nvidia.com/news/rubin-platformai-supercomputer. [9] Amol Deshpande, Zachary Ives, and Vijayshankar Raman. 2007. Adaptive query processing. Found. Trends Databases 1, 1 (Jan. 2007), 1–140. [10] Harish Doraiswamy, Vikas Kalagi, Karthik Ramachandra, and Jayant R. Haritsa. 2023. A Case for Graphics-Driven Query Processing. Proc. VLDB Endow. 16, 10 (jun 2023), 2499–2511. https://doi.org/10.14778/3603581.3603590 [11] Mark Harris and RAPIDS Development Team. 2020. [FEA] Make cudf::size_type 64-bit. GitHub Issue #3958, https://github.com/rapidsai/cudf/ issues/3958. RAPIDS cuDF repository. Accessed: 2026-05-19. [12] Dong He, Supun C Nakandala, Dalitso Banda, Rathijit Sen, Karla Saur, Kwanghyun Park, Carlo Curino, Jesús Camacho-Rodríguez, Konstantinos Karanasos, and Matteo Interlandi. 2022. Query processing on tensor computation runtimes. Proc. VLDB Endow. 15, 11 (July 2022), 2811–2825. https://doi.org/10.14778/ 3551793.3551833 [13] Sven Hepkema, Azim Afroozeh, Charlotte Felius, Peter Boncz, and Stefan Manegold. 2025. G-ALP: Rethinking Light-weight Encodings for GPUs. In Proceedings of the 21st International Workshop on Data Management on New Hardware (DaMoN ’25). Association for Computing Machinery, New York, NY, USA, Article 11, 10 pages. https://doi.org/10.1145/3736227.3736242 [14] Stefan Heule, Marc Nunkesser, and Alexander Hall. 2013. HyperLogLog in practice: algorithmic engineering of a state of the art cardinality estimation algorithm. In Proceedings of the 16th International Conference on Extending Database Technology (Genoa, Italy) (EDBT ’13). Association for Computing Machinery, New York, NY, USA, 683–692. https://doi.org/10.1145/2452376.2452456 [15] Kijae Hong, Kyoungmin Kim, Young-Koo Lee, Yang-Sae Moon, Sourav S Bhowmick, and Wook-Shin Han. 2024. Themis: A GPU-Accelerated Relational Query Execution Engine. Proc. VLDB Endow. 18, 2 (Oct. 2024), 426–438. https://doi.org/10.14778/3705829.3705856 [16] Yu-Ching Hu, Yuliang Li, and Hung-Wei Tseng. 2022. TCUDB: Accelerating Database with Tensor Processors. In Proceedings of the 2022 International Conference on Management of Data (Philadelphia, PA, USA) (SIGMOD ’22). Association for Computing Machinery, New York, NY, USA, 1360–1374. https: //doi.org/10.1145/3514221.3517869 [17] Zezhou Huang, Krystian Sakowski, Hans Lehnert, Wei Cui, Carlo Curino, Matteo Interlandi, Marius Dumitru, and Rathijit Sen. 2025. GPU Acceleration of SQL Analytics on Compressed Data. arXiv:2506.10092 [cs.DB] https://arxiv.org/abs/ 2506.10092 [18] Marko Kabić, Shriram Chandran, and Gustavo Alonso. 2025. Maximus: A Modular Accelerated Query Engine for Data Analytics on Heterogeneous Systems. Proc. ACM Manag. Data 3, 3, Article 187 (June 2025), 25 pages. https://doi.org/10.1145/ 3725324 [19] Marko Kabić, Bowen Wu, Jonas Dann, and Gustavo Alonso. 2025. Powerful GPUs or Fast Interconnects: Analyzing Relational Workloads on Modern GPUs. Proc. VLDB Endow. 18, 11 (Sept. 2025), 4350–4363. https://doi.org/10.14778/3749646. 3749698

Eiger: An Efficient Library for GPU-based Data Analytics

[20] Tim Kaldewey, Guy Lohman, Rene Mueller, and Peter Volk. 2012. GPU join processing revisited. In Proceedings of the Eighth International Workshop on Data Management on New Hardware (Scottsdale, Arizona) (DaMoN ’12). Association for Computing Machinery, New York, NY, USA, 55–62. https://doi.org/10.1145/ 2236584.2236592 [21] Tomas Karnagel, René Müller, and Guy M. Lohman. 2015. Optimizing GPU-accelerated Group-By and Aggregation. In ADMS@VLDB. https://api. semanticscholar.org/CorpusID:5017248 [22] Gregory Kimball, Zoltán Arnold Nagy, Devavret Makkar, Daniel Bauer, and Chengcheng Jin. (last accessed) 2026. Accelerating Large-Scale Data Analytics with GPU-Native Velox and NVIDIA cuDF. [Online] Available from: https://developer.nvidia.com/blog/accelerating-large-scale-data-analyticswith-gpu-native-velox-and-nvidia-cudf/. [23] Donald E Knuth, James H Morris, Jr, and Vaughan R Pratt. 1977. Fast pattern matching in strings. SIAM journal on computing 6, 2 (1977), 323–350. [24] Artem Kroviakov, Petr Kurapov, Christoph Anneser, and Jana Giceva. 2024. Heterogeneous Intra-Pipeline Device-Parallel Aggregations. In Proceedings of the 20th International Workshop on Data Management on New Hardware (Santiago, AA, Chile) (DaMoN ’24). Association for Computing Machinery, New York, NY, USA, Article 3, 10 pages. https://doi.org/10.1145/3662010.3663441 [25] Yinan Li, Bailu Ding, Ziyun Wei, Lukas M. Maas, Momin Al-Ghosien, Spyros Blanas, Nicolas Bruno, Carlo Curino, Matteo Interlandi, Craig Peeper, Kaushik Rajan, Surajit Chaudhuri, and Johannes Gehrke. 2025. Scaling GPU-Accelerated Databases Beyond GPU Memory Size. Proc. VLDB Endow. 18, 11 (Sept. 2025), 4518–4531. https://doi.org/10.14778/3749646.3749710 [26] Haotian Liu, Bo Tang, Jiashu Zhang, Yangshen Deng, Xinying Zheng, Qiaomu Shen, Xiao Yan, Dan Zeng, Zunyao Mao, Chaozu Zhang, Zhengxin You, Zhihao Wang, Runzhe Jiang, Fang Wang, Man Lung Yiu, Huan Li, Mingji Han, Qian Li, and Zhenghai Luo. 2022. GHive: A Demonstration of GPU-Accelerated Query Processing in Apache Hive. In Proceedings of the 2022 International Conference on Management of Data (Philadelphia, PA, USA) (SIGMOD ’22). Association for Computing Machinery, New York, NY, USA, 2417–2420. https://doi.org/10.1145/ 3514221.3520166 [27] Jigao Luo, Nils Boeschen, Muhammad El-Hindi, and Carsten Binnig. 2026. PystachIO: Efficient Distributed GPU Query Processing with PyTorch over Fast Networks & Fast Storage. arXiv:2512.02862 [cs.DB] https://arxiv.org/abs/2512.02862 [28] Clemens Lutz, Sebastian Breß, Steffen Zeuch, Tilmann Rabl, and Volker Markl. 2020. Pump Up the Volume: Processing Large Data on GPUs with Fast Interconnects. In Proceedings of the 2020 ACM SIGMOD International Conference on Management of Data (Portland, OR, USA) (SIGMOD ’20). Association for Computing Machinery, New York, NY, USA, 1633–1649. https://doi.org/10.1145/3318464. 3389705 [29] Clemens Lutz, Sebastian Breß, Steffen Zeuch, Tilmann Rabl, and Volker Markl. 2022. Triton Join: Efficiently Scaling to a Large Join State on GPUs with Fast Interconnects. In Proceedings of the 2022 International Conference on Management of Data (Philadelphia, PA, USA) (SIGMOD ’22). Association for Computing Machinery, New York, NY, USA, 1017–1032. https://doi.org/10.1145/3514221.3517911 [30] Vasilis Mageirakos, Joel André, Marko Kabić, Bowen Wu, Yannis Chronis, and Gustavo Alonso. 2026. To GPU or Not to GPU: Vector Search in Relational Engines. arXiv:2605.15957 [cs.DB] https://arxiv.org/abs/2605.15957 [31] Hubert Mohr-Daurat, Xuan Sun, and Holger Pirk. 2023. BOSS - An Architecture for Database Kernel Composition. Proc. VLDB Endow. 17, 4 (Dec. 2023), 877–890. https://doi.org/10.14778/3636218.3636239 [32] Thomas Neumann and Michael J. Freitag. 2020. Umbra: A Disk-Based System with In-Memory Performance. In Conference on Innovative Data Systems Research. https://www.cidrdb.org/cidr2020/papers/p29-neumann-cidr20.pdf [33] NVIDIA. (last accessed) 2026. cuCollections. [Online] Available from: https: //github.com/NVIDIA/cuCollections/tree/dev. [34] Oracle. (last accessed) 2026. SQL Tuning Guide (Chapter 10: Optimizer Statistics Concepts). [Online] Available from: https://docs.oracle.com/en/database/oracle/ oracle-database/26/tgsql/optimizer-statistics-concepts.html. [35] Tsuyoshi Ozawa and Kazuo Goda. 2026. Data Path Fusion in GPU for Analytical Query Processing. arXiv:2605.10511 [cs.DB] https://arxiv.org/abs/2605.10511 [36] Johns Paul, Bingsheng He, Shengliang Lu, and Chiew Tong Lau. 2019. Revisiting Hash Join on Graphics Processors: A Decade Later. In 2019 IEEE 35th International Conference on Data Engineering Workshops (ICDEW). 294–299. https://doi.org/ 10.1109/ICDEW.2019.00008 [37] Pedro Pedreira, Orri Erling, Konstantinos Karanasos, Scott Schneider, Wes McKinney, Satya R Valluri, Mohamed Zait, and Jacques Nadeau. 2023. The Composable Data Management System Manifesto. Proc. VLDB Endow. 16, 10 (June 2023), 2679–2685. https://doi.org/10.14778/3603581.3603604 [38] PostgreSQL. (last accessed) 2026. PostgreSQL 18 Documentation (Chapter 14.2: Statistics Used by the Planner). [Online] Available from: https://www.postgresql. org/docs/current/planner-stats.html. [39] Raghu Ramakrishnan and Johannes Gehrke. 2003. Database management systems (3 ed.). McGraw-Hill New York. [40] RAPIDS. (last accessed) 2026. cuDF: A GPU DataFrame Library. [Online] Available from: https://github.com/rapidsai/cudf.

[41] RAPIDS. (last accessed) 2026. GQE: GPU Query Engine. https://github.com/ rapidsai/gqe. [42] Viktor Rosenfeld, Sebastian Breß, Steffen Zeuch, Tilmann Rabl, and Volker Markl. 2019. Performance Analysis and Automatic Tuning of Hash Aggregation on GPUs. In Proceedings of the 15th International Workshop on Data Management on New Hardware (Amsterdam, Netherlands) (DaMoN’19). Association for Computing Machinery, New York, NY, USA, Article 8, 11 pages. https://doi.org/10.1145/ 3329785.3329922 [43] Ran Rui and Yi-Cheng Tu. 2017. Fast Equi-Join Algorithms on GPUs: Design and Implementation. In Proceedings of the 29th International Conference on Scientific and Statistical Database Management (Chicago, IL, USA) (SSDBM ’17). Association for Computing Machinery, New York, NY, USA, Article 17, 12 pages. https: //doi.org/10.1145/3085504.3085521 [44] Ravi Sethi and J. D. Ullman. 1970. The Generation of Optimal Code for Arithmetic Expressions. J. ACM 17, 4 (Oct. 1970), 715–728. https://doi.org/10.1145/321607. 321620 [45] Anil Shanbhag, Samuel Madden, and Xiangyao Yu. 2020. A Study of the Fundamental Performance Characteristics of GPUs and CPUs for Database Analytics. In Proceedings of the 2020 ACM SIGMOD International Conference on Management of Data (Portland, OR, USA) (SIGMOD ’20). Association for Computing Machinery, New York, NY, USA, 1617–1632. https://doi.org/10.1145/3318464.3380595 [46] Anil Shanbhag, Bobbi W. Yogatama, Xiangyao Yu, and Samuel Madden. 2022. Tile-based Lightweight Integer Compression in GPU. In Proceedings of the 2022 International Conference on Management of Data (Philadelphia, PA, USA) (SIGMOD ’22). Association for Computing Machinery, New York, NY, USA, 1390–1403. https://doi.org/10.1145/3514221.3526132 [47] Xuri Shi, Kai Zhang, X. Sean Wang, Xiaodong Zhang, and Rubao Lee. 2026. RayDB: Building Databases with Ray Tracing Cores. Proc. VLDB Endow. 19, 1 (Jan. 2026), 43–55. https://doi.org/10.14778/3772181.3772185 [48] Panagiotis Sioulas, Periklis Chrysogelos, Manos Karpathiotakis, Raja Appuswamy, and Anastasia Ailamaki. 2019. Hardware-Conscious Hash-Joins on GPUs. In 2019 IEEE 35th International Conference on Data Engineering (ICDE). 698–709. https://doi.org/10.1109/ICDE.2019.00068 [49] Evangelia A. Sitaridi and Kenneth A. Ross. 2016. GPU-accelerated string matching for database applications. The VLDB Journal 25, 5 (Oct. 2016), 719–740. https: //doi.org/10.1007/s00778-015-0409-y [50] Wenbo Sun, Asterios Katsifodimos, and Rihan Hai. 2023. An Empirical Performance Comparison between Matrix Multiplication Join and Hash Join on GPUs. In 2023 IEEE 39th International Conference on Data Engineering Workshops (ICDEW). 184–190. https://doi.org/10.1109/ICDEW58674.2023.00034 [51] Diego G. Tomé, Tim Gubner, Mark Raasveldt, Eyal Rozenberg, and Peter A. Boncz. 2018. Optimizing Group-By and Aggregation using GPU-CPU Co-Processing. In ADMS@VLDB. https://api.semanticscholar.org/CorpusID:52895287 [52] Transaction Processing Performance Council. 2022. TPC Benchmark H (Decision Support) Standard Specification. Technical Report. Transaction Processing Performance Council (TPC). https://www.tpc.org/TPC_Documents_Current_ Versions/pdf/TPC-H_v3.0.1.pdf Version 3.0.1. [53] David Wendt and Gregory Kimball. (last accessed) 2026. Mastering String Transformations in RAPIDS libcudf. [Online] Available from: https://developer.nvidia. com/blog/mastering-string-transformations-in-rapids-libcudf/. [54] Bowen Wu, Wei Cui, Carlo Curino, Matteo Interlandi, and Rathijit Sen. 2025. Terabyte-Scale Analytics in the Blink of an Eye. Proc. VLDB Endow. 19, 2 (Oct. 2025), 141–155. https://doi.org/10.14778/3773749.3773754 [55] Bowen Wu, Dimitrios Koutsoukos, and Gustavo Alonso. 2025. Efficiently Processing Joins and Grouped Aggregations on GPUs. Proc. ACM Manag. Data 3, 1, Article 39 (Feb. 2025), 27 pages. https://doi.org/10.1145/3709689 [56] Bobbi Yogatama, Weiwei Gong, and Xiangyao Yu. 2025. Scaling your Hybrid CPUGPU DBMS to Multiple GPUs. Proc. VLDB Endow. 17, 13 (Feb. 2025), 4709–4722. https://doi.org/10.14778/3704965.3704977 [57] Bobbi Yogatama, Yifei Yang, Kevin Kristensen, Devesh Sarda, Abigale Kim, Adrian Cockcroft, Yu Teng, Joshua Patterson, Gregory Kimball, Wes McKinney, Weiwei Gong, and Xiangyao Yu. 2025. Rethinking Analytical Processing in the GPU Era. arXiv:2508.04701 [cs.DB] https://arxiv.org/abs/2508.04701 [58] Bobbi W. Yogatama, Weiwei Gong, and Xiangyao Yu. 2022. Orchestrating data placement and query execution in heterogeneous CPU-GPU DBMS. Proc. VLDB Endow. 15, 11 (July 2022), 2491–2503. https://doi.org/10.14778/3551793.3551809 [59] Yichao Yuan, Advait Iyer, Lin Ma, and Nishil Talati. 2025. Vortex: Overcoming Memory Capacity Limitations in GPU-Accelerated Large-Scale Data Analytics. Proc. VLDB Endow. 18, 4 (May 2025), 1250–1263. https://doi.org/10.14778/3717755. 3717780 [60] Haitao Zhang, Ran Pang, Yuanyuan Zhu, Hao Zhang, Congli Gao, Ming Zhong, Jiawei Jiang, Tieyun Qian, and Jeffrey Xu Yu. 2025. TQEx: Tensor-based Query Engine Enhanced by Bridging the Gap. Proc. ACM Manag. Data 3, 6, Article 370 (Dec. 2025), 27 pages. https://doi.org/10.1145/3769835

Related documents

• TrajectoryDB: A New Database for Agent Trajectories#673620
• Preserving contextual information in cultural heritage metadata through multidimensional knowledge graphs#673621
• X-DigCheck: Co-Evolving Application Profiles and Knowledge Graphs, Demonstrated on the RTI Documentation of Rupe Magna#673622
• Open Tabular Insight Extraction: Where Do We Stand, and Where Should We Go?#673623
• Protocol effects on feature-based hardware-Trojan detection across Trust-Hub families#673624