Latest Posts (20 found)

Warnock: Harnessing GPU Geometry Amplification for Vector Graphics

Warnock: Harnessing GPU Geometry Amplification for Vector Graphics Mark J. Kilgard and Brian Barran HPG'26 We live in a strange world where the primary use case for GPUs is LLMs, the secondary use case is 3D graphics, and one has to perform contortions to accelerate 2D graphics with a GPU. This paper presents a technique to use modern GPU hardware features to offload a portion of 2D graphics rendering that is typically performed on a CPU. There are three hard problems associated with high-quality vector graphics: winding rules, curves, and antialiasing. Say you have given a polygon like so: and you would like to know if a particular point is contained within the polygon. The orthodox definition of the odd-even winding rule is: Draw a ray starting at and pointing in any direction. Count the number of polygon edges that the ray intersects. If the number of intersections is odd, then is inside of the polygon. For example, the ray in the diagram below intersects with the polygon 3 times, and so the ray origin is considered to be inside of the polygon. The flexibility to cast rays in any direction allows for a variety of vector graphics rasterization algorithms. CPU-based iterative algorithms typically cast rays horizontally. The idea is that the algorithms steps vertically one scan line at a time and incrementally tracks the set of edges which intersect the scan line. Data parallel approaches on GPUs can cast rays in a different direction. Here is a common way to do it (also used by this paper): Pick an arbitrary point (somewhere near the center of the polygon) and draw a triangle fan emanating from that point. A triangle fan is like those collapsible hand-held fans with ridges. Here is a triangle fan for the running example, with each triangle in a unique (partially transparent) color: Here is the key: to determine if is covered by the triangle, we will logically cast a ray with an origin of , and a direction that points directly away from the center of the triangle fan, like so: Putting it all together, the way to do this on a GPU is to use an intermediate 2D buffer (i.e., a stencil buffer) which holds an integer for every pixel. This buffer starts out with a value of 0 at each pixel. Then instruct the GPU to rasterize each triangle in the triangle fan. For each pixel that is covered by a triangle, increment the associated integer. When all triangles in a fan have been rendered, then a pixel is considered to be covered if the stencil buffer contains an odd value at that pixel. GPUs are extremely fast at rasterizing triangles (i.e., determining which pixels are covered by a triangle), so this technique works well. Most interesting vector graphics contains curves in addition to lines. A curve can be approximated by a series of lines, but this approximation is necessarily resolution dependent. If you zoom in on a curve, then more line segments are required to accurately represent it. The Warnock paper proposes to use geometry amplification hardware to decompose curves into sequences of line segments in real time. Somewhat older GPUs support a feature called tessellation , while newer GPUs support mesh shading . The idea is that the CPU can generate an abstract representation of an object, along with a program that that runs on the GPU each frame to convert from that representation to a variable number of triangles. The variable number part is the key, this GPU-side program can determine how finely to approximate a curve based on runtime parameters like how large the curve will be on screen. This paper uses polar filling on the GPU. The idea is that each triangle in the triangle fan is separated a maximum angle. Fig. 5. illustrates polar filling for a specific shape. Note how some areas of the light green shape are represented with few lines, while curvy areas are represented with many lines. Source: https://dl.acm.org/doi/10.1145/3820012 Section 3 of the paper goes into the nitty gritty of the math behind polar filling. In the domain of vector graphics antialiasing means to not make a binary decision about each pixel: “is the pixel covered by the shape or not?”. Rather some pixels on the edge of a shape can be considered to be partially covered. This is a hard problem for 2D vector graphics because the entire shape must be considered. Antialiasing cannot be resolved on a per-triangle basis. For example, if a given pixel is 50% covered by one triangle and 50% covered by another, should that pixel be considered to be 100% covered? It depends on the relative overlap of the two triangles. This paper doesn’t present any advances in this space, and instead relies on hardware antialiasing features to evaluate the stencil buffer at multiple locations per pixel. Table two contains performance results against industrial strength 2D graphics libraries, and against a prior NVIDIA implementation: Source: https://dl.acm.org/doi/10.1145/3820012 Warnock does well against industrial libraries… but Direct2D doesn’t get beat too badly. I’m particularly proud of this as I spent a few years of my life optimizing the Direct2D implementation. Modern GPUs have ray tracing support, which can be hijacked for some interesting purposes. Since 2D vector graphics rasterization is fundamentally defined in terms of rays, I wonder if ray tracing hardware could be another approach to 2D vector graphics rendering. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views

Honey, I Shrunk the Headers With Flow.ZIP

Honey, I Shrunk the Headers With Flow.ZIP Authors SIGCOMM'26 This paper presents a crazy brilliant technique for reducing the overheads associated with network packets. This was my first introduction to MPLS, which is deceptively simple. Fig. 1 shows the cost of header bytes in via “simulation of a FatTree topology and representative production storage workload”. The y-axis shows average flow completion time (FCT), which is the time between the first and last packets of a flow. The x-axis shows how loaded the network is. Source: https://dl.acm.org/doi/10.1145/3789240.3829104 The summary is that there are workloads where the cost of transmitting packet headers is non-trivial. Section 2 of the paper uses a mathematical model as further evidence. This problem is exacerbated by network virtualization techniques that add additional levels of encapsulation. If packet headers and encapsulation are the problem, how can it be solved? Encapsulate network packets by slapping on another header (an MPLS header to be precise)! Fig. 2 shows the layout of a !32-bit! MPLS header: identifies a path through the network. represents a traffic class for prioritization. indicates if this is the last MPLS header in a stack of MPLS headers represents time to live (like the IP TTL field). The paper proposes using MPLS for packet compression. Compression occurs near the sender, decompression occurs near the receiver. The idea is that for a given flow (i.e., connection) most bits in the IP/TCP/UDP header are identical for all packets in the flow. For example, say the sender is a cloud VM. The VM could send a normal packet (with bloated headers) to a SmartNIC. The NIC could then remove all header bits which are constant for the flow (e.g., addresses, ports, protocol) and replace them with a single 32-bit MPLS header. As the packet travels through the network to its destination, all switches/routers along the way would use the MPLS header for routing. Once the packet arrives at the destination SmartNIC, the NIC would remove the MPLS header and restore the original headers. Fig. 5 illustrates the compression scheme: Source: https://dl.acm.org/doi/10.1145/3789240.3829104 The original L3/L4 header (green) is replaced with a 32-bit MPLS header (blue) and the compression residue (orange). The compression residue holds bits which are not constant for the flow (e.g., payload length). This scheme relies on a lot of network hardware (e.g., NICs, switches) being in on the trick. The paper proposes a controller which configures the MPLS mappings in all of the network hardware. A typical switch supports an MPLS table with 16K entries. The controller gathers statistics about all flows and chooses 16K flows with the most bang for the buck. The remaining long tail of flows does not use MPLS compression. Fig. 8 shows flow completion times for 4 different workloads. Flow.ZIP is this paper. ROHC is an alternative technique called Robust Header Compression . Dangling Pointers This technique seems to be at odds with packet spraying, which is a hot topic these days. I wonder if there is a way to support both header compression and packet spraying. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Source: https://dl.acm.org/doi/10.1145/3789240.3829104 The summary is that there are workloads where the cost of transmitting packet headers is non-trivial. Section 2 of the paper uses a mathematical model as further evidence. This problem is exacerbated by network virtualization techniques that add additional levels of encapsulation. MPLS to the Rescue If packet headers and encapsulation are the problem, how can it be solved? Encapsulate network packets by slapping on another header (an MPLS header to be precise)! Fig. 2 shows the layout of a !32-bit! MPLS header: identifies a path through the network. represents a traffic class for prioritization. indicates if this is the last MPLS header in a stack of MPLS headers represents time to live (like the IP TTL field).

0 views
Dangling Pointers 1 months ago

Presto: A Match-Action TCP Stack for the Terabit Era

