AMD GPUs for Scientific Computing: Reviewing the Software Ecosystem & Developer Experience

AMD GPUs for Scientific Computing Reviewing the Software Ecosystem & Developer Experience Matthew Leach (Research Software Engineer) Goals ●Briefly introduce how GPUs differ from CPUs ●Summarise the state of the AMD ecosystem ●Introduce primitives library/code samples ●Discuss my experience going from CUDA to HIP/ROCm 2 Talk Overview 1. Scientific Computing a. What is it? b. Motivation for GPUs 2. HIP/ROCm (vs CUDA) a. Programming model b. Support c. Features d. Tooling e. Libraries 3. Primitives Library 4. Porting a. Experience b. Automated tools 3 1. Scientific Computing 4 Scientific Computing - Intro For this talk, scientific computing refers to: ●Simulations/models ●Compute-heavy applications ○AI ○Parameter searching 5 Scientific Computing - Motivation for GPUs Why GPUs? ●GPUs are massively parallel - increased performance if the problem can be structured in the right way ●Able to solve larger problems ●Able to perform more iterations e.g. parameter sweeping/fitting Typically well suited to problems which perform the same operation on multiple data items e.g. ●Neural nets ●Agent-based models ●Some FEM formulations 6 Scientific Computing - Developing for GPUs ●Kernels are functions which are run on the device (GPU). Each thread will execute the code in the kernel ○Written in C++ ○Also use device API to interact with GPU specific-functionality e.g. accessing shared memory, atomic operations ●Host API is used to send/receive data to/from GPU, launch kernels ●Program with parallel algorithms to maximise performance ○Lots of serial algorithms have parallel equivalents - often involve iterations (as weʼll see later!) ○Sometimes challenging to map algorithms to a parallel equivalent 7 2. HIP/ROCm (vs CUDA) 8 HIP ●HIP stands for Heterogeneous-computing Interface for Portability ●Default programming model for building applications for AMD GPUs ●Lightweight API wrapper which can interface with the AMD or NVIDIA backends ●Very closely mirrors the CUDA API ●Goal is to be platform agnostic 9 HIP/ROCm - Documentation AMDʼs HIP/ROCm high-level documentation is comprehensive and well structured. It is divided into: ●an introductory section ●installation instructions ●a series of how-to guides covering AI, HPC and advanced features ●a section of ʻConceptualʼ guides ●finally, a reference section These are accompanied by code examples and a full repo of sample programs. API references can be more sparse, but can often be understood via looking at their CUDA equivalents 16 HIP/ROCm - Features ●__host__/__device__/__global__ qualifiers ●Hints for block size ●Pinned memory ●Memory fences e.g. __threadfence() ●Sync functions e.g. __syncthreads(), __syncthreads_and() ●High resolution timer access from kernels ●Atomic operations ●Warp shuffles/balloting ●Streams/concurrent execution/graph API 17 HIP/ROCm - CUDA features that arenʼt supported ●Independent thread scheduling - all AMD threads within a warp execute in lockstep ●Warp matrix types/functions - AMD cards donʼt have tensor cores ●There might be more! 18 HIP/ROCm - Tooling ●Compiler - hipcc ○Calls clang or nvcc depending on the target ●Profiling ○ROCm ships with the rocprofiler-sdk library ○rocprofv3 for collecting raw counter info & tracing ○rocprof-sys collects systems info - host/device/communication ○rocprof-compute performs roofline/baseline analysis of kernels ●Debugging ○ltrace ○ROCgdb - can force synchronous kernel execution to identify actual fault locations ●Porting ○HIPIFY - more on this later 19 HIP/ROCm - Libraries 20 CUDA Library hip Library roc Library Comment cuBLAS hipBLAS rocBLAS Basic Linear Algebra Subroutines cuBLASLt hipBLASLt Linear Algebra Subroutines, lightweight and new flexible API cuFFT hipFFT rocFFT Fast Fourier Transfer Library cuSPARSE hipSPARSE rocSPARSE Sparse BLAS + SPMV cuSOLVER hipSOLVER rocSOLVER Lapack library AmgX rocALUTION Sparse iterative solvers and preconditioners with algebraic multigrid Thrust rocThrust C++ parallel algorithms library CUB hipCUB rocPRIM Low Level Optimized Parallel Primitives cuDNN MIOpen Deep learning Solver Library cuRAND hipRAND rocRAND Random Number Generator Library NCCL RCCL Communications Primitives Library based on the MPI equivalents RCCL is a drop-in replacement for NCCL HIP/ROCm - AI AMD offer dedicated guidance on using ROCm for AI in their documentation ●Installation guides & docker images offered for ○PyTorch ○TensorFlow ○JAX ○verl ○Stanford Megatron-LM ○DGL ○Megablocks ○Taichi 21 ●Portability? Yes ●Performant portability - it depends ○Different hardware architecture ○E.g. NVIDIA provide heavily optimised deep learning primitive implementations (cudnn) specialised for each of their cards - no reason the same performance couldnʼt be reached with the AMD cards, but requires someone doing the implementation. ○For some applications performance may be comparable without significant effort HIP/ROCm - Performance Portability 22 HIP/ROCm - Uptake/Momentum ●Experiencing strong growth in both support & community ●Usage in many top tier supercomputers ○Exascacle - El Capitan, Frontier ●Improving support for consumer & pro grade cards, but cards donʼt stay supported for that long, historically issues with ROCm not being available for new cards upon release ●ROCm 7.0 coming - According to AMD, offering up to 3x perf for inference & AI training ●AMD developing ʻTheRockʼ which is a “open source build platform for HIP and ROCm”. Still in early preview. Aims to simplify setup, support multiple platforms 23 3. GPU Primitives Library 24 GPU Primitives Library - Intro ●Equivalent CUDA and HIP/ROCm implementations of fundamental GPU operations ○Searches ○Sorts ○Reductions ●CMake infrastructure so it can easily be built for AMD or NVIDIA gpus ○Also includes benchmarking infrastructure ●Useful to ○Those new to parallel algorithms ○Those new to GPU programming ○Developers looking to compare CUDA and HIP/ROCm code ●Also includes two more involved projects ○FEM solver ○SPH solver 25 Porting from CUDA - Automation HIPIFY offers analysis and automated conversion of CUDA code to HIP code, e.g. 32 Porting from CUDA - Automation Some manual effort may still be required due to API differences, but HIPIFY will point this out for you. Watch out for other differences e.g. warp size 33 The Future ●AMD News - 7.0 coming ○Closer alignment with CUDA API to reduce manual portion of porting effort ●Alternatives like ○ZLUDA - aiming to be a drop in replacement for CUDA which can run on AMD hardware, has had some legal issues ○SCALE - superset of CUDA, compiles for AMD without using any of the CUDA kit 34 Get in touch: m.leach@sheffield.ac.uk https://rse.shef.ac.uk/ 35