Engineering Papers⌕ Search

SEARCH · Engineering Papers

Results for “Kernel”

Search indexed NASA NTRS and DOE OSTI research on propulsion, heat transfer, battery materials and energy systems. Follow report and document links to the original sources.

Quote a phrase for an exact phrase match. Source license links do not imply unrestricted reuse.

At least 1,009 records · Page 56

MemHC: An Optimized GPU Memory Management Framework for Accelerating Many-body Correlation

The many-body correlation function is a fundamental computation kernel in modern physics computing applications, e.g., Hadron Contractions in Lattice quantum chromodynamics (QCD). This kernel is both computation and memory intensive, involving a series of tensor contractions, and thus usually runs on accelerators like GPUs. Existing optimizations on many-body correlation mainly focus on individual tensor contractions (e.g., cuBLAS libraries and others). In contrast, this work discovers a new optimization dimension for many-body correlation by exploring the optimization opportunities among tensor contractions. More specifically, it targets general GPU architectures (both NVIDIA and AMD) and optimizes many-body correlation’s memory management by exploiting a set of memory allocation and communication redundancy elimination opportunities: first, GPU memory allocation redundancy: the intermediate output frequently occurs as input in the subsequent calculations; second, CPU-GPU communication redundancy: although all tensors are allocated on both CPU and GPU, many of them are used (and reused) on the GPU side only, and thus, many CPU/GPU communications (like that in existing Unified Memory designs) are unnecessary; third, GPU oversubscription: limited GPU memory size causes oversubscription issues, and existing memory management usually results in near-reuse data eviction, thus incurring extra CPU/GPU memory communications.

97 MATHEMATICS AND COMPUTING↗

An MLIR-based Compiler Flow for System-Level Design and Hardware Acceleration

The generation of custom hardware accelerators for applications implemented within high-level productive programming frameworks requires considerable manual effort. To automate this process, we introduce \sodaopt, a compiler tool that extends the MLIR infrastructure. \sodaopt automatically searches, outlines, tiles, and pre-optimizes relevant code regions to generate high-quality accelerators through high-level synthesis. \sodaopt can support any high-level programming framework and domain-specific language that interface with the MLIR infrastructure. By leveraging MLIR, \sodaopt solves compiler optimization problems with specialized abstractions. Backend synthesis tools connect to \sodaopt through progressive intermediate representation lowerings. \sodaopt interfaces to a design space exploration engine to identify the combination of compiler optimization passes and options that provides high-performance generated designs for different backends and targets. We demonstrate the practical applicability of the compilation flow by exploring the automatic generation of accelerators for deep neural networks operators outlined at arbitrary granularity and by combining outlining with tiling on large convolution layers. Experimental results with kernels from the PolyBench benchmark show that \sodaopt high-level optimizations improve execution delays of synthesized accelerators up to 60x. We also show that for the selected kernels, our solution outperforms the current of state-of-the art in more than 70% of the benchmarks and provides better average speedup in 55% of them.

Bohm Agostini, Nicolas↗

QASMBench: A Low-Level Quantum Benchmark Suite for NISQ Evaluation and Simulation

The rapid development of quantum computing (QC) in the NISQ era urgently demands a low-level benchmark suite and insightful evaluation metrics for characterizing the properties of prototype NISQ devices, the efficiency of QC programming compilers, schedulers and assemblers, and the capability of quantum system simulators in a classical computer. In this work, we fill this gap by proposing a low-level, easy-to-use benchmark suite called QASMBench based on the OpenQASM assembly representation. It consolidates commonly used quantum routines and kernels from a variety of domains including chemistry, simulation, linear algebra, searching, optimization, arithmetic, machine learning, fault tolerance, cryptography, and so on, trading-off between generality and usability. To analyze these kernels in terms of NISQ device execution, in addition to circuit width and depth, we propose four circuit metrics including gate density, retention lifespan, measurement density, and entanglement variance, to extract more insights about the execution efficiency, the susceptibility to NISQ error, and the potential gain from machine-specific optimizations. Applications in QASMBench can be launched and verified on several NISQ platforms, including IBM-Q, Rigetti, IonQ and Quantinuum. For evaluation, we measure the execution fidelity of a subset of QASMBench applications on 12 IBM-Q machines through density matrix state tomography, comprising 25K circuit evaluations. In addition we also compare the fidelity of executions among the IBM-Q machines, the IonQ QPU and the Rigetti Aspen M-1 system.

