Evaluating Application Characteristics for GPU Portability Layer Selection

This paper presents a study analyzing representative heterogeneous applications from major High Energy Physics experiments to identify key characteristics that influence performance across various GPU portability layers, thereby guiding developers in selecting the most suitable technology for their specific needs.

The Gist

Imagine you are a chef trying to cook a massive banquet. You have three different types of high-powered ovens in your kitchen: one made by NVIDIA, one by AMD, and one by Intel. Each oven cooks food differently, uses different knobs, and requires different recipes to work at its best.

If you write your recipe specifically for the NVIDIA oven (using a language called CUDA), you can't just put that same recipe into the AMD or Intel ovens. You'd have to rewrite the whole thing. This is a problem because you don't always know which oven you'll have in your kitchen tomorrow.

To solve this, the paper discusses "portability layers." Think of these as universal translators or smart adapters. They let you write one master recipe that the translator converts into the specific language each oven understands. The paper looks at several of these translators (like Kokkos, SYCL, OpenMP, and Alpaka) to see which one is the best fit for different kinds of cooking tasks.

Here is what the authors found when they tested these translators with real "recipes" from high-energy physics experiments (like those used to study subatomic particles):

1. The "Startup Time" Problem

Turning on a GPU (the oven) isn't instant. It takes a few milliseconds to wake up and get ready.

• The Issue: Some translators are slow to start the cooking process. For example, Kokkos can add a significant delay when using AMD ovens. If your cooking task is very short (like boiling an egg for 10 seconds), and the translator takes 5 seconds just to start the stove, you've wasted half your time.
• The Lesson: If your tasks are tiny and quick, avoid translators that make the startup slow.

2. The "Crowded Kitchen" Problem

In a real physics lab, the GPU isn't working alone. It's part of a larger system where many people (threads) are trying to use the oven at the same time.

• The Issue: Some translators are bad at handling crowds. Kokkos, for instance, has a rule that says, "Only one person can talk to the oven at a time," which causes a traffic jam if multiple chefs try to launch tasks simultaneously. SYCL is a bit inconsistent; sometimes it lets everyone cook at once, and sometimes it forces them to wait in line, depending on which version of the translator you are using.
• The Lesson: If your application needs many people working at once, you need a translator that knows how to manage a busy kitchen without locking the doors.

3. The "Toolbox Compatibility" Problem

Physics recipes often use special tools (libraries like ROOT or Eigen) that help with math and data.

• The Issue: Some of these tools don't play nice with the translators. For example, a popular math tool called Eigen often breaks when used with the NVIDIA compiler, which many translators rely on. Also, trying to use two different compilers (one for the CPU and one for the GPU) in the same project is like trying to build a house with two different sets of blueprints that don't match—it makes the construction (building the software) a nightmare.
• The Lesson: Before picking a translator, check if your favorite tools will fit inside it.

4. The "Furniture Arrangement" Problem

GPUs love simple, flat layouts. They prefer data to be arranged like a neat row of boxes. However, physics data often comes in complex, messy shapes (like a pile of different-sized suitcases).

• The Issue: Translators try to fix this mess by wrapping the data in special containers. While this makes the code portable, it adds "overhead"—like putting every single item in a suitcase before moving it, even if you only need to move one sock. This slows things down. Also, none of the translators are great at handling "jagged" data (rows of different lengths), which is very common in physics.
• The Lesson: If your data is complex and messy, the translator might slow you down trying to tidy it up.

5. The "Specialized Tools" Problem

Sometimes you need a specific tool, like a Random Number Generator (RNG) or a Fast Fourier Transform (FFT).

• The Issue: Each oven manufacturer has their own super-fast, specialized version of these tools. The universal translators often don't include these specialized versions, or they use their own slower versions. While you can force the translator to use the oven's native tool, it breaks the "portability" because that tool only works on that specific oven.
• The Lesson: If you rely heavily on these specific tools, you might have to choose between speed (using the oven's native tool) or portability (using the translator's generic tool).