Presto: A Match-Action TCP Stack for the Terabit Era Rajath Shashidhara, Antoine Kaufmann, and Simon Peter SIGCOMM'26 This paper presents Presto, a Goldilocks implementation of the TCP protocol. It is efficient and yet does not require fixed-function TCP-specific networking hardware. The paper is a tour-de-force in the way it isolates the specific problems that make TCP processing hard to pipeline, and describing clever solutions to these problems. The Reconfigurable Match-Action Table architecture one specific flavor of programmable network accelerator. Here are two previous paper summaries that reference the RMT architecture. At its core, the RMT architecture is a feed-forward pipeline through which network packets flow. Each pipeline stage has a content addressable memory, and a limited amount of compute. The hard part about mapping an application to the RMT architecture is that there is very limited communication between pipeline stages. Network packets flow forward through the pipeline. The one escape hatch is the pipeline can decide to recirculate a packet, which can cause information to be sent from the tail of the pipeline to the front. This paper which, implements a key-value store with RMT leans heavily on this recirculation. Mapping the various steps in TCP protocol handling onto the RMT architecture requires distributing the state associated with a connection across the RMT pipeline. The size of per-connection state at each pipeline stage is fixed. The hardest TCP feature to map onto RMT is segment reassembly. Segment reassembly is the task of tracking and handling received segments (i.e., packets), which may arrive out of order. The receive side of a TCP connection must track the start and end of a window of packets that may be accepted. For example, if the packet with sequence number 4 has been processed, and the window size is 10, then the sender is free to send packets [5, 6, …, 15]. The paper describes three segment reassembly designs, I’ll illustrate one (OOO-1) here. Fig. 4 illustrates a continuous stream of packets with monotonically increasing sequence numbers. is the lowest sequence number of packets that have not yet been received (i.e., the start of the TCP window). defines the end of the TCP window. and define a contiguous set of packets that have been received and are in the TCP window. Note that this design happily accepts these packets. Source: https://dl.acm.org/doi/10.1145/3789240.3829111 Fig. 3 illustrates the 4 pipeline stages that implement TCP receive window tracking. Note that each of the 4 state variables described above is tracked in a different pipeline stage. For example, say that and , and . This means that the next expected sequence number is 4, and no packets in the TCP window have arrived. Say that packet 6 arrives next. Presto will accept this packet and set and . If packet 5 arrives next, then will be set to 5. Finally, when packet 4 arrives, will be set to 4. At this moment (ooo-head-1 is equal to next-seq), the packets 4, 5, and 6 can be sent down the pipeline. This is accomplished with recirculation: a dummy packet is injected into the pipeline which flows through all stages and updates state variables as expected. Source: https://dl.acm.org/doi/10.1145/3789240.3829111 Results Fig. 9 shows throughput vs latency curves for Presto and TAS (a software TCP stack based on kernel bypass): Source: https://dl.acm.org/doi/10.1145/3789240.3829111 Fig. 10 shows power consumption: Source: https://dl.acm.org/doi/10.1145/3789240.3829111 Dangling Pointers It is a shame that Intel has discontinued the Tofino chips. The literature shows that the RMT architecture is flexible enough to efficiently implement a wide range of applications (e.g., key-value store, TCP protocol acceleration). Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 1 months ago

Understanding Host Network Stack Latency

Understanding Host Network Stack Latency Tianyu Zuo, Jaehyun Hwang, Ao Tang, Rachit Agarwal, and Qizhe Cai SIGCOMM'26 Think of this paper like the infamous: Latency Numbers Every Programmer Should Know , but for the Linux networking stack in particular. The authors measure end-to-end P99.9 tail latency for TCP communication. In other words: the time it takes to both send and receive a single TCP packet, in the slowest .1% cases. To do this, the authors instrumented the Linux kernel to break down where time is spent. The latency for a vanilla case is 19 microseconds . Of that, most of the time is spent in: Interrupt handling IP protocol processing TCP protocol processing CPU scheduling Switching from interrupts to busy polling reduces latency to 14 microseconds . Latency jumps to 1139 microseconds (!) when 48 threads run on a single core (each using different connection). The authors find a few reasons for this: The first is a bug in the Linux CPU scheduler. When a network packet is received some of the work occurs in a softIRQ (similar to a DPC on Windows). Say a packet arrives for thread A while thread B is running. The kernel will steal the CPU from B to execute the softIRQ work on behalf of A . In theory, this design is OK, however the Linux scheduler charges this time to thread B! This issue has been known for some time , and Linux includes a fix (which is disabled by default). The authors of this paper found that there is a bug in the fix that causes it to not achieve the desired result. The authors implemented a patch, which resolves the issue. Enabling this off-by-default feature and applying the patch reduces P99.9 latency from 1139 microseconds to 415 microseconds . After dipping their toes into the water of modifying the Linux scheduler, the authors tried a more aggressive change: scheduling threads based on the number of packets sent/received rather than elapsed time. This networking-specific scheduling policy reduced multi-threaded P99.9 latency to 276 microseconds . NICs support interrupt moderation , as a way to trade latency for host CPU time spent handling interrupts. The idea is that the NIC sends an interrupt for a batch of received packets, rather than for a single received packet. I think of it like Nagle’s algorithm for hardware interrupts. Dynamic interrupt moderation (DIM) adjusts interrupt moderation parameters to find an ideal balance between latency and overhead. The authors find that the default DIM algorithm on their NIC produces inconsistent results depending on the thread count. Tweaking the DIM algorithm further reduced multi-threaded P99.9 latency to 215 microseconds . Fig. 7 shows how the insights from this paper can improve Redis latency: Source: https://dl.acm.org/doi/10.1145/3789240.3829115 The red Linux curve is the baseline The blue Linux + ACCa curve uses the fixed patch for the Linux scheduling issue The green Linux + PCSched curve uses a CPU scheduling policy based on the number of packets sent/received rather than wall-clock time consumed And the final purple line adds a tweaked DIM algorithm In the world of performance optimization, understanding the problem frequently takes much longer than fixing it. I wonder when LLMs will advance to the point where they can do this level of performance analysis. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Interrupt handling IP protocol processing TCP protocol processing CPU scheduling

0 views
Dangling Pointers 1 months ago

Analysis of Area and Power Overhead for a Synthesizable 4-Phase Bundled-Data Asynchronous RISC-V Processor

Analysis of Area and Power Overhead for a Synthesizable 4-Phase Bundled-Data Asynchronous RISC-V Processor Sean Jacobs, Mark Indovina, and Yashaswini Suresha GLVLSI'26 This paper describes the design of an asynchronous (i.e., there is no clock signal) RISC-V processor. The power and area results are impressive. I’ve only vaguely understood how asynchronous designs operate; this paper helped me understand the space a bit better. As far as I understand, there are two primary advantages of asynchronous circuits: Performance is not limited by infrequently used critical path. Power savings due to not having a clock signal #2 can be mitigated in synchronous designs with the use of clock gating. Inter-pipeline stage communication in this processor uses a 4-phase protocol. This protocol is illustrated in Fig. 2: Source: https://dl.acm.org/doi/10.1145/3787109.3815227 The 4 phases of the protocol are: The request signal transitions from low to high, indicating that all upstream stage outputs are ready to be consumed The acknowledge signal transitions from low to high, indicating that all upstream stage outputs have been consumed by the downstream stage The request signal transitions from high to low The acknowledge signal transitions from high to low The key building block that makes the 4-phase protocol implementable is the Muller C-element. A C-element has two inputs ( and ) and one output ( ). The output remains steady until both inputs are equal, at which point the output value is assigned to be equal to the first input (i.e., ). Fig. 3 illustrates how Muller C-elements and inverters can be used to construct a pipeline that uses the 4-phase protocol: Source: https://dl.acm.org/doi/10.1145/3787109.3815227 The triangles with circles on the input side are active-low inverters. Here is a narrative of how this works. Assume that the bottom ( ) input of the leftmost C-element is 0, the top ( ) input is 1, and the output ( ) is 0. When the input Req signal transitions to 1, it will cause the output of the leftmost C-element to transition to 1. This will send an acknowledgment backward to the unpictured upstream stage and send a request to the downstream (middle C-element) stage. The middle C-element will perform similarly (responding with an acknowledgment and propagating the request down the pipeline). When is set to 1, is set to 0, which will cause the left-most C-element output to return to zero when the unpictured upstream stage sets Req=0. The control path described above can be extended to support an arbitrary feed-forward pipeline as illustrated in Fig. 4: Source: https://dl.acm.org/doi/10.1145/3787109.3815227 The storage elements (pipeline registers) are D-latches (no clock inputs). The combinational logic clouds represent the logic that implements each stage of a RISC-V processor. The delay elements connected to each Muller C-element ensure that the control path does not assert the request signal until after the combinational logic has finished its job. An EDA tool is used to compute the worst possible delay through a logic cloud, and the corresponding delay value is set to be at least that long. The astute reader will realize that the delay elements eliminate one of the primary advantages of using an asynchronous design: the delays are set equal to the worst possible delay (i.e., the critical path). The advantage of this approach is that there is no area overhead added to the combinational logic. It is essentially the same logic that would exist in a synchronous design. Table 1 compares area and power numbers for this asynchronous processor (28nm, 4-stage) against a synchronous processor. Area is essentially unchanged; power is significantly lower. I’m not sure how much clock gating there is in the baseline implementation. Source: https://dl.acm.org/doi/10.1145/3787109.3815227 Dangling Pointers The performance numbers above should be taken with a grain of salt. As far as I can tell, the authors only tested the processor in simulation, they did not fabricate a chip in real life. The Muller C-elements are described with synthesizable RTL, not standard cells. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Performance is not limited by infrequently used critical path. Power savings due to not having a clock signal The request signal transitions from low to high, indicating that all upstream stage outputs are ready to be consumed The acknowledge signal transitions from low to high, indicating that all upstream stage outputs have been consumed by the downstream stage The request signal transitions from high to low The acknowledge signal transitions from high to low

0 views
Dangling Pointers 1 months ago

Vector Value Prediction with Element-wise Stride Compression