97 MATHEMATICS AND COMPUTING↗

Julia as a unifying end-to-end workflow language on the Frontier exascale system

We evaluate Julia as a single language and ecosystem paradigm powered by LLVM to develop workflow components for high-performance computing. We run a Gray-Scott, 2-variable diffusion-reaction application using a memory-bound, 7-point stencil kernel on Frontier, the US Department of Energy’s first exascale supercomputer. We evaluate the performance, scaling, and trade-offs of (i) the computational kernel on AMD’s MI250x GPUs, (ii) weak scaling up to 4,096 MPI processes/GPUs or 512 nodes, (iii) parallel I/O writes using the ADIOS2 library bindings, and (iv) Jupyter Notebooks for interactive analysis. Results suggest that although Julia generates a reasonable LLVM-IR, a nearly 50% performance difference exists vs. native AMD HIP stencil codes when running on the GPUs. As expected, we observed near-zero overhead when using MPI and parallel I/O bindings for system-wide installed implementations. Consequently, Julia emerges as a compelling high-performance and high-productivity workflow composition language, as measured on the fastest supercomputer in the world.

Godoy, William↗

IRIS-MASH: Efficient Multi-device Asynchronous Multi-Stream Heterogeneous Computing

In the rapidly evolving field of high-performance computing (HPC), effectively leveraging heterogeneous devices through asynchronous task programming is paramount. This paper presents a robust asynchronous task programming model tailored for a multi-device, multi-stream execution environment that incorporates a diverse array of heterogeneous computing units, including GPUs from various vendors and other accelerators. Current state-of-the-art task programming models provide methodologies to support asynchronous task executions, but they typically handle homogeneous devices using native programming languages, while support for heterogeneous devices is limited to frameworks like OpenCL. This gap presents significant challenges in abstracting heterogeneous devices to harness their true asynchronous capabilities effectively using their native programming languages. By implementing asynchronous task execution, our model significantly boosts the performance of tiled algorithm task graphs through overlapping data transfers with computation and enabling the simultaneous execution of multiple kernels. We integrate this approach into a heterogeneous Intelligent Runtime System (IRIS) and assess its performance using a suite of tiled algorithm benchmarks from the heterogeneous math kernels library (MatRIS) based on IRIS. Experimental results demonstrate a performance improvement ranging from 1.6 × to 2 × over IRIS without asynchronous support, and a notable 22% performance enhancement compared to established runtime systems such as StarPU and PaRSEC. This approach significantly improves computation efficiency of HPC workflows and provides a solid base for future exploration and development in the area of asynchronous task programming in heterogeneous systems.