6. The "Construction Time" and "Moving Day" Problems

• Building the Recipe: Some translators make the "cooking time" (compilation time) much longer. For huge projects, using certain translators can make the build process take hours instead of minutes.
• Moving the Kitchen: If you build your software for a specific oven (e.g., an NVIDIA V100), it might not work on a newer one (an NVIDIA A100). Some translators require you to build a separate version for every single type of oven you might encounter. This creates a massive logistical headache for distributing the software to different labs.

The Final Verdict

The paper concludes that there is no "perfect" translator.

• Kokkos is great for some things but struggles with concurrency and startup times on certain hardware.
• SYCL is powerful but can be inconsistent depending on the compiler version.
• OpenMP and others have their own strengths and weaknesses regarding how they handle memory and different hardware.

The Takeaway: You can't just pick a translator because it's popular. You have to look at your specific "recipe" (your application). If your code is short and fast, pick a translator with low startup time. If your code is complex and uses many tools, pick one that plays well with those tools.

The authors also note that these technologies are evolving rapidly, like new models of ovens coming out every year. What works best today might change tomorrow, so developers need to keep watching the landscape. In the future, new standards might make these choices easier, but for now, careful testing is the only way to find the right fit.

Technical Summary

Technical Summary: Evaluating Application Characteristics for GPU Portability Layer Selection

Problem Statement

High Energy Physics (HEP) computing is increasingly reliant on GPUs for performance, yet the hardware landscape is diversifying beyond NVIDIA to include AMD and Intel. While vendor-specific languages like CUDA offer high performance, they restrict deployment to specific hardware, reducing portability. Conversely, portability layers such as Kokkos, SYCL, HIP, OpenMP, Alpaka, and std::par promise execution across diverse CPU and GPU architectures. However, these layers are not functionally equivalent; they exhibit distinct characteristics regarding performance, concurrency support, and compatibility with existing codebases. The core problem addressed is the lack of a clear framework for developers to select the most appropriate portability layer based on specific application characteristics, kernel structures, and runtime environments.

Methodology

The High Energy Physics Center for Computational Excellence (HEP-CCE) conducted a comparative study using representative heterogeneous applications from major HEP experiments:

• ATLAS: FastCaloSim (parametrized calorimeter simulation).
• CMS: Patatrack (heterogeneous pixel reconstruction with 40+ kernels, multithreading, and memory pools) and P2R (single-kernel track reconstruction).
• DUNE: Wire-Cell Toolkit (rasterization, scatter-add, and FFT convolution kernels).
• Event Generators: GPU ports of Sherpa (Pepper) and Madgraph (utilizing auto-generated code).

These testbeds were ported from serial C++ to various portability technologies to identify code characteristics that significantly impact performance and compatibility. The study evaluated these applications across multiple dimensions, including kernel launch latency, concurrency handling, library compatibility, memory management, and compilation overhead.

Key Findings and Results

1. Kernel Launch Latency

Kernel launch is not instantaneous and can range from microseconds to tens of microseconds.

• Kokkos: Observed to add significant overhead (tens of microseconds), particularly on AMD GPUs, which already possess higher native launch latencies than NVIDIA or Intel.
• SYCL, HIP, OpenMP, Alpaka: Did not exhibit large additional penalties compared to native implementations.
• AMD Specifics: Calling rocRAND on AMD GPUs resulted in extreme first-call latencies (exceeding 1 second), suggesting that short-running kernels may suffer severe performance degradation if using portability layers that exacerbate launch overhead.

2. Concurrency and Thread Pools

HEP applications often run within larger frameworks utilizing multi-threading (e.g., Intel TBB).

• Kokkos: Demonstrated significant incompatibilities with external thread pools. Its serial backend uses locks that serialize calls, and its threads backend forbids external thread calls. Concurrent kernel launches are currently only viable with CUDA/HIP backends using architecture-specific APIs, limiting portability.
• SYCL: Implementations were inconsistent. Some serialized concurrent calls from different threads depending on the compiler manufacturer and version. AMD architectures appeared to lack a concurrent SYCL solution at the time of the study.