Vector Value Prediction with Element-wise Stride Compression Yanmeng Huang, Ling Yang, Yuanhu Cheng, Junhui Wang, Quan Deng, Junbo Tie, Yongwen Wang, Hai Zhong, and Libo Huang GLSVLSI'26 This paper makes an astute observation: the contents of vector registers are losslessly compressible. There are many hardware and architectural optimizations that can be derived from this fact. This paper presents one: prediction of vector register values. The point of value prediction is to expose more parallelism to a modern processor, so that more instructions can be executed in parallel. Much like branch prediction, value prediction speculatively breaks dependencies to allow the processor backend to process a less-constrained set of instructions. If the predicted value ends up being wrong, then the processor must throw away some work and start over. One way to implement value prediction is much like branch prediction, using the PC and branch history as the index into a table which contains a predicted value and a level of confidence that the predicted value is correct. This paper describes such a technique. SIMD instructions are an orthogonal way to expose more parallelism to a processor backend. Since value prediction and SIMD are orthogonal, why not combine them? Two reasons: The predictor must predict the values of all lanes of a vector register. If any one lane is a bad apple, it spoils the bunch. The data structures needed to make accurate vector-wide predictions would consume a lot of on-chip memory (storage would grow linearly with vector width) The core idea of this paper is illustrated in Fig. 1: Source: https://dl.acm.org/doi/full/10.1145/3787109.3815244 The left side of the figure shows raw (uncompressed) vector registers with varying element widths. The right side shows compressed versions. The gray rectangle represents the value stored in lane 0 (no compression). The orange rectangles hold the differences (i.e., strides ) between neighboring elements. The magnitude of a stride is typically small, so strides are stored with only a few bits per stride. Here is a concrete example: The key point here is that the base + strides representation consumes fewer bits than the raw representation. If the values in a particular register cannot be represented with the base + strides representation, then those values are not considered candidates for value prediction. Fig. 6 shows IPC improvements provided by vector value prediction across a set of benchmarks: Source: https://dl.acm.org/doi/full/10.1145/3787109.3815244 Dangling Pointers I didn’t see specific handling for floating point values in the paper. I wonder if a dedicated compression scheme for those types is warranted. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. The predictor must predict the values of all lanes of a vector register. If any one lane is a bad apple, it spoils the bunch. The data structures needed to make accurate vector-wide predictions would consume a lot of on-chip memory (storage would grow linearly with vector width)

0 views
Dangling Pointers 1 months ago

Modular GPU Programming with Typed Perspectives

Modular GPU Programming with Typed Perspectives Manya Bansal, Daniel Sainati, Joseph W. Cutler, Saman Amarasinghe, and Jonathan Ragan-Kelley PLDI'26 GPU programming is great, who doesn’t want to copy memory with thousands of threads? The trouble is that some operations require multiple threads to cooperate, and current GPU programming languages only have kludgy ways of expressing such cooperation. For example, see the fine print on the function from the CUB library (quoting from the paper which quotes from the CUB documentation): Computes a block-wide reduction for thread0 using the specified binary reduction functor. • The return value is undefined in threads other than thread0. • A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective’s temporary storage (e.g., temp_storage) is to be reused or repurposed This paper introduces the Prism language, which uses perspectives to bring some much-needed hierarchy to GPU programming Prism supports 3 coarse-grain levels of nesting: , , . indicates when code/data applies to pairs of threads. applies to a warp of threads. is used to indicate a group of 8 blocks. Grids are broader (i.e., coarser) than blocks, which are broader than threads. Also, is broader than (this only applies when the larger integer is divisible by the smaller one). Each statement is associated with a stack of perspectives, and the operator can be used to push an element on to the stack. For example: Similarly, local variables are annotated with a perspective, to indicate how they are shared among threads: Perspectives are the foundation upon which Prism builds higher level concepts and a type system that enable hierarchical parallelism to be expressed. This allows for code that is more modular and easier to reason about than the typical flat parallelism found in other languages. For example, Prism statically disallows a code operating at a broad perspective from reading the value of a variable declared with a narrow perspective (only read up is allowed). Similarly, local variables can only be written from a perspective that is the same as the variable’s perspective or broader ( write down is allowed). Fig. 22 compares various Prism implementations of a matrix multiply against , not too shabby. Source: https://dl.acm.org/doi/10.1145/3808290 Dangling Pointers The interface exposed by GPUs doesn’t natively support nested parallelism. I wonder if there is an opportunity for hardware to take advantage of hierarchically information that is available at the Prism level of abstraction. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 2 months ago

RNG: Flat Datacenter Networks at Scale

RNG: Flat Datacenter Networks at Scale Giacomo Bernardi, Ratul Mahajan, C. Seshadhri, Enrico Carlesso, Chinchu Merine Joseph, Saurabh Kumar, Pavan Manikonda, Luiza Popa, Randy Ram, Steven Robinson, Elizabeth Tennent arXiv This paper from Amazon describes an alternative to fat tree network topologies. In a fat tree, switches are organized into a strict set of levels (top-of-rack switches being the leaves). The tree is called fat because each switch at given level can be connected to multiple switches at a higher level. The paper describes the core problem of fat trees as capacity fungibility . Only a small set of network links are available to route traffic between any specific pair of switches. Fig. 1 compares a fat tree to an expander network (yellow circles represent top-of-rack switches): Source: https://arxiv.org/abs/2604.15261 The expander network has no visible structure, the connections are effectively random. Switches in this network use spraypoint routing to move packets to their destination. In spraypoint routing, all pairs of switches ( , ) are classified according to how packets destined for should be routed. Fig. 3 illustrates the classification scheme assuming that a packet is destined for switch : Source: https://arxiv.org/abs/2604.15261 Switches are classified as waypoint level 0 (WP 0 ) switches relative to . These switches are directly connected to . When a packet destined for arrives at a WP 0 switch, the packet is forwarded to . Switches are classified as waypoint level 1 (WP 1 ) switches relative to . These switches are not directly connected to but are directly connected to one or more WP 0 switches. When a packet destined for arrives at a WP 1 switch the packet is randomly forwarded to a WP 0 switch. The network can have any number of waypoint levels. Once the maximum waypoint level is exceeded, the remaining switches are classified as inner ring ( IR ) and outer ring ( OR ) switches relative to . Inner ring switches are directly connected to the last waypoint level; all remaining switches are classified as outer ring switches. Inner ring switches randomly forward packets to a switch in the last waypoint level. Outer ring switches randomly forward packets to a switch from the set of neighboring switches with the shortest path to . Random forwarding (spraying) is based on ECMP hashing, such that packets from the same flow will not be reordered. Spraypoint routing requires all switches to have a full view of the network topology, in order to classify all pairs of switches. Section 5 of the paper describes CPU and memory requirements for Spraypoint. The take-home message is that off-the-shelf switches can support Spraypoint. Fig. 14 shows the bottom line: Source: https://arxiv.org/abs/2604.15261 Spraypoint allows the number of switches in a network to be reduced by 2-5x relative to a fat tree topology. ECMP hashing at each switch implies that this design goes to some length to avoid reordering packets within a flow. Recent work has shown that it is possible to relax this requirement (either in the NIC or the switch ). I wonder how Spraypoint could be improved if it was coupled with hardware like Falcon. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 2 months ago

Bonsai: Compiling Queries to Pruned Tree Traversals

Bonsai: Compiling Queries to Pruned Tree Traversals Alexander J Root, Christophe Gyurgyik, Purvi Goel, Kayvon Fatahalian, Jonathan Ragan-Kelley, Andrew Adams, and Fredrik Kjolstad PLDI'26 File this under: “so elegant, why wasn’t this discovered sooner?” This paper describes a generic language (Bonsai) and compiler for efficient tree traversals. The language is structured such that a simple compiler can quickly generate efficient code. The examples from the paper fall into two buckets: SQL-like queries and spatial data structures. Here is some example code from the paper which describes a SQL-like operation on a set of points: A is a structure with two elements ( , and ). The function returns the minimum associated with any point that has an x value within the range . And here is an example from the paper which describes a ray tracing function: The function finds all triangles which intersect a given ray and returns the one that is the closest to the origin of the ray. The key point of the paper is that both types of functions can be implemented as tree traversals. A key design point that makes Bonsai feasible is metadata stored in the trees which represent sets. Tree metadata is specified separately from queries. The following snippet from the paper contains a tree data structure represented as an algebraic data type: A set of points is represented as a tree. Leaf nodes hold points. Interior nodes have pointers to and children, and four pieces of metadata ( ). The keyword is used to describe the meaning of the metadata. The expression means that the value of all x fields in all points contained in the subtree lies within the range . The expression means that the minimum value of the field in any point in the subtree is . The Bonsai compiler takes a query and a description of the tree metadata as input and generates C++ code to perform the query via traversing the tree. There are three key properties of this process. The first is that the generated code for all filter and reduction operations are fused together. In other words, there are no intermediate trees (i.e., sets) produced during query execution. Secondly, tree metadata is used to accelerate both filtering and reduction operations. For filtering, the generated code checks tree metadata to determine if all data elements in a subtree will be accepted by the filter, or if all elements will be rejected by the filter. If all elements will be rejected, then there is no need to traverse the subtree. If all elements will be accepted, then there is no need for further evaluation of the filter expression for the subtree. In this case, traversal into the subtree can be skipped if the root of the subtree has metadata containing a pre-reduced value ( in the example above). Fig. 4 shows IR corresponding to two filtering examples. In both examples, there are two types of leaf nodes in the tree. nodes contain a single value whereas nodes contain an array of values. Fig. 4a shows the IR for a single filter expression . Here is the code, annotated with some comments to describe the semantics: Fig. 4b shows the IR for the logical and of two filter expressions: : Code generation tracks a symbolic interval associated with each expression. An interval is represented by two expressions: one that evaluates to the lower bound of the expression and one that evaluates to the upper bound. This interval analysis is used to produce the code that implements and . Interval analysis is fast to compile but can produce code that suffers from false positives. Symbolic interval analysis is general enough to handle spatial filters (e.g., those used in ray tracing). Fig. 7 compares Bonsai to a state-of-the-art library ( FCPW ) for geometric queries: Source: https://dl.acm.org/doi/10.1145/3808256 Fig. 8 compares Bonsai to relational databases for range joins. The range join logically computes the Cartesian product of two relations, and then removes elements from the result which are not near each other according to Manhattan distance . Source: https://dl.acm.org/doi/10.1145/3808256 Dangling Pointers A logical extension of this process would be to automatically determine what tree metadata would be most useful for a set of queries. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 2 months ago

