Experience Report: Writing A Portable GPU Runtime with OpenMP 5.1

06/06/2021
by   Shilei Tian, et al.
0

GPU runtimes are historically implemented in CUDA or other vendor specific languages dedicated to GPU programming. In this work we show that OpenMP 5.1, with minor compiler extensions, is capable of replacing existing solutions without a performance penalty. The result is a performant and portable GPU runtime that can be compiled with LLVM/Clang to Nvidia and AMD GPUs without the need for CUDA or HIP during its development and compilation. While we tried to be OpenMP compliant, we identified the need for compiler extensions to achieve the CUDA performance with our OpenMP runtime. We hope that future versions of OpenMP adopt our extensions to make device programming in OpenMP also portable across compilers, not only across execution platforms. The library we ported to OpenMP is the OpenMP device runtime that provides OpenMP functionality on the GPU. This work opens the door for shipping OpenMP offloading with a Linux distribution's LLVM package as the package manager would not need a vendor SDK to build the compiler and runtimes. Furthermore, our OpenMP device runtime can support a new GPU target through the use of a few compiler intrinsics rather than requiring a reimplementation of the entire runtime.

READ FULL TEXT

page 1

page 2

page 3

page 4

research
01/28/2019

The OoO VLIW JIT Compiler for GPU Inference

Current trends in Machine Learning (ML) inference on hardware accelerate...
research
08/09/2023

__host__ __device__ – Generic programming in Cuda

We present patterns for Cuda/C++ to write save generic code which works ...
research
10/27/2021

JACC: An OpenACC Runtime Framework with Kernel-Level and Multi-GPU Parallelization

The rapid development in computing technology has paved the way for dire...
research
01/27/2022

Porting OpenACC to OpenMP on heterogeneous systems

This documentation is designed for beginners in Graphics Processing Unit...
research
10/09/2020

C for a tiny system

We have implemented support for Padauk microcontrollers, tiny 8-Bit devi...
research
08/28/2022

ECP SOLLVE: Validation and Verification Testsuite Status Update and Compiler Insight for OpenMP

The OpenMP language continues to evolve with every new specification rel...
research
03/20/2023

Runtime-Adaptable Selective Performance Instrumentation

Automated code instrumentation, i.e. the insertion of measurement hooks ...

Please sign up or login with your details

Forgot password? Click here to reset