Miniskar, Narasinga Rao [ORNL] (ORCID:000000018259↗

Fine-Grained Application Energy and Power Measurements on the Frontier Exascale System

The increasing complexity and power/energy demands of heterogeneous exascale systems, such as the Frontier supercomputer, present significant challenges for measuring and optimizing power consumption in applications. Current tools either lack the resolution to capture fine-grained power and energy measurements, fail to validate in-band measurements against out-of-band power sensors, or cannot integrate this information with application performance events in a scalable manner. This paper introduces a novel open-source performance toolkit that integrates extended PAPI components with Score-P plugins to enable in-band, fine-grained power and energy measurements, while also supporting validation using power meter measurements for both CPUs and GPUs. One key contribution is the ability to perform millisecond-level power and energy measurements for AMD MI250X GPUs, mapping them to application performance events within a single trace and measurement system that scales. Our toolkit combines coarse-grained measurements from cray_pm counters with high-resolution metrics from rocm_smi and RAPL, converting GPU instantaneous accumulated energy into power to capture both transient and steady-state power behavior, a capability often missed by out-of-band and monitoring tools. By mapping these metrics to specific application regions, developers can identify energy hotspots, address inefficiencies in GPU kernel execution, and validate in-band measurements against external measurements. We demonstrate the effectiveness of this approach through case studies using benchmarks such as GPU rocblas_sgemm, BLIS c_blas_dgemm, and rocHPL, highlighting the variability of the measurements and the impact of transient power spikes on kernel-level efficiency.

Hernandez Mendoza, Oscar [ORNL] (ORCID:00000002538↗

SSCA1-K1

SSCA1-K1 is a parallel implementation of kernel 1 of the SSCA1 benchmark suite released by the DARPA HPCS program. This kernel is able to run in parallel on a distributed shared memory system at extreme scales using OpenSHMEM.

Baker, Matthew↗

Comb

Comb is a communication performance benchmarking tool. It is used to determine performance tradeoffs in implementing communication patterns on high performance computing (HPC) platforms. At its core comb runs combinations of communication patterns with execution patterns, and memory spaces. The current set of capabilities Comb provides includes: o Configurable structured mesh halo exchange communication. o A variety of communication patterns based on grouping messages. o A variety of execution patterns including serial, openmp threading, cuda, cuda batched kernels, and cuda persistent kernels. o A variety of memory spaces including default system allocated memory, pinned host memory, cuda device memory, and cuda managed memory with different cuda memory advice.

Burmark, Jason↗

Sandia Cloud Forensics and Incident Response Platform v 1.0

As systems and devices become virtualized and deployed in the cloud, the hypervisor becomes an increasingly appropriate place to collect performance data, system state, system landscape, function calls, transaction traces, and other characteristics. We developed a method by which an introspection application may be coupled with a hypervisor to “reach into” the VM with minimal intrusiveness to collect data critical to the reconstruction of events, files, and operations. Such a capability is required to take advantage of the hypervisor as an instrumentation platform and to integrate that data with more traditional collection mechanisms. The concept of a VM serviced by a lightweight hypervisor is a relatively new paradigm for forensic practitioners. Traditional forensic techniques, based on assumptions that the filesystem was directly interacting with the hardware through an abstraction, afforded the forensic practitioner the assumption that there was nothing controlling the application below the filesystem. This is not the case when using virtualized technologies. Hypervisors can covertly monitor, introspect and interact with the guest in a transparent fashion. The problems of storage and collection of actionable data are exhausting. The current challenge is most hypervisors do not expose a useful application programming interface (API) at a sufficient level to do transparent, fine-grained and customizable introspection. Scalable VM instrumentation and introspection at an in-depth level requires fast handling of events, as well as direct access to VM state. Furthermore, deep introspection benefits greatly from the ability to gather data from the hardware during the VM’s exit to the hypervisor. All of this requires identical access to the system as the hypervisor itself; improper use of this ability could easily cause system instability. It is for this reason we believe that the hypervisor developers have been hesitant to grant this much control through their APIs. However, our approach leverages other means to collect and monitor the guests in a targeted fashion. Virtual machine introspection (VMI) is a technique used to monitor the runtime state of a system-level virtual machine. The runtime state can be may include processor registers, memory, disk, network, and any other hardware-level events A review of research literature and current VMI technologies exposed several limitations and tradeoffs in VMI approaches, including: the use of in-guest agents; kernel to user space transitions (dramatically slowing down processing); VMI tool pre-configuration requirements; hypervisor version lock-in or source code patching; reliance on operating system symbols; limited processor features due to hypervisor (even if the hardware could do more). To address these constraints, a VMI tool was envisioned to provide the cloud forensic capabilities while having as few of these limitations as possible. The Kernel-based Virtual Machine Introspection (KVMi) tool was developed for hypervisors on Intel’s x86-64 architecture. To meet performance, scoping and use-case demands, the follow criteria were applied to KVMi to support live, forensic data collection: • Shall not require in-guest agents. • Shall work with any recent version of modern hypervisors. • Shall introspect upon any guest (VM) running modern versions of Windows, Linux and OS X. • Shall have the ability to find and track operating system artifacts • Shall not require operating system symbol files. • Shall be able to fully handle VM-exits, bypassing execution of the hypervisor if necessary, to facilitate new features the hypervisor may not support. • Shall be compilable/loadable on a running system with standard build tools.

Stout, WilliamM.S.↗

LOGAN: High-Performance X-Drop Pairwise Alignment on GPU (LOGAN) v1.0

Pairwise sequence alignment is one of the most computationally intensive kernels in genomic data analysis, accounting for more than 90% of the run time for key bioinformatics applications. This method is particularly expensive for third-generation sequences due to the high computational expense of analyzing these long read lengths (1Kb-1Mb). Given the quadratic overhead of exact pairwise algorithms such as Smith-Waterman, for long alignments, the community primarily relies on approximate algorithms that search only for high-quality alignments and stop early when one is not found. In this work, we present the first GPU optimization of the popular X-drop alignment algorithm, named LOGAN. Results show that our high-performance multi-GPU LOGAN implementation achieves up to 181.6 GCUPS and speed-ups up to 6.6x and 30.7x using 1 and 6 NVIDIA Tesla V100, respectively, over the state-of-the-art software running on two IBM Power9 processors using 168 threads, with equivalent accuracy. We also demonstrate a 2.3x LOGAN speed-up versus ksw2, a state-of-art vectorized algorithm for sequence alignment implemented in minimap2. To highlight the impact of our work on a real-world application, we couple the LOGAN aligner with a many-to-many long-read alignment software called BELLA, and demonstrate that our implementation improves the overall BELLA runtime by up to 10.6x. Finally, we adapt the Roofline model for our optimized kernel and demonstrate that our implementation is near-optimal on the NVIDIA Tesla V100s.

Guidi, Giulia↗

CSBCOO SPMM: Compressed Sparse Block Coordinate format sparse matrix dense vector multiply benchmark (CSBCOO SPMM) v1.0

The purpose of this microbenchmark is to provide a means to explore different programming methodologies using a simple, but not trivial, mathematical kernel. The kernel is based on the matrix vector product in the MFDn nuclear configuration interaction code's LOBPCG eigensolver. This particular mathematical operation is a key element in the solution of systems of equations and block eigensolvers.

Cook, Brandon↗

Code for the manuscript "Mori-Zwanzig Modal Decomposition"

We would like to create an open source repository in LANL's github on code written in Julia, in which we implement and extend the data-driven Mori-Zwanzig method for extracting large-scale spatio-temporal structures from data, which we call MZMD. This method is an extension of Dynamic Mode Decomposition (DMD) in which Mori-Zwanzig memory kernels are included into the associated companion matrix. In the code we would like to release, we apply MZMD to a flow over a cylinder with Reynolds number 100 rather than the much larger data set used in the associated manuscript. DMD is used extensively in the fluid dynamics community mainly for extracting large scale spatio-temporal structures (patters) from flow data. This is useful for understanding the key mechanisms that generate certain complex dynamical process relevant in engineering design. In MZMD, we improve upon DMD by adding the Mori-Zwanzig memory kernels, and show this improvement is especially important in strongly nonlinear regions of the flow.

Woodward, Michael↗

The effect of atmospheric transmissivity on model and observational estimates of the sea ice albedo feedback

The sea-ice-albedo feedback (SIAF) is the product of the ice sensitivity (IS) – how much the surface albedo in sea-ice regions changes as the planet warms– and the radiative sensitivity (RS) – how much the top of atmosphere radiation changes as the surface albedo changes. We demonstrate that the RS calculated from radiative kernels in climate models is reproduced from calculations using the “approximate partial radiative perturbation” method that uses the climatological radiative fluxes at the top of atmosphere and the assumption that the atmosphere is isotropic to shortwave radiation. This method facilitates the comparison of RS from satellite-based estimates of climatological radiative fluxes with RS estimates across a full suite of coupled climate models and, thus, allows model evaluation of a quantity important in characterizing the climate impact of sea ice concentration changes. The satellite based RS is within the model range of RS that differs by a factor of two across climate models in both the Arctic and Southern Ocean. Observed trends in Arctic sea ice are used to estimate IS which, in conjunction with the satellite-based RS yields an SIAF of 0.16 ± 0.04 W m-2 K-1. This Arctic SIAF estimate suggests a modest amplification of future global surface temperature change by approximately 14% relative to a climate system with no SIAF. We calculate the global albedo feedback in climate models using model specific RS and IS and find a model mean feedback parameter of 0.37 W m-2 K-1 which is 40% larger than the IPCC AR5 estimate based on using RS calculated from radiative kernel calculations in a single climate model.

Arctic, Antarctic, climate change, Climate & Earth↗

Temperature Inversions below 1 km from a V-Band Scanning Radiometer at the North Slope of Alaska

A single-channel (56.7 GHz) scanning radiometer was deployed in August 2022 at the Atmospheric Radiation Measurement (ARM) North Slope of Alaska site near Utqiaġvik. The radiometer is designed to provide temperature profiles between 0 and 1 km every 5 min. Averaging kernels show that this single-channel radiometer, taking observations at 10 discrete elevation angles, yields approximately the same information as a seven-channel V-band radiometer scanning three elevation angles. The instrument is able to reproduce the occurrence of temperature inversions between the surface and 1 km and their strength showing a correlation of 0.85, bias of −0.6 K, and slope of 1.04 with respect to radiosondes. Uncertainty in the inversion base height varies from 50 m near the surface to ∼300 m above 0.4 km when compared with radiosondes. Conversely, the inversion top height is overestimated and has higher uncertainty due to the degrading effects of the averaging kernels on the vertical resolution of the retrievals. Thanks to the high temporal resolution of the retrievals, the diurnal cycle of boundary layer temperature was evaluated showing that the radiometer can capture some aspects of the boundary layer thermal structure. The present analysis provides an overview of the capabilities of this simple observing configuration for selected applications.

Atmospheric profilers↗

Acceleration of tensor-product operations for high-order finite element methods

This article is devoted to graphics processing unit (GPU) kernel optimization and performance analysis of three tensor-product operations arising in finite element methods. We provide a mathematical background to these operations and implementation details. Achieving close to peak performance for these operators requires extensive optimization because of the operators’ properties: low arithmetic intensity, tiered structure, and the need to store intermediate results during the kernel execution. We give a guided overview of optimization strategies and we present a performance model that allows us to compare the efficacy of these optimizations against an empirically calibrated roofline.

Świrydowicz, Kasia↗

Scalability of high-performance PDE solvers

Performance tests and analyses are critical to effective high-performance computing software development and are central components in the design and implementation of computational algorithms for achieving faster simulations on existing and future computing architectures for large-scale application problems. In this article, we explore performance and space-time trade-offs for important compute-intensive kernels of large-scale numerical solvers for partial differential equations (PDEs) that govern a wide range of physical applications. We consider a sequence of PDE-motivated bake-off problems designed to establish best practices for efficient high-order simulations across a variety of codes and platforms. We measure peak performance (degrees of freedom per second) on a fixed number of nodes and identify effective code optimization strategies for each architecture. In addition to peak performance, we identify the minimum time to solution at 80% parallel efficiency. The performance analysis is based on spectral and p-type finite elements but is equally applicable to a broad spectrum of numerical PDE discretizations, including finite difference, finite volume, and h-type finite elements.

97 MATHEMATICS AND COMPUTING↗

Parthenon—a performance portable block-structured adaptive mesh refinement framework

On the path to exascale the landscape of computer device architectures and corresponding programming models has become much more diverse. While various low-level performance portable programming models are available, support at the application level lacks behind. To address this issue, we present the performance portable block-structured adaptive mesh refinement (AMR) framework Parthenon, derived from the well-tested and widely used Athena++ astrophysical magnetohydrodynamics code, but generalized to serve as the foundation for a variety of downstream multi-physics codes. Parthenon adopts the Kokkos programming model, and provides various levels of abstractions from multidimensional variables, to packages defining and separating components, to launching of parallel compute kernels. Parthenon allocates all data in device memory to reduce data movement, supports the logical packing of variables and mesh blocks to reduce kernel launch overhead, and employs one-sided, asynchronous MPI calls to reduce communication overhead in multi-node simulations. Using a hydrodynamics miniapp, we demonstrate weak and strong scaling on various architectures including AMD and NVIDIA GPUs, Intel and AMD x86 CPUs, IBM Power9 CPUs, as well as Fujitsu A64FX CPUs. At the largest scale on Frontier (the first TOP500 exascale machine), the miniapp reaches a total of 1.7 × 10 13 zone-cycles/s on 9216 nodes (73,728 logical GPUs) at [Formula: see text] weak scaling parallel efficiency (starting from a single node). In combination with being an open, collaborative project, this makes Parthenon an ideal framework to target exascale simulations in which the downstream developers can focus on their specific application rather than on the complexity of handling massively-parallel, device-accelerated AMR.

97 MATHEMATICS AND COMPUTING↗