MatchBox: A Semantic Foundation for Data Plane Portability

MatchBox: A Semantic Foundation for Data Plane Portability Eric Hayden Campbell, Robert Zhang, Divyanshu Saxena, Aditya Akella, and Işıl Dillig PLDI'26 This paper identifies problems associated with subtle differences in match-action table implementations across hardware vendors and proposes some formalism to allow generic tooling to be implemented for match-action tables. A match-action table is a data structure present in some NICs and switches which specifies how packet-processing hardware should modify packets on the fly. This post on Gigaflow also discusses how a virtual switch can be implemented with match-action tables. Fig. 2 shows two match-action tables programmed in a hypothetical switch ( IP addresses in these tables are written in CIDR syntax ): Source: https://dl.acm.org/doi/10.1145/3808277 The syntax is clear but the semantics have an important footnote: switch S has a policy which causes packets that miss in the routing table to be dropped. Now imagine a switch ( ) from another vendor with a different policy: packets which cause a miss in the routing table are broadcast to all output ports. The task of porting the match-action tables from to in a way to preserves behavior is non-trivial. Fig. 3 contains a correct solution: Source: https://dl.acm.org/doi/10.1145/3808277 The difference is in the ACL table. The permissive allow entry has been replaced with two more restrictive entries that only allow packets that would not have triggered a miss in the routing table. One automated way to compute the contents of the ACL table for switch is an algorithm that looks at all pairs (i.e., the Cartesian product) of entries in the route and ACL table for switch , and then drops pairs that don’t make sense (e.g., the route and ACL addresses have no overlap). If you squint at that definition, it looks a lot like a relation join operation. In other words if you think of the ACL and routing tables as relations, then a way to compute the ACL table for switch would be to join the route and ACL tables for switch . The authors argue that relational algebra is almost the correct hammer for this nail, but not quite, and so this paper proposes Match Algebra . In Match Algebra, a match-action table can be thought of as a partial function which maps a valuation to an (action, valuation) pair. A valuation is just a set of bits (extracted from packet headers), and an action describes how a packet is to be modified. Another way to think of a match-action table is an ordered set of rules. A rule comprises: A guard, which defines the set of packets that match the rule An (action, valuation) pair Example guards could be: IPv4 packets with source address matching 10.2.1.0/24 UDP packets with a destination port of 15354 or 15355 An example (action, valuation) pairs could be: (replace the destination port, 15356) (pass the packet through unmodified, x) The set of rules is ordered such that if guards associated with multiple rules apply to a particular packet, the highest priority rule wins. The most important operators in the Match Algebra are: Sequential composition: modify the packet according to table , and use the result as input to table Preferential composition: Check to see if table has a rule matching the incoming packet, if so then apply that rule. Otherwise, use table to modify the packet. Parallel join: apply tables and to the input packet to generate two (action, valuation) pairs: , . Merge the results together. The merge operation requires that the results from and do not overlap (e.g., causes a destination port to be modified while causes destination address to be modified). The paper describes some other operators and presents formal semantics for Match Algebra. The paper describes which is a language (embedded in OCaml) which operates at the Match Algebra level of abstraction. Fig. 13 shows an example program which transforms match action tables with operators from Match Algebra: Source: https://dl.acm.org/doi/10.1145/3808277 Results The paper presents the following real-world use cases of : Transforming a set of match-action tables to an equivalent set with a different form (for porting rules between switches) Porting firewall rules between AWS, GCP, and Azure (which each have a different required structure) Translating system-wide eBPF tables to per-CPU eBPF tables Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Source: https://dl.acm.org/doi/10.1145/3808277 The syntax is clear but the semantics have an important footnote: switch S has a policy which causes packets that miss in the routing table to be dropped. Now imagine a switch ( ) from another vendor with a different policy: packets which cause a miss in the routing table are broadcast to all output ports. The task of porting the match-action tables from to in a way to preserves behavior is non-trivial. Fig. 3 contains a correct solution: Source: https://dl.acm.org/doi/10.1145/3808277 The difference is in the ACL table. The permissive allow entry has been replaced with two more restrictive entries that only allow packets that would not have triggered a miss in the routing table. One automated way to compute the contents of the ACL table for switch is an algorithm that looks at all pairs (i.e., the Cartesian product) of entries in the route and ACL table for switch , and then drops pairs that don’t make sense (e.g., the route and ACL addresses have no overlap). If you squint at that definition, it looks a lot like a relation join operation. In other words if you think of the ACL and routing tables as relations, then a way to compute the ACL table for switch would be to join the route and ACL tables for switch . The authors argue that relational algebra is almost the correct hammer for this nail, but not quite, and so this paper proposes Match Algebra . Match Algebra In Match Algebra, a match-action table can be thought of as a partial function which maps a valuation to an (action, valuation) pair. A valuation is just a set of bits (extracted from packet headers), and an action describes how a packet is to be modified. Another way to think of a match-action table is an ordered set of rules. A rule comprises: A guard, which defines the set of packets that match the rule An (action, valuation) pair IPv4 packets with source address matching 10.2.1.0/24 UDP packets with a destination port of 15354 or 15355 (replace the destination port, 15356) (pass the packet through unmodified, x) Sequential composition: modify the packet according to table , and use the result as input to table Preferential composition: Check to see if table has a rule matching the incoming packet, if so then apply that rule. Otherwise, use table to modify the packet. Parallel join: apply tables and to the input packet to generate two (action, valuation) pairs: , . Merge the results together. The merge operation requires that the results from and do not overlap (e.g., causes a destination port to be modified while causes destination address to be modified). Source: https://dl.acm.org/doi/10.1145/3808277 Results The paper presents the following real-world use cases of : Transforming a set of match-action tables to an equivalent set with a different form (for porting rules between switches) Porting firewall rules between AWS, GCP, and Azure (which each have a different required structure) Translating system-wide eBPF tables to per-CPU eBPF tables

0 views
Dangling Pointers 2 months ago

Breadcrumb Filters: Fast Fully Featured Filters

Breadcrumb Filters: Fast Fully Featured Filters Andrew Krapivin, Aaditya Rangarajan, Alex Conway, Martin Farach-Colton, Rob Johnson, and Prashant Pandey SIGMOD'26 This paper presents the design of a breadcrumb filter , which is a membership testing data structure . Unlike a Bloom filter , a breadcrumb filter supports operations like deletion and merging. Breadcrumb filters also have the nice property that most operations access a single cache line (most of the time). A breadcrumb filter is a fingerprinting filter . Each item in a set is represented by its fingerprint (i.e., hash). Say a filter contains 1024 cache lines, and each cache line has storage for items. To insert an item into the filter, compute a 16-bit fingerprint of the item. Decompose that into a 10-bit integer (the cache line index) and a 6-bit integer (the remainder). Use the cache line index to determine which cache line to access. Find an empty slot in that cache line and place the remainder bits of the fingerprint into the empty slot to represent the item. A breadcrumb filter builds on top of these mechanics by cleverly dividing the filter storage into two sections: the front and back yards. The front-yard represents the fast path: each filter operation will touch one front-yard cache line. The backyard is only used to handle cases where a front-yard cache line fills up. The paper hyphenates “front-yard” but writes “backyard” as one word. The following pseudo-code illustrates how an item is inserted into a breadcrumb filter: A lookup operation follows a similar structure: To delete an item from a breadcrumb filter, it is sufficient to delete the item’s fingerprint . That wasn’t obvious to me up front. Imagine two items have the same fingerprint (hash). Inserting them both causes the same fingerprint to be inserted twice. Now, when one of them is deleted, it suffices to delete one of the copies of the fingerprint in the breadcrumb filter. The trick with deletion is that deleting a fingerprint from a front-yard cache line can require promoting an item from an associated backyard cache line. The real magic with breadcrumb filters comes in how items are moved between the front-yard and backyard. If a front-yard cache line is found to be full during insertion, then one item is moved to the backyard. That item could be the one that is currently being inserted, or it could be an item that was previously placed into the front-yard. The policy is: move the item which has the greatest value of the remainder bits . For example, if two items (A, and B) have remainder bits of 23 and 7, then item A will be moved to the backyard before B. This enables lookup operations to avoid touching backyard cache lines. For example, say the item that is being searched for has remainder bits = 23, and the (full) front-yard cache line contains items with remainder bits = [12, 5, 34, 3], there is no need to search the backyard. The “34” in the front-yard implies that no item with remainder bits value less than can be in the associated backyard cache lines. Note that this policy requires that delete operations sometimes move items from the backyard to the front-yard. The other trick is a mapping between front-yard and backyard cache lines which enables promotion of a deleted item from backyard to front-yard. Say each front-yard cache line is associated with two backyard cache lines, but those backyard cache lines are each associated with many front-yard cache lines. A front-yard cache line index is represented with 10 bits: The two backyard cache line indices associated with that front-yard cache line are: The key here is that very little information needs to be stored in the backyard in order to allow mapping from a backyard cache line index to a front-yard cache line index. To map backyard index back to , all one needs to know is the value of bit (which is stored in the backyard cache line). When an item is deleted from a front-yard cache line, the two associated backyard cache lines are searched for an item to be promoted back to the front-yard. Metadata (e.g., the value of bit ) is used to ensure that items are promoted back to the front-yard cache line from whence they came. Fig. 9 compares the throughput of the breadcrumb filter (BCF*) against other filters: Source: https://dl.acm.org/doi/10.1145/3786629 Dangling Pointers This design agrees with many others that the best one can do is read a single cache line for each lookup. I don’t have a better solution in mind, but it seems painfully slow if the filter doesn’t fit in cache. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 3 months ago

