Compilers and More: A GPU and Accelerator Programming Model

By Michael Wolfe

December 9, 2008

Okay, maybe the title should be “Languages and More,” but I promise to talk about compilers further on below.

In recent columns, I’ve discussed parallel programming and programming GPUs and accelerators in particular. In May, I predicted that accelerator-based systems would dominate high performance computing, and suggested that an evolutionary approach to migrating applications from CPUs to accelerators was possible and appropriate. In September, I discussed in more detail the specifics of GPU hardware architecture, pointing out its strengths for high performance computing (lots of parallelism), as well as its weaknesses (limited to specific parallelism domains). In October, I showed what it takes to start porting a CPU program to a GPU, exposing some of the complexities of the interactions between the host and the GPU. The specific examples use NVIDIA’s very popular CUDA language, but I discuss OpenCL briefly as well (which should be about ready for public discussion about the time you read this). In my most recent column, I showed the details of optimizing a simple matmul kernel for a GPU, including testing various ways to organize it and vary the parallelism parameters.

If you read these, or are familiar with current approaches to programming accelerators, you are either discomforted by the complexities, or excited at the levels of control you can get. The low-level programming model in CUDA and OpenCL certainly has its place, though it’s not for the faint of heart. So, to go back to the first of those columns, can we come up with a different model of GPU and accelerator programming, one that retains most of the advantages of CUDA or OpenCL, but without requiring complete program rewrites, that can be applied to different target accelerators, and that retains the potential to develop and test in a more accessible environment? In short, a model that allows HPC programmers to focus on domain science instead of on computer science?

Architectural Model

Let’s start by looking at the features of the architecture that we want to use to advantage. Current GPUs are specific implementations of a programming model that works well for graphics problems. They support two levels of parallelism: an outer fully-parallel doall loop level, and an inner synchronous (SIMD or vector) loop level. Each level can be multidimensional (2 or 3 dimensions), but the domain must be strictly rectangular. The synchronous level may not be fully implemented with SIMD or vector operations, so explicit synchronization is supported (and required) across this level. No synchronization is supported between parallel threads across the doall level.

For those familiar with memory models, current GPUs implement a particularly weak model. In particular, they don’t support memory coherence between threads, unless those threads are parallel only at the synchronous level and the memory operations are separated by an explicit barrier. Otherwise, if one thread updates a memory location and another reads the same location, or two threads store a value to the same location, the hardware does not guarantee the results. You can’t say it gets the wrong answers, because such programs are defined as being in error. There is a software-managed cache on a GPU, and there are some hardware caches that can be used as well, but only in certain situations (and limited to read-only data).

The most significant characteristic is that the memory on the GPU or accelerator is separate from the host memory. The host can’t simply read or write to the accelerator memory because it’s not mapped into the virtual memory space of the host. Similarly, the accelerator can’t simply read or write to host memory; the host memory doesn’t support the bandwidth necessary for the accelerator, not to mention the need to support the virtual memory map on the accelerator.

The chips support parallelism on the order of hundreds of threads (today), but effective programs need parallelism on the order of thousands. This provides enough slack parallelism to tolerate long latency memory operations by thread switching, or multithreading, an idea pioneered by the venerable Denelcor HEP almost 30 years ago.

In summary, today’s GPUs look like an attached processor with its separate memory, that supports a multidimensional rectangular domain of parallelism, including doall and synchronous parallelism. We’d like a programming model that simplifies most of the hardware details, but gives experts finer levels of control. We probably can’t hide the distinction between the two levels of parallelism, but we’d like to avoid requiring the programmer to insert explicit synchronization as much as possible. It’s easy to map doall parallelism onto SIMD parallelism, but not the other way around, so we’d like to encourage programmers to program in a doall style when possible and appropriate. We probably can’t completely hide the distinction between host memory and accelerator memory, but the details of transferring data should be handled automatically.

