1. Bibliographic Information

• Title: Lessons Learned Migrating CUDA to SYCL: A HEP Case Study with ROOT RDataFrame
• Authors:
  • Jolly Chen (University of Amsterdam, CERN)
  • Monica Dessole (CERN)
  • Ana Lucia Varbanescu (University of Twente)
• Journal/Conference: The paper is formatted for an ACM Conference (Conference'17 template), indicating its intent for publication in a peer-reviewed conference.
• Publication Year: 2024
• Abstract: The paper addresses the challenge of analyzing petabytes of data from CERN's Large Hadron Collider (LHC). It focuses on ROOT, a C++ data analysis framework, and its high-level interface, RDataFrame. To leverage modern heterogeneous computing systems, the authors explore moving beyond RDataFrame's CPU-only parallelism. They detail their experience migrating a core High Energy Physics (HEP) operation, histogramming, from an initial CUDA implementation to the more portable SYCL standard. The study presents a rigorous performance comparison between the native CUDA version and two SYCL compilers (AdaptiveCpp and DPC++), identifying performance bottlenecks and the methods used to find them. The paper concludes with actionable insights for developers working with SYCL in large, complex codebases.
• Original Source Link: The paper is available as a preprint on arXiv.
  • arXiv Link: https://arxiv.org/abs/2401.13310v2
  • PDF Link: http://arxiv.org/pdf/2401.13310v2

2. Executive Summary

• Background & Motivation (Why):

  • Core Problem: CERN's LHC generates massive datasets (petabytes) that require highly efficient analysis. With the upcoming High Luminosity LHC (HL-LHC) upgrade, the data volume is expected to increase by a factor of 30, overwhelming existing analysis capabilities.
  • Importance & Gaps: The primary analysis tool, ROOT RDataFrame, traditionally supports only CPU-based parallelism. However, modern high-performance computing facilities are increasingly heterogeneous, incorporating powerful GPGPUs. While a direct CUDA implementation can leverage NVIDIA GPUs, it introduces vendor lock-in and requires maintaining a separate codebase, which is complex and unsustainable. There is a critical need for a portable, single-source solution that can target various accelerator architectures.
  • Innovation: This paper presents a pioneering case study of integrating SYCL, an open standard for heterogeneous computing, into the large and complex ROOT framework. It moves beyond theoretical performance comparisons by documenting the practical challenges and lessons learned during the migration of a real-world, performance-critical application (histogramming).

• Main Contributions / Findings (What):

  • CUDA to SYCL Migration Process: The paper provides a detailed account of porting an existing CUDA implementation of the RDataFrame histogramming operation to SYCL. It highlights specific challenges related to build system integration, compiler differences, and debugging.
  • Comparative Performance Analysis: It offers an extensive performance evaluation of the native CUDA implementation against two leading SYCL compilers, Intel's DPC++ and AdaptiveCpp, on an NVIDIA GPU. The analysis dissects performance into kernel execution, memory transfers, and API call overhead.
  • Identification of SYCL Performance Factors: The research identifies and quantifies the impact of several key factors on SYCL performance:
    1. The efficiency of the SYCL2020 reduction interface.
    2. The performance trade-offs between SYCL memory management models (buffers vs. device pointers).
    3. The significant overhead caused by Just-In-Time (JIT) compilation and how it is affected by compiler choices.
  • Actionable Developer Guidance: Based on their findings, the authors provide concrete, practical advice for developers aiming to use SYCL, particularly for migrating existing CUDA code or integrating SYCL into large C++ projects.

3. Prerequisite Knowledge & Related Work

• Foundational Concepts:

  • High Energy Physics (HEP): A field of physics that studies the fundamental particles and forces that constitute our universe. Experiments at particle accelerators like the LHC generate enormous amounts of data from particle collisions.

  • CERN & LHC: CERN is the European Organization for Nuclear Research, which operates the Large Hadron Collider (LHC), the world's largest and most powerful particle accelerator.

  • ROOT & RDataFrame: ROOT is a comprehensive, open-source C++ framework used extensively in HEP for data processing, analysis, storage, and visualization. RDataFrame is a modern, high-level interface within ROOT for analyzing data stored in a columnar format. It features lazy execution, where a graph of computations is built first and then executed in a single event loop, enabling optimizations.

  • Histogramming: A fundamental data analysis technique used to summarize the distribution of a dataset. It involves dividing the range of data values into a series of intervals (bins) and counting how many values fall into each bin. Image 3 (1.jpg) shows an example of a 1D histogram produced by ROOT.

    该图像为图表，展示了r2变量的直方图分布。横轴为r2取值范围，纵轴为对应的计数（count）。图中数据呈现近似正态分布，峰值集中在r2接近0处。右上角统计框给出了样本总数（Entries）、均值（Mean）和标准差（Std Dev）等统计信息。

  • GPGPU (General-Purpose computing on Graphics Processing Units): The use of a GPU, which is designed for massively parallel processing, to perform computations for general-purpose applications that would traditionally be handled by the CPU.

  • CUDA (Compute Unified Device Architecture): NVIDIA's proprietary parallel computing platform and programming model. It allows developers to use NVIDIA GPUs for GPGPU. Its major drawback is vendor lock-in, as CUDA code can only run on NVIDIA hardware.

  • SYCL: An open, royalty-free C++ abstraction layer for programming heterogeneous hardware (CPUs, GPUs, FPGAs, etc.) from different vendors. It enables a single-source programming model, where host (CPU) and device (accelerator) code can be written in the same C++ source file, promoting code portability and maintainability.

  • DPC++ (Data Parallel C++) & AdaptiveCpp: Two prominent implementations of the SYCL standard. DPC++ is Intel's implementation, built upon LLVM/Clang. AdaptiveCpp (formerly hipSYCL) is a community-driven implementation originating from Heidelberg University.

  • Reductions: A common parallel programming pattern where an operation (e.g., sum, max, min) is applied to a set of input values to produce a single result. This is fundamental for calculating statistics in the histogramming operation.

  • SYCL Memory Management:

    • Buffers & Accessors: A high-level memory management model where SYCL buffers encapsulate data and accessors are used within kernels to declare data access intentions (read, write). The SYCL runtime automatically manages data transfers and dependencies.
    • Unified Shared Memory (USM): A lower-level, pointer-based memory model that gives developers more explicit control over data allocation and movement, similar to cudaMalloc and cudaMemcpy.

  • JIT vs. AOT Compilation:

    • Ahead-Of-Time (AOT): Code is compiled into machine-specific binary (cubin for NVIDIA GPUs) for a target architecture before runtime. This is fast to load.
    • Just-In-Time (JIT): Code is compiled into an intermediate representation (PTX for NVIDIA) that is forward-compatible. At runtime, the GPU driver compiles this PTX code into a machine-specific cubin. This provides flexibility but incurs a significant one-time (or per-run) compilation overhead.

• Previous Works & Differentiation: The paper acknowledges other SYCL performance studies. However, it differentiates itself by:

  • Focusing on a real-world, large-scale scientific application (ROOT) rather than micro-benchmarks or smaller libraries.
  • Providing a detailed narrative of the migration process from an existing, optimized CUDA codebase, which exposes practical integration challenges (like build systems) that other studies might not cover.
  • Performing a direct, in-depth comparison of two major SYCL compilers (DPC++ and AdaptiveCpp) for the same task on the same hardware, revealing subtle but important differences in their runtime behavior and performance.

4. Methodology (Core Technology & Implementation)

The paper details the implementation of GPU-accelerated histogramming, starting from the original CPU logic, moving to a native CUDA version, and finally migrating to SYCL.

• Principles: The core idea is to offload the computationally intensive task of filling a histogram from the CPU to a GPU. RDataFrame's new bulk-processing API provides a natural unit of work (a "bulk" of events) to transfer and process on the accelerator.

• Steps & Procedures:

  1. Baseline CPU Histogramming: As shown in Image 1 (2.jpg), the standard process involves a loop that loads bulks of data. For each event in the bulk, the Histogram Class performs three steps:

     • Find Bin: Determine which histogram bin a given data point belongs to. For fixed-width bins, this is a simple arithmetic calculation. For variable-width bins, it requires a binary search (std::lower_bound).

     • Add Bin Content: Increment the counter for the identified bin, typically using an atomic operation in a parallel context.

     • Update Stats: Perform reduction operations to calculate statistical moments like mean and standard deviation.

       该图像是一个流程示意图，展示了直方图类（Histogram Class）中数据处理的步骤。从“Trigger Event Loop”触发事件循环开始，依次经过“Load Bulk”加载数据块，然后进入直方图类内部的“Find Bin”找箱、“Add Bin Content”添加箱内容、“Update Stats”更新统计三个步骤，最后经过“Final Bulk”阶段返回“Return Results”结果，显示了直方图数据处理的循环和阶段关系。

  2. Native CUDA Implementation: To move this to the GPU, the authors created CUDA-specific components. The workflow, illustrated in Image 2 (3.jpg), is as follows:

     • The CPU loop loads a bulk of event data.

     • The bulk is explicitly copied to the GPU's global memory.

     • Two types of CUDA kernels are launched:

       • Histogram Kernel (Listing 2): This kernel implements a two-stage filling process to reduce atomic contention. First, threads within a block use fast atomicAdd on a local copy of the histogram in shared memory. After all threads in a block synchronize (__syncthreads()), the partial results from shared memory are atomically added to the final histogram in global memory.
       • Reduction Kernel: A separate kernel performs transform-reduce operations to calculate the required statistics (sum of weights, sum of squared weights, etc.).

     • After all bulks are processed on the GPU, the final histogram result is copied back to the host CPU.

       该图像是一个流程示意图，展示了从触发事件循环开始到结果返回的处理步骤。流程依次包括加载批量数据、将数据批量复制到GPU、填充直方图的各个区间、汇总区间内容、更新直方图批量，最后将结果复制回主机。流程中“复制到GPU”和“复制结果回传”步骤用红色标注，核心计算部分用绿色标注，表示数据传输与计算的区分。

  3. Porting to SYCL: The CUDA code was migrated to SYCL with the goal of creating a single, portable codebase.

     • Kernel Translation: The CUDA kernels were translated almost one-to-one into SYCL kernel function objects (functors), as shown in Listing 3. The concepts of threads, blocks, and shared memory in CUDA map directly to work-items, work-groups, and local memory in SYCL.
     • SYCL2020 Reductions: Instead of a custom reduction kernel, the authors used the high-level sycl::reduction interface introduced in the SYCL2020 standard. This simplified the code for calculating statistics.
     • Compiler-Specific Challenges: A key finding was a behavioral difference in the reduction interface between compilers. DPC++ correctly combined the result with the initial value of the reduction variable, while AdaptiveCpp overwrote it, contradicting the SYCL2020 specification. This required a workaround (an extra kernel) for AdaptiveCpp.
     • Build System Integration: A major practical hurdle was integrating the SYCL compilers into ROOT's complex CMake build system. DPC++ requires setting the global C++ compiler, which conflicts with ROOT's dependencies. The authors developed a custom CMake function (add_sycl_to_root_target) to apply the SYCL compiler only to specific files, enabling a clean integration and easy switching between DPC++ and AdaptiveCpp.

5. Experimental Setup

• Hardware & Software:
  • System: A node from the DAS-6 cluster.
  • CPU: AMD EPYC 7402P 24-core Processor.
  • GPU: NVIDIA RTX A4000 (Ampere architecture, compute capability 8.6).
  • Compilers: GCC 12.2.1 (host), CUDA Toolkit 12.3 (nvcc), Intel DPC++ (commit dbee22), and AdaptiveCpp (commit 67cb7a).
• Datasets: Synthetic datasets were generated as ROOT RNTuple files. Each file contained a single column of double-precision floating-point numbers (double). The total number of events varied from 50 million to 1 billion. The values were uniformly distributed in the range [0, 1] to ensure all histogram bins were exercised. A 1D histogram with 1000 fixed-width bins was used.
• Evaluation Metrics:
  1. Total Runtime:
     • Conceptual Definition: Measures the real-world, end-to-end time taken to execute the histogramming action. It includes all overheads: data loading, CPU processing, GPU computation, and data transfers. It is measured using a monotonic clock to avoid issues with system time adjustments.
     • Mathematical Formula: Not applicable (direct time interval measurement).
     • Implementation: std::chrono::steady_clock.
  2. GPU Activity Breakdown:
     • Conceptual Definition: Provides a detailed breakdown of where time is spent during GPU execution. It is measured using the NVIDIA NSight Systems profiler, which traces CUDA API calls, kernel executions, and memory operations. This allows for pinpointing specific sources of overhead that are invisible to a simple wall-clock timer.
     • Mathematical Formula: Not applicable.
     • Components:
       • CUDA Kernels: Time spent executing code on the GPU.
       • CUDA Memory Operations: Time spent on data transfers (e.g., memcpy).
       • CUDA API Calls: Time the CPU spends invoking CUDA runtime functions (e.g., kernel launches, memory allocation, stream management). This represents CPU-side overhead.
• Baselines: The primary baseline is the native CUDA implementation, representing a highly optimized, vendor-specific solution. The DPC++ and AdaptiveCpp implementations are compared against this baseline and each other to evaluate the performance and overhead of the SYCL abstraction.

6. Results & Analysis

The analysis systematically investigates three major performance factors.

• SYCL2020 Reduction Performance (Section 4.3):

  • Insight 1: Workload per Work-Item: The authors tested how changing the number of elements processed sequentially by each work-item in a reduction kernel affects performance. As shown in Figure 4 (4.jpg), AdaptiveCpp's performance improved dramatically (1.4x speedup) when each work-item processed 4 elements instead of 1. This is because it amortizes the latency of accessing memory. In contrast, DPC++ showed marginal improvement, suggesting its default launch configuration or kernel implementation was already more optimal or behaved differently.

    该图像为条形堆叠图，展示不同实现（DPC++与ACPP）在不同每个工作项元素数（1, 2, 4, 8, 16）条件下的总运行时间（秒），细分多个CUDA内存操作、内核和API调用类别的时间开销。图中颜色代表不同CUDA操作，ACPP总体运行时间普遍高于DPC++，且随着每个工作项处理元素数增多，时间变化趋势明显。

  • Insight 2: Fusing Reductions: The authors compared launching one kernel for each statistic versus a single kernel that computes all statistics simultaneously. Figure 5 (5.jpg) demonstrates that fusing reductions into a single kernel yielded significant speedups for both implementations (1.9x for DPC++, 1.4x for AdaptiveCpp on 1B events). This reduces kernel launch overhead and improves memory access efficiency, as the input data is read only once.

    该图像为图表，展示了不同实现（Single DPC++、Multi DPC++、Single ACPP、Multi ACPP）在不同事件数量（50M、100M、500M、1B）条件下的总运行时间（秒），柱状图堆叠显示了CUDA内存操作（不同颜色）、核函数和API调用在总时间中的占比。图中通过颜色区分了各部分耗时，反映了不同实现和规模下性能及瓶颈差异。

• Buffers vs. Device Pointers (USM) (Section 4.4):

  • Overall Runtime: Figure 6 (6.jpg) shows the total end-to-end runtime. For both SYCL compilers, there was no significant performance difference between using high-level buffers (BUF) and low-level device pointers (PTR). However, a clear performance hierarchy emerged: CUDA > DPC++ > AdaptiveCpp. The performance gap between SYCL and CUDA widened as the number of events increased.

    该图像为性能比较折线图，横轴为事件数量（单位百万），纵轴为总耗时（秒）。图中展示了五种实现方式（DPC++ BUF、DPC++ PTR、AdaptiveCpp BUF、AdaptiveCpp PTR、CUDA）的性能随事件数增加的变化趋势。结果显示CUDA实现耗时最短，AdaptiveCpp实现耗时最长，DPC++表现介于两者之间。

  • GPU Activity Breakdown: The profiler results in Figure 7 (7.jpg) reveal a more nuanced story. The SYCL implementations incurred substantially more overhead from CUDA API calls than the native CUDA version. This overhead came from the SYCL runtime's management of streams and events. DPC++ created a very large number of CUDA streams (up to 192), while AdaptiveCpp created only 4. Figure 8 (8.jpg) details the API call overhead for DPC++, showing time spent on event synchronization and stream management. The authors note that AdaptiveCpp (for which a figure was discussed but not provided in the resources) spent nearly twice as much time on API calls as DPC++, particularly on memory deallocations. This indicates that while the end-to-end time for BUF and PTR was similar, their underlying CPU and GPU overhead profiles differed.

    该图像是柱状图，展示了不同实现（BUF DP C++、PTR DP C++、BUF AC PP、PTR AC PP、CUDA）在处理不同事件数量（50M、100M、500M、1B）时的总运行时间（秒）。图中柱状按颜色区分了CUDA内存操作、内核执行和API调用的不同部分，反映了各实现和事件规模下的性能表现及时间构成。

    该图像为柱状图，比较了不同实现（BUF DPC++和PTR DPC++）在不同事件数量（50M、100M、500M、1B）下CUDA API调用总时间（秒）。图中用不同颜色表示多种CUDA API调用类型，显示各调用在总时间中的占比和变化趋势，随着事件量增加，总调用时间明显上升，PTR DPC++实现时间普遍高于BUF DPC++。

• Just-In-Time (JIT) Compilation Overhead (Section 4.5):

  • Insight: The authors discovered that compiling the SYCL code for the wrong NVIDIA GPU architecture (e.g., sm_75 for a sm_86 card) led to massive performance degradation due to JIT compilation at runtime.

  • DPC++ Analysis: Figure 10 (10.jpg) shows that when targeting the incorrect architecture, DPC++ spent an enormous amount of time in the cuModuleLoadDataEx API call, which is responsible for loading and JIT-compiling PTX code. This resulted in a slowdown of up to 98x. Even when targeting the correct architecture, a non-negligible overhead of ~1 second remained, suggesting an issue with the CUDA driver's JIT cache.

    该图像为柱状图，展示了在不同计算能力（75、86）和事件数量（50M、100M、500M、1B）条件下，不同CUDA API调用的总时间（单位：秒）。图中各色块代表不同CUDA函数调用，蓝色（cuModuleUnload）占用时间最长，明显高于其他调用。整体来看，cuModuleUnload是主要性能瓶颈。

  • AdaptiveCpp Analysis: In contrast, Figure 11 (11.jpg) shows that AdaptiveCpp suffered a much smaller penalty (up to 4.9x slowdown) from the incorrect target. The investigation revealed this was due to a difference in kernel-module splitting granularity. DPC++ bundled all templated kernel instantiations into a single large module, which had to be JIT-compiled in its entirety, even though only one instantiation was used. AdaptiveCpp created a separate module for each kernel, so only the required kernel was loaded and compiled, drastically reducing the JIT overhead. This behavior in DPC++ could be changed with the -fsycl-device-code-split=per_kernel flag.

    该图像为柱状图，展示了不同目标计算能力（75和86）及不同事件数量（50M、100M、500M、1B）下，各类CUDA API调用耗时（单位：秒）的对比。图中用不同颜色区分了cudaEventRecord、cudaEventSynchronize、cudaLaunchKernel、cudaMalloc、cudaFree、cudaMemcpyAsync、cudaStreamCreateWithFlags和cudaStreamWaitEvent八种CUDA API调用的时间贡献。整体趋势显示，计算能力和事件数量增加时，总调用时间显著上升，且cudaLaunchKernel和cudaMemcpyAsync占用时间较多。

7. Conclusion & Reflections

• Conclusion Summary: The paper successfully demonstrates the process of migrating a core HEP analysis operation from CUDA to SYCL, providing a roadmap and highlighting key challenges. The performance analysis concludes that while SYCL offers the crucial benefit of portability, it does not yet match the performance of a native CUDA implementation for this workload. DPC++ performs significantly better than AdaptiveCpp and closer to the CUDA baseline. The overhead in SYCL implementations stems primarily from the runtime's management of the underlying CUDA backend (API calls, stream creation) and from JIT compilation, which can be mitigated with correct compiler flags. Despite the performance gap, SYCL is a promising path forward due to its portability and the high-level abstractions (like SYCL2020 reductions) that can make code more concise and maintainable.

• Limitations & Future Work:

  • The study was conducted only on NVIDIA hardware. The authors plan to extend their work to other accelerator architectures from Intel and AMD to fully validate SYCL's portability promise.
  • The analysis focused on a single, relatively simple RDataFrame action (a single 1D histogram). Future work will involve benchmarking more complex use cases with multiple actions and more intricate data dependencies.
  • The authors plan to migrate more RDataFrame actions to the GPU and investigate pipelining techniques to better overlap computation and data transfers.

• Personal Insights & Critique:

  • Strengths: This is an excellent piece of practical, engineering-driven research. Its value lies not in a groundbreaking new algorithm, but in the meticulous documentation of a real-world migration effort. The detailed performance debugging using NSight Systems is exemplary and provides deep insights into the behavior of SYCL runtimes. The actionable advice for developers is concrete and highly valuable.
  • Critique: The use of synthetic, uniformly distributed data is a necessary simplification to isolate the performance of the framework, but its performance characteristics might differ from real-world HEP data, which often has complex distributions and correlations. However, this is a minor point, as the study's primary goal was to evaluate the framework's overhead, not to optimize for a specific data distribution.
  • Overall Impact: This paper serves as a vital data point for the scientific computing community considering a move from CUDA to SYCL. It tempers the hype around portability with a realistic assessment of the current performance overheads and implementation complexities. It highlights that while SYCL is functionally ready, achieving performance parity with native frameworks requires careful attention to compiler flags, runtime behavior, and architecture-specific tuning. The findings are highly transferable to other domains migrating large C++ codebases to heterogeneous platforms.