RABIT: Efficient Range Queries with Bitmap Indexing

RABIT: Efficient Range Queries with Bitmap Indexing Junchang Wang, Fu Xiao, and Manos Athanassoulis SIGMOD'26 This paper presents an optimization for range filtering (e.g., ). File this under: “so crazy it might actually work”. This paper builds on the concept of a bitmap index . If the cardinality of a column is (i.e., there are distinct values in the column), then a bitmap index stores additional 1-bit columns along with the original column. These additional columns are called point bitvectors . The value in additional column is 1 for a particular row if the value in that row is the th distinct value. Here is an example column with 6 rows and 4 distinct values (0, 1, 2, 6). The header indicates that this is the value column. And here is the same column with 4 additional 1-bit wide rows attached. The header indicates columns that hold point bitvectors. To find all rows which have a value of 2, simply read the 3rd (i.e., ) bitmap column. Similarly, to find all rows which have a value of 6, simply read the 4th bitmap column. To efficiently support range queries, the authors of this paper propose adding even more columns that hold 1-bit values. These columns are called cumulative bitvectors . A cumulative bitvector holds the bitwise-or of a set of point bitvectors. In the example above, let’s create a cumulative bitvector (named ) for the values 0 and 1. The value of this column for a given row will be 1 if the value contained in the row is either 0 or 1. In other words: . Similarly, we can create a cumulative bitvector for the values 2 and 6 using the equation: . Here is the full turkey: A range query can be executed by reading one or more cumulative bitvectors and possibly a few point bitvectors. For example, to find all rows that have a value < 2, all one needs to do is read the value of . Cumulative bitvectors do not add much value in this toy example because each cumulative bitvector only aggregates data for two values, but you can see how this could work well with more aggregation. This trick can even be made to work for range queries that partially overlap with a cumulative bitvector. This whole scheme relies on the fact that point and cumulative bitvectors are highly compressible. This paper assumes the use of WAH compression . The executive summary of WAH compression is to divide each bitvector into words (e.g., 32 or 64 bit). One bit of each word is metadata that determines if the word is a or a . The remaining bits of a literal word contain raw (uncompressible bits). The remaining bits of a fill word contain a value and a length (run-length encoding). Fig. 9 compares throughput of this scheme (labeled GE in the figure) to other work. means each cumulative bitvector aggregates data for 20 point bitvectors. Performance looks good even for columns holding 100K distinct elements. Source: https://dl.acm.org/doi/10.1145/3769819 Table 3 compares the storage requirements for this scheme versus other indexing schemes that support range queries, which seems too good to be true. Source: https://dl.acm.org/doi/10.1145/3769819 Dangling Pointers I wonder if this idea could be generalized to other types of filtering, such as string operations (e.g., . Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 3 months ago

Accelerating Stream Processing Engines via Hardware Offloading

Accelerating Stream Processing Engines via Hardware Offloading Zhengyan Guo, Mingxing Zhang, Yingdi Shan, Kang Chen, Jinlei Jiang, and Yongwei Wu SIGMOD'26 This paper describes a trick to offload partitioning from CPU to NIC via clever use of RSS. The context of the paper is distributed systems for processing streaming queries, but the trick seems applicable to databases in general. Hash partitioning is a common divide-and-conquer technique to implement joins and aggregations. Here are some posts about papers that use this partitioning: SPID-Join: Skew-Resistant In-DIMM Joins Breaking Through the Memory Wall of OLTP Systems with PIM High-Performance Query Processing with NVMe Arrays: Spilling without Killing Performance Efficiently Processing Joins and Grouped Aggregations on GPUs RSS is a NIC feature whereby the NIC hashes select fields from incoming packet headers and uses the result to determine which CPU core to send the packet to. This enables efficient load balancing across CPU cores without reordering packets within a given flow (i.e., connection). Here is a previous post describing a clever way to extract more value out of RSS in cloud VMs: Enabling Fast Networking in the Public Cloud If you have many nodes cooperating to process a query, then the hash partitioning may span many nodes. For example, node A could hash the join/aggregation key of each row and then forward the row to either node B, C, or D, E depending on the hash value. This enables the join/aggregation work to be split across nodes B, C, D, and E. This is all fine and dandy from the perspective of node A. However, nodes B, C, D, and E likely have multiple CPU cores. How can one of these nodes execute their join/aggregation in parallel? The answer is recursive: partition the incoming rows again (using a different hash function) and have each CPU core process one of these smaller partitions. The paper focuses on the cost of partitioning the dataset, which can cost just as much as the partition join/aggregation step that follows it. The key insight is that the partition algorithm looks a lot like the RSS load balancing algorithm present in the NIC hardware. Here is the punchline: establish multiple network connections (using different ports) between node A and each of nodes B, C, D, and E. When node A partitions rows, it determines a specific connection (not node) to send each row to. This doesn’t improve performance at the sender, but it dramatically helps the receivers. Each receiver configures RSS such that all connections are spread across the CPU cores on the receiver. The NIC then distributes received packets to the appropriate CPU cores without any partitioning work on the receiving nodes. The one downside to this approach is load imbalances that occur due to data skew. If some join/aggregation keys are more common than others, then some CPU cores may be assigned more work than others. The paper proposes to dynamically monitor load imbalance at each receiver and reconfigure the RSS settings of the NIC to move connections off hot cores. Section 5 of the paper describes synchronization necessary to move a connection between cores in the middle of the query. This is a good mitigation, but as we’ve seen in this paper , RSS configuration is not uniformly exposed on cloud VMs. Fig. 8 has performance results across a number of benchmarks: Source: https://dl.acm.org/doi/10.1145/3769754 Dangling Pointers The solution is great, but asymmetric. I wonder if there is a way to get similar benefits on at the sending node (send side scaling)? Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 3 months ago

F3: The Open-Source Data File Format for the Future