But our programming model shouldn’t focus on the details of today’s GPUs as the ultimate accelerator architecture. One can envision accelerators with mostly (or only) synchronous parallelism (like the Clearspeed CSX700 accelerator processor), or with mostly doall parallelism (like the Tilera TILE64 chip). Future accelerators may share physical and/or virtual memory with the host, and may support a stronger memory model with richer synchronization methods. Software and hardware cache architectures are likely to change rapidly. A robust programming model should express parallelism broadly enough that compilers and tools can map an application onto future generations of accelerators as well as it does onto today’s GPUs. In fact, a successful model should be able to map applications onto a multicore X64 processor, where the SSE instructions implement the synchronous parallelism, and the doall parallelism is mapped across cores. From the available details, this model would even map well onto Intel’s proposed Larrabee chip. There will be work to tune the performance for each architecture, both in the tools and even at the application level, but the parallelism model needs to be reasonably portable.

Programming Model

How should we implement an accelerator-targeted programming model? Three options immediately come to mind: library, language, or directives. If you look at the array of parallel programming choices (all intended to make parallel programming easy), they span all three options.

Library-based solutions are attractive for many such problems; they are easy to port and can be independent of processor or compiler vendor. The MPI communication library for large system communication is one well-known example. It’s often easier to create (and modify) a standard for a library than for a language.

Language-based solutions expose the semantics in the language, allowing compilers or other tools to analyze and optimize the program. Co-Array Fortran, which is (currently) part of the next (allegedly minor) revision of the Fortran standard, exposes MPI-like parallelism and communication in the language, similar in some respects to Unified Parallel C (UPC). A compiler for Co-Array Fortran might be able to discover that a data copy from one image (thread) to another in a loop could be vectorized, given the appropriate support in the communication layer; such analysis in an MPI program is left entirely to the programmer. However, languages are expensive to implement, typically change quite slowly, and mistakes are hard to remedy once the standard is written.

A directive-based approach has some of the advantages of language-based solutions, in that directives expose the semantics to the compiler and other tools, allowing intelligent analysis and optimization. Such an approach also allows a program to be developed and tested on platforms that don’t support the directives, since the base language is unchanged. OpenMP is a widely available, successful parallel programming model based on directives to describe the parallel regions of the program.

Getting good performance on today’s accelerators depends on selecting a region that has enough work to amortize the overhead of moving data between the host and accelerator. This is one instance of the more general problem of selecting a region that has enough compute intensity to amortize the data traffic across the memory hierarchy, be it separate memory or multilevel caches. Some day, we may trust compilers to make this determination automatically, but not yet. So let me propose a model that borrows strategies from OpenMP, since I’m the PGI representative to the OpenMP group. I’ll propose directives in C and Fortran programs to delineate the regions of the program (loops) that should be accelerated (compiled for the GPU or other accelerator). Since the architecture model uses regular rectangular domains, I’ll propose using parallel loops as the primitive parallel operation.

The keys to tuning are minimizing and perhaps optimizing the data traffic between the host and accelerator, and selecting a schedule for the parallelism. In many cases, a compiler can analyze the nested loops and determine the input and output data sets, so it can manage the data traffic automatically. However, we should never trust automatic analysis to solve all our performance problems, so we’ll need directives or clauses to modify or augment the analysis.

As for scheduling, we saw in my previous column that there can be many possible schedules for even the simplest of parallel loops. Recent academic research in this area depends on doing more or less what I did by hand: generating many versions of the program and running each of them, then choosing the best one from the bunch; see Shane Ryoo’s PhD dissertation (University of Illinois, 2008), and joint work from Professors Ramanujam and Sadayappan (Louisiana State University and The Ohio State University) as good examples. Such an approach is valid for research, or when searching for a good algorithm for a highly tuned library, but inappropriate for a compiler. Instead, we will depend on the compiler to determine a reasonably good schedule (as we do when we use automatic parallelization and vectorization today), again with directives or clauses to modify or augment the decisions.

It’s important that a programmer be able to control any compiler optimization decision here; the difference between good and bad performance is quite dramatic, and at least in the immediate future, any compiler decision will be made with only partial information. However, to support this requires that the compiler tell the programmer what decisions it has made, and hopefully why, so the programmer knows whether it’s appropriate to step in and make a change.

