EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA

EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA August 2025 AUTHOR(S): Mohamad Khaled Charaf American University of Beirut SUPERVISOR(S): Andrea Bocci - Aurora Perego CERN openlab Report x/2025 ABSTRACT This report details the porting and optimization of the CLUE clustering algorithm for the CMS experiment onto an Altera Agilex 7 FPGA. The primary goal was to explore the feasibility of using a SYCL-based High-Level Synthesis workflow within the alpaka performance portability framework to accelerate high-energy physics workloads on spatial architectures. Through an iterative process of analyzing compiler reports and refactoring code, the CLUE kernels were optimized to mitigate memory and data dependencies. The final design is synthesized to operate at 480 MHz, with most critical loops achieving a high-throughput Initiation Interval (II) of 1. Static analysis reveals the implementation is memory-centric, consuming 30% of the FPGA’s on-chip RAM while making minimal use of logic and DSP resources. Although a hardware fault with the accelerator card precluded final performance validation on physical hardware, the synthesis and simulation results establish a robust and highly-optimized baseline, demonstrating the viability of this approach for future heterogeneous computing at CMS. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 1 CERN openlab Report x/2025 TABLE OF CONTENTS 1 INTRODUCTION 3 2 BACKGROUND: FPGAs AND SYCL 3 2.1 WhatisanFPGA?.................................. 3 2.2 WhyFPGAsatCMS? ................................ 4 2.3 WhatisSYCL?.................................... 4 2.4 The Target FPGA: Altera Agilex 7 . . . . . . . . . . . . . . . . . . . . . . . . . 5 2.5 KeyFPGAConcepts................................. 5 2.5.1 Pipelining and Maximum Frequency (fMAX) ................ 5 2.5.2 Datapath and Occupancy . . . . . . . . . . . . . . . . . . . . . . . . . . 6 2.5.3 Scheduling................................... 6 2.5.4 Leveraging Parallelism . . . . . . . . . . . . . . . . . . . . . . . . . . . . 7 3 THE CLUE ALGORITHM 8 3.1 WhatisCLUE?.................................... 8 3.2 The Different Kernels of CLUE . . . . . . . . . . . . . . . . . . . . . . . . . . . 9 3.2.1 Spatial Indexing and Histogram Creation . . . . . . . . . . . . . . . . . . 9 3.2.2 Local Density (ρ)Calculation ........................ 9 3.2.3 Nearest Higher (δ)Calculation........................ 9 3.2.4 Seed Identification and Cluster Expansion . . . . . . . . . . . . . . . . . 11 4 METHODOLOGY AND IMPLEMENTATION 11 4.1 FPGA Development Workflow and Device Selection . . . . . . . . . . . . . . . . 11 4.2 Choosing the Right Parallelism Model . . . . . . . . . . . . . . . . . . . . . . . 12 4.3 Task Graph of the CLUE Implementation . . . . . . . . . . . . . . . . . . . . . 13 4.4 Memory Management and Data Transfer with SYCL USM . . . . . . . . . . . . 13 4.5 ChallengesEncountered ............................... 15 5 Kernel Optimization and Performance Tuning 15 5.1 Optimizing compute_histogram: On-Chip Memory and Caching . . . . . . . . . 16 5.2 Optimizing CalculateDensity and CalculateDistanceToHigher: Refactoring DataDependencies .................................. 16 5.3 Optimizing FindClusters: Moving Counters On-Chip . . . . . . . . . . . . . . 17 5.4 Addressing AssignClusters: A Challenging Control-Flow Problem . . . . . . . 18 5.5 Reducing Memory Traffic with a Pipe-Based Design . . . . . . . . . . . . . . . . 18 6 RESULTS AND ACHIEVEMENTS 19 6.1 Functional Implementation of CLUE Kernels . . . . . . . . . . . . . . . . . . . . 19 6.2 FPGA Resource Utilization Analysis . . . . . . . . . . . . . . . . . . . . . . . . 20 6.3 Performance Characterization and Bottlenecks . . . . . . . . . . . . . . . . . . . 20 7 CONCLUSION 21 8 REFERENCES 21 EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 2 CERN openlab Report x/2025 1 INTRODUCTION The unprecedented data volumes from the CERN Large Hadron Collider (LHC) and its upcoming High-Luminosity (HL-LHC) upgrade present a significant computational challenge for the Compact Muon Solenoid (CMS) experiment. To meet the tight, millisecond-level execution time budget for event reconstruction, CMS has actively explored heterogeneous computing platforms. This effort began with the adoption of NVIDIA GPUs in 2017, leading to a substantial increase in event processing throughput. However, the proliferation of different hardware architectures, including GPUs from various vendors (NVIDIA, AMD, Intel) and other accelerators like FPGAs,introduced a new challenge: code portability. Developing and maintaining separate codebases for each architecture duplicates effort and introduces bugs. To address this, CMS has adopted alpaka, a C++ template library that enables performance portability, allowing a single source code to run on multiple accelerators without extensive rewrites. This project continues that effort by exploring the integration of the Intel oneAPI toolchain for FPGA programming into the alpaka framework. By porting the CLUE clustering algorithm to an FPGA, we aim to investigate the feasibility and performance of this approach, contributing to CMS’s long-term heterogeneous computing strategy. Traditionally, FPGAs are programmed using a Hardware Description Language (HDL) like Verilog or VHDL. However, a recent trend is to use High-Level Synthesis (HLS), which allows developers to use higher-level languages like C++. This approach can significantly reduce design time and increase portability, which aligns perfectly with the goals of this project. 1 2 BACKGROUND: FPGAs AND SYCL 2.1 What is an FPGA? Unlike conventional CPUs and GPUs, which have a fixed Instruction Set Architecture (ISA), Field-Programmable Gate Arrays (FPGAs) are spatial architectures. This means that a program is mapped directly onto a configurable hardware fabric, creating a physical dataflow pipeline. Different parts of the program execute on different regions of the device simultaneously, enabling extreme parallelism and compelling energy efficiency. While an applicationspecific integrated circuit (ASIC) generally offers higher performance, its immutable design requires significant time and financial investment. FPGAs provide a more flexible and costeffective alternative, as they can be reprogrammed for new applications. As can be seen in Figure 1, an FPGA consists of a grid of programmable logic blocks, often called Adaptive Logic Modules (ALMs), alongside specialized blocks for Digital Signal Processing (DSP) and Random Access Memory (RAM). These components are connected via configurable routing interconnects. A simplified ALM contains a lookup table (LUT), which can implement any arbitrary Boolean logic, and a register, which is the most basic storage element that synchronizes data to a clock signal. DSP blocks are hardwired to perform arithmetic operations like multiplication and addition efficiently, while RAM blocks provide dense, on-chip data storage. 1The source code for this project is publicly available at: https://github.com/MohamadKhaledCharaf/ CLUE-SYCL EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 3 CERN openlab Report x/2025 (a) High-level architectural view showing the grid of programmable blocks. (b) Simplified view of a Register Figure 1: Views of an FPGA’s internal structure.[3] 2.2 Why FPGAs at CMS? The CMS experiment constantly seeks new ways to accelerate its data processing. This project evaluates FPGAs as a compelling choice for accelerating specific, computationally intensive algorithms within the High Level Trigger (HLT) reconstruction software. While the HLT processes a data rate that is significantly reduced by the initial trigger stages, it must execute complex algorithms within a latency budget on the order of seconds. The inherent parallelism of FPGAs allows them to perform these demanding tasks with high throughput and predictable latency, making them an ideal solution for offloading specific workloads from the HLT’s CPU farm to improve the overall system efficiency. 2.3 What is SYCL? SYCL is an open-standard, C++-based programming model that enables code to be written once and deployed across various hardware accelerators, including CPUs, GPUs, and FPGAs. It provides an abstraction layer that allows developers to write platform-agnostic code, which is then compiled for a specific target architecture. The Intel oneAPI DPC++ compiler, an implementation of the SYCL standard, provides specific extensions and features for FPGA development. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 4 CERN openlab Report x/2025 Figure 2: Block diagram of the BittWare IA-840f board featuring the Intel Agilex 7 AGF027 FPGA. [1] 2.4 The Target FPGA: Altera Agilex 7 At the heart of the BittWare IA-840f (see Figure 2) is the Intel®Agilex™7 AGF027 FPGA. According to the official device overview [2], this FPGA provides a rich set of programmable resources, including 2,696K Logic Elements (LEs), 3,840 high-performance 18x19 DSP blocks, and 51.3 Mb of embedded M20K RAM. In addition, the device features numerous hardened intellectual property (IP) blocks to support high-speed interfaces, providing native support for a PCIe Gen4 ×16 interface. Complementing the on-chip resources, the IA-840f board features a high-performance memory subsystem configured according to the Board Support Package (BSP). The system utilizes two 512-bit wide DDR4 channels that support memory interleaving with a 4 kB stride to maximize throughput. According to the optimization report, the theoretical maximum bandwidth the BSP can deliver is 42.7 GB/s. However, the report notes that a more realistic peak bandwidth is approximately 90% of this value (around 38.4 GB/s) due to overhead from the interconnect and memory controller. The BSP further characterizes the performance with a maximum read bandwidth of 42.7 GB/s and a write bandwidth of 21.6 GB/s. Furthermore, a separate host-shared memory region accessible over the PCIe interface delivers up to 30 GB/s, enabling efficient CPU–FPGA communication. 2.5 Key FPGA Concepts 2.5.1 Pipelining and Maximum Frequency (fMAX) In a digital circuit, the maximum clock frequency (fMAX) is limited by the propagation delay of signals through combinational logic between two consecutive registers. The longest path, known as the critical path, determines the speed of the entire circuit. Pipelining is a technique used to increase fMAX by inserting additional registers into the critical path. This breaks a long combinational path into shorter, faster stages, allowing the circuit to run at a higher clock frequency. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 5 CERN openlab Report x/2025 2.5.2 Datapath and Occupancy Adatapath is the network of registers and logic that performs computations. The occupancy of a datapath refers to the proportion of its pipeline stages that contain valid data at any given time. Unoccupied stages, known as bubbles, are analogous to no-operation (no-op) instructions for a CPU and reduce throughput. While deeper pipelining increases fMAX and peak throughput as illustrated by the example in Figure 3, it also increases latency, as more clock cycles are required to fill the pipeline initially. This trade-off is crucial: for applications processing large, continuous streams of data, the initial latency is easily amortized, and the gains in throughput are significant. Figure 3: An example of a datapath where a pipeline register is inserted to break a long combinational logic path into two shorter stages, enabling a higher fMAX [3] 2.5.3 Scheduling Scheduling is the process of assigning each operation in an FPGA design to a specific clock cycle. The goal is to create a deeply pipelined architecture that allows multiple operations to execute concurrently, maximizing throughput. Dynamic Scheduling The Intel DPC++ Compiler creates a dynamically scheduled, pipelined datapath. In this model, an operation does not pass data to the next stage until the successor signals it is ready. This is achieved using handshaking control logic with ‘valid‘ and ‘stall‘ signals. This mechanism is essential for handling variable-latency operations (e.g., memory access) without creating pipeline bubbles. Clustering the Datapath (Figure 4)To reduce the area overhead from handshaking logic, the compiler groups fixed-latency operations into clusters. This approach confines handshaking interfaces to the cluster boundaries, simplifying the control logic. Cluster Types The compiler supports two cluster types: •Stall-Enable Cluster (SEC): Propagates the ‘stall‘ signal to all internal stages in parallel. It is area-efficient but can only remove leading bubbles from the pipeline. •Stall-Free Cluster (SFC): Includes an exit FIFO that allows the cluster to continue processing even when stalled downstream. It is more effective at removing bubbles but requires more area. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 6 CERN openlab Report x/2025 Figure 4: Illustration of datapath clustering. Fixed-latency operations are grouped, and handshaking logic is only required between clusters and variable-latency units.[3] 2.5.4 Leveraging Parallelism FPGAs achieve high performance by exploiting parallelism on a massive scale. Instead of executing a sequence of instructions like a CPU, FPGAs create a physical circuit customized for the algorithm. This enables two primary forms of parallelism: pipelining (temporal parallelism) and vectorization (spatial parallelism). Pipelining As discussed previously, pipelining breaks a computation into a series of smaller stages. The key to parallelism is that new data can enter the first stage of the pipeline on every clock cycle, even before the first piece of data has finished its entire journey. As illustrated in Figure 5, this approach keeps all stages of the hardware datapath occupied simultaneously, maximizing occupancy and throughput. The compiler applies this principle to loops, aiming for an Initiation Interval (II) of 1, where a new loop iteration can begin every single clock cycle. Figure 5: Illustration of a pipelined datapath. New data (t0, t1, t2...) enters the pipeline on subsequent clock cycles, keeping the hardware fully occupied.[3] EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 7 CERN openlab Report x/2025 Vectorization To further increase throughput, the compiler can vectorize the hardware. As shown in Figure 6, vectorization creates multiple identical copies of the entire pipelined datapath. These parallel pipelines can then process multiple data items simultaneously. For example, three pipelines can process three data items in the time it takes one pipeline to process one, effectively tripling the throughput. This powerful technique comes at the direct cost of a proportional increase in FPGA resource usage, as the hardware logic is physically duplicated on the chip. Figure 6: Illustration of a vectorized datapath. The pipeline is duplicated to process multiple data items in parallel, multiplying throughput at the cost of additional FPGA area.[3] 3 THE CLUE ALGORITHM 3.1 What is CLUE? CLUE (CLUstering of Energy) is a fast, fully parallelizable density-based clustering algorithm developed for the high-occupancy environments of the CMS High Granularity Calorimeter (HGCAL) [6]. Inspired by the CFSFDP algorithm [5], CLUE is optimized for scenarios where the number of clusters is much larger than the average number of points per cluster. Its core feature is a fixed-grid spatial index for efficient neighborhood queries, which enables linear scalability. The algorithm is decomposed into several highly parallelizable kernels, making it well-suited for heterogeneous architectures like FPGAs and GPUs. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 8 CERN openlab Report x/2025 5. Memory Deallocation (free_device)Finally, once all processing is complete, the memory allocated on the FPGA is explicitly released using sycl::free. This is a critical step to prevent memory leaks on the device and return the accelerator’s resources to the system. This meticulous management of memory ensures that the application is both robust and performant. The deallocation requires the original pointer and the queue for context: sycl::free(d_points.x, q); The pointer d_points.x identifies the memory to be released. The queue qis again necessary to tell the SYCL runtime which device’s memory manager is responsible for this pointer and must perform the deallocation. 4.5 Challenges Encountered Porting the CLUE algorithm to an FPGA using SYCL presented several challenges: •Report-driven Optimization: Achieving an efficient spatial pipeline is highly dependent on the DPC++ compiler. A significant part of the work involved analyzing compiler reports to identify and mitigate performance bottlenecks, such as loops with a high initiation interval (II), which directly impacts throughput. •Data Dependencies: While FPGAs handle pipeline dependencies well, a naive implementation can still lead to stalls. Kernels had to be carefully structured to minimize data hazards and maximize datapath occupancy. •Resource Utilization: A critical aspect of FPGA design is managing hardware resources. Designs that overutilize any single resource type (e.g., ALMs,DSPs, or RAM) exceeding 90% can be difficult for the compiler to place and route, which often degrades the final operating frequency (fMAX). However, given the substantial capacity of the target Agilex 7 FPGA, this did not prove to be a significant constraint for implementing the CLUE algorithm. The kernels fit comfortably within the available resources, leaving significant capacity to spare. 5 Kernel Optimization and Performance Tuning Achieving high throughput on an FPGA requires creating a deeply pipelined hardware design where the main processing loops can initiate a new iteration every clock cycle. This ideal state is known as an Initiation Interval (II) of 1. The initial implementation of the CLUE kernels, while functionally correct, suffered from high II values due to various data and memory dependencies identified in the Intel DPC++ Loop Analysis report. This section details the stepby-step optimizations applied to each kernel to reduce their II and improve overall performance. A crucial general optimization was the consistent use of compiler hints. The attribute [[intel::kernel_args_restrict]] was applied to kernel definitions to inform the compiler that pointer arguments do not alias (overlap in memory). This resolves compiler conservatism and prevents it from serializing memory accesses unnecessarily. Furthermore, explicitly casting USM pointers to sycl::ext::intel::device_ptr within the kernel body was vital. This tells the compiler that the data resides exclusively on the device, eliminating ambiguity and preventing the generation of costly, stall-prone logic to handle potential host-side access. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 15 CERN openlab Report x/2025 5.1 Optimizing compute_histogram: On-Chip Memory and Caching The compute_histogram kernel’s primary bottleneck was a severe loop-carried memory dependency. The loop iterates through points and increments a counter for the corresponding tile. Initially, these counters resided in off-chip DDR4 memory. •Problem: The read-modify-write operation on an off-chip memory counter created a long-latency Read-After-Write (RAW) hazard. Each loop iteration had to wait for the previous iteration’s memory write to complete before it could read the updated counter value. This resulted in an extremely high II, reported as 1157 when compiled with a Board Support Package (BSP) that models realistic memory latencies. •Solution 1 (On-Chip Memory): The first and most impactful optimization was to move the tile-size counters from off-chip DDR4 to on-chip Block RAM (BRAM). This was achieved by declaring a local array, int tileSizes[NLAYERS][T::nTiles];, within the single_task kernel. This dramatically reduced the memory access latency, dropping the loop II from over 1000 to just 2. Figure 13: Loop Carried Dependancy with II=2 •Remaining Bottleneck (II=2): Although much improved, the II remained at 2. The Loop Analysis report confirmed that the read-modify-write cycle to the on-chip memory still constituted a loop-carried dependency that took more than a single clock cycle to resolve as can be seen in Figure 13. •Solution 2 (On-Chip Caching): To achieve the ideal II of 1, the "On-chip Memory Cache Technique" described in the Intel FPGA handbook is necessary [3]. This technique breaks the dependency by creating a small, register-based cache for recently accessed counter values. A read-modify-write operation on a register can complete in a single cycle. By servicing memory requests from this cache, the dependency chain is broken. While a manual attempt was made using the [[intel::ivdep]] attribute to assert no dependency, the most robust solution involves using Intel’s provided caching classes, which is a clear path for future work. 5.2 Optimizing CalculateDensity and CalculateDistanceToHigher: Refactoring Data Dependencies The density and distance calculation kernels both feature a triply nested loop structure to iterate over neighboring points. The initial implementation suffered from a classic loop-carried data dependency that serialized the outer loops. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 16 CERN openlab Report x/2025 •Problem: An accumulator variable (e.g., rhoi in CalculateDensity) was updated in the innermost loop. Because each iteration of an outer loop depended on the final accumulated value from all iterations of its inner loops, the compiler could not pipeline them. The outer loop had to stall until the entire nested structure below it completed, leading to poor hardware utilization and high latency. •Solution (Dependency Refactoring): The solution was to decouple the computations by introducing local accumulator variables at each loop level. For example, in CalculateDensity, variables rho2 and rho1 were introduced for the inner and middle loops, respectively. The innermost loop accumulates into rho2, which is then added to rho1 at the end of the middle loop. Finally, rho1 is added to the main accumulator rhoi at the end of the outer loop. This refactoring breaks the dependency chain, allowing the compiler to pipeline all three nested loops with an II of 1. The attribute [[intel::initiation_interval(1)]] was added to explicitly guide the compiler towards this goal. The optimized code structure for CalculateDensity is shown below: 1for(int i=0; i <numberOfPoints ; ++i){ 2float rhoi = 0.; 3// ... setup for point i ... 4[[intel::initiation_interval(1)]] 5for(int xBin =...){ 6float rho1 = 0; 7[[intel::initiation_interval(1)]] 8for(int yBin =...){ 9// ... 10 float rho2 = 0; 11 [[intel::initiation_interval(1)]] 12 for(int binIter = 0; binIter <binSize; ++binIter){ 13 // ... calculation ... 14 rho2 += ... ; 15 } 16 rho1 += rho2; 17 } 18 rhoi += rho1; 19 } 20 d_points.rho[i] =rhoi; 21 } The same principle was applied to the CalculateDistanceToHigher kernel for the deltai and nearestHigheri variables, successfully enabling full pipelining. 5.3 Optimizing FindClusters: Moving Counters On-Chip The FindClusters kernel had a similar issue to compute_histogram. Its main loop determines if points are seeds or followers and appends them to corresponding lists. •Problem: The counters for the number of seeds and the number of followers for each point were stored in off-chip memory. Modifying these counters created a high-latency memory dependency, resulting in an initial II of 3296. •Solution: As with the histogram kernel, the solution was to move these counters into on-chip memory. A local variable seeds_size and a local array follower_counters[] EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 17 CERN openlab Report x/2025 were declared inside the kernel. This change drastically reduced the dependency latency, bringing the II down to 2. Like the histogram kernel, this could be further optimized to II=1 using an on-chip caching strategy. 5.4 Addressing AssignClusters: A Challenging Control-Flow Problem The final kernel, AssignClusters, which traverses the DAG of followers from each seed, presents the most significant optimization challenge due to its complex and data-dependent control flow. •Problem 1 (Loop-Carried Stack Dependency): The kernel uses a local stack to perform a Depth-First Search (DFS) traversal of the follower graph. Pushing and popping from this stack within a while loop creates a strong loop-carried dependency on the stack pointer and its contents, making pipelining impossible. •Problem 2 (Dynamic Inner Loop): The inner loop iterates over the followers of the current point being processed. The number of followers is data-dependent and only known at runtime. The compiler cannot create a statically scheduled pipeline for an outer loop when the inner loop has a variable, unpredictable trip count. This forces the outer loop to wait for the inner loop to fully complete, preventing pipelining. •Proposed Solution (Architectural Redesign): A fundamental redesign is required to parallelize this kernel effectively on an FPGA. A promising approach is to use a producerconsumer pattern with SYCL pipes. A producer kernel could traverse the initial list of seeds, pushing pairs of {point, clusterID} into a hardware pipe. Multiple instances of a consumer kernel could then read from this pipe in parallel. Each consumer would process a point, assign its followers the same cluster ID, and push new {follower, clusterID} pairs back into the pipe. This architecture would transform the serial DFS traversal into a parallel, pipelined breadth-first-style traversal, better suited to the FPGA’s spatial architecture. However, before attempting this significant architectural change, we planned to first try a simpler experiment: fixing the inner loop’s dynamic trip count to its maximum value of 32 and observing the performance once the physical board becomes available. 5.5 Reducing Memory Traffic with a Pipe-Based Design A primary performance bottleneck in many FPGA applications is the latency and bandwidth limitations of accessing off-chip DDR4 memory. A powerful technique to mitigate this is to create a streaming architecture using SYCL pipes. Pipes allow data to flow between kernels directly using on-chip resources, following the principle of “load once, process many, store once.” The ideal implementation would be a fully-pipelined design where data is loaded from DDR4 a single time, a producer kernel streams it through a chain of consumer kernels connected by pipes, and a final kernel writes the results back. This would maximize on-chip data reuse and hide memory latency. However, this ideal architecture is not fully compatible with the CLUE algorithm. Kernels like CalculateDensity and CalculateDistanceToHigher perform data-dependent neighborhood searches. For each point they process, they must look up data from specific, non-sequential tiles in the spatial index, which resides in global memory. This requirement for random access to a large data structure prevents a purely passive, stream-based processing model. As a pragmatic compromise, a hybrid approach can be explored. While a fully-piped design is infeasible, the main outer loop that is present in most kernels, which iterates over all input EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 18 CERN openlab Report x/2025 points, can be structured as a streaming pipeline. In this model, illustrated in Figure 14, a producer kernel could iterate through the points and push them into a pipe. The subsequent kernel would act as a consumer, reading from that pipe to get its work item. Although this consumer kernel still needs to perform its own random-access lookups to global memory, this architecture organizes the high-level data flow more efficiently and aligns better with the FPGA’s spatial nature. Figure 14: Top: A traditional, memory-bound architecture where each kernel repeatedly accesses slow global memory. Bottom: An optimized streaming architecture using fast, on-chip SYCL pipes to reduce memory traffic to a single read and write, boosting performance.[3] 6 RESULTS AND ACHIEVEMENTS This section presents the results from porting and optimizing the CLUE algorithm for the Intel Agilex 7 FPGA. Our efforts focused on achieving functional correctness, analyzing resource utilization, and characterizing the performance of the final design based on the Intel DPC++ compiler’s static analysis and simulation reports. 6.1 Functional Implementation of CLUE Kernels The primary achievement of this project was the successful porting and optimization of the core CLUE computational kernels to SYCL for FPGAs. We successfully implemented and verified the functionality of: •Kernel 1: compute_histogram:The spatial indexing kernel correctly partitions the input points. •Kernel 2: kernel_calculate_density:The density calculation produces results consistent with the original CPU implementation. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 19 CERN openlab Report x/2025 •Kernel 3: CalculateDistanceToHigher:The nearest-higher separation calculation was also successfully ported. •Kernel 4 & 5: FindClusters and AssignClusters:The cluster finding and assignment kernels were implemented. The complex control-flow of AssignClusters remains a topic for future architectural redesign once the board becomes available. 6.2 FPGA Resource Utilization Analysis Static analysis reports generated by the Intel DPC++ compiler provide a detailed breakdown of the hardware resources required by our synthesized kernels. The total resource footprint for the five CLUE kernels (the “Kernel System”) is summarized in Table 1. Table 1: Total FPGA resource utilization for the complete CLUE Kernel System. Resource Used Utilization of Device (%) ALMs 81,003 4% Registers 162,300 3% RAM Blocks (M20K) 3,951 30% DSP Blocks 37 <1% The analysis reveals that the implementation is heavily memory-centric. The kernels consume a significant portion of the on-chip RAM (30%), which is expected given the optimization strategy of moving counters and data structures from off-chip DDR4 to on-chip memory. In contrast, the utilization of general-purpose logic (ALMs at 4%) and DSP blocks (<1%) is very low. This indicates that the CLUE algorithm is dominated by memory access and simple comparisons rather than complex arithmetic computations. 6.3 Performance Characterization and Bottlenecks Hardware Testing Status Deployment and performance measurement on the physical BittWare IA-840f card were precluded by a hardware fault discovered during setup. A support ticket has been opened with the vendor to resolve the issue. Consequently, all performance results presented in this report, including operating frequency and loop throughput, are based on the compiler’s static analysis and cycle-accurate simulation reports, not physical hardware measurements. Loop Throughput Analysis The iterative optimization process described in Section 5 was highly successful, with the final design synthesized to operate at a frequency of 480 MHz. The compiler’s Loop Analysis report, summarized in Table 2, shows that nearly all performance-critical loops achieved the ideal Initiation Interval (II) of 1. The two exceptions, in ComputeHistogram and FindClusters, achieved an II of 2 due to the remaining memory dependency on on-chip RAM, as discussed in the optimization section. The loops in AssignClusters remain unpipelined due to their challenging control-flow, marking a clear target for future architectural redesign. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 20 CERN openlab Report x/2025 Table 2: Loop Analysis summary for the optimized CLUE kernels. Kernel Loop Location Pipelined II Bottleneck / Notes FindClusters FindClusters.B2 (line 365) Yes 2 Memory dependency FindClusters.B3 (line 387) Yes 1 - AssignClusters AssignClusters.B1 (line 442) No n/a Out-of-order inner loop AssignClusters.B3 (line 449) No n/a Loop unpipelined AssignClusters.B4 (line 456) Yes 1 - CalculateDensity CalculateDensity.B1 (line 227) Yes 1 - CalculateDensity.B2 (line 236) Yes 1 - CalculateDensity.B5 (line 239) Yes 1 - CalculateDensity.B7 (line 246) Yes 1 - ComputeHistogram ComputeHistogram.B2 (line 195) Yes 2 Memory dependency ComputeHistogram.B3 (line 198) Yes 1 - ComputeHistogram.B5 (line 199) Yes 1 - CalculateDistanceToHigher CalculateDistanceToHigher.B1 (line 283) Yes 1 - CalculateDistanceToHigher.B2 (line 294) Yes 1 - CalculateDistanceToHigher.B5 (line 298) Yes 1 - CalculateDistanceToHigher.B7 (line 306) Yes 1 - 7 CONCLUSION This project successfully established a design and optimization workflow for accelerating the CMS CLUE reconstruction algorithm on an Altera Agilex 7 FPGA using modern SYCL. While a board-level fault precluded final performance validation on physical hardware, the synthesis and simulation results are highly encouraging, validating the feasibility of this approach and establishing a robust baseline for future development. The key achievement of this work was the practical application of targeted FPGA optimization techniques. By analyzing compiler reports, we identified and resolved critical performance bottlenecks. High-latency memory dependencies were mitigated by moving counters to onchip RAM, and loop-carried data dependencies were resolved by refactoring accumulator logic. These efforts resulted in a design that is synthesized to operate at 480 MHz, with most critical loops achieving an ideal Initiation Interval of 1, as confirmed by static analysis. The resource utilization reports further confirmed that the CLUE implementation is memory-centric, consuming 30% of the FPGA’s on-chip RAM while making minimal use of logic and DSP resources. This exploration has paved a clear path for future work, including the implementation of advanced on-chip caching to achieve a perfect II=1 on all memory-bound loops and an architectural redesign of the challenging cluster assignment kernel. Overall, this project validates the potential of FPGAs as a viable platform for accelerating key components of LHC data processing. The use of high-level synthesis with SYCL, within a portability library like alpaka, provides a sustainable path forward for harnessing the unique advantages of FPGAs in a heterogeneous computing environment. Acknowledgements I would like to extend my sincere gratitude to my supervisors, Andrea Bocci and Aurora Perego, for their invaluable guidance, support, and mentorship throughout this project. I am deeply thankful to the entire CERN OpenLab team for providing this incredible summer student opportunity and for fostering a stimulating research environment. 8 REFERENCES EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 21 CERN openlab Report x/2025 References [1] BittWare, a Molex company. IA-840f PCIe FPGA Card Featuring Intel Agilex FPGA. Accessed: 2024-09-01. 2024. url:https://www.bittware.com/products/ia-840f/. [2] Intel Corporation. Agilex™7 FPGAs and SoCs Device Data Sheet: F-Series and I-Series. Tech. rep. DS-1060. Document ID 683301. Available at: https://cdrdv2-public.intel. com/670300/ag_datasheet-683301-670300.pdf. Intel Corporation, Aug. 2025. [3] Intel Corporation. Intel®oneAPI DPC++/C++ Compiler FPGA Optimization Guide. Document Number: 2024.1. Intel Corporation. May 2024. url:https://www.intel. com / content / www / us / en / develop / documentation / oneapi - fpga - optimization - guide/top.html. [4] James Reinders et al. Data Parallel C++: Programming Accelerated Systems Using C++ and SYCL. 2nd ed. Apress, 2023. isbn: 978-1-4842-9690-5. doi:10.1007/978-1-48429691-2. [5] Alex Rodriguez and Alessandro Laio. “Clustering by fast search and find of density peaks”. In: Science 344.6191 (2014), pp. 1492–1496. [6] Marco Rovere et al. CLUE: A Fast Parallel Clustering Algorithm for High Granularity Calorimeters in High Energy Physics. 2020. arXiv: 2001.09761 [physics.ins-det]. [7] Marco Rovere et al. Visualisation of the main steps of the CLUE algorithm. CERN Document Server. 2020. url:https://cds.cern.ch/record/2709269/plots. EXPLORING FPGA ACCELERATION FOR CMS RECONSTRUCTION USING ALPAKA 22