F3: The Open-Source Data File Format for the Future Xinyu Zeng, Ruijun Meng, Martin Prammer, Wes McKinney, Jignesh M. Patel, Andrew Pavlo, and Huanchen Zhang SIGMOD'26 F3 is a file format for columnar data (e.g., Parquet ) that is designed to be efficient and extensible. The optimizations make sense, the extensibility mechanism is ingenious , dangerous , fascinating. The key assumption made by this paper is that the hardware and software will continue to improve. It is hard to argue with that. The trouble is that interoperable formats like Parquet take a snapshot of the state-of-the-art and freeze it in a specification. Some innovations that are invented after the format is frozen are incompatible with existing formats because they require a different data layout. Section 1 of the paper refers to many examples related to compression, indexing, and filtering. The goal of F3 is to be general enough to allow seamless incorporation of future innovations without changing the F3 spec nor F3 decoder implementations. Fig. 2 illustrates an F3 file: Source: https://dl.acm.org/doi/10.1145/3749163 A file consists of a metadata and a set of row groups. A specific row group contains data for all columns and a subset of rows. F3 contains incremental improvements over existing columnar formats, for example: F3 metadata supports random access, which is important for operations that only need to access a smaller percentage of all columns. F3 decouples file I/O from a row group storage. The rows associated with a given column in a row group are further subdivided into , which are actually stored. This allows row groups to be sized for efficient row-group level filtering, while the IO unit size is tuned to minimize working set while also amortizing the fixed costs associated with file I/O. F3 allows flexible . Each IO unit can contain a dedicated dictionary, or multiple IO units can share a dictionary. Columns with low cardinality will benefit from smaller dictionary scopes, whereas columns with high cardinality do better with larger dictionary scopes. The stand-out feature for F3 is the yellow block in the block. The idea is that an F3 file can contain within it the WebAssembly code needed to decode the encoded values in an IO unit. If someone invents a brilliant new encoding method that works well with some data sets, they can ship the decoder right along with the data set. Storage of the WASM code shouldn’t be too much of an issue, because the storage cost is amortized across all row groups. The big questions are performance and security. Section 6.2 has some comments on this. In theory, the WASM specification is air-tight, and a bug-free implementation should be able to securely run arbitrary WASM code in-process. WASM also supports performance optimizations like parallel compilation and SIMD instructions. Something I don’t see in the paper is a discussion about how filtering interacts with WASM decoding. I suppose the extensibility could only be used for decoding, and filtering could be hard coded into F3, but that seems against the extensible spirit of F3. Fig. 11 shows the working set reduction from decoupling IOUnit size from row group size: Source: https://dl.acm.org/doi/10.1145/3749163 Table 3 shows how flexible dictionary scopes allow one to trade encoding time for compression ratio (lower relative CR numbers mean smaller files on disk): Source: https://dl.acm.org/doi/10.1145/3749163 Fig. 15 quantifies WASM overhead by comparing decode time for hard coded F3 decoder implementations vs the same algorithms expressed in WASM: Source: https://dl.acm.org/doi/10.1145/3749163 Fig. 16 shows potential savings associated with using WASM extensibility to implement a custom decoder from the literature. Source: https://dl.acm.org/doi/10.1145/3749163 Dangling Pointers I wonder how well WASM decoders can be implemented on other hardware architectures. Is WASM the ideal language for expressing this, or convenient standard that already exists? Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Source: https://dl.acm.org/doi/10.1145/3749163 A file consists of a metadata and a set of row groups. A specific row group contains data for all columns and a subset of rows. F3 contains incremental improvements over existing columnar formats, for example: F3 metadata supports random access, which is important for operations that only need to access a smaller percentage of all columns. F3 decouples file I/O from a row group storage. The rows associated with a given column in a row group are further subdivided into , which are actually stored. This allows row groups to be sized for efficient row-group level filtering, while the IO unit size is tuned to minimize working set while also amortizing the fixed costs associated with file I/O. F3 allows flexible . Each IO unit can contain a dedicated dictionary, or multiple IO units can share a dictionary. Columns with low cardinality will benefit from smaller dictionary scopes, whereas columns with high cardinality do better with larger dictionary scopes.

0 views
Dangling Pointers 3 months ago

Enabling Packet Spraying over Commodity RNICs with In-Network Support

Enabling Packet Spraying over Commodity RNICs with In-Network Support Xiangzhou Liu, Wenxue Li, Zihao Wang, and Kai Chen EUROSYS'26 This paper proposes changes to top-of-rack (ToR) switch hardware to enable packets from a single flow to utilize many network paths ( Falcon offers a similar benefit via changes to the NIC rather than the switch). The Falcon approach is robust but more invasive. The switch-based approach from this paper is a more incremental change. The sweet spot for packet spraying is a data center that has a large number of network paths compared to network flows (i.e., connections). In such an environment, there is an incentive to spray the packets associated with one flow across multiple paths. The trouble with packet spraying is that packets will commonly arrive out-of-order. The system has to be able to distinguish the out-of-order case from genuine packet loss. Section 2.2. of the paper describes three techniques for handling packet loss: PFC - the trouble is that this has scalability limits, and only addresses packet loss due to buffer overflows (not angels flying down and flipping your bits) Timeouts - the trouble is that practical timeout values have to be large Selective Repeat Selective repeat is a feature of modern RDMA NICs which is similar to the bitmaps tracked by Falcon hardware. The idea is that a receiving NIC tracks an expected sequence number (ePSN) for each flow. If a packet arrives with a sequence number greater than (but not too much greater than) the ePSN, the NIC accepts it and records this fact in a per-flow bitmap. The receiving NIC then sends a NACK to the sender, requesting that the sender resend the packet corresponding to the ePSN. In the out-of-order case, this NACK is unnecessary as the expected packet will arrive soon enough. The question asked by this paper is: can one easily modify switch hardware to filter the unnecessary NACKs? The core idea proposed by this paper is that the NICs and switches agree on the number of paths for a particular flow. Sending NICs use a packet’s sequence number to determine which path to use (e.g., . The switch can then track information for each path associated with each flow. When a NACK arrives at a switch, the switch can drop the NACK as necessary (thus avoiding unnecessary retransmissions). For example, say there is a single flow mapped to 4 paths. Packets with PSNs [0, 4, 8, 12, …] will travel over path 0. Packets with PSNs [1, 5, 9, 13, …] will travel over path 1. If packets arrive at the receiving NIC in this order: [0, 1, 5 , 4] the NIC will send a NACK when it receives packet 5. However, the switch will drop that NACK because it “knows” that no packet has been received out of order with respect to its flow. Fig. 12 has simulated performance results for collective operations common in AI workloads. is the work described in this paper. CCT is a measure of how long the collective operation took. Source: https://dl.acm.org/doi/10.1145/3767295.3803588 Dangling Pointers This feels like an engineering solution to a business problem of how to get many NIC vendors to align on a packet spraying solution. I suspect there are many applications where it would not be too difficult to introduce multiple flows. For example, in a machine learning workload, the weights/deltas associated with layer could be assigned to flow . This would increase network path utilization without any hardware changes. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. PFC - the trouble is that this has scalability limits, and only addresses packet loss due to buffer overflows (not angels flying down and flipping your bits) Timeouts - the trouble is that practical timeout values have to be large Selective Repeat

0 views
Dangling Pointers 4 months ago

LightDSA: Enabling Efficient DSA Through Hardware-Aware Transparent Optimization

LightDSA: Enabling Efficient DSA Through Hardware-Aware Transparent Optimization Yuansen Wang, Teng Ma, Yuanhui Luo, Dongbiao He, Zheng Liu, and Yunpeng Chai EUROSYS'26 This paper describes performance characteristics of the Intel DSA hardware accelerator, and software techniques to maximize performance when using DSA. My takeaway is: the DSA supports a variety of convenience features, but each one is so expensive that you are better off adding software complexity to avoid these paths. The Data Streaming Accelerator ( DSA ) is a hardware accelerator in recent Intel chips. It can implement simple memory operations like , , , and CRC generation. Fig. 2 contains a high-level diagram of the DSA architecture. Source: https://dl.acm.org/doi/10.1145/3767295.3769356 Operations are written into work queues as 64-byte descriptors. A work descriptor (WD) describes a single operation, whereas a batch descriptor (BD) references many work descriptors. A batch is the fundamental unit of control (the DSA signals the CPU when a batch has completed). The DSA contains multiple engines and arbiters to spread work across the engines. Source and destination buffers accessed by the DSA do not need to be pinned into memory, the DSA can handle page faults. The DSA supports demand faulting via the page request service (PRS). This enables the DSA to send an interrupt to the OS (via the IOMMU) requesting the OS to resolve the fault. This paper reports a similar finding to a previous paper looking at the PCIe page request interface: demand faulting is convenient, but slow. The LightDSA authors recommend that software forcibly fault in pages before submitting descriptors. The DSA supports operations that require unaligned reads and writes of data from/to DRAM but the authors find that 64-byte aligned accesses are much faster. Fig. 8 has some numbers, even 32-byte alignment is expensive (compare the light green and dark green bars). Source: https://dl.acm.org/doi/10.1145/3767295.3769356 The authors recommend having software ensure that all writes performed by the DSA are 64-byte aligned. Software can do this by executing the operation for the first few bytes of each task, up until the destination buffer is 64-byte aligned. Like many HW/SW interfaces, the DSA writes both result data and metadata to memory. Result data is associated with each work descriptor, while completion metadata is associated with each batch descriptor. Metadata is read by software to learn when an operation completes. In such a scheme, it is important that software observes the batch metadata write after the result writes have completed. If the metadata write can land first, then software may try to read the result buffer before it has actually been updated. The DSA supports multiple traffic classes (TC). As with discrete PCIe accelerators, writes from the DSA associated with the same traffic class will land in host memory in order (these are posted writes ). However, writes associated with different TCs may be reordered. Here is a previous paper that describes performance problems with reordering. Section 3.8 of the DSA architecture specification describes two choices that software developers have. Either they should configure work descriptors and batch descriptors to use the same traffic class, or they should configure the DSA to enforce ordering via the (readback) flag. When that flag is set, the DSA will ensure that all result writes have landed in host memory by issuing a read request to read the most recently written result data back to the DSA, waiting for the response to come back, and then issuing metadata writes associated with the batch descriptor. Discrete PCIe devices can use the same trick to enforce ordering across traffic classes. Fig. 7 shows the performance cost of using this feature: Source: https://dl.acm.org/doi/10.1145/3767295.3769356 My takeaway is that DSA users should ensure that work and batch descriptors use the same traffic class, to avoid having to invoke this slow read-back path. Because the DSA contains multiple engines, tasks can complete in a different order than the order in which they are submitted. This is fine in itself, but the authors note that software must take care to efficiently support allocating work and batch descriptors in light of this. Time spent bookkeeping to handle out-of-order completion is overhead that adds up for small tasks. The solution proposed by this paper is for software to maintain two batch descriptor lists (free, and busy). When software needs to recycle descriptors from the busy list to the free list, it checks most (but not all) batch descriptors in the busy list to see if the hardware has completed the batch. This is in contrast to an approach which simply checks to see if the oldest-submitted batch has completed. The paper finds that it is optimal for the recycling process to ignore the 25 most recently submitted batch descriptors but check the completion status of all other outstanding batches. Figs. 12 and 13 compare the performance you can expect to see from using DSA naively versus using the techniques described in this paper (LightDSA). My takeaway is that the DSA is powerful, but only if you use it carefully. Source: https://dl.acm.org/doi/10.1145/3767295.3769356 Dangling Pointers I suspect the elevator pitch for DSA is something like: “just re-compile your existing C/C++ code and all of the memcpy/memcmp time will be optimized out”. It seems like DSA falls short of that. I wonder if the elevator pitch would be better realized if application code was written in other languages (like an explicitly pipeline parallel language). Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 4 months ago