So let me propose two directives. The first delineates an accelerator region, with optional clauses to control the data movement between host and accelerator memory. Borrowing liberally from OpenMP, I’ll propose a #pragma acc prefix for C directives, and !$acc prefix in Fortran. In C, I’ll describe an acceleration region as:

    #pragma acc region     {        /* loops to be accelerated go here */     } 

Fortran doesn’t have structured blocks (yet), so we’ll use region and end region directives:

   !$acc region        ! loops to be accelerated go here    !$acc end region 

Compare these to the OpenMP parallel regions. I propose optional clauses to tell the compiler what data needs to be copied into the region, from host to accelerator, what data needs to be copied out, and what data is local to the region; local data corresponds roughly to OpenMP private data. Compiler analysis is often able to determine the in/out/local data automatically.

The second directive is used to describe the mapping of parallel loops onto the hardware parallelism, what I called the schedule earlier. This corresponds roughly to the OpenMP loop directive, which describes the work-sharing pattern of parallel loops. It’s probably easiest to explain with a familiar example; in my most recent column, I showed several versions of matmul in CUDA with different schedules. The first (and simplest) version would be written (in Fortran) using these directives as:

   !$acc region       !$acc do parallel       do j = 1, m         do k = 1, p           !$acc do parallel, vector(32)           do i = 1, n             a(i,j) = a(i,j) + b(i,k)*c(k,j)           enddo         enddo       enddo    !$acc end region 

The loop directives do two things: the first is to tell the compiler about loop-level parallelism, augmenting its analysis. The second is to tell the compiler how to schedule or map the loop-level parallelism onto the hardware. In this loop, both the i and j loops exhibit doall parallelism, but we want to map the stride-1 i loop onto the synchronous (vector) parallelism in strips of size 32, using doall parallelism between the strips. We expect compilers to issue a warning message if a programmer inserts a do parallel directive on a loop that compiler analysis shows is in fact not parallel. Compare this code for clarity with the actual CUDA kernel.

This isn’t intended to be a user guide, tutorial, even a formal proposal, but I hope to convince you that a directive-based approach is feasible in the short term, and can address many of the problems programmers will face when porting large applications for use on host+GPU platforms in particular, and host+accelerators in general.

This model does use reasonably sophisticated compiler analysis, but nothing that hasn’t been implemented in commercial parallelizing compilers for many years. In this example, the compiler must take the following steps:

  • Determine what data is input to the region; for this loop, the input data is a(1:n,1:m), b(1:n,1:p), c(1:p,1:m), and the loop limits.
  • Determine what data is output to the region; this is simply a(1:n,1:m).
  • Determine what data is local to the region, which is empty (except perhaps for the loop counters). Classical data flow and array region analysis solves all three of these problems.
  • Determine which loops can run in parallel, augmented by information in the directives. For this loop, the j and i loops are completely parallel; the k loop requires a sum reduction, which is less efficient but could still be parallelized.
  • Determine the loop schedule; in this example, the schedule is specified by the directives. Without the loop directives, the compiler would have to search among the possible schedules and select a best one; note to academics: this is still a fertile area for continued research.
  • Generate code for the accelerator. For the most part, this is a classical compiler problem, and well known methods apply. On a target like the NVIDIA GPU, optimizing for the software-managed cache adds some complexity, but such problems have been addressed on past machines as well.
  • Generate host code to move data to the accelerator, launch the accelerator kernel(s), and move results back from the accelerator.

Final Words

Will adoption and use of directives such as these make GPUs more generally applicable? These directives may make GPUs more accessible, but there are still serious limitations to the parallelism GPUs support. The restrictions include rectangular domains, two levels of parallelism, limited synchronization, and a weak memory model (in the formal sense). This makes it unlikely that anyone will be porting unstructured mesh code or dynamic pointer-chasing data structures to a GPU anytime soon.

