Nowadays, query processing in column-store database systems is highly tuned to the underlying (co-)processors. This approach works very well from a performance perspective, but has several shortcomings from a conceptual perspective. For example, this tuning introduces high implementation as well as maintenance cost and one implementation cannot be ported to other (co-)processors. To overcome that, we developed a column-store speci c abstraction layer for hardwaredriven vectorization based on the Single Instruction Multiple Data (SIMD) parallel paradigm. Thus, we are able to implement vectorized query operators in a hardware-oblivious manner, which can be specialized to di erent SIMD instruction set extensions of modern x86-processors. To soften the limitation to x86-processors, we describe our vision to integrate GPUs in our abstraction layer by interpreting GPUs as virtual vector engines in this paper. Moreover, we present some initial evaluation results to determine a reasonable virtual vector size. We conclude the paper with an outlook on our ongoing research in that direction.
Analytical database queries typically access a small
number of columns, but a high number of rows and are, thus most
e ciently answered by column-store database systems [
For this reason, vectorization based on the Single Instruction Multiple Data (SIMD) parallel paradigm has established itself as a core query optimization technique in column32nd GI-Workshop on Foundations of Databases (Grundlagen von Datenbanken), September 01-03, 2021, Munich, Germany.
Copyright © 2021 for this paper by its authors. Use permitted under
Creative Commons License Attribution 4.0 International (CC BY 4.0).
store systems [
Besides vectorization, hardware also shifts from
homogeneous x86-processors towards heterogeneous systems with
di erent computing units (CU) [
Our Contribution and Outline. In this paper, we describe our vision and present some initial steps to integrate GPUs in our abstraction layer by interpreting GPUs as virtual vector engines. Generally, GPUs use a Single Instruction Multiple Thread (SIMT) execution model which can be also interpreted as SIMD combined with multi-threading. Based on that, GPUs seem like a perfect match for our TVL to soften the limitation to x86-processors. Thus, our contributions are the following in this paper: 1. We start with an introduction in our SIMD abstraction layer TVL as well as with an architectural description of NVIDIA GPUs in Section 2. 2. Then, we present our general idea of SIMDization of GPUs in Section 3. To interpret a GPU as virtual vector engine, we have to determine a reasonable virtual vector size, which can be most e ciently processed in parallel. The vector size is important because it is an (a) Template Vector Library (TVL) (b) NVIDIA GPU Architecture essential part of SIMD and an integral part of our TVL. For that, we present some TVL-oriented experiments.
Based on the results, we derive a virtual vector size. 3. In Section 4, we summarize our lessons learned and describe our ongoing activities in that direction. 4. We close the paper with related work in Section 5 and a short summary in Section 6. 2.
In this section, we introduce our Template Vector Library as SIMD abstraction layer for column-stores. Moreover, we brie y describe the GPU architecture and execution model. 2.1
Vectorization is a state-of-the-art query optimization
technique in in-memory column-stores, because all recent
x86processors o er powerful SIMD extensions [
SIMD provides data parallelism by executing a single
instruction on multiple data elements simultaneously [
However, these SIMD extensions are increasingly diverse in
terms of (i) the number of available vector instructions, (ii)
the vector length, and (iii) the granularity of the bit-level
parallelism, i.e., on which data widths the vector
instructions are executable [
Our abstraction approach follows a separation of concerns concept as shown in Fig. 1(a). On the one hand, it o ers hardware-oblivious but column-store speci c primitives, which are similar to intrinsics. The primitives are derived from state-of-the-art vectorized columnar query operators.
We organized these primitives in seven self-descriptive
classes like load/store (L/S) or an arithmetic class for
a better organization including a uni ed interface per
class [
From an implementation perspective, our abstraction concept is realized as a header-only library, where the hardwareoblivious primitives abstract from SIMD intrinsics. These primitives are generic functions representing a uni ed interface for all SIMD architectures. In addition to the primitives, we introduced generic data types: base t: The base type can be any scalar type. vector t: The vector type contains one or more values of
the same base type. mask t: A mask is a scalar value, which is large enough to
store one bit for each element in a vector.
Using the provided primitives and the data types, we can implement columnar query operators in a hardwareoblivious way. For the hardware-conscious mapping, we use template metaprogramming requiring hardware-conscious implementations for all primitives and for all underlying SIMD extensions. This function template specialization has to be implemented, whereby the implementation depends on the available functionality of the SIMD extension. In the best case, we can directly map a TVL primitive to a SIMD intrinsic. However, if the necessary SIMD intrinsic is not available, we are able to implement an e cient workaround in a hardware-conscious way. This implementation is independent of any query operator and must be done only once for a speci c SIMD extension.
Figure 2 illustrates how a hardware-oblivious vectorized operator using our TVL looks like and how it can be called. We show a simple aggregation (summation) operator consisting of four TVL primitives: set1: lls a vector register with a given value load: loads multiple consecutive data elements into a vector
register add: executes an element-wise addition on data elements in
two vector registers hadd: executes an horizontal addition on data elements in
one vector register The aggregation operator assumes that the number of data elements is a multiple of the number of elements per vector and works as follows: One vector register called resVec is
lled with zeros. Afterwards, the operator iterates over the input column in and loads a number of consecutive data elements into a second vector register called dataVec. Then, the data elements in both vector registers are added elementwise and the result is stored in resVec. When all elements of the input column have been processed, the horizontal aggregation hadd is carried out to determine the nal sum result.
To specialize this hardware-oblivious operator implementation during query compile-time, we use three template parameters called processingStyle (ps): (i) the vector extension (e.g., SSE, AVX, NEON, or scalar), (ii) the vector size in bit, and (iii) the base data type with bit granularity (e.g., int8, int64, oat). The de nition of ps is shown in Figure 2, which is used to call the aggregation-operator. // Calling the operator. using ps = t v l : : a v x 2 < t v l : : v 2 5 6 < u i n t 6 4 _ t > > ; // for example size_t count = 1024; uint64_t * array = g e n e r a t e _ d a t a ( count ) ; uint64_t sum = agg<ps> ( array , elemCount ) ;
Graphics Processing Units (GPUs) are increasingly used
for large-scale query processing in database systems [
Figure 1(b) depicts a simpli ed architecture of an NVIDIA GPU. Generally, a modern GPU consists of (i) a large global main memory with a memory bandwidth of up to 1.2 TB/s and (ii) a number of compute units called Streaming Multiprocessors (SMs). Each SM has a number of simplistic cores, a xed set of registers, and shared memory. This shared memory serves as scratchpad and can be accessed by all cores in the SM. Moreover, the GPU has an on-chip L2 cache, which is shared across all SMs and optionally, each SM may have a local L1 cache. The number of SMs, and cores per SM, the size of global memory, the size of the L2 cache, etc. varies across GPU products.
The execution model of GPUs is called Single
Instruction Multiple Threads (SIMT), whereas SIMT is very similar
to SIMD. While multiple data are processed by a single
instruction in SIMD, multiple threads are processed by a single
instruction in lock-step. That means, each thread in SIMT
executes the same instruction, but on di erent data. While
a thread switch is very costly on CPUs, GPUs can handle
thread switching with more ease. GPUs feature a fast thread
switching based on groups of 32 threads called warps. The
warp scheduler issues instructions to warps available on an
SM beyond the number of physical cores to hide latency
[
For the GPU implementation, general purpose parallel
programming models such as CUDA [
Our overall vision is to fully support GPUs in our TVL as shown in Figure 1(a). For that, we have to create a GPU-speci c hardware-conscious implementation of the hardware-oblivious interface. To achieve that, we interpret GPUs as virtual vector engines and call this interpretation SIMDization. For an optimal SIMD processing, we have to nd a reasonable vector size. This vector size is virtual because we want to interpret regular arrays as virtual vector registers. For example, on Intel x86-processors, the di erent SIMD instruction set extensions either have a vector length of 128-, 256- or 512-bit. If we assume 64-bit per data element, we can simultaneously process 2 elements in a vector register of the size 128-bit. The wider the vector registers, the more data can be processed in parallel. We are not aware of any work that has ever determined a vector size for GPUs. 3.1
To determine a reasonable virtual vector size for our SIMDization approach for GPUs, several experiments were executed. A rst set of experiments were conducted on an NVIDIA RTX Quadro 8000 GPU. This GPU has a global main memory size of 48 GiB, a memory bandwidth of 672 GB/s, an L2 cache with a size of 6 MiB, and 72 SMs. Each SM has 64 cores resulting in a total of 4,608 cores.
For the hardware-conscious implementation, we can distinguish three di erent main groups of hardware-oblivious SIMD primitives across all TVL classes: (i) load/store primitives, (ii) element-wise primitives, and (iii) horizontal primitives. Element-wise primitives are characterized by the feature that they do not introduce dependencies between the elements of the same vector register, e.g., element-wise arithmetic, comparisons, or boolean logic. In contrast to that, horizontal primitives do not treat the elements of a vector independently. An example is the horizontal addition. As shown in our example sum-aggregation operator (cf. Figure 2), di erent primitives from all groups are used to implement a vectorized query operator. In this operator, we use an element-wise and a horizontal add primitive.
In our rst experiments, we evaluated the element-wise and horizontal addition on the GPU. For that, we implemented simple vectorized CUDA kernels and executed these kernels using di erent vector sizes. For the horizontal addition, we use the function cub::DeviceReduce from the CUDA SDK 11.2 as most e cient implementation as foundation. Since our GPU has 4,608 cores and usually more threads than cores are used for good performance, we investigated vector sizes in the range from 256 KiB to 1 GiB. In terms of number of elements, we evaluated the range from 213; ::; 227 number of 64-bit data elements. For each vector size, di erent CUDA con gurations with blocks and threads per block are possible and we tested a large number of di erent con gurations in a systematical way.
Element-wise Addition: This kernel was performed on two input columns A and B. Each column was a sequence of unsigned randomly generated 64-bit integer values with a size of 1 GiB. The result was written back into column B. Within a loop, our vectorized kernel is called until the whole columns have been processed. Figure 3 shows the best achieved throughput over all con gurations for each vector size. As we can see, small vector sizes negatively a ect the performance. For a vector size of 256 KiB, a throughput of 248.87 GiB/s is achieved. The best performance is gained by using a vector size of 1 GiB. This results in a throughput of 530.99 GiB/s. A vector size of 8 MiB is 3.8% slower than a vector size of 1 GiB. We conclude, higher vector sizes lead to higher throughput and in the best case, the vector size corresponds to column size.
Horizontal Additon: This kernel adds all elements in a vector register together and the kernel is based on the cub::DevideReduce function being shipped as part of the CUDA SDK. For our evaluation, we generated di erent columns where the column size corresponds to the vector size. Again, the data elements are 64-bit unsigned integer values. Figure 4 shows the best achieved throughput over all congurations for each vector size. Similar to the element-wise addition, higher vector sizes lead to higher throughput.
Aggregation Operator: Based on these evaluations, we could draw the conclusion that the best vector size should be the column size. To validate our hypothesis, we evaluated the sum-aggregation operator consisting of both primitives as next. For this evaluation, we generated a single column with a size of 1 GiB and varied the vector size from 256 KiB to 1 GiB. As illustrated in Figure 5, we obtain a completely di erent result. The throughput increases up to a vector size of 2 MiB. Then, the throughput decreases and stabilizes at a low level. The main di erence to our previous experiments is that a single vector register or array is now the main driver of the processing. As shown in Figure 2, one vector register resVec is lled with zeros at the beginning of the aggregation operators. Afterwards, we repeatedly load into a second vector dataV ec of the column and conduct an element-wise add between both. The result is stored in resVec vector. This resVec vector is frequently accessed, and thus, should be kept in cache.
To evaluate the cache- tting in more detail, we slightly modi ed our experiment. We generated a new input column of size 1,5 GiB containing 64-bit unsigned integers and varied the vector size according to the on-chip L2 cache size: 1=4, 1=3, 1=2, 1, 2 times of the cache size. Figure 6 depicts the resulting throughputs. As we can see, we obtain the best performance when our vector size is a third of the L2 cache size. Larger vector sizes lead to a lower throughput. 3.2
As shown above, we are able to determine a reasonable virtual vector size providing the best performance for our vectorized aggregation on the GPU. In this section, we present evaluation results comparing the vectorized aggregation with the native CUDA aggregation. For that, we generated various columns|sequences of 64-bit integer values|with increasing sizes from 2 MiB to 8 GiB. For the vectorized aggregation, we applied the best performing vector size of 2 MiB in all experiments. For the native CUDA aggregation, we used the function cub::DevideReduce from the CUDA SDK as already done for our horizontal aggregation.
The results are depicted in Figure 7. The throughput of the native CUDA aggregation is slightly higher as for our vectorized aggregation. This is especially true for large columns, which is not particularly surprising. Nevertheless, the result is promising to specialize our hardware-oblivious vectorized query operators to GPUs and to obtain a reasonably good performance. We hope to increase the throughput of our vectorized execution with additional GPU-speci c optimization techniques as discussed in Section 4. 3.3
We also executed all our experiments on a second NVIDIA GPU namely an NVIDIA GTX 1070 Ti. This GPU provides a global main memory size of 8 GiB, a memory bandwidth of 256 GB/s, an L2 cache of 2 MB size, and 19 SMs. Each SM has 128 cores resulting in a total of 2,432 cores. In general, we observed a behavior similar as with the NVIDIA Quadro RTX 8000. In particular, we executed the cache- tting experiment to determine the best virtual vector size for the aggregation operator. In contrast to the RTX 8000 GPU, we obtain the best performance when our vector size is a fourth of the L2 cache (512 KiB) size as shown in Figure 8.
Based on the previous section, we conclude that SIMD processing using a virtual vector model is generally possible on a GPU. Choosing the right vector size and con guration is critical to achieve a good performance. Sub-optimal congurations reduce the performance by more than one order of magnitude. Overall, we obtain good performance which is slower than native CUDA implementations. For the aggregation operator using the ideal vector size, our approach is 37.4% slower in the worst case for 8 GiB data size, and 10.8% slower in the best case for 8 MiB data size than a native CUDA approach. By validating our approach on a di erent GPU, we have also shown that our conclusions are more generally applicable beyond a single GPU model. Moreover, our hardware-oblivious query operators can be specialized to SIMD extensions as well as to GPUs in a uni ed way.
To summarize, our results are promising and our ongoing research in that direction will focus on the following aspects: (1) Virtual Vector Size: To extend our work, we are looking forward to evaluate other query operators. We seek to explore if optimal vector sizes and con gurations depend on the query operators.
(2) Implementation: Completing the GPU TVL requires to implement all primitives. To optimize the hardwareconscious implementations, we want to explore the usage of shared memory, registers, and persistent caching.
(3) Optimization: Besides an optimal hardwareconscious implementation for GPUs, we want to investigate more mapping strategies for a broader optimization. Currently, the TVL has a 1:1 mapping of hardwareoblivious primitives to hardware-conscious implementations. As shown in our experiments, the CUDA native aggregation outperforms the vectorized approach. To improve the performance of our vectorized approach, an idiom-based mapping strategy could be helpful. By replacing a vectorized computation by a semantically equivalent but more performant code at query compile-time, a speedup can possibly be achieved. For that, it is necessary to identify often used idioms that lose performance by vectorization and replace them by idiomatic implementations. 5.
To address the portability of code across heterogeneous
computing units, OpenCL [
In a recent work, Shanbhag et al. [
In this paper, we evaluated the integration of GPUs in our SIMD abstraction layer TVL by interpreting GPUs as virtual vector engines. By conducting a number of experiments, we have shown that our approach is promising. While the observed throughput does not outperform CUDA native implementations, it allows developers to use CUDA-based fast GPU primitives without requiring knowledge of GPU implementations. Our vectorized approach still achieves reasonable performance that is not an order of magnitude slower than native CUDA implementations. However, tuning the primitives and operators by choosing the right virtual vector size and con guration is critical for achieving good performance.
This work was partly funded by the German Research Foundation (DFG) within the RTG 1907 (RoSI).