Accelerating Transactional Execution via Processing-In-Memory

Accelerating Transactional Execution via Processing-In-Memory André Lopes, Daniel Castro, and Paolo Romano EUROSYS'26 This paper describes a way to implement OLTP for a processing-in-memory architecture. As with other academic research, it uses UPMEM ( here are two summaries of papers that rely on UPMEM). Something I found surprising in this work is conflicts can cause transactions to abort, even if all transactions only access data in the same UPMEM bank. A UPMEM DIMM is like a DRAM DIMM, but each bank contains a multi-threaded in-order core which can access data from the bank it is co-located with. This paper calls these processors DPUs (some other papers call them IDPs). The only way for two DPUs to communicate with each other is for the host CPU to read data from one bank and write it into another. The system described in this paper is called PIM-TIDE. It assumes that transactions come pre-sliced into computational graphs comprising subtransactions. A single subtransaction only accesses data within one bank (and thus executes on a specific DPU). Users of PIM-TIDE do not need to know exactly which data words a subtransaction will access, but they do need to be able to restrict a subtransaction to only access data from a specific bank. This works well on TPC-C style transactions where most of the database can be partitioned on . The host CPU groups transactions into batches and sends work to the DPUs one batch at a time. Within a batch, transactions are categorized into two groups: Local transactions execute entirely within a DPU Distributed transactions execute across multiple DPUs All subtransactions associated with distributed transactions within a batch are assigned a unique sequence number, which determines the order in which the subtransactions will commit. Local transactions are not preassigned to a commit order. DPUs first process all distributed subtransactions in a batch and then execute all local subtransactions. Algorithm 3 illustrates PIM-TIDE’s concurrency control scheme for dealing with intra-DPU conflicts between subtransactions. is stored in fast on-chip memory and is indexed by a hash of the word address. If a transaction aborts, state can be rolled back and the transaction is retried. All transactions assigned to a DPU will commit eventually. Inter-DPU conflicts are handled by deterministic concurrency control (i.e., the pre-assigned sequence numbers). Source: https://dl.acm.org/doi/10.1145/3767295.3803621 Results Fig. 3 compares performance of PIM-TIDE vs a CPU baseline for a mix of TPC-C transactions, wow that is a significant speedup. I believe all transactions/subtransactions are written by hand in C code that is compatible with UPMEM DPUs. Source: https://dl.acm.org/doi/10.1145/3767295.3803621 Dangling Pointers TPC-C is easy to partition; I wonder how well PIM-TIDE does on workloads that are not as partitionable. Also, this scheme doesn’t seem to allow for interactive transactions, how important are those in the real world? Thanks for reading Dangling Pointers! Subscribe for free to receive new posts. Local transactions execute entirely within a DPU Distributed transactions execute across multiple DPUs

0 views
Dangling Pointers 4 months ago

Yield Not Thy Core

Yield Not Thy Core Achilles Benetopoulos, Peter Alvaro, Andi Quinn, and Robert Soule EUROSYS’26 This paper describes a solution to the placement problem in distributed systems. If you model a computation as a directed graph, how do you optimally distribute the graph among a set of cooperating computers? The authors propose a dynamic placement system and implement it in Magpie . One common solution to the placement problem is to ship data over the network. For example, a set of compute nodes could access data via network requests to a separate set of nodes running Redis servers. At the opposite end of the spectrum, code can be shipped over the network. The canonical example is expressing computation as a SQL query which is sent to the node(s) that hold the relevant data. Magpie proposes a more fluid solution, where both code and data can move dynamically. In Magpie, an object represents data that is operated on. What makes Magpie objects unique is that pointers to data stored in an object are encoded as tuples. This allows Magpie to dynamically move objects around the system without invalidating pointers. The downside of this approach is that it prevents traditional libraries (that rely on raw pointers) from being used in user code. Magpie assumes a high degree of inter-object locality, so any given object is stored by exactly one node (i.e., a single object is never split between multiple nodes). User code is expressed in terms of nanotransactions and epics . A nanotransaction runs to completion on a single node and accesses a pre-specified set of objects. The Magpie runtime ensures that all objects accessed by a given nanotransaction are resident on a single node before executing the nanotransaction. The code for a nanotransaction is simple, because there is no need to query data over the network, and there is no need to deal with locking. If a hazard is present between two nanotransactions, they will execute serially. In Magpie, nanotransactions are written in Rust. An epic is a computation graph where each vertex is a nanotransaction and each edge is a data dependency. In contrast to nanotransactions, a single epic can be distributed across multiple nodes. Magpie schedules nanotransactions once all data dependencies are satisfied. Conflicts between concurrently running epics are handled via snapshot isolation . Any particular epic has a consistent view of each object and may abort in the event of a conflict. Scheduling and data movement are implemented hierarchically. A worker node can locally determine if it has ownership of all dependencies required for a nanotransaction. If this is the case, then the worker node executes the transaction immediately. Otherwise, the worker node uses a local ownership cache to try to determine if another node has all required dependencies and communicates with that node if possible. Failing that, scheduling is performed by a global orchestration node. Fig. 9 compares Magpie to memcached executing a workload that involves a user-specified read-modify-write operation: Source: https://dl.acm.org/doi/10.1145/3767295.3803616 Magpie is able to offer a lower latency because it is able to ship the entire read-modify-write operation to the server that holds the relevant data, rather than requiring multiple roundtrips. Some applications may benefit from being able to indicate that an object is rarely changed and thus can be distributed among multiple nodes at the same time. Thanks for reading Dangling Pointers! Subscribe for free to receive new posts.

0 views
Dangling Pointers 4 months ago

Pipeline Parallel Decompression