Can this programming model be adapted to make parallel programming easy in general? I’ve argued that parallel programming is difficult, and always will be, regardless of the programming model, and I’m not backing down. To reiterate, this directive model is intended to make accelerator programming accessible, so programmers can focus on algorithms and performance, not on syntax and other trivialities.

This proposed style of parallel programming isn’t universal, but it does address a significant segment of the parallel community. The model is portable, across GPUs, across accelerators, even to multicore CPUs, though we need to develop the compilers. Moreover, it’s nicely incremental; you can use these directives to accelerate parts of your program without having to undertake a whole rewrite, and, as with OpenMP, you can still build and test your application on the host by ignoring the directives altogether.

—–

Michael Wolfe has developed compilers for over 30 years in both academia and industry, and is now a senior compiler engineer at The Portland Group, Inc. (www.pgroup.com), a wholly-owned subsidiary of STMicroelectronics, Inc. The opinions stated here are those of the author, and do not represent opinions of The Portland Group, Inc. or STMicroelectronics, Inc.

—–

UPDATE: This article was original published during our SC08 coverage on November 20. Due to recent developments in heterogeneous compilers by PGI and CAPS Enterprise, as well as the ratification of the OpenCL standard, we felt it was worth another look. — Editor

Subscribe to HPCwire's Weekly Update!

Be the most informed person in the room! Stay ahead of the tech trends with industy updates delivered to you every week!

Why HPC Storage Matters More Now Than Ever: Analyst Q&A

September 17, 2021

With soaring data volumes and insatiable computing driving nearly every facet of economic, social and scientific progress, data storage is seizing the spotlight. Hyperion Research analyst and noted storage expert Mark No Read more…

GigaIO Gets $14.7M in Series B Funding to Expand Its Composable Fabric Technology to Customers

September 16, 2021

Just before the COVID-19 pandemic began in March 2020, GigaIO introduced its Universal Composable Fabric technology, which allows enterprises to bring together any HPC and AI resources and integrate them with networking, Read more…

What’s New in HPC Research: Solar Power, ExaWorks, Optane & More

September 16, 2021

In this regular feature, HPCwire highlights newly published research in the high-performance computing community and related domains. From parallel programming to exascale to quantum computing, the details are here. Read more…

Cerebras Brings Its Wafer-Scale Engine AI System to the Cloud

September 16, 2021

Five months ago, when Cerebras Systems debuted its second-generation wafer-scale silicon system (CS-2), co-founder and CEO Andrew Feldman hinted of the company’s coming cloud plans, and now those plans have come to fruition. Today, Cerebras and Cirrascale Cloud Services are launching... Read more…

AI Hardware Summit: Panel on Memory Looks Forward

September 15, 2021

What will system memory look like in five years? Good question. While Monday's panel, Designing AI Super-Chips at the Speed of Memory, at the AI Hardware Summit, tackled several topics, the panelists also took a brief glimpse into the future. Unlike compute, storage and networking, which... Read more…

AWS Solution Channel

Supporting Climate Model Simulations to Accelerate Climate Science

The Amazon Sustainability Data Initiative (ASDI), AWS is donating cloud resources, technical support, and access to scalable infrastructure and fast networking providing high performance computing (HPC) solutions to support simulations of near-term climate using the National Center for Atmospheric Research (NCAR) Community Earth System Model Version 2 (CESM2) and its Whole Atmosphere Community Climate Model (WACCM). Read more…

ECMWF Opens Bologna Datacenter in Preparation for Atos Supercomputer

September 14, 2021

In January 2020, the European Centre for Medium-Range Weather Forecasts (ECMWF) – a juggernaut in the weather forecasting scene – signed a four-year, $89-million contract with European tech firm Atos to quintuple its supercomputing capacity. With the deal approaching the two-year mark, ECMWF... Read more…

Why HPC Storage Matters More Now Than Ever: Analyst Q&A

September 17, 2021

With soaring data volumes and insatiable computing driving nearly every facet of economic, social and scientific progress, data storage is seizing the spotlight Read more…