3. External Library and Compiler Compatibility

• Eigen: Many versions are incompatible with NVIDIA's nvcc, affecting portability layers that rely on it (Kokkos, Alpaka). It is also currently incompatible with OpenMP offload.
• ROOT: Cannot be fully compiled with NVIDIA's nvc++. Hybrid compilation (using g++ and nvc++) is possible but complicates build rules and requires careful memory allocation management (GPU data must be allocated by nvc++).
• Build Complexity: Supporting multiple portability layers in one project requires complex build rules (e.g., #ifdef usage, file separation) to avoid header conflicts between layers like Kokkos and CUDA.

4. Data Structures and Memory Transfers

• Object-Oriented vs. Flat Layouts: Complex, object-oriented data structures common in HEP map poorly to GPUs, which prefer flat layouts or Structures of Arrays (SoA).
• Portability Overhead: Layers like Kokkos and SYCL provide abstractions (e.g., Kokkos::Views, sycl::buffers) that increase allocation and transfer overhead, especially for small objects.
• Initialization: Default memory initialization in Kokkos::Views can be a performance bottleneck; this can be mitigated using Kokkos::ViewAllocateWithoutInitializing.
• Unified Shared Memory (USM): Automatic USM transfers (used by std::par) were observed to be slower than explicit transfers and may not trigger at optimal times.
• Jagged Arrays: None of the APIs natively support multidimensional vectors of varying sizes (jagged arrays), requiring manual crafting, dimension reduction, or padding.

5. RNGs, FFTs, and Atomics

• Native Libraries: Vendor-specific libraries (cuRAND, rocRAND, oneMKL) offer tuned performance but lack portability. While Kokkos provides cross-architecture RNG/FFT implementations, they may not match native performance.
• HEP-CCE Contributions: The group developed a header-only library wrapping device-native RNGs with a uniform API and hardware-native backends for oneMKL to allow SYCL to call native implementations across vendors.
• Atomics: Atomic operations map poorly to GPU hardware. Performance varies wildly by compiler and hardware (e.g., atomic<int> on NVIDIA with OpenMP/clang13 was 35x slower than CUDA). Integer atomics generally underperform compared to float/double atomics.

6. Compilation Time

• Compiler Overhead: nvc++ is generally 2x–3x slower than g++. Kokkos adds a 10%–20% overhead.
• Template Heavy Code: Madgraph, utilizing heavily templated auto-generated code, experienced compilation times with SYCL's icpx backend that were over 1000x slower than other backends for large kernels.

7. Runtime Provisioning

• Kokkos: Requires building with specific hardware identifiers. A binary built for one GPU (e.g., V100) may not run on another (e.g., A100). Supporting shared libraries and multiple architectures creates a complex matrix of build configurations, challenging binary distribution.
• Alpaka: Requires separate compilation for each backend compiler.
• SYCL/OneAPI: Multi-platform binaries are emerging but often limited to specific architectures or require compile-time identification.
• OpenMP: Some implementations (LLVM Clang) allow listing multiple architectures at compile time and dynamically selecting the backend at runtime.

Significance and Conclusion

The paper concludes that no single GPU portability solution is optimal for all scenarios. The choice of portability layer must be driven by a careful analysis of the application's specific characteristics, including kernel runtime, concurrency requirements, data structure complexity, and external library dependencies.

The authors emphasize that the ecosystem is rapidly evolving. Both portability layers (Kokkos, SYCL, etc.) and native compilers (clang, nvcc, hipcc) are improving, with trends moving toward increased feature support and optimized performance. The paper advocates for frequent monitoring of these technologies to ensure the best match between codebases and portability layers. Additionally, the paper notes that emerging C++ standards (such as P2300 for asynchronous execution) may simplify these choices in the future, though broad hardware support remains years away.

Ultimately, selecting a portability layer without fully understanding the code's nature, external dependencies, and runtime environment risks significant issues in both performance and deployment flexibility.