This isn’t a paper summary, but rather a description of a hobby experiment I’ve been hacking on ("research quality" code). This quote (attributed to either Anonymous or David Clark ) originally referred to networking, but applies to parallel programming as well: There is an old network saying: Bandwidth problems can be cured with money. Latency problems are harder because the speed of light is fixed—you can’t bribe God. Standard "cured with money" parallelization techniques (e.g., shared-nothing architectures, data parallelism) try to minimize cross-core communication. These hammers are great for hitting nails labeled: "improve throughput by throwing more cores at the problem”. Not everything is a nail. Important problems which cannot be solved with this kind of approach include: Parallel network packet processing in cases where load balancing schemes like RSS do not apply Parallel transaction processing when there is high contention between transactions Parallel encryption of a single stream of data Pipeline parallelism has the potential to provide "bribing God” solutions to some of these problems. A potential additional benefit that pipeline parallelism brings to the table is better usage of CPU caches because of a smaller working set. For example, if 8 cores cooperate to process 1 input file, the working set (input data, output data, intermediate data structures) is potentially 8 times smaller than the case where each core processes a separate input file. This caching advantage also applies to instruction caches, as pipeline parallelism distributes the computational steps of an algorithm across cores. Pipeline parallelism has some major drawbacks: Fine-grain synchronization/communication Load imbalance The purpose of this experiment is to put some numbers on the costs and benefits in a real-world application ( DEFLATE decompression). DEFLATE decompression is hard to parallelize because of two tight feedback loops: The position of encoded token in the input stream is not known until token is decoded (because input data is encoded with a variable length code). The output generated by a match (i.e., length & distance tuple) cannot be computed until some amount of previous output has been generated (because a match references previously generated output) A Negative Nancy might view these as problems, but a Positive Pipeliner views them as a guide for how to decompose the algorithm into pipeline stages. The general technique is to dedicate a pipeline stage to each of these feedback loops and whittle them down to be as tight as possible. The design I’ve landed on has three pipeline stages: , , and . The stage computes the length of each encoded token. It simply reads the next 13 bits from the input stream and uses them as an index into a lookup table. The inner loop looks like this: Note that in contrast to non-pipelined implementations, the only thing this code (and the lookup table) are concerned with is finding the length of each token, everything else is dealt with in another pipeline stage. Each iteration of this loop runs in about 8 clock cycles, and the lookup table fits in the L1 cache. The CPU cannot run multiple iterations of this loop in parallel due to the tight dependency chain. The input to the lookup stage is the encoded bits associated with each input token ( in the code above). These bits are used to perform another lookup (in a larger lookup table, stored in the L2 cache) which results in much more information about each token. Optimizing this stage is easy, because it doesn’t contain any tight feedback loops. The CPU can process multiple loop iterations in parallel, which enables it to hide the latency of accessing the L2. If necessary, it would be easy to split this pipeline stage into two. The inner loop looks like this: The structure contains metadata about the input token (literal value and/or information about a match). This data structure does not contain the exact distance associated with the match, the variables named deal with that detail from the DEFLATE spec. The stage writes literals and matches to the output buffer. This code leans on the CPU store-to-load forwarding hardware to deal with match operations which must read data that was recently produced. Each iteration of the inner loop performs a word-sized write of literal data, plus a 32B read and write to read and write match data. Actual store-to-load forwarding is rare, as most match distances are large. The Silesia Corpus contains commonly used files to benchmark compression algorithms. has English text with short matches whereas contains data dumps with longer matches. is an optimized library which can decompress roughly 2-3x faster than the standard . The following chart shows baseline performance on in a shared-nothing architecture where each CPU core decompresses a separate input file. There is one data point for each core count (1, 2, …, 8). As you would expect, throwing more cores at the problem improves throughput, at the cost of slight latency increase. If you want a more interesting tradeoff of throughput vs. latency, you have to bribe God. For example, say you are writing a decompression application. If the user requests a bulk decompression of 100 files, then the optimal choice may assign each file to a CPU core. But if the user requests to decompress a single file, then you would prefer to decompress using multiple CPU cores. And here is the same chart with the 3-stage pipeline implementation added in orange (compare it to the third blue dot from the left for a 3-core vs 3-core comparison): For a 37% cost in throughput, you get a 2x reduction in latency. Here is the chart for , which shows a similar story. Data-parallel throughput saturates at 6 cores. Pipeline parallelism allows a 2.6x latency reduction at the cost of 14% throughput. Dangling Pointers I think there is room for language/runtime support to improve performance of pipeline parallel algorithms on multicore CPUs (by reducing load imbalance). is bound by the chase stage, whereas is bound by the output stage. The programmer could supply multiple implementations of the pipeline (with some compiler help to reduce code duplication), and the runtime could dynamically switch between them depending on which stage is the bottleneck. High level synthesis tools are capable of automatic pipelining. Such techniques could be used to automatically generate many pipeline implementations for the runtime to choose between. The description above leaves out a few implementation details regarding the lookup tables. Because the lookup table data is spread across two cores (i.e., pipeline stages), there is enough room to store data for 2 Huffman tokens (2 literals, or a full match). This provides a large speedup compared to traditional implementations that store all data in the caches of a single core. Because the stage is throughput bound rather than latency bound, it can afford to access the lookup table via a layer of indirection. The 13 input bits are used to lookup a index, and that index is used to access the final data in another lookup table. The second lookup table has fewer entries, but each entry is larger. This reduces the total working set. This design leans heavily on CPU branch prediction. The code snippets shown earlier are for the common cases, with branches used to implement uncommon cases (e.g., a single encoded token that is wider than 13 bits). As long as those cases are rare, branch prediction does a great job of keeping the inner loops humming. An interesting puzzle arose during this experiment. I found that performance could swing widely (~10%) based on where the operating system located stacks of the various threads. The stack address would change from run to run because of ASLR . A little to offset the stack by a small amount would resolve this issue. It seems to be an important consideration when trying to maximize usage of the L1 cache. Subscribe now Parallel network packet processing in cases where load balancing schemes like RSS do not apply Parallel transaction processing when there is high contention between transactions Parallel encryption of a single stream of data Fine-grain synchronization/communication Load imbalance The position of encoded token in the input stream is not known until token is decoded (because input data is encoded with a variable length code). The output generated by a match (i.e., length & distance tuple) cannot be computed until some amount of previous output has been generated (because a match references previously generated output)

0 views
Dangling Pointers 4 months ago

SG-IOV: Socket-Granular I/O Virtualization for SmartNIC-Based Container Networks

SG-IOV: Socket-Granular I/O Virtualization for SmartNIC-Based Container Networks Chenxingyu Zhao, Hongtao Zhang, Jaehong Min, Shengkai Lin, Wei Zhang, Kaiyuan Zhang, Ming Liu, and Arvind Krishnamurthy ASPLOS'26 SR-IOV is a PCIe feature that enables a single device to expose multiple virtual functions, each of which appears as a separate device. This can be used to securely share one hardware device among multiple virtual machines. For containerized workloads, one would think that SR-IOV could be used to expose a virtual NIC to each container. This would save CPU cycles as network virtualization would be handled by the NIC rather than software. The trouble is that SR-IOV doesn’t scale to high container counts (they top out on the order of 100 of virtual functions per physical NIC). This paper introduces Socket-Granular I/O Virtualization (SG-IOV) which enables NIC virtualization at the socket level . The authors have an implementation working on NVIDIA BlueField-3. The key assumption that SG-IOV makes is that container networking uses stream sockets (e.g., TCP) rather than datagram sockets (e.g., UDP). In other words, software running inside a container wants to reliably send a stream of bytes rather than a stream of packets. Streams are transmitted through Warp Pipes , which are simply ring buffers (comprising a base address, head and tail pointers). Warp Pipes can be stored in host memory or NIC local memory. A socket is associated with one or more dedicated Warp Pipes. A server can have thousands of Warp Pipes, because the low-level NIC hardware (which may have scalability limits) doesn’t directly interact with Warp Pipes. Updates to head and tail pointers are not performed directly via memory writes, instead they are communicated through a Cross-FIFO . A single message in a cross-FIFO contains three fields: Ring buffer ID (which Warp Pipe to update) Which pointer to update (head or tail) The new value of the head or tail pointer There are multiple implementations of the Cross-FIFO interface. For example, the authors use a PCIe-based implementation for host→NIC communication, and a RDMA based implementation for NIC→NIC communication. Each cross-FIFO uses limited NIC resources; therefore many sockets share a single cross-FIFO, which is OK because the messages they contain are coarse-grain (e.g., increment tail pointer by 16KiB). Say an application in container A on host 1 wants to send 8KiB of data to an application in container B on host 2. The flow looks like this: Host CPU (host 1): The application calls the standard socket API The payload is copied into the Warp Pipe associated with the socket A message is enqueued into a PCIe cross-FIFO, indicating that the head pointer should be incremented by 8KiB ARM CPU running on the NIC (host 1): Dequeue the message from the cross-FIFO, update the head pointer Enqueue task descriptors to the low-level NIC hardware to read the payload data (i.e., local DMA) from the Warp Pipe and store it in the correct Warp Pipe on host 2 (i.e., RDMA); note that these task descriptors can configure the NIC to perform appropriate network virtualization After the data has been transmitted, enqueue a message to the NIC on host 2 (via a RDMA-based cross-FIFO) to update the head pointer for the Warp Pipe on host 2 ARM CPU running on the NIC (host 2): Dequeue the message from the cross-FIFO Update the local head pointer Enqueue a message to the host CPU (host 2) via a PCIe-based cross-FIFO, indicating that the head pointer should be incremented Host CPU (host 2): Dequeue the message from the PCIe-based cross-FIFO Send data to the application when it calls And a similar sequence would update tail pointers. The key design principles of SG-IOV are: All state is tracked by software running on the host CPU and ARM CPUs running on the NIC All low-level NIC hardware is used in a stateless manner, and is multiplexed across many sockets A benefit of this design is that it is flexible enough to handle the loopback case efficiently. Section 8 of the paper shows that SG-IOV enables the use of hardware-accelerated network virtualization (saving host CPU cycles) while offering latency and bandwidth that is competitive with other container-based network virtualization systems. Figs. 17 and 18 show latency and bandwidth numbers for a few benchmarks: Source: https://dl.acm.org/doi/10.1145/3779212.3790218 Dangling Pointers It seems like one downside of this system is excessive memory usage. If each socket has dedicated large ring buffers, then DDIO may not be effective (see here , here , and here for other papers that discuss this problem). Subscribe now Ring buffer ID (which Warp Pipe to update) Which pointer to update (head or tail) The new value of the head or tail pointer The application calls the standard socket API The payload is copied into the Warp Pipe associated with the socket A message is enqueued into a PCIe cross-FIFO, indicating that the head pointer should be incremented by 8KiB Dequeue the message from the cross-FIFO, update the head pointer Enqueue task descriptors to the low-level NIC hardware to read the payload data (i.e., local DMA) from the Warp Pipe and store it in the correct Warp Pipe on host 2 (i.e., RDMA); note that these task descriptors can configure the NIC to perform appropriate network virtualization After the data has been transmitted, enqueue a message to the NIC on host 2 (via a RDMA-based cross-FIFO) to update the head pointer for the Warp Pipe on host 2 Dequeue the message from the cross-FIFO Update the local head pointer Enqueue a message to the host CPU (host 2) via a PCIe-based cross-FIFO, indicating that the head pointer should be incremented Dequeue the message from the PCIe-based cross-FIFO Send data to the application when it calls All state is tracked by software running on the host CPU and ARM CPUs running on the NIC All low-level NIC hardware is used in a stateless manner, and is multiplexed across many sockets

0 views