Cerebras Brings Its Wafer-Scale Engine AI System to the Cloud

September 16, 2021

Five months ago, when Cerebras Systems debuted its second-generation wafer-scale silicon system (CS-2), co-founder and CEO Andrew Feldman hinted of the company’s coming cloud plans, and now those plans have come to fruition. Today, Cerebras and Cirrascale Cloud Services are launching... Read more…

AI Hardware Summit: Panel on Memory Looks Forward

September 15, 2021

What will system memory look like in five years? Good question. While Monday's panel, Designing AI Super-Chips at the Speed of Memory, at the AI Hardware Summit, tackled several topics, the panelists also took a brief glimpse into the future. Unlike compute, storage and networking, which... Read more…

ECMWF Opens Bologna Datacenter in Preparation for Atos Supercomputer

September 14, 2021

In January 2020, the European Centre for Medium-Range Weather Forecasts (ECMWF) – a juggernaut in the weather forecasting scene – signed a four-year, $89-million contract with European tech firm Atos to quintuple its supercomputing capacity. With the deal approaching the two-year mark, ECMWF... Read more…

Quantum Computer Market Headed to $830M in 2024

September 13, 2021

What is one to make of the quantum computing market? Energized (lots of funding) but still chaotic and advancing in unpredictable ways (e.g. competing qubit tec Read more…

Amazon, NCAR, SilverLining Team for Unprecedented Cloud Climate Simulations

September 10, 2021

Earth’s climate is, to put it mildly, not in a good place. In the wake of a damning report from the Intergovernmental Panel on Climate Change (IPCC), scientis Read more…

After Roadblocks and Renewals, EuroHPC Targets a Bigger, Quantum Future

September 9, 2021

The EuroHPC Joint Undertaking (JU) was formalized in 2018, beginning a new era of European supercomputing that began to bear fruit this year with the launch of several of the first EuroHPC systems. The undertaking, however, has not been without its speed bumps, and the Union faces an uphill... Read more…

How Argonne Is Preparing for Exascale in 2022

September 8, 2021

Additional details came to light on Argonne National Laboratory’s preparation for the 2022 Aurora exascale-class supercomputer, during the HPC User Forum, held virtually this week on account of pandemic. Exascale Computing Project director Doug Kothe reviewed some of the 'early exascale hardware' at Argonne, Oak Ridge and NERSC (Perlmutter), while Ti Leggett, Deputy Project Director & Deputy Director... Read more…

Ahead of ‘Dojo,’ Tesla Reveals Its Massive Precursor Supercomputer

June 22, 2021

In spring 2019, Tesla made cryptic reference to a project called Dojo, a “super-powerful training computer” for video data processing. Then, in summer 2020, Tesla CEO Elon Musk tweeted: “Tesla is developing a [neural network] training computer called Dojo to process truly vast amounts of video data. It’s a beast! … A truly useful exaflop at de facto FP32.” Read more…

Berkeley Lab Debuts Perlmutter, World’s Fastest AI Supercomputer

May 27, 2021

A ribbon-cutting ceremony held virtually at Berkeley Lab's National Energy Research Scientific Computing Center (NERSC) today marked the official launch of Perlmutter – aka NERSC-9 – the GPU-accelerated supercomputer built by HPE in partnership with Nvidia and AMD. Read more…

Esperanto, Silicon in Hand, Champions the Efficiency of Its 1,092-Core RISC-V Chip

August 27, 2021

Esperanto Technologies made waves last December when it announced ET-SoC-1, a new RISC-V-based chip aimed at machine learning that packed nearly 1,100 cores onto a package small enough to fit six times over on a single PCIe card. Now, Esperanto is back, silicon in-hand and taking aim... Read more…

Enter Dojo: Tesla Reveals Design for Modular Supercomputer & D1 Chip

August 20, 2021

Two months ago, Tesla revealed a massive GPU cluster that it said was “roughly the number five supercomputer in the world,” and which was just a precursor to Tesla’s real supercomputing moonshot: the long-rumored, little-detailed Dojo system. “We’ve been scaling our neural network training compute dramatically over the last few years,” said Milan Kovac, Tesla’s director of autopilot engineering. Read more…

CentOS Replacement Rocky Linux Is Now in GA and Under Independent Control

June 21, 2021

The Rocky Enterprise Software Foundation (RESF) is announcing the general availability of Rocky Linux, release 8.4, designed as a drop-in replacement for the soon-to-be discontinued CentOS. The GA release is launching six-and-a-half months after Red Hat deprecated its support for the widely popular, free CentOS server operating system. The Rocky Linux development effort... Read more…

Google Launches TPU v4 AI Chips

May 20, 2021

Google CEO Sundar Pichai spoke for only one minute and 42 seconds about the company’s latest TPU v4 Tensor Processing Units during his keynote at the Google I Read more…

Intel Completes LLVM Adoption; Will End Updates to Classic C/C++ Compilers in Future

August 10, 2021

Intel reported in a blog this week that its adoption of the open source LLVM architecture for Intel’s C/C++ compiler is complete. The transition is part of In Read more…

AMD-Xilinx Deal Gains UK, EU Approvals — China’s Decision Still Pending

July 1, 2021

AMD’s planned acquisition of FPGA maker Xilinx is now in the hands of Chinese regulators after needed antitrust approvals for the $35 billion deal were receiv Read more…

Leading Solution Providers

Contributors

Hot Chips: Here Come the DPUs and IPUs from Arm, Nvidia and Intel

August 25, 2021

The emergence of data processing units (DPU) and infrastructure processing units (IPU) as potentially important pieces in cloud and datacenter architectures was Read more…

Julia Update: Adoption Keeps Climbing; Is It a Python Challenger?

January 13, 2021

The rapid adoption of Julia, the open source, high level programing language with roots at MIT, shows no sign of slowing according to data from Julialang.org. I Read more…

10nm, 7nm, 5nm…. Should the Chip Nanometer Metric Be Replaced?

June 1, 2020

The biggest cool factor in server chips is the nanometer. AMD beating Intel to a CPU built on a 7nm process node* – with 5nm and 3nm on the way – has been i Read more…

HPE Wins $2B GreenLake HPC-as-a-Service Deal with NSA

September 1, 2021

In the heated, oft-contentious, government IT space, HPE has won a massive $2 billion contract to provide HPC and AI services to the United States’ National Security Agency (NSA). Following on the heels of the now-canceled $10 billion JEDI contract (reissued as JWCC) and a $10 billion... Read more…

Quantum Roundup: IBM, Rigetti, Phasecraft, Oxford QC, China, and More

July 13, 2021

IBM yesterday announced a proof for a quantum ML algorithm. A week ago, it unveiled a new topology for its quantum processors. Last Friday, the Technical Univer Read more…

Intel Launches 10nm ‘Ice Lake’ Datacenter CPU with Up to 40 Cores

April 6, 2021

The wait is over. Today Intel officially launched its 10nm datacenter CPU, the third-generation Intel Xeon Scalable processor, codenamed Ice Lake. With up to 40 Read more…

Frontier to Meet 20MW Exascale Power Target Set by DARPA in 2008

July 14, 2021

After more than a decade of planning, the United States’ first exascale computer, Frontier, is set to arrive at Oak Ridge National Laboratory (ORNL) later this year. Crossing this “1,000x” horizon required overcoming four major challenges: power demand, reliability, extreme parallelism and data movement. Read more…

Intel Unveils New Node Names; Sapphire Rapids Is Now an ‘Intel 7’ CPU

July 27, 2021

What's a preeminent chip company to do when its process node technology lags the competition by (roughly) one generation, but outmoded naming conventions make it seem like it's two nodes behind? For Intel, the response was to change how it refers to its nodes with the aim of better reflecting its positioning within the leadership semiconductor manufacturing space. Intel revealed its new node nomenclature, and... Read more…

  • arrow
  • Click Here for More Headlines
  • arrow
HPCwire