Applied Science, Faculty of
Electrical and Computer Engineering, Department of
DSpace
UBCV
Dashti, Mohammad K.
Doctor of Philosophy - PhD
University of British Columbia
In an era of hardware diversity, the management of applications' allocated memory is a complex task that can have significant performance repercussions. Non-uniformity in the memory hierarchy, along with heterogeneity and asymmetry of chip designs, make the costs of memory accesses unpredictable if the allocated memory is not managed carefully. Poor memory allocation and placement on symmetric, non-uniform memory access (NUMA) server systems can cause interconnect link congestion and memory controller contention, which can drastically impact performance. Furthermore, asymmetric systems consisting of CPU and GPU cores suffer similarly depending on how memory is allocated and what types of cores are accessing it. Furthermore, modern chip designs are integrating compute units into the memory modules to bring compute closer to memory and eliminate the high costs of transferring data. Hence, accessing memory efficiently is becoming increasingly challenging as systems evolve toward heterogeneity. Our contribution is a detailed analysis and insight into software and hardware memory management techniques on modern symmetric server-class and asymmetric heterogeneous processors consisting of integrated CPU and GPU cores, and to recently commercially available Processing In Memory (PIM) systems. In particular, we examine three types of systems that are affected by memory access bottlenecks. First, we look at NUMA systems, then integrated CPU-GPU systems, and finally PIM systems. We analyze these systems and suggest possible solutions to mitigate memory access issues. For NUMA systems, we propose a holistic memory management algorithm that intelligently distributes memory pages to reduce congestion and improve performance. Then, we provide a detailed analysis of memory management methods on integrated CPU-GPU systems focusing on performance and functionality trade-offs. Our goal is to expose the performance impact of memory management techniques and assess the viability and advantage of running applications with complex behaviors on such integrated CPU-GPU systems. Finally, we examine PIM systems with a case study of image decoding, which is a critical stage for many deep learning applications. We show that the extreme compute scalability of PIM systems can be utilized to accelerate image decoding with performance potential that can surpass CPUs and GPUs.
Towards Understanding and Mitigating Memory-AccessChallenges in Computing SystemsbyMohammad K. DashtiBSc, Old Dominion University, 2004MSc, King’s College London, 2006MSc, Simon Fraser University, 2012A THESIS SUBMITTED IN PARTIAL FULFILLMENTOF THE REQUIREMENTS FOR THE DEGREE OFDoctor of PhilosophyinTHE FACULTY OF GRADUATE AND POSTDOCTORALSTUDIES(Electrical and Computer Engineering)The University of British Columbia(Vancouver)August 2022c©Mohammad K. Dashti, 2022The following individuals certify that they have read, and recommend to the Fac-ulty of Graduate and Postdoctoral Studies for acceptance, the dissertation entitled:Towards Understanding and Mitigating Memory-Access Challenges inComputing Systemssubmitted by Mohammad K. Dashti in partial fulfillment of the requirements forthe degree of Doctor of Philosophy in Electrical and Computer Engineering.Examining Committee:Alexandra Fedorova, Professor, Electrical and Computer Engineering, UBCSupervisorMargo Seltzer, Professor, Computer Science, UBCCo-supervisorSathish Gopalakrish, Associate Professor, Electrical and Computer Engineering,UBCUniversity ExaminerChen Feng, Assistant Professor, School of Engineering, UBCUniversity ExaminerPascal Felber, Professor, Computer Science, University of NeuchatelExternal ExaminerAdditional Supervisory Committee Members:Matei Ripeanu, Professor, Electrical and Computer Engineering, UBCSupervisory Committee MemberMieszko Lis, Associate Professor, Electrical and Computer Engineering, UBCSupervisory Committee MemberiiAbstractIn an era of hardware diversity, the management of applications’ allocated mem-ory is a complex task that can have significant performance repercussions. Non-uniformity in the memory hierarchy, along with heterogeneity and asymmetry ofchip designs, make the costs of memory accesses unpredictable if the allocatedmemory is not managed carefully.Poor memory allocation and placement on symmetric, non-uniform memoryaccess (NUMA) server systems can cause interconnect link congestion and mem-ory controller contention, which can drastically impact performance. Furthermore,asymmetric systems consisting of CPU and GPU cores suffer similarly dependingon how memory is allocated and what types of cores are accessing it. Furthermore,modern chip designs are integrating compute units into the memory modules tobring compute closer to memory and eliminate the high costs of transferring data.Hence, accessing memory efficiently is becoming increasingly challenging as sys-tems evolve toward heterogeneity.Our contribution is a detailed analysis and insight into software and hardwarememory management techniques on modern symmetric server-class and asymmet-ric heterogeneous processors consisting of integrated CPU and GPU cores, andto recently commercially available Processing In Memory (PIM) systems. In par-ticular, we examine three types of systems that are affected by memory accessbottlenecks. First, we look at NUMA systems, then integrated CPU-GPU systems,and finally PIM systems. We analyze these systems and suggest possible solutionsto mitigate memory access issues.For NUMA systems, we propose a holistic memory management algorithmthat intelligently distributes memory pages to reduce congestion and improve per-iiiformance. Then, we provide a detailed analysis of memory management methodson integrated CPU-GPU systems focusing on performance and functionality trade-offs. Our goal is to expose the performance impact of memory management tech-niques and assess the viability and advantage of running applications with com-plex behaviors on such integrated CPU-GPU systems. Finally, we examine PIMsystems with a case study of image decoding, which is a critical stage for manydeep learning applications. We show that the extreme compute scalability of PIMsystems can be utilized to accelerate image decoding with performance potentialthat can surpass CPUs and GPUs.ivLay SummaryComputer systems are becoming increasingly complex. To maintain high perfor-mance and energy efficiency, software and hardware designers are adopting designswith many computing cores with complex memory hierarchies. This complexityaffects the performance of applications running on these systems as these appli-cations are mostly written without regard to the underlying hardware. This workprovides an analysis of these systems and examines how applications are affectedwhen run on these kinds of hardware. The work proposes solutions that alleviateperformance degradation and discusses challenges for adapting applications to runon such complex hardware.vPrefaceThe research work in this dissertation spans multiple previously published workscompleted both individually and in collaboration. Chapters 2 and 4 contain workdone in collaboration. Chapters 3 and 5 present my individual contributions.Chapter 2 is a modified version of previously published work. It is c©2013ACM, reprinted with permission from:• Mohammad Dashti, Alexandra Fedorova, Justin Funston, Fabien Gaud, Re-naud Lachaize, Baptiste Lepers, Vivien Quema, and Mark Roth. 2013. Traf-fic Management: A Holistic Approach to Memory Placement on NUMASystems. In Proceedings of the 18th International Conference on Architec-tural Support for Programming Languages and Operating Systems (ASPLOS13). ACM, New York, NY, USA, 381394.https://doi.org/10.1145/2451116.2451157The work in this PhD dissertation extends my previous work on NUMA sys-tems [27, 38]. This early work is distinct from the contribution presented in thisdissertation: The MSc thesis [38] 1) examined thread scheduling on NUMA nodesbased on classifying application threads by their memory access aggressiveness. 2)contained a brief overview of four main mechanisms that would be used in the Car-refour algorithm, 3) included analysis of the technique for sampling page accessesusing hardware counters.The work described in this PhD dissertation builds on this foundation, wherewe assume that the placement of threads on cores is already performed. We developthe Carrefour algorithm, which is concerned with the placement of memory pagesacross nodes through the application of the different page placement techniquespreviously discussed.viSince this chapter contains collaborative work, it is essential to clearly specifymy contributions. I started the initial analysis and implementation of the Carrefouralgorithm examining the transfer of pages across NUMA nodes. I later designedand implemented the module responsible for collecting metrics from the Perfor-mance Monitoring Unit (PMU!). Baptiste Lepers and Fabien Gaud later improvedthe module. Fabien Gaud, Justin Funston, Baptiste Lepers and I continued theexploration of NUMA aspects and concluded that memory traffic congestion ismore crucial to optimize than locality. I conducted several of the experiments inthis work and was the lead investigator of AutoNUMA and other Linux NUMApatches providing my insights and comparing its performance to our algorithm. Iwrote the sections of the papers relevant to my contributions. All co-authors in-cluding myself provided technical and editorial advice.Chapter 3 is a modified version from my previously published paper. It isc©2017 ACM, reprinted with permission from:• Mohammad Dashti, and Alexandra Fedorova. Analyzing memory manage-ment methods on integrated CPU-GPU systems. ACM SIGPLAN Interna-tional Symposium on Memory Management (ISMM), 2017.https://doi.org/10.1145/3092255.3092256I was the lead researcher, investigator, designer, and implementer of all thetechniques in this paper. I conducted all the experiments and wrote all the chaptersin this paper. My supervisor provided technical and editorial advice.Chapter 4 is a modified version of previously published work. It is c©2021USENIX, reprinted with permission from:• Joel Nider, Craig Mustard, Andrada Zoltan, John Ramsden, Larry Liu, Ja-cob Grossbard, Mohammad Dashti, Romaric Jodin, Alexandre Ghiti, JordiChauzi, Alexandra Fedorova, A Case Study of Computing-in-Memory inoff-the-Shelf Systems, to appear in USENIX Annual Technical Conference(ATC), 2021.https://www.usenix.org/system/files/atc21-nider.pdfviiI was the lead investigator of GPU and PIM performance comparisons. I im-plemented the GPU versions of snappy compression, snappy decompression, andhyper-dimensional computing. I performed all of the GPU experiments and as-sisted in writing and editing the paper.Chapter 5 Summarizes my analysis of image decoding on PIM systems. I con-ducted the research to investigate the viability of offloading image decoding toPIM systems. I improved and re-implemented a previous PIM image decodingalgorithm which was originally written by Daniel Ng. I performed all the experi-ments and wrote the whole manuscript (except for the JPEG introduction subsec-tion which was written by Daniel Ng and edited by me) with technical and editorialguidance from my supervisors.viiiTable of ContentsAbstract . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . iiiLay Summary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . vPreface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . viTable of Contents . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . ixList of Tables . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . xiiList of Figures . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . xiiiAcknowledgments . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . xvi1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 11.1 Summary of Contributions . . . . . . . . . . . . . . . . . . . . . 52 Memory Management on NUMA Systems . . . . . . . . . . . . . . . 72.1 Background . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 82.2 Traffic Congestion on Modern NUMA Systems . . . . . . . . . . 102.3 Design and Implementation . . . . . . . . . . . . . . . . . . . . . 142.3.1 The Mechanisms . . . . . . . . . . . . . . . . . . . . . . 142.3.2 The Algorithm . . . . . . . . . . . . . . . . . . . . . . . 152.3.3 Implementation . . . . . . . . . . . . . . . . . . . . . . . 182.4 Evaluation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 222.4.1 Testbed . . . . . . . . . . . . . . . . . . . . . . . . . . . 23ix2.4.2 Single-Application Workloads . . . . . . . . . . . . . . . 242.4.3 Multi-Application Workloads . . . . . . . . . . . . . . . 292.4.4 Overhead . . . . . . . . . . . . . . . . . . . . . . . . . . 302.4.5 Impact on Energy Consumption . . . . . . . . . . . . . . 342.4.6 Discussion: Hardware Support . . . . . . . . . . . . . . . 352.5 Related Work . . . . . . . . . . . . . . . . . . . . . . . . . . . . 362.6 Summary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 403 Analyzing Memory Management Methods on Integrated CPU-GPUSystems . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 413.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 423.2 Background . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 433.3 The NVIDIA/CUDA System . . . . . . . . . . . . . . . . . . . . 443.4 Evaluation of the NVIDIA/CUDA System . . . . . . . . . . . . . 483.4.1 Hardware/Software . . . . . . . . . . . . . . . . . . . . . 483.4.2 Benchmarks and Experiments . . . . . . . . . . . . . . . 493.5 AMD Kaveri APU . . . . . . . . . . . . . . . . . . . . . . . . . 593.5.1 OpenCL 1.2 and OpenCL 2.0 . . . . . . . . . . . . . . . 603.5.2 HSA: Heterogeneous System Architecture . . . . . . . . . 603.6 Evaluation of the AMD APU System . . . . . . . . . . . . . . . . 623.6.1 Hardware/Software . . . . . . . . . . . . . . . . . . . . . 623.6.2 Benchmarks and Experiments . . . . . . . . . . . . . . . 623.7 Related Work . . . . . . . . . . . . . . . . . . . . . . . . . . . . 713.8 Summary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 724 Analyzing an off-the-Shelf Processing-in-Memory System . . . . . . 734.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 734.2 Architecture . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 744.3 Evaluation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 774.3.1 Workloads . . . . . . . . . . . . . . . . . . . . . . . . . 774.3.2 Memory bandwidth scaling . . . . . . . . . . . . . . . . . 794.3.3 Data copy . . . . . . . . . . . . . . . . . . . . . . . . . . 834.3.4 DPU Speed . . . . . . . . . . . . . . . . . . . . . . . . . 84x4.3.5 Communication and control granularity . . . . . . . . . . 854.4 Summary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 855 Characterizing Image Decoding on Processing In Memory Systems 875.1 Introduction . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 875.2 Preprocessing and the CPU bottleneck . . . . . . . . . . . . . . . 885.2.1 Preprocessing in DNN Training . . . . . . . . . . . . . . 895.2.2 Preprocessing Primitives and Decoding . . . . . . . . . . 915.3 PIM-Decoding: Overview and Implementation . . . . . . . . . . 935.3.1 JPEG introduction . . . . . . . . . . . . . . . . . . . . . 935.3.2 PIM-Decoder . . . . . . . . . . . . . . . . . . . . . . . . 945.3.3 Tasklets . . . . . . . . . . . . . . . . . . . . . . . . . . . 955.4 Methodology . . . . . . . . . . . . . . . . . . . . . . . . . . . . 965.5 Evaluation . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 975.5.1 Bulk Image Decoding . . . . . . . . . . . . . . . . . . . 975.5.2 Performance Breakdown . . . . . . . . . . . . . . . . . . 985.6 PIM and Emerging Applications . . . . . . . . . . . . . . . . . . 1025.7 Related Work . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1025.8 Summary . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1036 Conclusion . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 104Bibliography . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 106xiList of TablesTable 2.1 NUMA traffic congestion effects . . . . . . . . . . . . . . . . 12Table 2.2 Statistics used by the Carrefour algorithm . . . . . . . . . . . . 15Table 2.3 Number of pages replicated, interleaved, and co-located . . . . 28Table 2.4 Carrefour performance improvement over Linux . . . . . . . . 29Table 2.5 Carrefour performance improvement for multiple applications . 33Table 3.1 Summary of allocation schemes . . . . . . . . . . . . . . . . . 48Table 3.2 Data TLB loads and Interrupts . . . . . . . . . . . . . . . . . 65Table 5.1 CPU, GPU, and PIM systems used in experiments . . . . . . . 96xiiList of FiguresFigure 1.1 A modern NUMA system, with four nodes and four cores pernode. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2Figure 1.2 A high level view of the NVIDIA integrated CPU-GPU systemused in this work. . . . . . . . . . . . . . . . . . . . . . . . . 3Figure 1.3 The Processing In Memory (PIM) architecture used in this work. 4Figure 2.1 Performance by thread and memory configuration . . . . . . . 11Figure 2.2 Streamcluster IC imbalance by memory configuration . . . . . 14Figure 2.3 Global decisions in Carrefour. . . . . . . . . . . . . . . . . . 16Figure 2.4 Carrefour evaluation, PARSEC/Metis . . . . . . . . . . . . . 25Figure 2.5 Carrefour evaluation, NAS . . . . . . . . . . . . . . . . . . . 26Figure 2.6 Load imbalance for selected benchmarks . . . . . . . . . . . 27Figure 2.7 Latency and locality for selected benchmarks . . . . . . . . . 28Figure 2.8 Carrefour evaluation, multi-application workloads . . . . . . 31Figure 2.9 Multi-application memory controller imbalance . . . . . . . . 32Figure 2.10 Multi-application interconnect imbalance . . . . . . . . . . . 32Figure 2.11 Multi-application memory latency . . . . . . . . . . . . . . . 32Figure 2.12 Multi-application local memory access ratio . . . . . . . . . . 33Figure 2.13 Effect of Carrefour on energy consumption . . . . . . . . . . 35Figure 3.1 sharedalloc usage example . . . . . . . . . . . . . . . . 48Figure 3.2 Write bandwidth under different memory management methods. 50Figure 3.3 Performance of Rodinia applications with the different globalmemory allocation methods. . . . . . . . . . . . . . . . . . . 51xiiiFigure 3.4 Accesses are cached in L2. (a) Each GPU thread accesses itsrespective array index. (b) All GPU threads access the samearray index. . . . . . . . . . . . . . . . . . . . . . . . . . . . 53Figure 3.5 Accesses are cached in L1. (a) Each GPU thread accesses itsrespective array index. (b) All GPU threads access the samearray index. . . . . . . . . . . . . . . . . . . . . . . . . . . . 55Figure 3.6 CPU-GPU concurrent benchmarks. Both the CPU and GPUwork on independent data. Default-conc (y-axis) is the defaultimplementation of the concurrent version of the benchmarkwhere data has to be explicitly copied between the CPU/GPU 56Figure 3.7 The locking scheme used in fine-grained sharing . . . . . . . 58Figure 3.8 Fine-grained sharing benchmark showing the performance ofhostalloc and sharedalloc . . . . . . . . . . . . . . . 58Figure 3.9 Memory system architecture of AMD Kaveri. Figure repro-duced from [93]. CP: Command Processor, GNB: GarlicNorth Bridge, UNB: Unified North Bridge. . . . . . . . . . . 61Figure 3.10 Write memory bandwidth for OpenCL 1.2 (coarse), OpenCL2.0 (SVM) and HSA (fine-grained). . . . . . . . . . . . . . . 63Figure 3.11 Write throughput as multiple threads concurrently update asingle memory location allocated as coarse (non-coherent) orfine (coherent) with HSA. . . . . . . . . . . . . . . . . . . . 66Figure 3.12 Latency of OpenCL 1.2, OpenCL 2.0, HSA kernel executiontime. Array size = 64 MB and each array entry is either notpadded, padded to 32 Bytes, or padded to 64 Bytes (cache linesize). (a) Kernel: memory[idx]++;. (b) Kernel: if(idx%2 == 0 ) memory[idx]++;. . . . . . . . . . . . . . . 67Figure 3.13 Slowdown of OpenCL 2.0 compared to OpenCL 1.2 as we in-crease the padding size of array entries. . . . . . . . . . . . . 68Figure 3.14 Runtime of Rodinia applications normalized to OpenCL 1.2 . 70Figure 4.1 DPU Architecture . . . . . . . . . . . . . . . . . . . . . . . . 76Figure 4.2 Memory bandwidth scales with memory capacity. . . . . . . . 79Figure 4.3 Snappy compression throughput . . . . . . . . . . . . . . . . 80xivFigure 4.4 Runtime breakdown for compression for dickens file . . . . . 81Figure 4.5 Snappy compression on GPUs . . . . . . . . . . . . . . . . . 81Figure 4.6 GPU kernel dimension effects on performance for compres-sion - dickens 32MB file . . . . . . . . . . . . . . . . . . . . 82Figure 4.7 HDC speedup over the host as file size increases . . . . . . . 83Figure 4.8 GPU HDC speedup over the host as file size increases . . . . 84Figure 4.9 Snappy decompression on the GPU relative to a single host CPU. 85Figure 5.1 Training throughput of Resnet18 with ImageNet-1K datasetwith preprocessing done on the CPU, GPU, and no preprocessing 90Figure 5.2 Training throughput of Alexnet with ImageNet-1K dataset withpreprocessing done on the CPU, GPU, and no preprocessing . 91Figure 5.3 Training throughput of Mobilenet with ImageNet-1K datasetwith preprocessing done on the CPU, GPU, and no preprocessing 92Figure 5.4 Breakdown of the major primitives used in preprocessing withDALI-CPU and DALI-GPU . . . . . . . . . . . . . . . . . . 92Figure 5.5 Time to decode a sample image with varying number of tasklets 95Figure 5.6 Decoding time for increasing number of images. Each JPEGimage is 2MB in size . . . . . . . . . . . . . . . . . . . . . . 98Figure 5.7 Decoding time for increasing number of images. Each JPEGimage is 311KB in size . . . . . . . . . . . . . . . . . . . . . 99Figure 5.8 Decoding time for increasing number of images. Each JPEGimage is 67KB in size . . . . . . . . . . . . . . . . . . . . . 100Figure 5.9 Breakdown of PIM-decoding for 2MB images . . . . . . . . . 100Figure 5.10 Breakdown of PIM-decoding for 311KB images . . . . . . . 101Figure 5.11 Breakdown of PIM-decoding for 67KB images . . . . . . . . 101xvAcknowledgmentsThis is the most important section of the whole dissertation. None of the con-tent of this dissertation would have been created without what I’m about to say inthis section. I used to read the Acknowledgement section in dissertations of pastPhD students from different universities in various places around the world to geta glimpse of how they were able to achieve what I thought was an impossible task.Because, let me tell you this, getting a PhD is anything but a smooth ride. In theacknowledgement section, students typically thank their supervisors, friends, col-leagues, and family members amongst other things. This is absolutely mandatoryas the people you thank definitely had a significant effect on you, but I assure youthat you probably also fought with all those people you mention sometime duringthe PhD and probably (out of frustration maybe) thought that they are partly thereason why you are struggling in your PhD! The real question is, would you thankthose people if you had not successfully completed your PhD? For me the answeris yes.My PhD journey took almost ten years! I quit the PhD two times during thisperiod and I thought that (more like I was certain that) it would be impossible forme to successfully complete the PhD. I reached the conclusion that it wouldn’t bethe end of the world if I didn’t complete the PhD. I started searching for peoplewho were successful who were also PhD drop-outs. I tried to convince myself thatleaving the PhD was a normal thing. And, I can guarantee you that it is definitelynormal and it is not the end of the world! However, for me, It didn’t feel right and Ialways felt that I had to get this done, but at the same time I felt it was impossible.The first person who made this feel less impossible was Sasha. Sasha has beenmy supervisor since January 2010 at SFU. Sasha pulled me back in after I decidedxvito leave the PhD the first time and injected some hope that kept me going. Whenthat hope drained, someone else came into the rescue as I was almost drowningin my PhD. This person is Margo, who also became my supervisor during the lastyear of my PhD at UBC. Margo, you know that I can’t thank you enough for whatyou did. I still remember the email that I sent you and how you accommodatedmy frustration and provided the perfect solution and then actively jumped in andhelped me down the stretch.Of course, my family was always supportive and made life meaningful out-side of the academic world. The first person in my family that I’m indebted to ismy mother. Thank you Mom for everything, and I know that thanking you in anacknowledgement section is infinitesimal and comes nowhere close to what youdeserve. I also have to thank my brothers for keeping me sane during the past tenyears or so (more like since I was born :P)! Thanks guys. Thanks to my Aunt andUncle (my aunt’s husband) for always being there for us. Thanks to my cousinswho are more like brothers to me. Thanks to all my family members who werethere for us.I also have to thank all my lab mates from both SFU and UBC. Special thanksgo to Tyler, Craig, and Svetozar. Thanks to all the lab mates who convinced mefor making the switch to Vim from Emacs. This change had a tremendous positiveeffect in my life ever since :PLast, but certainly not least is my wife. Thank you for always being there forme. Thanks for leaving everything behind and coming with me to Canada. It wouldbe impossible to have achieved this if you weren’t by my side all these years. I alsowant to thank my kids. You are the apple of my eyes and I love you!Now if you don’t have faith or don’t believe in the metaphysical world (depend-ing on your beliefs, this might be physical rather than metaphysical), don’t read thisparagraph :). Despite all the hardship that I’ve been through during my PhD (andevery hardship in life), I never lost faith in Ahlulbayt. They are the fortresses thatalways provide love, tranquility, and safety to those who sought refuge in them.My love for Imam Hussain kept me going and will keep me going no matter whatlife throws at me.xviiIn memory of my father and my grandmother (Yommah)xviiiChapter 1IntroductionAccessing memory efficiently has become exceedingly more challenging as sys-tems have become more heterogeneous. Due to power limitations and the break-down of Dennard scaling [39, 45], chip designers can no longer rely on the scalingdown of transistors and increasing of clock frequencies to increase performanceand reduce energy. The focus has turned to multicore chip designs where multi-ple symmetric or asymmetric cores occupy the chip area. Furthermore, systemsintegrating different types of accelerators are becoming widely available [54, 93].In symmetric designs, all cores are of the same type and there is usually fullcache coherency between the cores. Modern servers are implemented in this man-ner. The servers are built from several processor nodes, each containing a multicoreCPU and a local DRAM serviced by one or more memory controllers, as shownin figure 1.1. The nodes are connected into a single cache-coherent system by ahigh-speed interconnect. Physical address space is globally shared, so all corescan transparently access memory in all nodes. Accesses to a local node go througha local memory controller and accesses to remote nodes must traverse the intercon-nect and access a remote controller. Remote accesses typically take longer thanlocal ones, giving these systems a property of Non-Uniform Memory Access time(NUMA!). In Chapter 2, we show that poor memory placement can lead to memorytraffic congestion, which may heavily impact performance. The primary contri-bution in that chapter is a holistic method for memory placement that providessignificant performance improvement over existing methods.1Memory Node #1 Memory Node #2Memory Node #3 Memory Node #4Node 1L3 cacheC5L3 cacheC6C8C7L3 cacheC13L3 cacheC14C16C15Node 2Node 3 Node 4C1 C2C4C3C9 C10C12C11Figure 1.1: A modern NUMA system, with four nodes and four cores pernode.After studying memory management in symmetric server-class systems, wethen examine asymmetric systems. In asymmetric systems, there might be twoor more different types of processors integrated on the same chip or connectedthrough a Peripheral Component Interconnect (PCI) bus. One processor could bea general-purpose CPU while the other could be a GPU, FPGA, fixed-functionaccelerator, etc. One of the emerging hardware architectures conforming to theasymmetric design is to have an accelerator type of processor such as a GPU in-tegrated on the same die as a general purpose processor where they both share thesame physical memory (figure 1.2).The integrated CPU-GPU architecture can eliminate data transfer costs com-pared to traditional discrete systems where the accelerator sits on a separate chipconnected to the CPU through a PCI! bus. Such tight integration enables the CPUand the accelerator to share data without explicit copying. The elimination of datacopies can sometimes significantly improve performance compared to much more2 CPU 1 ARM A15 D$ 32K I$ 32K CPU 0 ARM A15 D$ 32K I$ 32K CPU 2 ARM A15 D$ 32K I$ 32K CPU 3 ARM A15 D$ 32K I$ 32K L2 $ (2MB) CPU Complex 64K Shared/L1$ 128K L2$ Kepler GPU 192 CUDA cores Memory Controller DRAM (2GB) 32K Registers Figure 1.2: A high level view of the NVIDIA integrated CPU-GPU systemused in this work.powerful discrete GPUs. For compute-intensive applications, a discrete GPU isgenerally better due to its superior compute capabilities. However, for applicationsthat need to transfer large amounts of data between the CPU and GPU, a less pow-erful integrated GPU, which consumes less energy can significantly outperform adiscrete GPU.Furthermore, the theoretical speedup potential for asymmetric chips comparedto symmetric ones is higher [59]. In other words, if there is enough parallelismin a given application, running this application on an asymmetric chip designedusing a specific number of transistors performs better compared to running it on asymmetric chip designed with the same transistor budget. However, on currentlyavailable integrated systems, there are a variety of memory management methodsand their internals are poorly analyzed and understood as the memory on these sys-tems is typically managed by a software framework such as OpenCL or CUDA, inaddition to a runtime library and a GPU driver. These frameworks offer a rangeof memory management methods that vary in ease of use, consistency guaranteesand performance. This introduces uncertainty as to which method a programmer3should use, and whether reduced performance is worth the improved programma-bility. Our main contribution in chapter 3 is a detailed analysis of memory man-agement methods for such integrated systems. We investigate how these differentmemory management methods behave in terms of performance and functionality,reveal the performance trade-offs that these memory management strategies make,and explain the technical reasons for these performance effects.Figure 1.3: The Processing In Memory (PIM) architecture used in this work.Furthermore, the Von Neumann bottleneck (or the memory wall) continues toproduce an imbalance between CPU and memory system performance. Hence,there has been a trend in system architectures where designers try to bring com-pute closer to memory in order to eliminate this bottleneck. Our last contributionin this dissertation is analyzing a recently available commercial product that imple-ments this Processing In Memory (PIM) architecture. Figure 1.3 shows an abstractrepresentation of the PIM system that we consider in this work. Such tight com-pute and memory integration presents challenges that we analyze and evaluate inchapter 4. Furthermore, in chapter 5, we examine image decoding as a case studyon PIM systems. Image classification in Deep Neural Network (DNN) training re-lies heavily on image decoding. Furthermore, the DNN training and preprocessingpipelines in regular CPU+GPU servers have been shown to be bottlenecked by thedecoding computations [78, 105], so we analyze such a critical workload and showthat using a PIM architecture has a significant potential for such problems. We also4discuss some of the limitations that are present on currently available PIM systems.1.1 Summary of Contributions• In the first part of this thesis (chapter 2), we present our collaborative workto show that poor memory placement on NUMA systems can lead to mem-ory traffic congestion, which may heavily impact performance. We dis-cover that memory traffic congestion is more crucial to optimize than lo-cality and present a holistic method, Carrefour, for memory placement thatprovides significant performance improvement over existing methods. Weimplemented Carrefour in Linux and obtained significant performance im-provements relative to the default kernel, as well as significant improvementscompared to NUMA-aware patchsets available for Linux.• In the second part of the thesis (chapter 3), we try to answer the followingquestion: When applications run on integrated CPU/GPU systems, where ex-plicit memory copy is not necessary, are performance advantages easily ob-tained or are there programmability challenges to obtaining high-performingmemory accesses on these systems? Our contribution of this second partis a detailed analysis of memory management methods on such integratedsystems. We study their viability for accelerating different types of appli-cations and provide a clear view of the current landscape of such systems.We show that to conceal low-level details of memory management, hardwareand software manufacturers sacrifice performance, and our work is the firstto measure the impact that programmability and ease of use have on per-formance. We analyze the common memory management methods of themost widely used software frameworks for heterogeneous systems: CUDA,OpenCL 1.2, OpenCL 2.0, and HSA, on NVIDIA and AMD hardware. Wefocus on performance and functionality trade-offs, with the goal of expos-ing their performance impact and assess the viability of running applicationswith complex behaviors on such integrated CPU-GPU systems.• In the third part of the thesis (chapters 4, and 5), the first contribution is ourcollaborative work on studying processing in memory characteristics on a5currently available commercial product from UPMEM. We show that somememory-bound applications can be limited by the conventional memory hi-erarchy, and PIM can alleviate these limitation and improve performance forthese applications. We also compare the performance of this system withother CPU and GPU systems and demonstrate some of the current limita-tions of PIM.The second contribution is presented in chapter 5 where I present a casestudy of image decoding on PIM systems. This case study expands on theoutcomes of chapter 4 and presents a PIM-decoder implementation and eval-uation that demonstrate the capabilities of PIM systems for such a complextask. We show that using such memory slots for extra compute capabilitiescan be a viable alternative to systems which only consists of CPUs and GPUsas the core compute units. We also discuss some of PIM’s limitations andthe need for assessing applications’ suitability for running on PIM systems.6Chapter 2Memory Management on NUMASystems1NUMA systems are characterized by Non-Uniform Memory Access times, whereaccessing data on a remote node takes longer than accessing data on a local one.NUMA hardware has been built since the late 80’s, and the operating systemsdesigned for it were optimized for access locality. They co-located memory pageswith the threads that accessed them to avoid the cost of remote accesses. In contrastto older systems, modern NUMA hardware has much smaller remote wire delays,so remote access costs per se are not the main concern for performance, as we showin this work. Instead, congestion on memory controllers and interconnects, causedby memory traffic from data-intensive applications, hurts performance a lot more.Because of that, memory placement algorithms must be redesigned to target trafficcongestion. This requires an arsenal of techniques that go beyond optimizing lo-cality. In this chapter we describe Carrefour, an algorithm that addresses this goal.We implemented Carrefour in Linux and show that it significantly improves perfor-mance relative to the default kernel and to other NUMA-aware patchsets availablefor Linux. Carrefour never hurts performance by more than 3% when memoryplacement cannot be improved. We present the design of Carrefour, the challengesof implementing it on modern hardware, and draw insights about hardware supportthat would help optimize system software on future NUMA systems.1This chapter is a modified version of a previously published work [37].72.1 BackgroundModern servers are built from several processor nodes, each containing a multi-core CPU and a local DRAM serviced by one or more memory controllers (seeFigure 1.1). The nodes are connected into a single cache-coherent system by ahigh-speed interconnect. Physical address space is globally shared, so all corescan transparently access memory in all nodes. Accesses to a local node go througha local memory controller; accesses to remote nodes must traverse the interconnectand access a remote controller. Remote accesses typically take longer than lo-cal ones, giving these systems the property of Non-Uniform Memory Access time(NUMA).It is well understood that optimal performance on NUMA systems can beachieved only if we place threads and their memory in consideration of the sys-tem’s physical layout. For instance, previous work on NUMA-aware memoryplacement focused on maximizing locality of accesses, where the placement ofmemory pages is done such that data accesses are satisfied from a local node when-ever possible [19, 28]. That was done to avoid very high costs of remote memoryaccesses. Contrary to insights from previous work, we discover that on modernNUMA systems remote wire delays, which are delays resulting from traversing agreater physical distance to reach a remote node, are not the most important sourceof performance overhead. However, congestion on interconnect links and in mem-ory controllers, which results from high volume of data flowing across the system,can dramatically hurt performance. This motivates the design of new NUMA-aware memory placement policies.To make these statements concrete, consider the following facts. On NUMAsystems circa 1990, the time to access data from a remote node took 4-10 timeslonger than accessing from a local node [99]. On NUMA systems that are builttoday, remote wire delays add at most 30% to the cost of a memory access [29]. Formost programs, this latency differential alone would not have a substantial impacton performance. However, fast modern CPUs are able to generate memory requestsat very high rates. Massive data traffic creates congestion in memory controllerqueues and on interconnects. When this happens, memory access latencies canbecome as large as 1000 cycles, from a normal latency of only around 200. Such a8dramatic increase in latencies can slow down data-intensive applications by morethan a factor of three. Fortunately, high latencies can be avoided or substantiallyreduced if we carefully place memory pages on nodes to avoid traffic congestion.In response to the changes in hardware bottlenecks, we approach the problemof thread and memory placement on NUMA systems from an entirely new per-spective. We look at it as the problem of traffic management. Our algorithm, calledCarrefour2, places threads and memory so as to avoid traffic hotspots and preventcongestion in memory controllers and on interconnect links. This is akin to trafficmanagement in the context of city planning: popular residential and business hubsmust be placed so as to avoid congestion on the roads leading to these destinations.The mechanisms used in our algorithm: e.g., migration and replication of mem-ory pages, are well understood, but the algorithm itself is new. Our algorithmmakes decisions based on global observations of traffic congestion. Previous algo-rithms optimized for locality, and relied on local information, e.g., access patternof individual pages. We found that in order to effectively manage congestion onmodern systems we need an arsenal of techniques that go beyond optimizing local-ity. While locality plays a role in managing congestion (when we reduce remoteaccesses, we reduce interconnect traffic), alone it is not sufficient to achieve thebest performance. The challenge in designing Carrefour was to understand how tocombine different mechanisms in an effective solution for modern hardware.Implementing an effective NUMA-aware algorithm on modern systems presentsseveral challenges. Modern systems do not have the same performance monitoringhardware that was present (or assumed) on earlier systems. Existing instructionsampling hardware cannot gather the profiling data needed for the algorithm withthe desired accuracy and speed. We had to navigate around this problem in ourdesign. Furthermore, the memory latencies that we are optimizing are lower thanon older systems, so we can tolerate less overhead in the algorithm.We implemented Carrefour in Linux and evaluated it with several data-centricapplications: k-means clustering, facial recognition, map/reduce, and others. Car-refour significantly improves performance of these applications compared to otherexisting scheduling methods. When memory placement cannot be improved Car-2Carrefour– (French) intersection, crossroads.9refour never hurts performance by more than 3%. Existing NUMA-aware patchesfor the Linux kernel perform less reliably and, in general, fall short of improve-ments achieved with Carrefour.2.2 Traffic Congestion on Modern NUMA SystemsWe first demonstrate that the effects of traffic congestion are more substantial thanthose of wire delays, motivating why memory placement algorithms must be re-designed. We report data from two sets of experiments. In the first set, our goalis to measure the effects of wire delays only. We run applications in two con-figurations: local-memory and remote-memory. Under local-memory, the threadand its data are co-located on the same NUMA node; under remote-memory, thethread runs on a different node than its data. To ensure that wire delay is the dom-inant performance factor, we had to avoid congestion on memory controllers andinterconnects, so we run one application at a time and use only one thread in eachapplication. We do not include applications with CPU utilization less than 30%,because memory performance is not their main bottleneck. The experiments arerun on a system described in Section 2.4 (Machine A). We use applications fromthe NAS, PARSEC and map/reduce Metis suites, also described in Section 2.4.Figure 2.1(a) shows relative completion time under remote-memory vs. local-memory configuration. The performance degrades by at most 20% under remote-memory, which is consistent with at most 30% difference in local-vs-remote mem-ory latencies measured in microbenchmarks [29].In the second set of experiments, we want to observe traffic congestion, sowe run each application with as many threads as there are cores. We demonstratehow performance varies under two memory placement policies on Linux, as theyinduce different degrees of traffic congestion. The first policy is First-touch (F): thedefault policy where the memory pages are placed on the node where they are firstaccessed. The second policy is Interleaving (I), where memory pages are spreadevenly across all nodes. Although these are not the only possible policies and notnecessarily the best policies, comparing them illustrates the salient effects of trafficcongestion.Figure 2.1(b) shows the absolute difference in completion time between first-10 0 20 40 60 80 100BT CG DC EP FT IS LU MG SP UA bodytrack facesim fluidanimate streamcluster swaptions x264 kmeans matrixmult PCA wrmem Per for mance di ff er ence (%)(a) Performance difference for single-thread versions of applications between local andremote memory configurations. 0 20 40 60 80 100BT (F)CG (F)DC (F)EP (-)FT (F)IS (I)LU (F)MG (F)SP (F)UA (F)bodytrack (-)facesim (I)fluidanimate (-)streamcluster (I)swaptions (-)x264 (I)kmeans (I)matrixmult (-)PCA (I)wrmem (F)Per for mance di ff er ence (%)(b) Absolute performance difference for multi-thread versions of applications betweenFirst-touch (F) and interleaving (I).Figure 2.1: Performance difference of applications depending on the threadand memory configuration.11Streamcluster PCABest (I) Worst (F) Best (I) Worst (F)Local access ratio 25% 25% 25% 33%Memory latency 476 1197 465 660Mem-ctrl. imbalance 8% 170% 5% 130%IC: imbalance, (avg) 22% (59%) 85% (33%) 20% (48%) 68% (31%)L3MPKI 16.85 16.89 7.35 7.4IPC 0.29 0.15 0.52 0.36Table 2.1: NUMA traffic congestion effectstouch and interleaving. The policy that performed the best is indicated in paren-thesis next to the application name; a “-” is shown when the application performsequally well with either policy. We observe that the differences are often muchlarger than what we can expect from wire delays alone. For Streamcluster, a k-means clustering application from PARSEC, the performance varies by a factor oftwo depending on the memory placement policy.To illustrate that these differences are due to traffic congestion, we show in Ta-ble 2.1 some supporting data for the two applications, Streamcluster and PCA (amap/reduce application):Local access ratio: The percent of all memory accesses sourced from a local node.Memory latency: The average number of cycles to satisfy a memory request fromany node.Memory controller imbalance: The standard deviation of the load across allmemory controllers, expressed as percent of the mean. Load is measured as thenumber of requests per time unit.Average interconnect (IC) usage: The utilized interconnect bandwidth as percentof total, averaged across all links.Interconnect (IC) imbalance: The standard deviation of utilization across thelinks as percent of mean utilization.L3MPKI: The number of last-level (L3) cache misses per thousand instructions.IPC: The number of instructions per cycle.The data in Table 2.1 leads to several curious observations. First, we see thatlocality of memory accesses either does not change regardless of the memory man-12agement policy, or decreases under the better performing policy. For Streamclus-ter, most of the memory pages happen to be placed on a single node under first-touch (because a single thread initializes them at the beginning of the program).Under interleaving the pages are spread across all nodes, but since the threads ac-cess data from all four nodes, the overall access ratio is about the same in bothconfigurations. For PCA, interleaving decreases the local access ratio and yet in-creases performance. So the first surprising conclusion is that better locality doesnot necessarily improve performance.And yet, the IPC substantially improves (2× for Streamcluster and 41% forPCA), while the L3 miss rate, as well as L1 and L2 miss rates, remain unchanged.The explanation emerges if we look at the memory latency. Under interleaving, thememory latency reduces by a factor of 2.48 for Streamcluster and 1.39 for PCA.This effect is entirely responsible for performance improvement under the betterpolicy. The question is, what is responsible for memory latency improvements? Itturns out that interleaving dramatically reduces memory controller and intercon-nect congestion by alleviating the load imbalance and mitigating traffic hotspots.Rows 5, 6 in Table 2.1 show significant reductions in imbalance under interleaving,and Figure 2.2 illustrates these effects visually for Streamcluster. So even with-out improving locality (we even reduce it for PCA), we are able to substantiallyimprove performance. Yet, existing NUMA-aware algorithms disregarded trafficcongestion, optimizing for locality only. Our work addresses this shortcoming.Although the two selected applications performed significantly better under in-terleaving, this does not mean that interleaving is the only desired policy on modernNUMA hardware. In fact, as Figure 2.1(b) shows, many NAS applications fared alot worse with interleaving. In the process of designing the algorithm we learnedthat a range of techniques — interleaving, page replication and co-location — mustbe judiciously applied to different parts of the address space depending on globaltraffic conditions and page access patterns. So the challenge in designing a goodalgorithm is understanding when to apply each technique, while navigating aroundthe challenges of obtaining accurate performance data and limiting the overhead.131%1%1%97%25% 25%25% 25%Figure 2.2: Traffic imbalance under first-touch (left) and interleaving (right)for Streamcluster. Nodes and links bearing the majority of the trafficare shown proportionately larger in size and in brighter colors. Thepercentage values show the fraction of memory requests destined foreach node. The figure is drawn to scale.2.3 Design and ImplementationWe begin by describing the mechanisms composing the algorithm: page co-location,interleaving, and replication. Then we explain how they fit together.2.3.1 The MechanismsPage co-location is when we place a physical page on the same node as the threadthat accesses it. Co-location works well for pages that are accessed by a singlethread or by threads co-located on the same node.Page interleaving is about evenly distributing physical pages across nodes.Interleaving is useful when we have imbalance on memory controllers and inter-connect links, and when pages are accessed by many threads. Operating systemsusually provide an interleaving allocation policy, but only give an option to enableor disable it globally for the entire application. We found that interleaving worksbest when judiciously applied to parts of the address space that will benefit from it.Page replication is about placing a copy of a page on several memory nodes.Replication distributes the pressure across memory controllers, alleviating traffichotspots. An added bonus is eliminating remote accesses on replicated pages.When done right, replication can bring very large performance improvements. Un-fortunately, replication also has costs. Since we keep multiple copies of the same14Global statisticsMC-IMB Memory controller imbalance (as defined in Section 2.2)LAR Local access ratio (as defined in Section 2.2)MAPTU Memory (DRAM) accesses per time unit (microsecond)Per-application statisticsMRR Memory read ratio. Fraction of DRAM accesses that are readsCPU% Percent CPU utilizationPer-page statisticsNumber of accesses The number of sampled data loads that fell in that pageAccess type Read-only or read-writeTable 2.2: Statistics collected for the algorithm.page, we must synchronize their contents, which is like running a cache coherencyprotocol in software. The costs can be very significant if there is a lot of fine-grained read/write sharing. Another potential source of overhead is the synchro-nization of page tables. Since modern hardware walks page tables automatically,page tables themselves must be replicated across nodes and kept in sync. Finally,replication increases the memory footprint. We should avoid it for workloads withlarge memory footprints for fear of increasing the rate of hard page faults.2.3.2 The AlgorithmOur memory management algorithm has three components: measurement, globaldecisions and page-local decisions. The measurement component continuouslygathers various metrics (Table 2.2) that later drive page placement decisions. Globaland per-application metrics are collected using hardware counters with very lowoverhead. Per-page statistics are collected via instruction-based sampling (IBS) [44],which can introduce significant overheads at high sampling rates. Section 2.3.3 de-scribes how we keep the overheads at bay. Global decisions are based on system-wide traffic congestion and workload properties which determine what mecha-nisms to use. Page-local decisions examine access patterns of individual pagesto decide their fate.15MAPTU > 50 ?Enable CarrefourDisable CarrefourYesNoStep 1MRR > 95%&&Free RAM ≥ 1 −_ ?Enable replicationDisable replicationYesNoStep 2MC_IMB > 35% ?Enable interleavingDisable interleavingYesNoStep 3LAR < 80% ?Enable co-locationDisable co-locationYesNoStep 4Figure 2.3: Global decisions in Carrefour.Global DecisionsThe global decision-making process is outlined in Figure 2.3.Step 1: We only want to run Carrefour for applications that generate substan-tial memory traffic. Other applications would not be affected by memory placementpolicies, so there is no reason to subject them to sampling overhead. This decisionis driven by the application’s memory access rate (MAPTU – see Table 2.2). Car-refour is enabled for applications with the MAPTU above a certain threshold. TheMAPTU threshold is to be determined experimentally and the right setting mayvary from system to system. We found the threshold of 50 MAPTU worked wellon all hardware we evaluated, and the performance was not very sensitive to itschoice. To determine the right MAPTU threshold on a system different from ours,we recommend running a benchmark suite under different NUMA policies, notingwhich applications are affected and using the lowest observed MAPTU from thoseexperiments.Once we decided whether there is sufficient memory traffic to justify runningCarrefour, we need to decide which of the available mechanisms, replication, inter-leaving and co-location, should be enabled for each application given its memoryaccess patterns. The goal here is to choose the most beneficial techniques and avoidassociated overhead. The next three steps take care of this decision.Step 2: We decide whether it is worthwhile to use replication. Replicationrisks introducing significant overhead if it exhausts RAM (and causes additionalhard page faults) or requires frequent synchronization of pages across nodes (seemore discussion in Section 2.3.3). To avoid the first peril, we conservatively en-16able replication only if there is sufficient free RAM to replicate the entire residentset. That is, the fraction of free RAM must be at least 1− 1NUM NODES . This isa conservative threshold, because not all pages will be replicated, and not all res-ident pages will be accessed frequently enough to generate significant page faultoverhead if evicted. Evaluating the trade-off between replication benefit and po-tentially increased page-fault rate was outside the scope of the work. This requiresworkloads that both benefit from replication and have very large memory-residentsets, which we did not encounter in our experiments.To avoid the overhead associated with the synchronization of page contentacross nodes, we do not replicate pages that are frequently written. An applicationmust have the memory read ratio (MRR) of at least 95% in order for its memorypages to be considered for replication3. The setting of this parameter can have avery significant effect on performance. While we found that the performance wasnot sensitive when we varied the parameter in the range of 90-99%, it is alwayssafe to err on the high side.Step 3: We decide whether to use interleaving. Interleaving improves perfor-mance if we have large memory controller imbalance. We enable interleaving ifmemory controller imbalance is above 35% but found that the performance wasnot highly sensitive to this parameter. Applications that benefit from interleavingusually begin with a large imbalance.Step 4: We decide whether or not to enable co-location. Co-location willbe triggered only for pages that are accessed from a single node, so it will notexacerbate the imbalance if memory-intensive threads are evenly spread acrossnodes, which is the case since threads are bound to cores. Therefore, we enableco-location if the local access rate is slightly less than ideal (LAR < 80%). Perfor-mance is not highly sensitive to this parameter; we observed that if this parameteris completely eliminated from the algorithm and co-location is always enabled thenthe largest performance impact is only a few percent.Although we expect that optimal settings for the parameters used in the algo-3MRR is approximated as fraction of L1 refills from DRAM in modified state, because there is nohardware counter that provides this quantity precisely per core, as opposed to per-node. Similarly,due to hardware counter limitations described in Section 2.3.3, it is very difficult to measure theMRR per page. That is why we use the MRR for the entire application.17rithm vary from one system to another, we found that we did not need to adjust thesettings when we moved between the two experimental systems used in our eval-uation. Although our systems had the same number of nodes and both used AMDCPUs, they differed in the number of cores per node, the cache-coherency protocol(broadcast vs. directory-based), and one had a higher interconnect throughput thanthe other. Therefore, it is possible that the algorithm parameters settings are ratherstable across all but drastically different systems.Page-local DecisionsCarrefour makes page-local decisions depending on the mechanisms enabled: e.g.,pages are only considered for replication if replication is enabled for that applica-tion. The following explanation assumes that all three mechanisms are enabled.To decide the fate of each page, we need at least two memory-access samplesfor that page. If the page was accessed from only a single node we migrate it tothat node. We do not migrate the thread, because we assume that threads are boundto cores. If the page is accessed from two or more nodes, it is a candidate for eitherinterleaving or replication. If the accesses are read-only, the page is replicated.Otherwise it is marked for interleaving. To decide where to place a page markedfor interleaving, we use two probabilities: Pmigrate and Pnode. Pmigrate determinesthe likelihood of migrating the page away from the current node. Pmigrate is theMAPTU of the current node as the fraction of MAPTU on all nodes, so the higherthe load on the current node relative to others, the higher the chance that we willmigrate a page. Pnode gives us the probability of migrating a page to a particularnode, and it is the complement of Pmigrate for that node, so Carrefour will migratethe page to the least loaded node.2.3.3 ImplementationWe implemented Carrefour in the Linux kernel 3.6.0. Carrefour measures the se-lected performance indicators, and with periodicity of one second, makes decisionsregarding page placement and resets statistic counters. To a large extent, Car-refour relies on well-understood mechanisms in the Linux kernel, such as physicalpage migration. The non-trivial aspects of the implementation were understanding18how to accomplish fast and accurate sampling of memory accesses and navigatingaround the overheads of replication. We describe how we overcame these chal-lenges in the two sections that follow.Fast and Accurate Memory Access SamplingA crucial goal of the algorithm is to quickly and accurately detect memory pagesthat cause the most DRAM accesses, and accurately estimate the read/write ratioof those pages. To that end, we used Instruction-Based Sampling (IBS): hardware-supported sampling of instructions available in AMD processors. Intel processorssupport similar functionality in the form of PEBS: Precise Event-Based Sampling.IBS can be configured to deliver instruction samples at a desired interval, e.g., af-ter expiration of a certain number of cycles or micro-ops. Each sample containsdetailed information about the sampled instruction, such as the address of the ac-cessed data (if the instruction is a load or a store), whether or not it missed inthe cache and how long it took to fetch the data. Unfortunately, every deliveredsample generates an interrupt, so processing samples at a high rate becomes verycostly. Other systems that relied on IBS performed off-line profiling [69, 86], sothey could tolerate much higher overhead than what would be acceptable in ouronline algorithm.After experimenting with IBS on our systems, we found that for most appli-cations the sampling interval of 130,000 cycles incurs a tolerable overhead of lessthan 5%. The desired sampling rate can be trivially derived for new systems: itamounts to experimenting with different sampling rates and settling for the onethat generates acceptable runtime overhead.Our initial decision was to filter out all the samples that did not generate aDRAM access. However, we found that such filtering produced too few samples.Even memory-intensive workloads access DRAM only a few times for every thou-sand instructions. That, combined with a low IBS sampling frequency, gave us asampling rate of less than 10−6%, and made it difficult to generate a sufficient num-ber of samples. Furthermore, filtering samples that did not access DRAM causedus to miss the accesses generated by the hardware prefetcher. These accesses arenot part of any instruction so they will not be tagged by IBS. For prefetch-intensive19applications, we obtain a very small number of samples and a distorted read-writeratio.To address this problem, we used two solutions. First, we adapted the samplingrate. When the program begins to run, we sample it at a relatively high rate of onceper 65,000 cycles. If after this measurement phase we take fewer than ten actionsin the algorithm (an action is any change in page placement), we switch to a muchlower rate of 1/260,000 cycles. Otherwise we continue sampling at the high rate.The second solution was, when filtering IBS samples, to retain not just the datasamples that accessed DRAM but those that hit in the first-level cache as well.First-level cache loads include accesses to prefetched data, so we avoid prefetcher-related inaccuracy. On the one hand, considering cache accesses can introduce“noise” in the data, because we could be sampling pages that never access DRAM.On the other hand, Carrefour is activated for memory-intensive applications only,and for them, there is a higher correlation between the accesses that hit in the cacheand those that access DRAM.With these two solutions combined, we were able to successfully identify thepages that are worth replicating, while this was nearly impossible prior to intro-ducing these solutions. For example, for Streamcluster we used to be able to detectonly a few percent of the pages that are worth replicating, but with these solutionsin place, we were able to identify 100% of them4.However, even though performance became much better (we were able to speedup Streamcluster by 26% relative to the default kernel), we were still far from the“ideal” manual replication, which sped it up by more than 2.5×. To approach idealperformance, we had to mitigate the overheads of replication, which we describenext.ReplicationReplication has overhead from the following three sources. First, there is the initialset-up cost and slightly more expensive page faults. Modern hardware walks pagetables automatically, so a separate copy of a page table must be created for eachnode. Page faults become slightly more costly, because a new page table entry4Streamcluster holds shared data in a single large array, so it is trivial to detect which data isworth replicating and implement a manual solution to use as the performance upper-bound.20must be installed on every node. To avoid these costs when we are not likely tobenefit from replication, we avoid replication unless the applications has at least afew hundred pages marked for replication5.The second source of overhead comes from additional hard page faults if weexceed the physical RAM capacity by replicating pages. As explained earlier, weavoided this overhead by conservatively setting the free memory threshold whenenabling replication.The final and most significant source of overhead stems from the need to syn-chronize the contents of replicated pages when they are written. This involves aphysical page copy and is costly. Before explaining how we avoid this overheadwe provide a brief overview of our implementation of replication.In Linux, a process address space is represented by a mm struct, which keepstrack of valid address space segments and holds a pointer to the page table, which isstored as a hierarchical array. Since modern hardware walks page tables automati-cally, we cannot modify the structure of the page table to point to several physicallocations (one for each node) for a given virtual page. Instead, we must maintain aseparate copy of the page table for each node and synchronize the page tables whenthey are modified, even for virtual pages that are not replicated. Linux dictates thatthe page table entry (PTE) be locked when it is being modified. We do not makeany changes to this locking protocol. The only difference is that we designate onecopy of the page table the master copy and lock the PTE only in the master copywhile installing the corresponding PTEs into all other replicas.When a page is replicated, we create a physical copy on every memory nodethat runs threads from the corresponding application. We install a different virtual-to-physical translation in each node’s page table. We write-protect the replicatedpage, so when any node writes that page we receive a page protection fault. Tohandle this fault, we read-protect the page on all nodes except the faulting one, andenable writing on the faulting node. If another node accesses that page, we mustcopy the new version of the page to that node, enable the page for reading andprotect it from writing.We refer to all the actions needed to keep the pages synchronized as page col-5We use the threshold of at least 500 pages. Performance is not highly sensitive to this parameter.21lapses. Collapses are extremely costly, and would occur if we replicate a page thatis write-shared. Even with very infrequent writes (e.g., one in 1000 accesses), thecollapse overhead could be prohibitively high. With limited capabilities of IBS, weare unable to detect read/write ratio on individual pages with sufficient accuracy.That is why we use the application-wide MRR and disable replication for the en-tire application if the MRR is low. Furthermore, we monitor collapse statistics ofindividual pages and disable replication for any page that generated more than fivecollapses.With these optimizations, as well as those described in Section 2.3.3, we wereable to avoid replication costs for the applications we tested and approached within10% the performance of manual replication for Streamcluster.2.4 EvaluationIn this section, we study the performance of Carrefour on a set of benchmarks.The main questions that we address are the following:1. How does Carrefour impact the performance of applications, includingthose that cannot benefit from its heuristics?2. How does Carrefour compare against existing heuristics for modern NUMAhardware?3. How well does Carrefour leverage the different memory placement mecha-nisms?4. What is the overhead of Carrefour?To assess the performance of Carrefour, we compare it to three other configu-rations:Linux: A standard Linux kernel with the default first-touch memory allocationpolicy.Manual interleaving: A standard Linux kernel with the interleaving policy man-ually enabled for the application.AutoNUMA: A Linux patchset [34] that aims to automatically place memorypages on nodes based on where they are accessed.22The rest of the section is organized as follows: we first describe our experi-mental testbed. Next, we study single-application scenarios, followed by work-loads with multiple co-scheduled applications. We then detail the overhead of Car-refour and conclude with a discussion on additional hardware support that wouldimprove Carrefour’s operation.2.4.1 TestbedWe used two different machines for the experiments:Machine A has four 2.3GHz AMD Opteron 8385 processors with 4 cores in each(16 cores in total) and 64GB of RAM. It features 4 nodes (i.e., 4 cores and 16GBof RAM per node) interconnected with HyperTransport 1.0 links.Machine B has four 2.6GHz AMD Opteron 8435 processors with 6 cores in each(24 cores in total) and 64GB of RAM. It features 4 nodes (i.e., 6 cores and 16GBof RAM per node) interconnected with HyperTransport 3.0 links.We used Linux kernel v3.6 for all experiments. For the AutoNUMA config-uration, we used AutoNUMA v27 and disabled PMD scan because it decreasesperformance on all applications we measured6.We used the following set of applications: the PARSEC benchmark suitev2.1 [25], the FaceRec facial recognition engine v5.0 [24], the Metis MapReducebenchmark suite [74] and the NAS parallel benchmark suite v3.3 [21]. PARSECapplications run with the native workload. From the available workloads in NASwe chose those with the running time of at least ten seconds. We excluded appli-cations whose CPU utilization was below 33%, because they were not affected bymemory management policies. We also excluded applications using shared mem-ory across processes (as opposed to threads) or memory-mapped files, becauseour replication mechanism does not yet support this behaviour. For FaceRec, weused two kinds of workloads: a short-running one and a long running one (namedFaceRecLong in the remainder of the paper) to show that Carrefour is able to suc-cessfully handle short workloads (less than 4s on machine B when running withCarrefour). Each experiment was run ten times. Overall, we observed a stan-6We also tested AutoNUMA v28. The performance results are very similar. However, we ob-served a significantly higher standard deviation (up to 18% on machine A and 23% on machine B)which makes profiling very difficult.23dard deviation between 1% and 2% for the Linux, Manual interleaving and Car-refour configurations. AutoNUMA has a more significant standard deviation, up to9% on machine A and 13% on machine B.2.4.2 Single-Application WorkloadsPerformance comparison. Figures 2.4 and 2.5 show the performance improve-ment relative to default Linux obtained under all configurations for machine A andmachine B. Performance improvement is computed as:De f aultLinuxtime−SystemtimeDe f aultLinuxtime∗100%,where System can be either Carrefour, Manual interleaving or AutoNUMA.We can make two main observations. First, Carrefour almost systematicallyoutperforms default Linux, AutoNUMA, and Manual interleaving, sometimesquite substantially. For instance, when running Streamcluster or FaceRecLong,we observe that Carrefour is up to 12% faster than Manual interleaving, up to 62%faster than AutoNUMA, and up to 73% faster than default Linux on machine B.Second, we observe that, unlike other techniques, Carrefour never performs sig-nificantly worse than default Linux: the maximum performance degradation overLinux is below 3%. In contrast, AutoNUMA and Manual interleaving cause per-formance degradations of up to 17% and 10% respectively. Besides, we notice thatManual Interleaving has a very irregular impact on performance. While it doesfairly well for PARSEC, NAS applications suffer significant performance degrada-tion (up to 38%) when run with manual interleaving.IS is an exception among the NAS benchmarks: Manual interleaving improvesits performance, while Carrefour does no better than default Linux. This is becauseIS suffers from very short imbalance bursts that we are not able to correct due tolimited sampling accuracy achievable with low overhead using existing hardwarecounters. Section 2.4.6 discusses hardware support that would help us address thisproblem.To understand the reasons for the performance impact of the different policies,we study in detail a set of applications whose performance is improved the most by24-20 0 20 40 60 80 100BodytrackFacesimFluidanimateStreamclusterSwaptionsx264FaceRecFaceRecLongKmeansMatrixmultiplyPCAWrmemPerformance improvement with respect to Linux (%)AutoNUMAManual interleavingCarrefourMachine A 0 30 60 90BodytrackFacesimFluidanimateStreamclusterSwaptionsx264FaceRecFaceRecLongKmeansMatrixmultiplyPCAWrmemPerformance improvement with respect to Linux (%)AutoNUMAManual interleavingCarrefourMachine BFigure 2.4: PARSEC/Metis: AutoNUMA, Manual interleaving and Car-refour vs.Default Linux.Carrefour. We present several metrics: the load imbalance on memory controllers(Figure 2.6(a)), the load imbalance on interconnect links (Figure 2.6(b)), the aver-age memory latency (Figure 2.7(a)) and the local access ratio (Figure 2.7(b)).We draw the following observations. First, Carrefour much better balances theload on both memory controllers and interconnect links than Linux and AutoN-UMA. Not surprisingly, Manual interleaving is also good at balancing the load.Nevertheless, we observe in Figure 2.7(a) that Carrefour induces lower averagememory latencies than Manual interleaving, which explains its better performance.To understand why Carrefour reduces memory latencies we refer to Figure 2.7(b),which shows that Carrefour not only balances the load on memory controllers andinterconnect links, but also often induces a much higher ratio of local memoryaccesses than other techniques. This is a consequence of Carrefour’s judiciously25-40-30-20-10 0 10BT CG DC EP FT IS LU MG SP UAPer for mance impr ovement wi th respect t o Li nux (%) AutoNUMAManual interleavingCarrefourMachine A-40-20 0 20 40 60BT CG DC EP FT IS LU MG SP UAPer for mance impr ovement wi th respect t o Li nux (%) AutoNUMAManual interleavingCarrefourMachine BFigure 2.5: NAS: AutoNUMA, Manual interleaving and Carrefour vs. De-fault Linux.applying the right techniques (interleaving, replication or co-location) in placeswhere they are beneficial. Interleaving mostly balances the load; replication andco-location, in addition to balancing the load, improve the local access ratio. Bet-ter locality improves latencies in two ways: it avoids remote wire delays and, moreimportantly, decreases congestion on the interconnect links.We also study MG, which is a multi-grid solver, as a representative exampleof the NAS applications, for which Carrefour does not bring significant improve-ments over default Linux (but still performs better than AutoNUMA and Manualinterleaving in most cases). MG has a low imbalance and a good local access ratioto begin with, so manual interleaving has a negative impact on such workloads, sig-nificantly decreasing the local access ratio and as a result stressing the interconnect.Looking inside Carrefour. To better understand the behavior of Carrefour, weshow in Table 2.3 the number of replicated pages, the number of interleaved pages,26 0 20 40 60 80 100 120 140 160Facesim Streamcluster FaceRec FaceRecLong PCA MG SPLoad imbal anceon memor y cont rol l er s (%)LinuxAutoNUMAManual interleavingCarrefour(a) Memory controllers 0 20 40 60 80Facesim Streamcluster FaceRec FaceRecLong PCA MG SPLoad imbal anceon int er connect l i nks (%)LinuxAutoNUMAManual interleavingCarrefour(b) Interconnect linksFigure 2.6: Load imbalance for selected single-application benchmarks (ma-chine A).and the number of co-located pages for the chosen benchmarks. These numbersprovide a better insight into how Carrefour manages the memory. We observe thatall three memory placement mechanisms are in use, and that most applicationsrely on two or three techniques. As discussed previously, MG does not suffer fromtraffic congestion, so Carrefour does not enable any technique for this application.A legitimate question we can ask is whether Carrefour always selects the besttechnique. In Table 2.4, we report the performance improvement over Linux ob-tained when running a full-fledged Carrefour and when running a reduced versionof Carrefour enabling only one technique at a time. We observe that Carrefour sys-tematically selects the best technique. It is also interesting to remark how differenttechniques work together and how the numbers provided here echo those in Ta-ble 2.3. We notice that, for all the studied applications except MG and SP, thecombination of several techniques employed by Carrefour outperforms any single27 0 200 400 600 800 1000 1200Facesim Streamcluster FaceRec FaceRecLong PCA MG SPAvg lat ency (nbCycl es/ req)LinuxAutoNUMAManual interleavingCarrefour(a) Average memory latency 0 20 40 60 80 100 120Facesim Streamcluster FaceRec FaceRecLong PCA MG SPRat io of l ocalmemor y accesses (%) LinuxAutoNUMAManual interleavingCarrefour(b) Local memory access ratioFigure 2.7: DRAM latency and locality for selected single-applicationbenchmarks (machine A).No. replicated pages No. interleaved pages No. migrated pagesFacesim 0 431 10.1kStreamcluster 25.4k 14.5k 858FaceRec 4k 3 1.3kFaceRecLong 4.1k 5 1.4kPCA 31k 33 41.3kMG 0 0 1SP 0 305 1.7kTable 2.3: Number of memory pages that are replicated, interleaved and co-located on single-application workloads (machine A).28Carrefour Replication Interleaving Co-locationFacesim 42% -3% 0% 37%Streamcluster 66% 60% 33% 20%FaceRec 40% 38% 26% 1%FaceRecLong 45% 42% 20% 1%PCA 33% 32% 17% 15%MG -2% -2% -2% -2%SP 8% -1% -7% 8%Table 2.4: Performance improvement over Linux when running Car-refour and the three different techniques individually on single-application workloads (machine A).technique, even when a given technique has a dominant impact (e.g., for Stream-cluster). The slight performance degradation for MG corresponds to the monitoringoverhead of Carrefour.2.4.3 Multi-Application WorkloadsWe next study how Carrefour behaves in the context of workloads with multiple ap-plications that are co-scheduled on the same machine. The goal is to assess whetherCarrefour is able to work on complex access patterns and to make the distinctionbetween the diverse requirements of different applications. We consider severalscenarios based on some of the applications previously studied in Section 2.4.2: (i)MG + Streamcluster, (ii) PCA + Streamcluster, and (iii) FaceRecLong + Stream-cluster. We chose these scenarios because they exhibit a nuance of patterns, whichrequire combining several memory placement techniques to achieve good perfor-mance.Each application is run with half as many threads as the number of cores (i.e.,8 threads on machine A, 12 on machine B). With two applications, the workloadoccupies all the available cores. The threads of each application are clustered onthe same node, so each application uses all the cores on two of the four nodes on amachine.Figure 2.8 shows, for each workload, the performance improvement with re-spect to Linux for AutoNUMA, Manual interleaving and Carrefour on machine A29and machine B. We observe that Carrefour always outperforms AutoNUMA andManual interleaving, by up to 62% and 36% respectively. Besides, Carrefour alsooutperforms default Linux, while Manual interleaving hurts MG with a 25% slow-down. Overall, the results obtained with Manual interleaving are closer to the onesof Carrefour compared to the other setups.The reason why Manual interleaving performs relatively well in these scenar-ios is because, with each application using two domains, there is a lot less cross-domain traffic than in the single-application case. Hence there are fewer problemsthat need to be fixed and there is a smaller discrepancy between the performanceof different memory management algorithms.To explain the results, we show the same detailed metrics as in Section 2.4.2:load imbalance on memory controllers, load imbalance on interconnect links, av-erage DRAM latency and ratio of local DRAM accesses in Figures 2.9, 2.10, 2.11and 2.12 respectively. The metrics are aggregated for the two applications of eachworkload. As was the case with the single-application workloads, we see thatCarrefour systematically improves the latency of the studied workloads for tworeasons: a more balanced load on memory controllers and interconnect links aswell as an improved locality for DRAM accesses. Note that, for each workload,the depicted latencies are averaged over the two applications. This explains thesmall latency variations between Linux, AutoNUMA and Manual Interleaving inthe case of MG + Streamcluster.Finally, we show in Table 2.5 the performance improvement for each applica-tion when the effects of only one of the techniques are enabled. We observe thatthe multi-applications workloads perform as well or better with an arsenal of tech-niques used in Carrefour rather than with any single technique alone (especiallyfor PCA + Streamcluster).2.4.4 OverheadCarrefour incurs CPU and memory overhead. The first source of CPU overhead isthe periodic IBS profiling. To measure CPU overhead, we compared performanceof Carrefour with Linux on those applications where Carrefour does not yield anyperformance benefits. We observed the overhead between 0.2% and 3.2%. The30-20 0 20 40 60 80MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterPer for mance impr ovement wi th respect t o Li nux (%)MG Streamcluster PCA Streamcluster FaceRecLong StreamclusterAutoNUMAManual interleavingCarrefourMachine A-20 0 20 40 60 80MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterPer for mance impr ovement wi th respect t o Li nux (%)MG Streamcluster PCA Streamcluster FaceRecLong StreamclusterAutoNUMAManual interleavingCarrefourMachine BFigure 2.8: Multi-application workloads: AutoNUMA, Manual interleavingand Carrefour vs. Default Linux.31 0 20 40 60 80 100MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterLoad imbal anceon memor y cont rol l er s (%)LinuxAutoNUMAManual interleavingCarrefourFigure 2.9: Multi-application workloads: load imbalance on memory con-trollers (machine A). 0 20 40 60 80MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterLoad imbal anceon int er connect l i nks (%)LinuxAutoNUMAManual interleavingCarrefourFigure 2.10: Multi-application workloads: load imbalance on interconnectlinks (machine A). 0 100 200 300 400 500 600 700MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterAvg lat ency (nbCycl es/ req)LinuxAutoNUMAManual interleavingCarrefourFigure 2.11: Multi-application workloads: average memory latency (ma-chine A).32 0 20 40 60 80 100MG+ StreamclusterPCA+ StreamclusterFaceRecLong+ StreamclusterRat io of l ocalmemor y accesses (%)LinuxAutoNUMAManual interleavingCarrefourFigure 2.12: Multi-application workloads: local memory access ratio (ma-chine A).Carrefour Replication Interleaving Co-locationMG2% / 71% 2% / 73% -5% / 17% -1% / 6%+ StreamclusterPCA24% / 57% 18% / 57% 8% / 5% 14% / 1%+ StreamclusterFaceRecLong53% / 71% 53% / 71% 12% / 9% 12% / 4%+ StreamclusterTable 2.5: Multi-application workloads: performance improvement overLinux when running Carrefour and the three different techniques indi-vidually (machine A).adaptive sampling rate in Carrefour is crucial to keeping this overhead low. Asecond and potentially significant source of CPU overhead is replication, if weperform a lot of collapses. A single collapse costs a few hundred microsecondswhen it occurs in isolation. Parallel collapses can take a few milliseconds, becauseof lock contention. That is why it is crucial to avoid collapses and other synchro-nization events by disabling replication for write-intensive workloads, as is done inCarrefour.The first source of memory overhead is the allocation of data structures to keeptrack of profiling data. This overhead is negligible: e.g., 5MB on Machine A with64GB of RAM. Carrefour’s data structures are pre-allocated on startup to avoidmemory allocation during the runtime. We limit the number of profiled pages to30,000 to avoid the cost of managing dynamically sized structures. The second33source of memory overhead is memory replication. When enabled, replication in-troduces a memory footprint overhead of 400MB (353%), 60MB (210%), 60MB(126%) and 614MB (5%) for Streamcluster, FaceRec, FacerecLong and PCA re-spectively.2.4.5 Impact on Energy ConsumptionIt was observed that remote memory accesses require significantly more energythan local ones [35]. Since Carrefour can either decrease or increase the num-ber of remote memory accesses, we were interested in evaluating its impact onenergy consumption7. We show the results for selected applications from single-application workloads on Machine A. We report the increase in energy consump-tion as well as the increase in completion time of all configurations compared todefault Linux in Figure 2.13. Completion time increase is computed here as:Systemtime−De f aultLinuxtimeDe f aultLinuxtime∗100%.We observe that there is a strong relationship between the completion time andthe energy consumption: if the completion time is decreased, the energy consump-tion is also decreased proportionally. As a result, Carrefour saves up to 58% ofenergy. When no traffic management is needed, Carrefour on its own has a lowimpact on energy consumption (e.g., 2% on MG).More generally, we found that the increase of remote memory accesses haslittle or no impact on the global energy consumption of the machine. For example,Manual interleaving drops the local access ratio of MG from 97% to 25% andthus proportionally increases the number of remote accesses. However, the energyconsumption increase is slightly lower that the completion time increase, whichindicates that the extra energy overhead of remote memory accesses have no strongimpact on overall energy consumption.7We used IPMI which gives access to the current global power consumption on our servers.34-70-60-50-40-30-20-10 0 10 20 30 40 50 60 70AutonumaManual Int.CarrefourAutonumaManual Int.CarrefourAutonumaManual Int.CarrefourAutonumaManual Int.CarrefourAutonumaManual Int.CarrefourAutonumaManual Int.CarrefourAutonumaManual Int.CarrefourI ncr ease in wi th respect t o Li nux (%)Energy consumptionCompletion timeSPMGPCAFaceRecLongFaceRecStreamclusterFacesimFigure 2.13: Increase in completion time and energy consumption for se-lected single-application benchmarks with respect to Linux (machineA). Lower is better.2.4.6 Discussion: Hardware SupportWe have shown that a traffic management system such as Carrefour can bring sig-nificant performance benefits. However, the challenge in building Carrefour wasthe need to navigate around the limitations of the performance monitoring units ofour hardware as well as the costs of replicating pages. We draw some insights onthe features that could be integrated into future machines to further mitigate theoverhead and improve accuracy, efficiency and performance of traffic managementalgorithms.First, Carrefour would benefit from hardware profiling mechanisms that sam-ple memory accesses with high precision and low overhead. For instance, it wouldbe useful to have a profiling mechanism that accumulates and aggregates page ac-cess statistics in an internal buffer before triggering an interrupt. In this regard, theAMD Lightweight Profiling [42] facility seems a promising evolution of profilinghardware8, but we believe the hardware should go even further, and not only accu-mulate the samples but be configured to aggregate them according to user needs,to reduce the number of interrupts even further.Second, Carrefour would benefit from dedicated hardware support for memoryreplication. We believe that there should be interfaces allowing the operating sys-8Unfortunately, Lightweight Profiling is only available on some AMD processors and we werenot able to evaluate it in this work.35tem to indicate to the processor which pages to replicate. The processor would thenbe in charge of replicating the pages on the nodes accessing it and maintaining con-sistency between the various replicas (in the same way as it maintains consistencyfor cache lines). Given that maintaining consistency between frequently writtenpages is costly, we believe that such processors should also be able to trigger aninterrupt when a page is written too frequently. The OS would then decide to keepthe page replicated or to revert the replication decision.This hardware support can be made a lot more scalable than cache coherencyprotocols, because it is controlled by the OS. Since the OS is armed with betterhardware profiling, it will invoke it only for pages that perform little write sharing.So the actual synchronization protocol would be triggered infrequently.2.5 Related WorkWe explain how Carrefour relates to different works on multicore systems. First,we review systems that maximize data locality. Second, we contrast Carrefour withprevious contention-aware systems. Third, we consider application-level tech-niques to mitigate contention on data-sharing applications. Finally, we discusstraffic characterization observations for modern NUMA systems.NUMA-aware thread and memory management policies were proposed for ear-lier research systems [28, 31, 70, 99] as well as in commercial OS. Their maindifference from our work is that their goal was to optimize locality. However,on modern systems, the main performance problems are due to traffic congestion.Our algorithm is the first one that meets the goal of mitigating traffic congestion.Among the above-mentioned works, the one most related to Carrefour is the systemby Verghese et al. [99] for early cache-coherent NUMA machines, which leveragespage replication and migration mechanisms. Their system relies on assumptionsabout hardware support that do not hold on currently available machines (e.g., pre-cise per-page access statistics). Thus, Carrefour’s logic is more involved, as it ismore difficult to amortize the costs of the monitoring and memory page placementmechanisms. The authors noticed that locality-driven optimizations could, as a sideeffect, reduce the overall contention in the system. However, their system does notsystematically address contention issues. For instance, shared written pages are not36taken into account, whereas Carrefour uses memory interleaving techniques whenthere is contention on such pages. Moreover, the load on memory controllers isignored when making page replication/migration decisions.Similarly to earlier NUMA-aware policies, Solaris and Linux focus primarilyon co-location of threads and data, but to the disadvantage of data-sharing work-loads, replication is not supported [19]. Linux provides the option to interleaveparts of the address space across all memory nodes, but the decision when to invokethe interleaving is left to the programmer or the administrator. Solaris supports thenotion of a home load group, such that the thread’s memory is always allocated inits home group and the thread is preferentially scheduled in its home group. This,again, favours locality, but does not necessarily address traffic congestion.The AutoNUMA patches for Linux [91] also implement locality-driven opti-mizations, using two main heuristics. First, threads migrate toward nodes holdingthe majority of the pages accessed by these threads. Memory residence is deter-mined by page fault statistics. Second, pages are periodically unmapped from aprocess address space and, upon the next page fault, migrated to the requestingnode. As shown in the evaluation section, this approach yields irregular results.We attribute this limitation to the following sources of overhead: local thread/pagemigration decisions that do not take data sharing patterns nor access frequenciesinto account (thus leading to page bouncing or useless migrations) nor memorycontroller or interconnect load (thus leading to memory load imbalance/conges-tion), continuous overhead due to the scanning/unmapping of page-table entriesand the corresponding soft page faults. In contrast, Carrefour makes global dataplacement decisions based on precise traffic patterns and adjusts the monitoringoverhead based on the observed contention level.Locality-driven optimizations for data-sharing applications were addressed ina study that dynamically identified data-sharing thread groups and co-located themon the same node [97]. However, that solution was for a system with a centralized(UMA) memory architecture. Thus, it only studied the benefits of thread placementfor improved cache locality and did not address the placement of memory pageson multiple memory nodes.Zhou and Demsky [106] investigated how to distribute memory pages to cacheson a many-core Tilera processor, to implement an efficient garbage collector. The37authors tried various policies but found that maximizing locality was the best ap-proach for their system. This contrasts with the systems Carrefour is targeting,where reducing congestion is more important than improving locality. Also, theauthors used a different method for monitoring page access patterns that relies onsoftware-serviced TLB misses, which is not possible on x86.Several studies addressed contention issues in the memory hierarchy. Someof these works were designed for UMA systems [66, 75, 107] and are inefficienton NUMA systems, because they fail to address or even accentuate issues suchas remote access latencies and contention on memory controllers and on the inter-connect links [26]. Other works have been specifically designed for NUMA sys-tems but only partially address contention issues. The N-Mass thread placementalgorithm [72] attempts to achieve good DRAM locality while avoiding cache con-tention. However, it does not address contention issues at the level of memory con-trollers and interconnect links. Two studies [20, 73] have shown the importance oftaking memory controller congestion into account for data placement decisions, butthey did not provide a complete solution to address multi-level resource contention.The most comprehensive work to date on NUMA-aware contention management isthe DINO scheduler [26], which spreads memory intensive threads across memorydomains and accordingly migrates the corresponding memory pages. However,DINO does not address workloads with data sharing between threads or processes,which require identifying per-page memory access patterns and making the appro-priate data placement decisions.When introducing a new resource-management policy in the OS, it is worthasking whether a similar or better effect could be achieved by restructuring theapplication. In our context, it is important to consider the so-called no-sharingprinciple of application design. The key idea behind no-sharing is that the datamust be partitioned or replicated between memory nodes, and a thread needing toaccess data in a different domain than its own either migrates to the target domainor asks the thread running in that domain to perform the work on its behalf, insteadof fetching the data over the memory channels [23, 30, 48, 77, 85, 90, 94]. Whilethe no-sharing architecture was primarily motivated by the need to avoid locking, itcould similarly help reduce the amount of traffic sent across the interconnect, andthus alleviate the traffic congestion problem.38Unfortunately, no-sharing architectures are not a universal remedy. First of all,they trade-off data accesses for messages or thread migration; the trade-off is onlyworth making if the size of the data used in a single operation is much larger thanthe size of the message or the state of the migrated thread [30]. Second, adoptinga no-sharing architecture often requires significant changes to the application (andto the OS, if the application is OS-intensive). A good illustration of the potentialchallenges can be gleaned from two studies that converted a database system tothe no-sharing design. The first study took the path of least resistance and sim-ply replicated and/or partitioned the database among domains, adding a message-routing layer on top [90]. While this worked well for small read-mostly workloads,for large workloads replication had very significant memory overhead (unaccept-able because of increased paging), and partitioning required a priori knowledgeof query-to-data mapping, which is not a reasonable assumption in a general case.A solution that overcame these limitations, DORA [85], required a very signifi-cant restructuring of the database system, which could easily amount to millions ofdollars in development costs for a commercial database.Our goal was to address scenarios where adopting a no-sharing architecture isnot feasible either for technical reasons or for practical considerations. Providingan OS-level, rather than an application-level, solution allows us to address manyapplications at once. Understanding the limitations of the OS-level solution anddetermining what optimizations can be done only at the level of an application isan open research question.A study characterized the performance of emerging “scale-out” workloads onmodern hardware [46]. The authors observed that there is little inter-thread datasharing, and the utilization of the off-chip memory bandwidth is low. As a result,they argue that memory bandwidth on existing processors is unnecessarily high.Our findings do not agree with this observation. Although it is also true that theworkloads we consider perform little fine-grained data sharing, they still stress thecross-chip interconnect, because they access a large working set, which is spreadacross the entire NUMA memory space. The authors of the scale-out study re-ported a very low bandwidth utilization (<10%), even for database workloads.In contrast, our measurements show a utilization as low as 45% in some cases.These differences could be due to the authors of this study used a system with only39two chips and 12 cores in their experiments. On larger systems, more threads aremaking requests to remote memories, so there is greater pressure on bandwidth.Furthermore, we found that low bandwidth utilization is not necessarily a healthysymptom. In our experiments, performance dropped steeply even as bandwidthutilization went from 10% to 30%. The interconnect became the bottleneck even ata fraction of the total available bandwidth, memory requests are not spread evenlyin time; they are bursty. Burstiness causes requests to clash on the link even if theoverall bandwidth is not exceeded. In summary, we conclude that contrary to thesuggestion made in [46], it is too early to reduce the bandwidth of cross-chip inter-connects on large multicore systems, especially if they are used for running largedata-centric workloads.2.6 SummaryIn this chapter we discussed the design and implementation of Carrefour, whichis a memory management algorithm for NUMA systems that manages traffic onmemory controllers and interconnects. Earlier NUMA-aware memory manage-ment policies aimed to mitigate the cost of remote wire delays, which is no longerthe main bottleneck on modern systems. Carrefour’s design was motivated bythe evolution of modern NUMA hardware, where traffic congestion plays a muchlarger role in performance than wire delays.System design principles embodied in Carrefour are important not only fortoday’s systems, but also for future hardware. The amount of memory bandwidthper core is projected to decrease in the future [30], so managing traffic congestionwill be as crucial as ever.40Chapter 3Analyzing Memory ManagementMethods on IntegratedCPU-GPU Systems1Heterogeneous systems that integrate a multicore CPU and a GPU on the samedie are becoming more common. In the previous chapter, we examined memorymanagement on systems that consist of multiple multi-core processors connectedtogether in a non-uniform memory arrangement. In this chapter, we focus on sys-tems that integrate two types of processors (multi-core CPUs and a GPU) on thesame chip. On these systems, both the CPU and GPU share the same physicalmemory as opposed to using separate memory dies. Although integration elimi-nates the need to copy data between the CPU and the GPU, arranging transparentmemory sharing between the two devices can carry large overheads. Memory onCPU/GPU systems is typically managed by a software framework such as OpenCLor CUDA, which includes a runtime library, and communicates with a GPU driver.These frameworks offer a range of memory management methods that vary in easeof use, consistency guarantees and performance. In this study, we analyze someof the common memory management methods of the most widely used softwareframeworks for heterogeneous systems: CUDA, OpenCL 1.2, OpenCL 2.0, and1This chapter is a modified version of a previously published work [36].41HSA, on NVIDIA and AMD hardware. We focus on performance/functionalitytrade-offs, with the goal of exposing their performance impact and simplifying thechoice of memory management methods for programmers.3.1 IntroductionCopying data over the Peripheral Component Interconnect (PCI) bus often limitsperformance on discrete GPU systems. Integrated CPU-GPU systems address thisproblem by placing the CPU and the GPU on the same die (Fig. 1.2), with the addedadvantage of reduced area and power footprint. In integrated CPU-GPU systems,both the CPU and the GPU share the physical memory of the system. Therefore,it might appear that this tight integration allows for seamless data sharing. As weexplain in this chapter, this is not the case. There is still a lot of software andhardware restrictions on CPU-GPU integrated systems that forbid complete andseamless unification of memory.Software frameworks such as CUDA and OpenCL provide a number of mem-ory management methods to simplify programmability. For example, one methodmight allocate a memory area that can be seamlessly accessed from either the CPUor the GPU without the need to manually translate the pointers or copy data; an-other might provide a memory area that can be used for atomic accesses. Thesememory management methods are hidden in the runtime library and the GPUdriver. Their implementation details are poorly understood, and the effects on per-formance remain somewhat of a mystery. For example, to enable seamless sharingof memory between the CPU and the GPU, a memory buffer could be configureduncacheable or access-protected by the driver, which means that an interrupt willbe handled every time the memory is accessed. Although many of these designchoices are justifiable from the standpoint of programmability, they can introducelarge performance overheads.The variety of memory management methods and poor understanding of theirinternals introduce the uncertainty as to which method a programmer should use,and whether reduced performance is worth the improved programmability. In thiswork we investigate how these different memory management methods behave interms of performance and functionality. Our main contribution is revealing the per-42formance trade-offs that these memory management strategies make and explain-ing the technical reasons for these performance effects, which required thoroughanalysis and some reverse-engineering effort due to lack of documentation. Weuse an NVIDIA system with the CUDA programming framework and an AMDAPU with the OpenCL 1.2, OpenCL 2.0 and HSA programming frameworks. Webelieve that our choice of platforms and frameworks covers a wide spectrum ofcurrent and future integrated CPU-GPU systems.3.2 BackgroundGPUs are massively parallel processors containing thousands of cores that hidelatency by having thousands of instructions in-flight to ensure that there will alwaysbe something executing on the available cores. If, for example, a thread runningon one GPU core issues a memory instruction and the data is not immediatelyavailable, the thread will be de-scheduled and a new thread will start executingimmediately. On the contrary, a typical CPU consists of tens of processing coresthat are more complex because they implement architectural techniques, such asdeep out-of-order pipelines, branch predictors and multi-level cache hierarchies,aiming to reduce the latency of a sequential instruction stream.Applications to be accelerated on a GPU typically have sequential and parallelregions. Sequential regions do not benefit from running on the GPU, since itsspecialty is massively parallel SIMT (Single Instruction Multiple Threads) code.Parallel regions can either be offloaded to the GPU, or continue running on theCPU. The portion of code that runs on the GPU is called a kernel. In CUDAterminology, each kernel is composed of a number of blocks. The blocks are furtherbroken down into warps, which consist of 32 threads each. In other words, a GPUkernel consists of blocks of warps of threads. In the OpenCL and HSA models theblocks are called workgroups that are subdivided into wavefronts, each consistingof 64 threads or work items. Determining how to best distribute the parallel regionsbetween the CPU and the GPU can be complex as it depends on multiple factorssuch as the offloading overhead, the cost of data transfer, the acceleration potentialon the GPU, and memory requirements.On current systems, interacting with the GPU occurs through a programming43framework, which typically consists of a programming language (or an extensionto a language), a runtime, and a driver. The two most widely used program-ming frameworks are CUDA and OpenCL. Furthermore, AMD introduced HSA,a lower-level framework that utilizes hardware features such as shared pageablememory and user-space work queues [47]. All frameworks have APIs that allowspecifying how to allocate memory and how to launch kernels on the GPU.At the hardware level, GPUs schedule warps (or wavefronts) rather than in-dividual threads for execution. GPUs follow a single-instruction multiple-thread(SIMT) paradigm. Hence, a group of 32 (or 64) threads execute instructions inlockstep. Divergence between the individual threads is allowed and tracked bythe GPU so that possible re-convergence can be resumed at a later point in execu-tion. This type of divergence is referred to as thread divergence. There is someoverhead with divergence, so it is best for all threads to follow the same executionpath. Another type of divergence is referred to as memory divergence. It is bestfor the threads within a warp to access consecutive words in memory. This way ofaccessing memory allows the GPU to coalesce these memory accesses into a sin-gle access which loads or stores all of the consecutively accessed data. If, however,each thread tries to access a random memory location, the memory unit of the GPUwill generate multiple memory requests and performance can degrade.Although this study targets platforms that integrate CPUs and GPUs, it is ap-plicable to any system in which two types of heterogeneous units need to com-municate with each other. Instead of relying on simulators, we chose to presentour analysis on real systems employing different software models and implement-ing different CPU-GPU coherency schemes. Future systems ranging from mobiledevices to exascale computing are expected to house such integrated chips to ac-commodate a variety of applications [101].3.3 The NVIDIA/CUDA SystemThe NVIDIA system used in our study includes a Kepler GPU with 192 coresand an ARM quad-core CPU; further details are provided in Section 3.4. Thereis no hardware coherency between the CPU and the GPU caches; software musttake care to keep the data consistent. The programming framework used on this44platform is CUDA. We discover that memory management methods offered byCUDA sometimes apply restrictions on how the data is cached and whether it canbe accessed concurrently by the CPU and the GPU. We now explain how thesemethods work.cudaMalloc: The first and most commonly used method of managing mem-ory is to allocate it on the host (CPU) using malloc() and on the device (GPU)using cudaMalloc(). This method does not utilize shared virtual memory, sothe CPU uses a different virtual address space than the GPU. Therefore, two dif-ferent buffers need to be allocated with different pointers pointing to each buffer.Data allocated on the host cannot be accessed from the device and vice versa, butit can be copied back and forth between the two buffers using cudaMemcpy();the DMA copy engine will perform the copying from one location on the uni-fied physical memory to another. It is up to the programmer to explicitly invokecudaMemcpy() when the data is needed. As a result, the resulting code mayend up cluttered with allocations and copy commands and can be difficult to debugand maintain. We will refer to this allocation scheme throughout the chapter as thedefault scheme.cudaHostAlloc: On integrated systems, data allocated withcudaHostAlloc() is placed in an area that is accessible from both theCPU and the GPU without the need to perform explicit memory copies. Fromour experiments, we found that the data is allocated on a device-mapped region/dev/nvmap; nvmap is the GPU memory management driver on integratedNVIDIA systems. This region physically resides in the system’s main memory,shared by the CPU and the GPU, but has different access properties than other vir-tual regions. Our investigations revealed that this region is configured uncacheableon the CPU and the GPU. As a result, any read or write from or to that memoryregion goes directly to main memory, bypassing caches. We will refer to thisallocation scheme as hostalloc.cudaMallocManaged: Managed memory was introduced in CUDA 6.0 toeliminate the need for explicit data copying between the CPU and the GPU byproviding a unified shared address space across the CPU and GPU. This unifiedaddress space allows seamless pointer sharing across both processing units. Amanaged memory buffer is allocated via cudaMallocManaged(), and simi-45larly to hostalloc, there is no need to explicitly copy data between the host anddevice. Furthermore, just like hostalloc, the data is mapped to a /dev/nvmapregion. In contrast to hostalloc, data is cached in CPU caches, and in the GPUL2 caches. The price of unrestricted caching is that the memory cannot be accessedconcurrently by the CPU and the GPU. Once the host launches a GPU kernel, it isnot allowed to access data in the same managed region as the one accessed by theGPU until the kernel completes; an attempt to do so will cause a SEGFAULT. Ifthe host wishes to access a managed area after the GPU kernel completes, it mustissue an explicit CUDA API synchronization call, which flushes caches to ensure aconsistent view of the data. This restriction is not an issue for programs where theCPU and the GPU execute sequentially, but can be a limitation for programs wherethe CPU and the GPU concurrently access data in the same allocated region (evenif the actual data is disjoint). We refer to this allocation scheme as managed.These memory management schemes represent different points on the spec-trum between the functionality and ease of use and restrictions on memory accesses(which leads to performance degradation, as we will show later). While the defaultscheme places no restrictions on memory accesses, it is the most difficult to use forthe programmer, because a pointer allocated on the CPU cannot be accessed on theGPU (and vice versa), so data must be manually copied if sharing between the de-vices is desired. The managed scheme obviates the need for copying, but preventsconcurrent access to data by the CPU and the GPU, so that caches can be flushedon kernel invocation boundary. This is a coarse consistency management method,which ensures that the programmer sees consistent copies of shared data at theexpense of restrictions when the memory can be accessed. Finally, hostallocmakes another choice with respect to consistency management: it does not restrictconcurrent access, but “shields” the programmer from accidentally reading incon-sistent data by disabling caching.We observe that there is another point in this spectrum that is not covered bythe available memory management methods: allowing the programmer seamlessmemory sharing without concurrency or caching restrictions, but requiring themto manually flush caches when consistency management is required. This is sim-ilar to what is in use on a modern multi-socket systems where barriers and cacheflushes are inserted to ensure consistency. To explore this point, we introduce a46new configuration scheme that we call sharedalloc.sharedalloc: Like hostalloc, it allows concurrent accesses to datafrom the CPU and the GPU, but at the same time allows caching, like the managedmethod. To implement sharedalloc we modified the Tegra Linux kernel codeto introduce a new system call, gpu hook(), that hooks into the GPU driver andallows us to control caching and explicitly flush caches to maintain data consis-tency.We control caching by setting two flags that are passed as arguments toioctl(). These flags make data cacheable on the CPU and the GPU. We dis-covered how to use these flags by tracing calls to ioctl() and examining the(limited) publicly available source code from NVIDIA. Once these flags are set,we can simply use cudaHostAlloc() to allocate memory without caching re-strictions.sharedalloc does not modify the existing memory consistency model usedon Nvidia GPUs [18, 71]. We place the burden on the programmer to decide whenit is best to flush the GPU caches to maintain consistency across the CPU and theGPU. The CUDA runtime already uses this cache flushing when using some of thebuilt-in methods. For example, when calling cudaDeviceSynchronize, theGPU caches are flushed. Our implementation simply allows the programmer toflush the caches whenever they deem appropriate for their application. Of course,there is a risk of having inconsistent views from the CPU and GPU if the cachesare not appropriately flushed. Therefore, the programmer must ensure that the GPUcache is flushed when they want the data to be correctly read from the CPU. Weshowcase practical cases for how to use sharedalloc later in this chapter.In our implementation of sharedalloc, the programmer flushes the cachesby invoking gpu hook() manually. However, in programs adhering to a data-race-free synchronization model, the invocation of this function could be embeddedinto existing synchronization functions in the framework’s runtime. Therefore ourimplementation is a proof-of-concept and does not strictly rely on manual flushing.Figure 3.1 shows an example using sharedalloc. The difference between usingsharedalloc and hostalloc is the additional gpu hook() system calls.A summary of caching and concurrency trade-offs for the different memoryallocation schemes is shown in Table 3.1.47Allocation scheme CPU caching GPU caching CPU-GPU concurrencydefault cached cached difficult + copy overheadhostalloc not cached not cached easymanaged cached cached not possiblesharedalloc cached cached easyTable 3.1: Summary of allocation schemes1 // Initialize context 2 cudaFree(0); 3 // Inform Linux that next allocation will be sharedalloc 4 gpu_hook(INIT); 5 // Allocate data using the same syntax as cudaHostAlloc 6 cudaHostAlloc(data, size, flags); 7 // Work on data using CPU 8 cpu_work(data); 9 // Flush CPU caches to push data to main memory 10 gpu_hook(FLUSH_CPU); 11 // Launch kernel to manipulate data on the GPU 12 kernel<<<blocks, threads>>>(data, ...); 13 // Wait for kernel to finish 14 cudaDeviceSynchronize(); 15 // Flush GPU caches to push data to main memory 16 gpu_hook(FLUSH_GPU); Figure 3.1: sharedalloc usage example3.4 Evaluation of the NVIDIA/CUDA System3.4.1 Hardware/SoftwareWe performed our experiments on the NVIDIA Jetson TK1 board, a K1 SoC inte-grating a CPU and a GPU. The CPU is a ”4-Plus-1” 2.32GHz ARM quad-coreCortex-A15, and the GPU is a Kepler ”GK20a” GPU with 192 SM3.2 CUDAcores. There are 2GB of DDR3L DRAM running at 933MHz with 64-bit datawidth. All 192 CUDA cores share 32K registers, 64KB L1 cache/shared memory,and a 128KB L2 cache. Global memory is normally cached only in L2, unless theapplication is compiled with a flag enabling caching in L1. The cache line size is32 bytes. The system software is Linux For Tegra R21.4 with CUDA 6.5.In all of our experiments, we set the board in the ”performance” mode and48we maximize the CPU and GPU clock rates. We also maximize the system andmemory bus clocks.3.4.2 Benchmarks and ExperimentsWe run four types of tests to analyze the memory management behaviour on thisplatform:Bandwidth Experiments: We run microbenchmarks measuring read andwrite bandwidth from the CPU and the GPU under various memory managementschemes and highlight the performance differences.Rodinia Benchmarks: We port a number of applications from the Rodiniabenchmark suite [32] to use the memory management methods described in theprevious section. The original implementation of Rodinia uses the default method.GPU Caching Behaviour: We explore the behaviour of caching on the GPUusing microbenchmarks.Concurrent Workload Experiments: We analyze the performance of applica-tions where the CPU and GPU concurrently access portions of data from the sameallocated region. We wrote a microbenchmark and modified four applications fromRodinia to perform concurrent work by the CPU and the GPU.Bandwidth ExperimentsTo demonstrate the impact of caching restrictions on performance, we measured theread and write memory bandwidth achieved by either the CPU or the GPU with thedifferent memory management methods. For the CPU tests, a single CPU threadsequentially reads or writes a data array, which we vary in size. For the GPU tests,we launch many threads to read or write the array in parallel, each thread readinga single array entry. We vary the size of the array, and the number of threads varieswith it. For example, for an 8MB array, we have 220 8-byte entries, so we launch220 GPU threads to access these entries. We perform 10 million iterations over thearray in total, so caching will be beneficial where it is used and when the array fits.The reported bandwidth is the total number of accessed bytes divided by the runtime.Figure 3.2 shows the write bandwidth. We omit the read bandwidth numbers,490 2 4 6 8 10 12 14 16 18 20 8B 256B 8KB 16KB 32KB 64KB 128KB 256KB 512KB 1MB 2MB 4MB 8MB Write Bandwidth (GB/s) Array size GPU-default/managed/sharedalloc GPU-hostalloc CPU-default CPU-hostalloc Figure 3.2: Write bandwidth under different memory management methods.because they are almost identical. The legend reports the device performing theaccess and the memory management method. For example, the line labeled GPU-hostalloc refers to the GPU accessing a hostalloc’ed region; the line labelledCPU-hostalloc refers to the same region accessed by the CPU. (The CPU and GPUtests are run, separately, not concurrently.) We show a single line labeled GPU-default/managed/sharedalloc, because the performance of these memory manage-ment methods on the GPU is the same in this simple experiment, as expected;neither of them disables caching.We observe that caching restrictions imposed by the hostalloc methodhave a significant performance impact on both the CPU and the GPU. The CPUloses more than half the bandwidth when caching is disabled. For the GPU, themanaged, default and sharedalloc scheme are able to exceed the avail-able bandwidth of roughly 12GB/s when the data fits into the cache (the array sizeis smaller than 128KB), while hostalloc runs much slower due to the lack ofcaching. When the data no longer fits into the cache, all systems become limited bythe hardware throughput and deliver the same performance (notice the sharp dropat 128KB).The managed, default, and sharedalloc deliver the highest read/writebandwidth. The restriction on the CPU/GPU concurrency imposed by managedis not evident from this very simple test. However, many previous studies demon-50strated benefits of truly concurrent execution by the CPU and the GPU [22, 33, 61,76, 96]; for these applications, the restriction on CPU/GPU concurrency may be alimitation. We demonstrate a potential effect of these limitations with Concurrentexperiments later in this section.Rodinia Benchmarks0 1 2 3 4 5 Normalized overall slowdown Default hostalloc managed hostalloc-­‐cache sharedalloc 9.5 6.1 (a) Overall0 1 2 3 4 5 Normalized device slowdown Default hostalloc managed hostalloc-­‐cache sharedalloc 19.3 8.1 (b) Kernel-onlyFigure 3.3: Performance of Rodinia applications with the different globalmemory allocation methods.Figure 3.3 shows the slowdown under different memory management meth-ods, normalized to default for Rodinia benchmarks. Figure 3.3a shows the over-all runtime, Figure 3.3b shows the runtime of the GPU kernel only. In additionto the memory management schemes described in the previous section we also51experiment with hostalloc-cache. This is an intermediate option betweenhostalloc and sharedalloc that enables caching on the CPU, but not on theGPU.Due to caching restrictions, hostalloc causes performance degradation formost benchmarks; lud suffers the largest slowdown of 9.5×. The slowdown can beattributed to CPU caching, GPU caching, or both. We can deduce which effect isdominant by comparing hostalloc-cache, which only enables CPU caching,to sharedalloc, which also lifts restrictions on GPU caching. Additionally, wecan compare Figures 3.3a and 3.3b, which separate the overall runtime (affectedby CPU and GPU caching) from the runtime of the GPU kernel (affected by GPUcaching only). For example, the hostalloc version of kmeans does not sufferany slowdown on the GPU side, but its overall runtime is affected, leading us toconclude that the culprit is lack of caching on the CPU side. Particlefinder, ludand streamcluster suffer from both CPU and GPU caching restrictions. For manybenchmarks, the lack of GPU caching is the main source of overhead, and this iswhat the kernel runtime figure shows. We also verified that the actual costs of theallocation calls are negligible; it is the cost of accesses that have the major impacton performance.The managed scheme performs closely to the default, because both useidentical caching configurations. The exception is dwt2d, which actually useshostalloc and not default in the baseline implementation. Managed can pro-duce worse performance than default in scenarios where the cost of cache flushingthrough the synchronization call is high; with the default scheme the data is copiedby the programmer, so they have the option of cherry-picking what to copy, whilethe managed scheme always flushes caches on kernel launches to maintain con-sistency. However, for the applications presented here, the cost of the infrequentcache flushing and the cost of copying data between the CPU and the GPU is smallacross the board (not having to go through the PCIe bus certainly helps), so this ef-fect is not obvious. We will show later in the chapter that for concurrent CPU-GPUapplications the frequent cache flushing can have high overheads.Summary: For applications without concurrent CPU/GPU accesses to thesame data, the managed method provides equivalent performance to default withsubstantially better programmability. Hostalloc takes a toll on performance be-52cause it disables caching, and sharedalloc does not deliver any performancebenefits over managed because the applications analyzed here do not have anyCPU/GPU concurrency.GPU Caching BehaviourWe next examine the GPU caching restrictions implemented by hostalloc. Thedetails of its behaviour are not publicly available, so we infer them from simplemicrobenchmarks.0 20 40 60 80 100 120 140 160 180 1 2 4 8 16 32 64 128 1024 2048 4096 8192 Kernel /me (ms) Number of threads hostalloc no pad managed no pad hostalloc pad managed pad (a)0 50 100 150 200 250 300 350 1 2 4 8 16 32 64 128 1024 2048 4096 8192 Kernel 0me (ms) Number of threads hostalloc no pad managed no pad hostalloc pad managed pad (b)Figure 3.4: Accesses are cached in L2. (a) Each GPU thread accesses itsrespective array index. (b) All GPU threads access the same array index.Figure 3.4 shows the performance of a GPU kernel, launched with an increas-ing number of threads, that accesses an array managed either by hostalloc ormanaged. Each entry in the array is either padded to ensure that each entry is 32bytes (the size of the cache line in the L2 cache) or not padded. Figure 3.4a shows53a version of the kernel where each thread accesses its respective entry in the array(data[tid]++); Figure 3.4b shows a version where all threads access the firstentry (data[0]++).This experimental design enables us to observe how performance changes asmore and more threads access the same cache line. In the padded version whereeach thread accesses its own array element, each thread also accesses its owncacheline. In the unpadded version of the same experiment, several threads ac-cess the same cache line. In the experiment where all threads access the sameelement, the cache line is shared by all threads.We observe that as the number of threads accessing the same L2 cache lineincreases, performance gets worse with hostalloc, but is unaffected when themanaged scheme is used. Hostalloc provides stricter consistency between theGPU and the CPU, and it restricts caching on the GPU by disabling L2 caching al-together. This explains why writing to the same cache line slows down execution:there is more competition, so writing to the same cache line causes significant slow-down since every write has to go to main memory to ensure consistency. Hence,hostalloc is more appropriate for code that performs fine-grained data sharing.Figure 3.5 shows the results when data is cached in L1 caches in addition toL2. By default, global data are not cached in the L1 cache on Kepler class GPUs.However, this can be overridden using a compiler flag. The effects are less apparentwhen enabling L1 caching; however as the data set size increases and spills tothe L2 cache, the behavior becomes similar to the previous case. Furthermore,enabling the caching in L1 for global data can negatively affect applications whichrely on shared (local) memory. The L1 cache is used for caching shared memoryby default, so having global memory cached as well can reduce the effective usageof this cache.Concurrent Workload ExperimentsWe next demonstrate the utility of sharedalloc and hostalloc for work-loads that utilize both the CPU and the GPU concurrently. We use a microbench-mark and modify four Rodinia applications to have the CPU and GPU work col-laboratively on allocated data.540 5 10 15 20 25 30 35 40 45 50 1 2 4 8 16 32 64 128 1024 2048 4096 8192 Kernel 0me (ms) Number of threads hostalloc no pad managed no pad hostalloc pad managed pad (a)0 5 10 15 20 25 30 35 40 45 50 1 2 4 8 16 32 64 128 1024 2048 4096 8192 Kernel 0me (ms) Number of threads hostalloc no pad managed no pad hostalloc pad managed pad (b)Figure 3.5: Accesses are cached in L1. (a) Each GPU thread accesses itsrespective array index. (b) All GPU threads access the same array index.In the microbenchmark, we allocate a 128 MB integer array (32M entries)using the different allocation methods. The array is accessed 100 times in total.In the first iteration, the GPU increments the entries of the first half of the arrayand the CPU increments the entries of the second half. In the second iteration,the CPU works on the first half, and the GPU works on the second. The patterncontinues until all iterations are completed. Each iteration is executed concurrentlyby the CPU and the GPU. However, when the data is allocated using managed,the CUDA runtime does not allow the GPU and the CPU to access the allocatedregion concurrently. One way to overcome this limitation is to split the allocatedregion into multiple allocations, so that the CPU and the GPU work on differentallocations at any given time. While this would be feasible for simple programswhere the boundaries of data division are known a priori, for more complicated55scenarios where shared data is arbitrarily split across threads, such a solution wouldbe substantially more complicated to implement. Therefore, for the managedcase, we execute the GPU and the CPU portions in lockstep: i.e., the GPU kernelfinishes execution, then the CPU code begins to run. Under default, the data hasto be copied between the CPU and GPU in-between iterations since each iterationdepends on the values from the previous iteration.The default applications in Rodinia do not impose much interactions and shar-ing between threads running on the CPU and GPU. This has been explored in somedetails in a previous paper [53]. So, we modified a number of applications to in-troduce such interactions. In the modified Rodinia applications, the kernel workson half of the data by launching half as many GPU threads as in the original. Thesecond half of the data is processed by CPU threads (via OpenMP) concurrentlywith the GPU threads, except for the managed scheme where they run in lockstep.-­‐100 -­‐50 0 50 100 150 200 microbench cfd gaussian bfs kmeans % Improvement over default-­‐conc hostalloc managed sharedalloc Figure 3.6: CPU-GPU concurrent benchmarks. Both the CPU and GPU workon independent data. Default-conc (y-axis) is the default implementa-tion of the concurrent version of the benchmark where data has to beexplicitly copied between the CPU/GPUFigure 3.6 shows the performance under all memory management methodscompared to default. In terms of programmability, the default version is the mostcumbersome since data accesses can be non-trivial and figuring out which portionof the data to copy back and forth may be difficult. For managed, hostalloc,56and sharedalloc, the data is visible on both the CPU and GPU without theneed of explicit copies. sharedalloc improves performance by more than156% over the default method for the microbenchmark, and over 45% over themanaged scheme. For hostalloc, the main cause of the performance degra-dation is restricted caching. Although concurrency is possible with hostalloc,the overhead of not caching data on the CPU and GPU far exceeds concurrencybenefits. Comparing sharedalloc and managed, sharedalloc is superiorto managed because the CPU and the GPU are able to work concurrently. Cfd,and gaussian gain 25-50% performance relative to managed; bfs and kmeans gaina modest 4%.The applications presented in Figure 3.6 perform coarse-grained data sharing.That is, threads run a computation on disjoint sets of data and any data exchangesbetween threads (i.e., where one thread reads the data that another thread wrote) areinfrequent. In the following experiment we explore a fine-grained sharing scenarioand demonstrate the behaviors of the hostalloc and sharedalloc schemesin this case.To support a realistic scenario, we implemented a locking algorithm based onthe Peterson’s lock [87]. The algorithm uses three locks: a CPU-local lock (aTAS lock), a GPU-local lock (a TAS lock), and a shared Peterson’s lock. Threadsrunning on the CPU first compete for the CPU-local lock; threads running on theGPU will first compete for the GPU-local lock. The thread that acquired eachlock is called the leader thread. The leader threads on the CPU and the GPU willthen compete to acquire the shared Peterson’s lock. This arrangement is neededto avoid a deadlock on the GPU, which would occur if multiple threads within thewarp contended for the shared lock.Once the lock is acquired, the thread in the critical section accesses a shareddata region, allocated either with hostalloc or sharedalloc. Following thecritical section, the thread performs accesses to a non-shared region. We varythe number of accesses in the non-critical section relative to the critical section.Figure 3.7 illustrates the fine-grained scenario.While the shared region itself can be allocated using hostalloc orsharedalloc, the shared lock must be always allocated using hostalloc,because it requires strict consistency between the CPU and the GPU.57 CPU 1 CPU 0 CPU 2 CPU 3 L2 Cache CPU L2 Cache GPU Memory Controller GPU lock CPU lock Shared lock Warp(32 threads) Leader thread DRAM Shared CPU-­‐GPU region Figure 3.7: The locking scheme used in fine-grained sharing0 0.2 0.4 0.6 0.8 1 1.2 1.4 1.6 1.8 1 10 50 100 200 Execu/on /me (s) Number of accesses performed outside the cri/cal sec/on hostalloc sharedalloc Figure 3.8: Fine-grained sharing benchmark showing the performance ofhostalloc and sharedallocFigure 3.8 shows the results. When the number of accesses outside the criticalsection is small, hostalloc is superior to sharedalloc. Since hostalloc58does not cache the data, we don’t need to flush the caches after we release theshared lock since the latest copy of the data is already in main memory. How-ever, for sharedalloc, we need to flush the caches after releasing the lock.sharedalloc flushes the entire cache including the data that was actually mod-ified in the critical section, which causes higher overhead. However, as the numberof accesses outside the critical section increases (the critical section becomes rela-tively less important), this advantage of hostalloc evaporates, and the cachingadvantage of sharedalloc begins paying off as better performance.Summary: The managed memory management method is a significant im-provement in programmability relative to default, because it enables accessing thesame memory on the CPU and the GPU without explicit copy statements. How-ever, the downside is that the CPU and the GPU are unable to access the sameallocated region concurrently even if they access disjoint data. The hostallocmethod does not impose the concurrency restriction at the expense of disablingcaching on the CPU and the GPU. sharedalloc neither restricts caching norprevents concurrency, but needs to flush the caches when data written on one de-vice must be made visible on another. Cache flushing calls could be embeddedinto the CUDA API so there would be no negative impact on programmability.The downside of sharedalloc is in indiscriminate flushing of the entire cachewhen only a small portion of the data needs to be synchronized. This becomes evi-dent in fine-grained sharing scenarios, where hostalloc might be advantageousdepending on the size of the shared buffer and the number of accesses performedin the critical section.3.5 AMD Kaveri APUWe next focus on a different platform that implements a unified memory system.We evaluate an AMD A10-7850K Kaveri APU. There are several programmingframeworks available for the AMD APU. We briefly describe the programmingframeworks used on the AMD APU and then evaluate their performance.593.5.1 OpenCL 1.2 and OpenCL 2.0Memory management in OpenCL 1.2 is similar in programmability to the defaultscheme in CUDA. The CPU and the GPU do not share the same virtual addressspace. Therefore memory buffers have to be explicitly managed and communicatedbetween the two devices. A pointer on the CPU is inaccessible on the GPU andvice versa.OpenCL 2.0 makes memory management considerably simpler by introducingShared Virtual Memory (SVM). With SVM, the CPU and the GPU see a unifiedvirtual address space, similar to CUDA’s managed method.OpenCL 2.0 allows allocating a memory buffer either as coarse-grained orfine-grained. Coarse-grained buffers are allocated for access primarily on a sin-gle device (the CPU or the GPU), while fine-grained buffers are suitable for fine-grained sharing by the two devices. Coarse-grained buffers need to be mapped andunmapped using API calls, which act as synchronization points to ensure properaddress translation and consistency when one type of device accesses the data af-ter it might have been written by another. Fine-grained buffers do not need to bemapped or unmapped and are guaranteed to be consistent at kernel launch and ker-nel completion. Fine-grained buffers can also be declared with an atomic flag; inthat case, they would be kept consistent between the CPU and the GPU even duringkernel execution with the proper use of synchronization primitives.3.5.2 HSA: Heterogeneous System ArchitectureThe HSA framework is a low-level software layer that gives programmers finercontrol over the underlying integrated hardware. Similar to OpenCL 2.0, HSAemploys a unified virtual address space, where pointers can be seamlessly accessedfrom both the CPU and GPU.Using a special compiler, CL Offline Compiler [13], unmodified OpenCL ker-nels can be compiled to HSAIL (HSA Intermediate Language) and BRIG (a binaryrepresentation of HSAIL) which can then be loaded and run by the HSA runtime.However, the host-side source has to be modified, as OpenCL APIs are differentfrom those in HSA. The main modifications are related to the initialization of HSAand in launching the kernels.60On the Kaveri system, there are three different memory buses, named GARLIC,ONION and ONION+ (see Figure 3.9), and memory accesses on the GPU may takeone of these three paths [93]. If coherence with the CPU is not required as in thecase for coarse-grained buffers, memory requests take the GARLIC path whichleads directly to main memory without communication with the CPU. If kernel-granularity coherence with the CPU is required, as would be the case with fine-grained buffers, memory requests take the ONION path. If instruction-granularitycoherence is required (fine-grained buffers with atomics), the requests take theONION+ path. This path bypasses the L1/L2 caches on the GPU [93]. For theworkloads that we used in this chapter, both fine-grained and fine-grained-with-atomics memory allocations had identical performance, so we only present theanalysis for fine-grained. Fine-grained-with-atomics need to be used along withsynchronization primitives to ensure consistency. This would be similar to thefine-grained concurrent setting on the NVIDIA system. CPU 1 L1 CPU 0 L1 CPU 2 L1 CPU 3 L1 CPU GPU DRAM L2 $ (2MB) UNB L2 $ (2MB) CU L1 CU L1 CU L1 CU L1 CU L1 CU L1 CU L1 CU L1 GNB L2 $ (WB) CP Onion Onion+ Garlic Figure 3.9: Memory system architecture of AMD Kaveri. Figure reproducedfrom [93]. CP: Command Processor, GNB: Garlic North Bridge, UNB:Unified North Bridge.613.6 Evaluation of the AMD APU System3.6.1 Hardware/SoftwareWe perform our OpenCL and HSA experiments on an AMD A10-7850K KaveriQuad-Core 3.7 GHz system. Each pair of CPU cores share a 2MB L2 cache. Theintegrated GPU is an AMD Radeon R7 720 MHz Sea Islands family (GCN 1.1).It has eight Compute Units (CU). The GPU L2 cache is implemented as a physi-cally partitioned and distributed cache among the CUs. The total L2 cache size is512 KB with each cache slice 64-128 KB [12]. The OpenCL software stack usesthe proprietary FGLRX driver, while the HSA framework uses the Radeon OpenCompute Platform (ROCm).3.6.2 Benchmarks and ExperimentsFirst, we compare the performance of the memory management methods availablein OpenCL 1.2, OpenCL 2.0, and HSA using microbenchmarks. Then, we explorethe caching behaviour of GPU accesses. Finally, we study the performance ofRodinia benchmarks, which we ported to OpenCL 2.0 and HSA.The AMD platform offers us more data points to compare relative to theNVIDIA platform, because there are three programming frameworks and threetypes of memory allocation schemes (coarse-grained, fine-grained and fine-grainedwith atomics). To make the results easier to read and qualitatively comparable tothe other system, we focus on the following combinations of the frameworks andthe memory management methods.1. OpenCL 1.2, coarse-grained: The memory is allocated usingclCreateBuffer with the CL MEM READ WRITE flag. This method isqualitatively similar to the default scheme on the NVIDIA/CUDA platform.It is the most cumbersome in terms of programmability and provides nodata consistency between the devices and no sharing of virtual addresses.2. OpenCL 2.0, SVM, fine-grained: The memory is allocated usingclSVMAlloc specifying a fine-grained buffer. SVM is similar to CUDA’smanaged in that it provides a unified virtual address space for the CPU and62the GPU. Fine-grained buffers are kept consistent between the two devicesat kernel invocation boundaries at the expense of performance, similarly toCUDA managed. Unlike CUDA managed, this configuration does notrestrict concurrent access by the CPU and the GPU; however there is noguarantee on data consistency if the data is accessed concurrently by bothdevices2.3. HSA, fine-grained: This is similar to the above method but utilizes the HSAframework. The memory is allocated using the default system malloc call,which is a major portability advantage for HSA.From now on, we refer to the three chosen configurations simply as OpenCL1.2, OpenCL 2.0 and HSA, for brevity.Bandwidth Experiments0 2 4 6 8 10 12 8B 256B 8KB 16KB 32KB 64KB 128KB 256KB 512KB 1MB 2MB 4MB 8MB Bandwidth (GB/s) Array size OpenCL 1.2 (coarse) OpenCL 2.0 (SVM) HSA (fine-grained) Figure 3.10: Write memory bandwidth for OpenCL 1.2 (coarse), OpenCL 2.0(SVM) and HSA (fine-grained).We measure the read and write bandwidth under the three chosen configura-tions using the same test as in Section 3.4: we vary the size of the allocated array,2For truly concurrent and consistent accesses the programmer is expected to use fine-grainedbuffers with atomics along with synchronization primitives.63and each GPU thread reads or writes its own dedicated 8-byte entry. The num-ber of threads is, therefore, the same as the number of array entries. We verifiedboth experimentally and by reading the code and documentation that none of thememory management methods on the AMD system disable CPU caching or affectmemory access latency on the CPU, so we report the results of experiments on theGPU only.Figure 3.10 shows the results for the write bandwidth test. (We do not presentthe read bandwidth results, because they were almost identical to the write test.) Weobserve that the three different configurations deliver vastly different performance.For large buffer sizes, OpenCL 1.2 delivers the highest bandwidth, limited onlyby the speed of the hardware. HSA delivers roughly 10% lower throughput, andOpenCL 2.0 delivers little more than half the throughput of OpenCL 1.2. Forsmall buffer sizes, the performance differences are smaller, with HSA slightly onlytrailing the OpenCL versions.HSA enhances programmability over OpenCL 1.2 in two main ways: (1) HSAprovides a unified virtual address space between the CPU and the GPU, and (2)HSA configures the memory to be coherent between the CPU and the GPU. Bothof these programmer-friendly features come at a cost. Providing a unified virtualaddress space requires the GPU to perform address translation using CPU pagetables. As described in a previous article that includes several co-authors fromAMD [100], the GPU on the Kaveri system has a multi-level TLB hierarchy thatis accessed for address translation. Upon a TLB miss, a translation request is dis-patched to the IOMMU (Input/Output Memeory Managament Unit) that has itsown TLB hierarchy and walks CPU page tables upon a TLB miss. If the trans-lation is not found in the page table, IOMMU dispatches a page fault handlingrequest to the CPU. The address translation process has overhead; and this is partof the reason why HSA delivers lower throughput for large array sizes.The second cause of overhead is the cost of maintaining coherency betweenthe CPU and GPU. We couldn’t confirm that the platform implements hardwarefeatures enabling coherency, so this cost is either the actual cost of hardware co-herency or the cost of the software driver to functionally emulate hardware co-herency. We refer to this cost as functional hardware coherency. To isolate theoverhead of coherency from the address translation overhead, we perform a simple64experiment where multiple GPU threads concurrently update a single memory lo-cation. This way, the overhead of address translation is negligible. We compare twoconfigurations: (1) the memory is allocated (with HSA) as a fine-grained buffer,just like in Figure 3.10, and (2) the memory is allocated (with HSA) as a coarse-grained buffer. While the first configuration enables functional hardware coher-ence, the second one does not. Figure 3.11 shows the results and demonstratesthat functional hardware coherence, either implemented in software or hardware,incurs 15-25% overhead, which increases with the number of concurrent accesses.Returning to Figure 3.10, for small buffer sizes, the address translation over-head is small, so HSA trails the OpenCL implementations because of the functionalhardware cache coherency overhead. With larger buffer sizes address translationoverhead becomes dominant, so HSA benefits from the the functional hardwaresupport relative to OpenCL 2.0 (which handles address translation in software),but loses to OpenCL 1.2, which provides no unified virtual address space and thisincurs a smaller address translation overhead.Table 3.2 presents the number of data TLB loads and the number of interrupts,both measured on the CPU-side, for the experiment from figure 3.10 for the 1MBarray size. Indeed, OpenCL 1.2 performs a lot fewer TLB accesses than OpenCL2.0 and HSA; that is because it uses a coarse-grained allocation. OpenCL 2.0 andHSA perform a similar number of TLB loads if memory allocation is configuredas coarse-grained. What we find surprising is that OpenCL 1.2 generates the mostinterrupts and yet delivers the best performance. That could be explained by a highcost of address translation on the Kaveri platform (dissected in details in previ-ous work [100]); OpenCL 1.2 wins because it performs fewer of these expensiveoperations on a coarse-grained memory buffer.Metric OCL 1.2 (coarse) OCL 2.0 (SVM) HSA (fine)dTLB loads (billion) 17.7 28.2 25.9K interrupts/s 21.1 14.7 7.Table 3.2: Data TLB loads and InterruptsOpenCL 2.0 delivers the worst performance of all, which is attributed to theruntime/driver implementation. OpenCL 2.0 SVM implementation provides the650 200 400 600 800 1000 1200 1400 1600 1800 2000 128 512 1024 Throughput (K writes/ms) Concurrent GPU threads fine-grained coarse-grained Figure 3.11: Write throughput as multiple threads concurrently update a sin-gle memory location allocated as coarse (non-coherent) or fine (coher-ent) with HSA.convenience of a unified virtual address space and data consistency in software(runtime/driver). Profiling confirmed that a significant portion of the time spentby OpenCL 2.0 workloads is attributed to the driver. The GPU driver is respon-sible for producing a unified virtual address space on systems with no associatedhardware support. The driver allocates a GPU page table for a given range ofCPU virtual addresses and performs address translation. These driver activities in-cur performance overhead, which could explain the poor performance of OpenCL2.0. Furthermore, OpenCL 2.0 flushes caches once the kernel has finished runningto ensure that the CPU sees consistent data. That, and other software overheadis responsible for lower performance with this framework. Table 3.2 shows thatOpenCL 2.0 performs more TLB loads than both OpenCL 1.2 and HSA, and gen-erates almost twice as many interrupts as HSA, further pointing to driver overheadin this framework. Unfortunately, the OpenCL stack on our platform is proprietary,so we could not find further details on software overhead by examining the sourcecode.Summary: Similarly to our analysis of the NVIDIA platform, we observe thatprogrammability comes at a cost. Unified virtual address space requires performingaddress translation. Enabling the CPU and the GPU to see consistent copies of data66incurs overhead. OpenCL 2.0 pays a substantial cost for supporting these features,whereas HSA delivers them at a substantially lower overhead.Exploring GPU Caching Behaviour0 5 10 15 20 25 30 35 40 45 50 no padding padding (32 Bytes) padding (64 Bytes) Time (ms) OCL 1.2 OCL 2.0 HSA 101.12 (a)0 2 4 6 8 10 12 14 16 18 20 no padding padding (32 Bytes) padding (64 Bytes) Time (ms) OCL 1.2 OCL 2.0 HSA 99.56 (b)Figure 3.12: Latency of OpenCL 1.2, OpenCL 2.0, HSA kernel executiontime. Array size = 64 MB and each array entry is either not padded,padded to 32 Bytes, or padded to 64 Bytes (cache line size). (a)Kernel: memory[idx]++;. (b) Kernel: if(idx %2 == 0 )memory[idx]++;.To further understand the performance differences related to GPU caching,similar to our experiments with the NVIDIA platform, we varied the padding ofarray entries written by the threads. Unlike our NVIDIA experiments, we show theresults for a single array size (as opposed to showing the performance across ar-ray sizes), to highlight interesting performance effects that are unique to the AMDplatform. In this experiment, each thread writes its own eight-byte array entry in672.46 29.44 1.50 0 5 10 15 20 25 30 35 0 1 2 3 4 5 6 7 x slowdown to OCL 1.2 Padding in 8-byte increments OCL 2.0 Figure 3.13: Slowdown of OpenCL 2.0 compared to OpenCL 1.2 as we in-crease the padding size of array entries.a 64MB array. The entry is either not padded, padded to 32 bytes, or to 64 bytes.The size of the L2 cache line is 64 bytes, so we vary how many cache lines areaccessed. Figure 3.12a shows the results. Unlike the NVIDIA platform, the ”nopadding” case has much better performance across the board, because the GPUcan utilize memory coalescing to read/write each of the eight array entries with asingle memory access. This effect was not obvious on the NVIDIA platform, butis significant here.Once the entries are padded to 32 bytes, performance deteriorates for HSA andOpenCL 2.0 because the GPU can no longer coalesce memory accesses since eachcache line now contains only two entries. The second reason becomes apparent ifwe compare the OpenCL 2.0 bars in the figure. Except for the ”no padding” casewhich utilizes coalescing, whenever the cache line contains more than one entrythat might be accessed concurrently by more than one GPU thread, performancedrops drastically especially for OpenCL 2.0 (and to some extent for HSA).We believe that the observed caching behaviour is similar to the cache be-haviour we noticed with hostalloc memory on the NVIDIA platform. We donot see these effects with OpenCL 1.2 since the buffer is not shared between theCPU and GPU. With OpenCL 2.0, this software-triggered cache behaviour main-tains the consistency of SVM buffers and hence the drastic drop in performance.68HSA also seems to experience similar behaviour but to a lesser extent. As discussedearlier, HSA utilizes some hardware features which eliminate some software co-herency overheads. However, we notice that the 64-byte padding case seems toperform worse than OpenCL 2.0. As the number of cache lines used increases(which is the case with 64-byte array entries), any hardware mechanism is inferiorto simply flushing the cache.Figure 3.12b further demonstrates the case of performance deterioration forOpenCL 2.0 whenever coalescing cannot be utilized and the cacheline containsmultiple accessed entries. Since only even-numbered threads will write to the arrayentries, the ”no padding” case means that the cacheline contains 4 accessed entries,and hence the huge performance drop for OpenCL 2.0. Finally, figure 3.13 showsthe performance slowdown of OpenCL 2.0 compared to OpenCL 1.2 as we increasethe padding size (i.e. decrease the number of entries in each cacheline). Except forutilizing coalescing (the 0-padding case), the figure confirms that the more theentries accessed on the same cacheline, the larger is the slowdown.Rodinia BenchmarksTo further compare the memory management methods on the AMD system, weported four applications (BFS, Gaussian, kmeans, streamcluster) from their orig-inal implementation in OpenCL 1.2 to OpenCL 2.0 and HSA. For the OpenCL2.0 version, we allocate all the data used by the CPU and GPU as SVM fine-grainbuffers. Also, we remove all data copies since they become unnecessary once datais allocated as SVM-fine buffers. Programmability is increased as the lines of codeassociated with data copies and the multiple CPU and GPU allocations are elim-inated. For HSA, we port by using the HSA APIs for initialization and kernellaunching on the CPU side. The application kernels themselves are not modifiedfrom the implementations in OpenCL 1.2 and 2.0. We use the CL Offline Com-plier (CLOC) to generate a BRIG binary module that can be read using an HSAAPI call.Figure 3.14 shows the running times of the three implementations normalizedto the running time with OpenCL 1.2. Based on the analysis of microbenchmarksin the previous sections, the performance slowdowns when using OpenCL 2.0 and690 0.2 0.4 0.6 0.8 1 1.2 1.4 bfs gaussian kmeans streamcluster Normalized running <me OpenCL 1.2 OpenCL 2.0 HSA 3.67 Figure 3.14: Runtime of Rodinia applications normalized to OpenCL 1.2HSA are expected. For BFS and kmeans the slowdowns are mainly attributed toslower executions of the GPU kernel due to the effects highlighted earlier. Al-though the OpenCL 1.2 version actually does copy data between the CPU andGPU in this realistic test, in contrast to the previous section’s microbenchmarks,the copying is performed only once in each direction. The kernels are executedmultiple times, amortizing the overhead of data copying.For streamcluster, the overhead of HSA is higher than OpenCL 2.0, becausethe HSA version does not use local memory on the GPU, in contrast to the OpenCL1.2 and OpenCL 2.0 versions. We were not able to allocate local memory in HSA,so we used global memory instead. This caused the HSA version to have a higheroverhead than the OpenCL 2.0 version.The performance speedup of HSA compared to OpenCL 1.2 for Gaussian isbecause the initialization of HSA takes less time than the initialization of OpenCL(25 ms vs 300 ms). Since the application runs for only about 2 seconds, the initial-ization overhead is significant.Summary: Performance overheads shown with the microbenchmarks thatstress the memory system do not necessarily apply to real applications that per-form substantially more computation. Programmability does not always come ata higher cost. Exploring the overheads of programmability in applications withconcurrent access to data by the CPU and the GPU, such as those introduced in70previous work [22, 33, 61, 76, 96], would further elucidate this trade-off.3.7 Related WorkPrior to our work, Hestness, Keckler and Wood [57] analyzed the CPU and GPUmemory system behavior on a simulator. Their analysis focused on how the mem-ory behaviour of an application changes when it runs on the CPU vs on the GPU;they find significant differences and attribute them to the differences in memoryhardware. Our work studies the cost of memory accesses in different program-ming frameworks and analyses how programmability enhancements may affecttheir overhead.A second paper by Hestness [57] also analyzed the impact of hardware cachecoherency on applications [58]. In contrast to our work, they used a simulatedsystem. Similarly to our experience, they found that many existing GPGPU ap-plications do not benefit much from the physically unified memory, because theywere written at the time when this convenience was not available on GPU systems.To accommodate their study, they modify a number of applications to take advan-tage of the unified physical memory and cache coherence, and show that adaptingthe software is key to capturing the benefits of the hardware.Implementing a full-blown CPU-like cache coherency on integrated CPU-GPUsystems has practical limitations, so prior work looked at alternative software/hard-ware methods to implement partial coherency [15]. The authors relied on a simu-lator to evaluate their scheme which eliminates the need for CPU-GPU hardwarecoherency. Our evaluation on a real system showed that software-only designs mayincur large overheads.Similar to our performance analysis of the AMD platform, Mukherjee char-acterize the benefits of the new features of OpenCL 2.0 and HSA [81]. Our workfocuses mainly on the cost of memory accesses under the different allocation meth-ods on multiple frameworks.Sun et al present a benchmark suite that explores the collaborative executionpatterns on the CPU and GPU using an AMD APU [96]. In our work, we eval-uate the performance of existing platforms from a memory management scope.We briefly touch on concurrency workloads but mainly focus on memory access71performance.3.8 SummaryWe presented an analysis of memory management methods on real integrated CPU-GPU systems. We analyzed the cost of programmability enhancements, such as theunified virtual address space and data consistency. We see two potential avenuesfor future work. (1): Hardware improvements to decrease the overhead of virtualaddress translation and cache coherency. (2): Improvements to software frame-works that discretely use these expensive features. However, to enable meaningfulanalysis in these future studies, there is a need for applications that truly rely on andexercise the features such as the unified virtual address space and cache coherency.72Chapter 4Analyzing an off-the-ShelfProcessing-in-Memory System14.1 IntroductionThe memory wall has plagued computer systems for decades. Computer scien-tists have pursued the ideas of in-memory and near-memory computing, aiming tobring computation closer to data. Yet most of the proposed hardware never made itpast simulation or proprietary prototypes, so some questions about this technologycould not be answered. In this chapter, we examine a recent commercially availableProcessing In Memory (PIM) system from UPMEM [40] to better understand itsstrengths and limitations. Most in-memory and near-memory solutions proposedin the past required specialized hardware that was incompatible in some way withcurrently deployed systems [16, 17, 51, 64, 67]. The hardware we evaluate wasdesigned specifically to be used as a drop-in replacement for conventional DRAM,which has imposed some limitations that are not present in many of the simulatedarchitectures.UPMEM’s DRAM DIMMs include general-purpose processors, called DRAMProcessing Units (DPU). Each 64MB slice of DRAM has a dedicated DPU, socomputing resources scale with memory size. We examine a number of applications1This chapter is a modified version of a previously published work [83]73that require high memory bandwidth and whose computational needs increase withthe size of the data. We observed that, indeed, application throughput scaled withdata size, but scaling was not always the best that this hardware could achieve. Themain reasons were the difficulty of accessing data located inside DPU-equippedDRAM from a host CPU without making a copy, the limited processing power ofthe DPUs, and the granularity at which the DPUs are controlled.The focus of this chapter is on understanding the strengths and limitations ofPIM when it is integrated into an off-the-shelf system. After examining and diag-nosing the main features of the PIM system, we delve into simple application casestudies where we compare the performance of PIM to that of a CPU-only systemand a GPU system. The current status of PIM hardware is similar to those of accel-erators where parallel regions of applications are offloaded to the parallel computeunits. Therefore, comparing PIM to other parallel systems gives a better overviewof where such hardware stands in the realm of accelerators.4.2 ArchitectureMany PIM architectures were proposed in the past [16, 17, 43, 49, 50, 52, 62, 80,92, 95, 102, 104], ranging from logic gates embedded in each memory cell up tomassive arrays of parallel processors that are located close to the memory. Theyall share the goal of overcoming the limitations of memory bandwidth but differvastly in design. Nguyen et al. [82] introduce a classification system for variousdesigns based on the distance of the compute units from the memory. We studyspecific hardware from the COMPUTE-OUT-MEMORY-NEAR (COM-N) model.This model includes compute logic that is located outside of the memory arraysbut inside the memory package (same or different silicon die). This is sometimescalled CIM (compute-in-memory) but is more widely known as PIM (processing-in-memory). After our study was concluded, Samsung announced the productionof similar PIM hardware (called FIMDRAM) that also falls into the COM-N model[68]. While we could not perform an in-depth comparison, some similarities areapparent from the published design. These common design points are likely toappear in new designs as well, and the lessons we learned from studying UPMEM-hardware will be generally applicable to this class of PIM.74Organization In the architecture we study, DRAM is organized into ranks (tworanks per DIMM), and each rank contains eight memory chips. Each DRAM chipincludes eight DRAM Processing Units (DPUs). Each DPU is coupled with a64MB slice of DRAM. There are 64 DPUs per rank and 128 DPUs per 8GB mem-ory DIMM. Since the DPUs and DRAM are built on the same silicon die, intro-ducing the DPUs comes at the cost of some DRAM capacity. With this design, thesum of DPUs and DRAM is limited by the silicon area on the die and the processdensity (number of transistors per unit area). Our experimental system has 36GBof DPU DRAM and 576 DPUs.DPU Capabilities A DPU is a simple general-purpose processor. UPMEM’sdesign only supports integer operations (no floating-point hardware). Each DPUhas exclusive access to its slice of DRAM over a private bus. The more totalDRAM there is, the more DPUs there are. In other words, we get an additionalunit of computation, and an additional bus, with each additional unit of memory– computing capacity and intra-DIMM bandwidth scale with the size of memory.Figure 4.1 shows a schematic view of a single DPU and its periphery. Before aDPU can compute on the data, it must copy it from DRAM to a 64KB privateSRAM buffer, called the working memory. The copy is performed with an explicit,blocking DMA instruction. The DMA engine can copy up to 2KB in a single oper-ation but the copy time increases linearly with the size. To hide DMA latency, eachDPU has 24 hardware threads, called tasklets, that can be scheduled for executionsimultaneously. Because it is an interleaved multithreading (IMT) design [103],only one tasklet can advance at each cycle. When a tasklet is blocked on a DMAoperation, other tasklets can still make progress.DPU Speed Processing capabilities of DPUs are well below that of a commonhost CPU. DPUs have a simple, in-order design, and are clocked at anywhere from266MHz (in our experimental system), to 500MHz (projected in commercial offer-ings). This was necessary to fit within the power envelope of commodity systems.DPU communication In DRAM there are no communication channels be-tween individual slices of memory dedicated to DPUs. Thus, DPUs cannot sharedata, other than by copying it via the host. Communication between DPUs on dif-ferent chips would require additional pins in the package, which would change theform factor from standard DRAM. Creating a communication bus between DPUs75Figure 4.1: DPU Architectureon the same chip is theoretically possible, but severely restricted by the numberof available metal layers. Therefore, we cannot expect any PIM designs in theCOM-N model to have any DPU-to-DPU communication until these problems canbe solved. This constraint implies that programs running on DPUs must shard dataacross DPUs ahead of the execution and use algorithms that don’t require datasharing.Interleaving At the chip level, interleaving dictates the ordering of bytes asthey are written into DRAM. In the DIMMs in our test system, each DRAM chipconnects to the DDR4 bus by an 8-bit interface. 8 chips together form the 64-bit wide bus. When a 64-byte cache line is committed to memory, the data isinterleaved across DRAM chips at byte granularity, the first byte going into chip 0,the second byte into chip 1, etc. So when the write is done, the first chip receivesbytes 0, 8, 16, ..., 56. Interleaving at this level is used to hide read and write latencyby activating multiple chips simultaneously. Interleaving is transparent to the hostsince the data is reordered appropriately in both directions (i.e., during reads andwrites). But a DPU, which can access data only in its own slice, sees only whatis written to that particular chip (i.e., every nth byte of the data). This propertymakes it difficult to access the data by the host CPU and the DPU in the samelocation. Unless a program running on the DPU can be modified to operate onnon-contiguous data, we must make a separate copy, transposing it in the process,76so that each DPU receives a contiguous portion of the data.Transposition UPMEM’s SDK counteracts the interleaving by transposing thedata. While the cost of the transposition is small (the current SDK provides anefficient implementation that uses SIMD extensions), the cost of data copy, whichrequires going over the DRAM bus, is not. For applications that can place theirdataset into the DPU memory and compute on it using only DPUs, the one-timecost of the copy can be negligible, but for applications where the data is short-livedor both the DPU and host CPU must access it, frequent copying will stress thememory channel – the very bottleneck this architecture aims to avoid.Control granularity Accessing multiple chips in each bus cycle also means itis natural to schedule DPUs across those chips as a group. With the current SDKimplementation, a full rank must be scheduled as a single unit. We can only send acommand to an entire rank of DPUs and cannot launch a DPU or read its memorywhile another DPU in the same rank is executing. Similarly, the host cannot accessany memory in a rank if any DPU in that rank is executing. In practice, this re-sults in treating DPU-equipped memory as an accelerator, rather than part of mainRAM. The host typically keeps programs and data in the “normal” DRAM (i.e.,host DRAM), copies data for processing into the DPU DRAM, and copies back theresults.4.3 Evaluation4.3.1 WorkloadsConsidering DPU architecture properties we chose programs that exhibit threeproperties: 1) their computational needs scale with the size of the data, 2) theycan be easily parallelized, and 3) the parallel computations share minimal data.The SDK for the DPU architecture includes a C compiler and libraries withfunctions to load and launch DPUs, copy and transpose the data, etc [14]. Whilewe could write the code in a familiar language, the main challenges had to dowith memory management (as the DPU needs to copy data from DRAM into theworking memory prior to processing it), and with efficiently controlling ranks ofDPUs.77Snappy Snappy is a light-weight compression algorithm designed to processdata faster than other algorithms, such as gzip, but at the cost of a lower compres-sion ratio [56]. Snappy divides the original file into blocks of a fixed size (typically64KB) and compresses each block individually. The original implementation isdesigned for sequential (i.e., single-threaded) processing. The file format concate-nates compressed blocks head to tail, without any separation marker. The result isthat to decompress block n, the program must first decompress blocks 0 .. n−1.To enable parallelism, we changed the compressed file format by prepending eachblock with its compressed size. This small change enables the host program toquickly locate the start of each compressed block without having to decompressthe previous blocks. We also modified the header by adding the decompressedblock size. That enables us to experiment with different block sizes. Snappy isbandwidth-intensive because of the small amount of processing, especially duringdecompression. The majority of the time is spent copying literal strings or previ-ously decoded strings. We consider both snappy compression and decompressionas the algorithms exhibit different behaviours.Hyperdimensional computing Hyperdimensional Computing (HDC) [63] isused in artificial intelligence to model the behaviour of large numbers of neurons. Itrelies on vectors with high dimensionality (at least 10,000 dimensions), called hy-pervectors, to classify data and find patterns. Our HDC application performs clas-sification of raw electromyography (EMG) signals into specific hand gestures [88].We built upon a prior reference implementation [79]. The classification works bycomputing the Hamming distance between a previously encoded input hypervectorand previously trained vectors. Our implementation distributes the raw signal dataamong DPUs, which then encode it into hypervectors and perform the classifica-tion.Our baseline for performance comparison is a single host CPU2 since our focusin this chapter is to understand when the DPU architecture overcomes the von Neu-mann bottleneck and when it does not. In chapter 5, we examine a more complexapplication and compare performance to multiple CPU cores and a GPU.2Our host machine is a commodity server with an Intel(R)Xeon(R) Silver 4110 CPU @ 2.10GHzwith 64GB of conventional DDR4 DRAM.784.3.2 Memory bandwidth scalingIn the DPU architecture, each additional slice of DRAM gets its own DPU alongwith the working memory and the DMA channel. Herein lies its main strength:scaling computing capacity and intra-DIMM bandwidth with the size of mem-ory. To demonstrate, we ran a simple program where all DPUs iterate over the datain DRAM, copying every byte from their respective DRAM slices into the workingmemory. As shown in Fig. 4.2, the data processing bandwidth increases with thesize of the data, reaching roughly 200GB/s for 36GB worth of DRAM (approxi-mately 350MB/s per DPU). For a system with 128GB of DRAM and DPUs clockedat 500MHz (which is the target for commercial release), the aggregate bandwidthattained by all DPUs would reach 2TB/s.Figure 4.2: Memory bandwidth scales with memory capacity.For applications whose compute and bandwidth needs grow with the size of thedata, scaling of compute and bandwidth resources with memory has a direct andpositive impact on performance. E.g., Fig. 4.3 shows the throughput of Snappycompression as the file size (and the number of DPUs sharing the work) increases.We use the optimal number of DPUs for each file size – i.e., using more DPUsyields no additional speedup. The input data resides in host DRAM; the exper-iment copies the data to DPU DRAM, allocates and launches the DPUs to runcompression, and copies the compressed data back to host DRAM. As the file sizegrows so does the throughput – a direct result of compute capacity and internal79Figure 4.3: Snappy compression throughputbandwidth scaling with additional memory.Breaking down the runtime, we observe that the execution time on DPUs (Runin Fig. 4.4) remains roughly the same even as we increase the file size, becausemore DPUs are doing the work. Again, this is the effect of scaling compute re-sources with data size. The DPU execution time does increase for 512MB and 1GBfile sizes, because as we reach the limits of our experimental hardware, each DPUgets more data to process. The non-scalable components of runtime in Fig. 4.4,Copy In, Copy Out and Setup, are discussed in §4.3.3.A final experiment illustrating the strength of scaling involves a comparisonwith a GPU3, an accelerator that shares similarities with DPUs (massive paral-lelism, high memory bandwidth), but does not have the property of scaling com-putational resources with memory. Fig. 4.5 shows the throughput relative to a hostCPU (same as in the DPU experiments) of snappy compression as the file sizegrows. Contrary to the DPU architecture, the GPU scaling is flat and the through-put does not surpass that of the host CPU by more than 5×. This coincides with theinternal memory bandwidth of the GPU, which is approximately 5× higher than3We used an RTX 2070 GPU with 8GB of GDDR6 and 448GB/s global memory bandwidth.There are 36 SMs, with 64 CUDA cores per SM, clocked at 1.4GHz, There is a 64KB L1 cache perSM, and 4MB L2 cache shared by all SMs.80Figure 4.4: Runtime breakdown for compression for dickens fileFigure 4.5: Snappy compression on GPUsthat of the CPU.Where each DPU has a dedicated working memory and a DMA channel, allGPU cores within an SM (streaming multiprocessor) share an L1 cache. This data-intensive workload with a random access pattern caused thrashing in the small64KB L1 cache. Memory fetching became the bottleneck and as a result, we ob-tained the best performance by launching kernels consisting of a single thread percompute block (and hence one thread per SM) to reduce contention for the cache.The input file fit in the GPU memory, so there were no fetches over the PCIe bus81Blocks / Threads per blockTime (s)0.001.002.003.004.000.000.100.200.308192 / 14096 / 22048 / 41024 / 8512 / 16256 / 32128 / 6464 / 12832 / 25616 / 512Total Time (s)Time per Block (s)Figure 4.6: GPU kernel dimension effects on performance for compression -dickens 32MB fileonce processing started. In other words, the GPU fell victim to a variant of the vonNeumann bottleneck.The RTX 2070 has 36 SM (streaming multiprocessors), each SM has 64 CUDAcores. Each SM has limited resources (registers, L1 cache, ) that are shared amongall the running threads. As mentioned, the best kernel grid dimension setup is toassign each 32K compression block to a single GPU thread inside the SM. In otherwords, it is best to have a single thread in the SM utilize all the available resourcesthan have multiple threads sharing the resources. So, although there are 64 avail-able cores in each SM, it is best to only have one core in this SM running and usingall the registers and L1 cache to itself! Having many threads increases resourcecontention and increases diversion among the threads since block compression isnot a uniform process among all the 32K blocks composing the file. Figure 4.6demonstrates this. This figure shows that the best performance for decoding thefile is to launch blocks containing a single thread that utilizes all the available re-sources. As we we increase the number of threads inside the block, performanceworsens since both the total time and the time to complete each block starts toincrease.824.3.3 Data copyAs we discussed in §4.2, the need for de-interleaving (transposing) data for DPUsmeans that it is difficult to avoid making copies when data needs to be accessed byboth DPUs and host CPUs. For applications where DPUs can be used exclusively,the data can be copied to DPU memory just once and reside there for the rest of theruntime, making copy overhead negligible. In programs where data is ephemeralor must be accessed by both the host CPU and DPUs, frequent copying will besubject to memory channel limitations.The compression runtime breakdown presented in the previous section(Fig. 4.4) illustrates this point. The DPU execution time remains roughly con-stant as data size increases, because each DPU has a dedicated DMA channel to itsDRAM slice. In contrast, the time to copy input data to DPU memory and copythe output back increases with the file size and is limited by the DRAM channel.Setup overhead in compression does not scale, because dividing up the workrequires sequential parsing of the input file. This is a property of the workload andimplementation, and not fundamental to the architecture. Copy overhead, on theother hand, is fundamental to the system and is the reason why the compressionthroughput in Fig 4.3 grows sublinearly with the increasing number of DPUs.Figure 4.7: HDC speedup over the host as file size increasesHDC is less affected by the copy-in overhead than Snappy and has negligible83copy-out cost due to the small dataset. Overall, copying takes about 10% of thetotal runtime, and thus HDC enjoys better scaling, as shown in Figure 4.7. Thisbetter scaling is also apparent when comparing the performance to the GPU. Fig-ure 4.8 shows the same experiment but run on the GPU. We ported the same versionof the application to CUDA. Only for the smallest file size does the GPU have abetter performance than PIM. HDC and other AI applications, where DPUs canrepeatedly perform inference on long-lived data that remains in DPU memory, aretherefore well-positioned to circumvent the memory channel bottleneck.File Size (MB)Speedup0.002.004.006.008.0010.0012.000 1 2 3Figure 4.8: GPU HDC speedup over the host as file size increases4.3.4 DPU SpeedAnother example where DPU processing power was less impressive compared toother available accelerators was snappy decompression. On the DPU, this work-load performed very similarly to compression. On the GPU, in contrast with com-pression, L1 cache was not the bottleneck: decompression works with smaller inputblock sizes and is largely write-intensive. Figure 4.9 shows decompression perfor-mance on the GPU. The GPU showed similar scaling trends as on the DPU andoutperformed the host CPU by up to 24×, while DPUs’ advantage was at most8×.84File size (MB)Speedup0.005.0010.0015.0020.0025.008 16 32 64 128 256 512 1024 2048dickensmozillanci saospamfilexmlFigure 4.9: Snappy decompression on the GPU relative to a single host CPU.4.3.5 Communication and control granularityCompliance with the DDR interface and software limitations in the current SDKmakes it inefficient to control DPUs one at a time: as long as a single DPU inthe rank is busy, all others are inaccessible. This means that the program mustdecide how it will divide the data between DPUs in a rank and launch them all atonce. While this is a trivial decision for embarrassingly parallel workloads, otherworkloads might under-utilize the hardware as a result. The fact that DPUs cannotcommunicate and steal work from one another makes matters worse.4.4 SummaryApplications can be limited by the conventional memory hierarchy, and addingmore CPUs may improve performance substantially. We have shown some real-world applications that can be accelerated by PIM by scaling memory bandwidthwith compute resources. PIM shares many attributes (such as massive parallelism)with other accelerators but the distribution of processing units inside the mem-ory gives it a unique advantage in certain cases. The architecture we evaluated85has some limitations that prevent exploiting the full capabilities of the hardware,which remain as challenges for future designs. Costly data transfers which mustbe performed in order to share data between the host CPU and PIM memory addssignificant overhead to processing time but can be mitigated by reusing data in-place. Despite the limitations, we find that PIM can be effective when applied tothe correct class of problems by scaling the processing power with the size of thedataset. Hence, in the next chapter, we explore such an application with detailedexamination.86Chapter 5Characterizing Image Decodingon Processing In MemorySystems5.1 IntroductionFor Processing-in-Memory (PIM) systems to be widely adopted as an alternativeto accelerators and other types of computing systems, software systems and ap-plications utilizing this new compute paradigm need to be further scrutinized. Inthis chapter, we examine an application that constitutes a critical component inmany domain fields including image manipulation, storage, and deep learning.In particular, data preprocessing pipelines in deep learning applications have be-come a major bottleneck directly affecting the overall performance of model train-ing [60, 78, 105]. Preprocessing tasks are normally executed by CPUs (and par-tially offloaded to GPUs) which should steadily feed the resulting transformed oraugmented data to GPUs. GPUs are then responsible for the core training compu-tations. To avoid data stalls where the GPUs sit idle waiting for preprocessed datato become available, the CPUs should continually feed the GPUs at a high enoughrate to keep the GPUs busy [78, 105]. Nvidia’s data loading library DALI [3]partly mitigates the shortcomings of the preprocessing pipeline by offloading some87preprocessing tasks to the GPUs. However, using GPUs for preprocessing meansthere would be less GPU compute and memory capacity available for training.In chapter 4, we demonstrated the potential of using DPUs for various prob-lems. Image preprocessing (and batch image decoding in particular) consist ofhugely parallel tasks, and DPUs provide abundant parallelism that can be utilizedfor such workloads. In this chapter we examine the preprocessing pipeline for DLapplications. We particularly focus on image classification where preprocessingconsists of multiple operations including image decoding, resizing, cropping, andothers. We identify image decoding to be the predominant component in imageclassification tasks. We provide an image decoding implementation on the DPUsand characterize its performance comparing it to other CPU and GPU implemen-tations.5.2 Preprocessing and the CPU bottleneckPrevious work has shown that the CPU is the bottleneck in preprocessing [78]. Thisbottleneck means that the CPU cores are unable to provide a steady stream of datato keep the GPUs busy performing the training computations. One possibility formitigating this bottleneck is dedicating more CPUs so that the parallel preprocess-ing work can be done on these CPUs and data can be fed more quickly to GPUs.However, it was also shown that for some workloads it may require more than 24CPU cores to partly alleviate this bottleneck. This cost might not be feasible asmany typical training servers will have a ratio of GPU to CPU cores of 1:3 [78].Therefore dedicating 24 CPU cores to eliminate the bottleneck could be extremelycostly or impractical. Furthermore, adding more CPUs could be difficult due tophysical constraints as there is limited space on the board and there are limits tohow many boards can be added in typical servers. As we showed in chapter 4,PIM systems have a compute scaling property that can be utilized for differentworkloads. JPEG decoding presents an interesting challenge as it contains a subtlecomputational intensity (the actual decoding work) in addition to memory trans-fer requirements (copying images). Furthermore, the computational capabilities ofthe PIM DPUs are far lower that their CPU or GPU counterparts. However, PIMwould be attractive as it augments existing systems with additional compute units88by simply plugging in memory modules. Therefore, in this chapter we evaluatewhether PIM, with its lower compute capabilities but higher scalability, is suitablefor such a complex workload.5.2.1 Preprocessing in DNN TrainingBefore starting our PIM evaluation, we examine the current landscape of DNNtraining as a motivating use case. Most modern DNN frameworks use custom andhigh performance data loading libraries. Nvidia’s DALI [3] is a state-of-the-artdata loading and preprocessing library compatible with many DNN frameworks.It consists of high-performance implementations of many preprocessing primitivesthat are GPU-accelerated. DALI also allows for running these primitives on theCPU, the GPU, or both. There are other data loading libraries in frameworks suchas PyTorch, TensorFLow, MxNet, and others. However, DALI has the best perfor-mance out of other competing data libraries [78].We study the throughput from multiple image training models to see how muchpreprocessing affects the overall performance. We compare the CPU and GPUversions of DALI when running these workloads. The actual training computationsare always done on the GPU; only the preprocessing tasks are the ones that areaffected by configuring DALI to use the CPU or GPU for preprocessing.In figures 5.1, 5.2, and 5.3 we show the training throughput from three com-monly used models in image classification: Resnet18, Alexnet, and Mobilenet us-ing the Imagenet-1k dataset [89]. We run these experiments on our GPU systemdescribed in table 5.1. The top flat yellow line in these figures represents the maxi-mum throughput that can be achieved if there was no bottleneck at all. We measurethe no-processing case in a similar way as was done in previous work, which usedsynthetically preprocessed data that is already loaded on the GPU [78]. In otherwords, this is the case if there was no overhead at all for performing preprocess-ing. We show later in section 5.2.2 that image decoding is the major contributor topreprocessing overhead. Once decoding overhead is eliminated, data can flow intothe GPUs for performing the actual training. Practically, however, the overhead ofpreprocessing (and decoding in particular) cannot be completely eliminated.In these experiments, we compare both a CPU and a GPU version of prepro-89# CPU coresTraining Throughput (samples/s)025050075010001 2 3 4 5 6 7 8 9 10 11 12preprocessing-CPU preprocessing-GPU no-preprocessingFigure 5.1: Training throughput of Resnet18 with ImageNet-1K dataset withpreprocessing done on the CPU, GPU, and no preprocessingcessing and verify previous conclusions which were made about the CPU being abottleneck for preprocessing [78, 105]. The trends in all studied DNN image clas-sification models are similar: there is a need for more compute power (more CPUcores) in order to approach the maximum achievable throughput. Using CPU coresalone for preprocessing is detrimental to performance as is conveyed by lookingat the blue lines (Preprocessing-CPU) in the figures. For some models, even us-ing 12 CPU threads (which is not practical on typical training servers [78]) is notenough to reach maximum throughput. This is why most modern frameworks relyon using GPU-accelerated preprocessing libraries such as DALI to alleviate suchoverhead [3, 78, 105]. However, even relying on GPU-accelerated DALI doesnot fully eliminate the preprocessing bottleneck as can be seen from the red lines(Preprocessing-GPU) in the figures.90# CPU coresTraining Throughput (samples/s)01000200030001 2 3 4 5 6 7 8 9 10 11 12Preprocessing-CPU Preprocessing-GPU no preprocessingFigure 5.2: Training throughput of Alexnet with ImageNet-1K dataset withpreprocessing done on the CPU, GPU, and no preprocessing5.2.2 Preprocessing Primitives and DecodingThere are four main tasks that comprise preprocessing for the models used in ourexperiments: CropMirrorNormalize, Resize, Flip, and Decode. CropMirrorNor-malize does three operations as indicated in its name: cropping the image to aspecific size used for training, mirroring the image, and normalizing it. Resize re-sizes the image to the dimensions required for training, and flip does an axis flip.Decode performs decompression of JPEG images, and is the most computationallydemanding out of all the preprocessing tasks as demonstrated in figure 5.4. Usingthe GPU mode for preprocessing in DALI, the GPU-accelerated version of theprimitives are executed very quickly; it is the decoding part that takes almost all ofthe preprocessing time. Using the CPU for preprocessing, decoding still consumesthe most time, but other primitives such as resizing are still costly.91# CPU coresTraining Throughput (samples/s)02004006008001 2 3 4 5 6 7 8 9 10 11 12preprocessing-CPU preprocessing-GPU Max speedFigure 5.3: Training throughput of Mobilenet with ImageNet-1K dataset withpreprocessing done on the CPU, GPU, and no preprocessingTime (s) per image0.00000.10000.20000.30000.40000.5000DALI-CPU DALI-GPUCropMirrorNorm Resize Flip DecodeFigure 5.4: Breakdown of the major primitives used in preprocessing withDALI-CPU and DALI-GPU925.3 PIM-Decoding: Overview and ImplementationIn chapter 4, we examined PIM systems in details and demonstrated their abilityand suitability for accelerating memory-bound workloads. We also showed thatsuch systems are able to massively scale computations across the DPUs which op-erate independently of each other. We now examine batch JPEG decoding (whichis a critical and computationally expensive step in DNN training workloads and inpreprocessing in particular) and analyze its suitability for running on a PIM system.5.3.1 JPEG introductionJPEG is a digital image encoding standard which uses lossy compression to storeimages [6, 7]. Neural networks frequently use JPEG images as training datasets,especially for image and object recognition tasks. Since JPEG stores image data inan encoded format, the data has to be decoded before it can be used for training.JPEG consists of four major steps: Decoding the Huffman bitstream [5], dequan-tization, Inverse Discrete Cosine Transform (IDCT), and color conversion fromYCbCr to RGB.JPEG decoding works as follows: first, the Huffman codes are derived from theHuffman Tables that are read from the markers (which are special locations in theJPEG header that mark the start of different segments). Using the Huffman codes,the encoded image data can be decoded bit by bit. Decoded data is grouped into 8x 8 pixel blocks called Minimum Coded Units (MCUs). Therefore, a 1920 x 1080gray-scale image, for example, has 32400 MCUs, and a colored image has 97200MCUs since it has 3 color channels. After decoding into MCUs, each MCU can beindividually dequantized. This is a simple step in which each pixel in each MCUis multiplied by its corresponding value in the Quantization Table read from themarkers. The dequantized MCUs now undergo inverse discrete cosine transform.Basically, each value in an MCU represents the sum of cosine functions oscillatingat different frequencies. This sum has to be transformed back to a finite sequenceof data points for each MCU to obtain the original values of the 3 color channelsfor each pixel. Lastly, the 3 color channels used by JPEGs are YCbCr. Therefore,these have to be converted to RGB values before they can be viewed as BMP filesor used for machine learning.935.3.2 PIM-DecoderPIM-decoder is implemented in C using the UPMEM SDK [14] 1 It is based ona previous implementation [4], but with major modifications allowing for a com-pletely asynchronous mode of operation when launching the decoding work on theDPUs and correctly synchronizing the CPU-DPU data transfers.One major shortcoming of the PIM architecture described in chapter 4 is itslack of floating point units. Although performing floating point operations ispossible using DPUs, they are extremely inefficient as they are emulated in soft-ware [98]. Therefore, it is critical that our implementation produces decoded im-ages that are of equal in term of accepted image similarity metrics such as peaksignal-to-noise ration PSNR [9] to those decoded with other methods such as tur-bojpeg or nvjpeg which are official ISO/ITU-T Reference implementations [8, 11].We verified that all of the images used in our experiments are not less than99.97% similar to images decoded with the turbojpeg reference implementationusing the PSNR metric. We also verified that the 0.03% difference does not givedecoding a performance advantage when producing the images. For example, Inour 2MB JPEG image (we describe the details in section 5.4), which has a com-pression quality of 99%, the DPU decoded image is only 0.03% different fromthe turbojpeg decoded image. So, if we have two images (that are encoded at twodifferent qualities) and we decode them with turbojpeg where their similarity isabout 0.03%, then the difference in the running time (if there is any) is how muchturbojpeg is supposedly losing performance from lower quality.We run this experiment to show how significant would the advantage be ifwe assume that there is a performance advantage from decoding an image whichwas encoded at lower quality. When running the test, we encode a raw imageat different JPEG qualities and compare similarities until we get a 0.03% differ-ence between two qualities. Then, we run turbojpeg decoding on the two imagesrecording the running time in each case. We find that if we have image X that isencoded at 96% quality and image Y that is encoded at 97% quality (we encodedthe JPEGs from a bmp file which was decoded by turbojpeg), the correspondingdecoded image similarity is 99.97% (the same 0.03% difference). The turbojpeg1Our PIM-Decoder implementation is publicly available on Github [1].94decoding running time for both images is almost exactly the same. We verified thatfor all images used in our experiments, the difference in running time for decodingto both lower and higher qualities is insignificant ranging from 0%-4%. Anotherparallel work to ours has also verified that images decoded using PIM systems areof similar quality to those decoded using the CPU [84]. We stress that our goalfrom implementing PIM-Decoder is not to beat the CPU at decoding. Rather, weshow that for this specific application, DPUs can decode images at good enoughspeed and quality without using any additional CPUs on the server board as we areonly utilizing the available memory slots.The PIM-Decoder allows decoding of images in different batch sizes. The waywe designed the decoder is to assign every image in the batch to a separate DPU.Therefore, each DPU is responsible for decoding a single image. Furthermore,since each DPU can only work on a 64MB portion of memory, the decoded imagesize cannot exceed this value. For image training datasets, this size is adequate andwe did not run into cases where we needed to decode to a larger image size.5.3.3 TaskletstaskletsTime (s)0.005.0010.0015.0020.001 2 3 4 5 6 7 8 9 10 11 12 13 14Figure 5.5: Time to decode a sample image with varying number of tasklets95Recall from Chapter 4 that tasklets are essentially threads running on PIM. Ourgoal is to use as many tasklets as are necessary to obtain the best performance fromDPUs. Figure 5.5 shows decode time as a function of the number of tasklets, allcooperatively decoding the same image. We run this test for different images withdifferent sizes and the same conclusions are reached. As we use more tasklets tocollaborate on the decoding work, decoding is completed more quickly. We use aminimum of eight tasklets in all of our evaluation experiments as this gives us thebest performance.5.4 MethodologySystem Total cores Memory Memory TDPcapacity bandwidthIntel Xeon 2x8 physical cores 128GB 71.53 GB/s 85 WSilver 4215 CPU (2x16 threads)Nvidia Geforce 64 SMs 8 GB (GPU) 448 GB/s 225 WRTX 2070 GPU (2304 CUDA cores) 16 GB (CPU)w/AMD Ryzen 51600 CPUUPMEM PIM 2556 DPUs 159GB 1.7 TB/s 383 WTable 5.1: CPU, GPU, and PIM systems used in experimentsFor all of our analysis and experiments discussed in section 5.5, we use threesystems shown in table 5.1. We consider two methods as the GPU and CPU base-lines that we compare our PIM-Decoding implementation to. The first method isthe software library from Nvidia called nvjpeg, which is proprietary and the cur-rent state of the art with any CUDA-capable Nvidia GPU [8]. This is the decodingmethod utilized in DALI for performing the GPU-accelerated decoding describedpreviously. We use the Nvidia RTX 2070 described in table 5.1 to obtain the per-formance results for this method. The CPU baseline method that we consider islibjpeg-turbo, which is a highly optimized CPU-only decoding method and is anofficial ISO/ITU-T Reference implementation [11]. We use the Intel Xeon CPUdescribed in the table for obtaining the corresponding CPU performance results.96For PIM-Decoding, we use a currently available commercial PIM system avail-able from UPMEM. The system contains 2,556 DPUs that are organized into 20double-rank DIMMs. Each DIMM contains 128 DPUs. The DPUs are furtherorganized into chips on these DIMMs, where every 8 DPUs are hosted on a sin-gle chip. The UPMEM DIMMs are installed on the same Intel Xeon CPU systemdescribed in table 5.1.Our main goal from the experiments run in section 5.5 is to determine whetheroffloading bulk image decoding to the DPUs is feasible and does not regress per-formance compared to running the same workloads on the CPU or GPU. Thesefindings can help direct future research involving DNN training software and hard-ware architectures. Furthermore, these PIM systems are fairly new in the marketand more commercial products will become available [10]. Presenting the pros andcons through running and analyzing real applications is timely.5.5 Evaluation5.5.1 Bulk Image DecodingIn these experiments, we compare the performance of JPEG decoding when usingthe CPU, GPU, and PIM (DPU 2). Figure 5.6 shows the decoding time as we in-crease the batch size of images. For the CPU baseline, there are three cases thatrange in the number of used CPU cores that are utilized for decoding images. Us-ing a single CPU core is specified by CPU-1, three cores by CPU-3, and six coresby CPU-6. Using multiple CPU cores demonstrate the parallelization capabilitiesof utilizing only the CPU for bulk image decoding. Each JPEG image is 2MB insize. From the figure, we see that for smaller batch sizes, decoding on the DPUsis inferior to other processors. The DPUs are simple processors with significantlylower compute capabilities compared to CPUs and GPUs. If we consider a singleimage, for example, the decoding time on the DPUs is slower than that of the CPUby a factor of 10-30x. It is only as we take advantage of the extreme compute andbandwidth capabilities of the DPUs that we start to gain performance surpassing2For simplicity, we sometimes refer to the whole UPMEM PIM system as DPU. It should be clearfrom context whether we refer to a single DPU or the the system as a whole.97# of imagesTime (s)020406064 128 256 512 1024DPU GPU CPU-1 CPU-3 CPU-6Figure 5.6: Decoding time for increasing number of images. Each JPEG im-age is 2MB in sizethat of GPUs and CPUs.The same trends, but to a lesser effect, are also observed when decoding batchesof smaller image sizes as demonstrated in figures 5.7 and 5.8. The performancegains from DPUs are not as pronounced as those in the larger image case is becauseof the larger copy to compute ratio on DPUs. We discuss in more details how thecost of copying data to and from DPUs dominates for these smaller image-sizecases in the next section.5.5.2 Performance BreakdownFigure 5.9 shows a breakdown of where the time is spent in the DPU for decodingdifferent batch sizes of images, where each image size is 2MB. This breakdowncorresponds to the performance shown previously in figure 5.6. A number of ob-servations can be made from the figure. First, the computation time (the actualdecoding work) on the DPUs remains constant as we decode more images. Sinceeach DPU is working independently on each image, we have perfect scaling with98# of imagesTime (s)010203064 128 256 512 1024DPU GPU CPU-1 CPU-3 CPU-6Figure 5.7: Decoding time for increasing number of images. Each JPEG im-age is 311KB in sizemore DPUs. In other words, the time to perform decoding computations is notaffected by how many DPUs we use and hence we can scale as many images tobe decoded on the available DPUs. Second, as we decode more images, two mainsources of overhead become more significant. The first source of overhead is aresult of having to copy the JPEG images from regular DRAM to the DPUs. Asexplained in chapter 4, data needs to be de-interleaved for it to be readable fromDPU memory. So, this cost of copying includes the time to transfer the images fromregular DRAM and de-interleave them on the DPU memory. The other source ofoverhead is the cost of copying data (decoded JPEGs) from the DPUs to regularDRAM. In the UPMEM machine on which we run our experiments on, there aretwo UPMEM DIMMs per memory channel, so each four ranks share a memorychannel. Therefore, the copying becomes costlier with more images since the totalcopy size will increase and because copying will not be done completely in parallel(since ranks share memory channels).99# of imagesTime (s)01234564 128 256 512 1024DPU GPU CPU-1 CPU-3 CPU-6Figure 5.8: Decoding time for increasing number of images. Each JPEG im-age is 67KB in sizeDPUs / # of imagesTime (ms)02500500075001000064 128 256 512 1024CPU to DPU copyInternal copy to mramsInternal copy to iraminternal copy to wramDPU to CPU copyDecode image on DPUsDPU setupFigure 5.9: Breakdown of PIM-decoding for 2MB images100DPUs / # of imagesTime (ms)020004000600064 128 256 512 1024CPU to DPU copyInternal copy to mramsInternal copy to iraminternal copy to wramDPU to CPU copyDecode image on DPUsDPU setupFigure 5.10: Breakdown of PIM-decoding for 311KB imagesDPUs / # of imagesTime (ms)050010001500200064 128 256 512 1024CPU to DPU copyInternal copy to mramsInternal copy to iraminternal copy to wramDPU to CPU copyDecode image on DPUsDPU setupFigure 5.11: Breakdown of PIM-decoding for 67KB images1015.6 PIM and Emerging ApplicationsPIM systems aim to eliminate long-latency data copying by bringing compute unitscloser to data in memory. However, in their current incarnations, available PIMsystems do not fully mitigate these data movement costs. Therefore, it is neces-sary to examine and analyze applications to decide whether or not PIM would beuseful. Recent work has shown that there are certain application characteristicssuch as parallelizability, vectorizability, data set sizes, and offload costs which willultimately determine whether an application is a good fit for PIM [41]. Further-more, some other solutions have opted for creating specialized in-memory hard-ware specifically designed for only a single application such as personal recom-mendation systems [64, 65]. Although domain-specific solutions can provide muchbetter performance, having general-purpose PIM systems may prove to be the bet-ter option if an increasing number of applications are shown to benefit from beingrun on these systems. The general-purpose PIM-suitability of several applicationswas demonstrated in previous work [55].5.7 Related WorkOur work in this chapter focused on mitigating the cost of decoding, which weidentified to be the most expensive step in preprocessing, by creating an image de-coding implementation on a processing in memory system. Our work was inspiredby the work of Mohan et al. [78] which analyzed both fetch and preprocessing datastalls. They provide a CPU-only solution for preprocessing stalls through the mem-oization of temporarily stored mini batches and coordination of these stored datato reduce the load on CPUs. Therefore, they introduce a solution that increasesthe preprocessing rate by relieving the CPUs from performing excessive prepro-cessing work. Although memoization relieves the CPUs from re-computations, itplaces an extra burden on memory resources. Furthermore, extra logic has to beput in place to determine when to discard old memoized data to make space fornew data. Our work only focuses on the major step of preprocessing, which is de-coding, and provides a PIM solution that completely eliminates the need for CPUsto perform image decoding.Other previous work also demonstrated the bottlenecks that exist in prepro-102cessing for deep learning pipelines [60]. Similar to our conclusion about decoding,they identified that the decoding step in computer vision workloads can create abottleneck that slows down the feeding of data to the GPUs for training. Further-more, previous work also demonstrated that the rate of preprocessing is a criticalfactor which needs to be considered to avoid creating data stalls in the trainingpipelines [105].In addition to further investigating PIM hardware for eliminating preprocess-ing bottlenecks, other approaches implement dedicated accelerators for performingcritical preprocessing stages such as image decoding. Nvidia, for example, has adedicated image decoding hardware on their newer GPUs [2].5.8 SummaryAs Processing in memory architectures continue to proliferate and become a morepromising alternative for eliminating data movement bottlenecks, it becomes im-perative to examine the suitability for applications to adopt this new computingparadigm. In this chapter we examine one such application in details and show thatdecoding a bulk of images using a currently available PIM system can be a viablealternative to performing such a task using CPUs or GPUs. Adding more CPUs toscale a workload’s performance might be difficult due to physical constraints as itmay not be feasible to add more CPUs or boards to the existing servers. However,using memory slots for extra compute with PIM, the compute capabilities of thesystem are expanded and PIM-suitable applications can utilize this along with theexisting CPU and GPU units. However, there are still limitations of current pro-cessing in memory systems that can hinder utilizing them for some applicationsor scenarios. Lighter workloads, such as smaller image batches, may not benefitfrom such an architecture as the cost of copying data to and from the DPUs candominate. Therefore, it is indispensable to examine the behavior of applications todiagnose their suitability to use PIM.103Chapter 6ConclusionComputer hardware is becoming more diverse. Cores of varying types and countsare being integrated on single chips, and system software and higher level softwareframeworks are responsible for realizing the possible performance gains from suchhardware complexity. This trend of increasing core counts and types will continueas computer architects can no longer rely on just increasing clock frequencies ofjust of a handful of cores. From a memory access perspective, data accessed byapplications running on these complex underlying systems will traverse sophisti-cated hierarchies and networks of different levels of caches, interconnects, memorycontrollers, and diverse memory buses. As we saw in this dissertation, software de-signs and implementations for managing memory accesses can have a substantialeffect on the performance of these applications.We answered several questions related to memory management on three typesof modern systems with complex designs. First, we showed that default systemsoftware design choices on NUMA systems lead to poor memory placement caus-ing memory traffic congestion, which may heavily impact performance. We dis-covered that memory traffic congestion is more crucial to optimize than localityand presented a holistic method, Carrefour, for memory placement that providessignificant performance improvement over existing methods.Then, we turned our focus to heterogeneous CPU-GPU systems, which inte-grate a multicore CPU and a GPU on the same chip sharing some of the memoryhierarchy. We showed that to conceal low-level details of memory management,104hardware and software manufacturers sacrifice performance. We showed that sys-tems having a unified memory bus do not necessarily mean that there would beseamless CPU-GPU collaboration as caching and other restrictions would still pre-vent complete integration. We further dissected the two integrated systems show-casing how to best use the available memory management schemes.Finally, we focused on Processing In Memory (PIM) systems as they are thenewest technology to address the Von Neumann bottleneck. We analyzed a cur-rently available PIM system and characterized the performance of image decodingas a case study on these systems. We showed that the perfect compute scalability inthese systems can be adapted to many types of applications. However, the currentPIM system from UPMEM lacks a lot of features to make it a truly Processing inMemory system: it is still mostly an accelerator in the way we offload to it dataand computations. Applications that have mostly all of their data resident in mem-ory would benefit most from current PIM systems and if data needs to be copied alot, we would suffer from the same memory wall problem that we would hope toeliminate from PIM.105Bibliography[1] PIM-Decoder.https://github.com/UBC-ECE-Sasha/PIM-JPEG/tree/DPU XFER ASYNC.→ page 94[2] Loading Data Fast with DALI and the New Hardware JPEG Decoder inNVIDIA A100 GPUs. https://developer.nvidia.com/blog/loading-data-fast-with-dali-and-new-jpeg-decoder-in-a100. → page 103[3] Fast AI Processing with NVIDIA DALI. https://developer.nvidia.com/blog/fast-ai-data-preprocessing-with-nvidia-dali.→ pages 87, 89, 90[4] JPEG Decoder.https://github.com/UBC-ECE-Sasha/PIM-JPEG/tree/daniel. → page 94[5] Huffman coding. https://en.wikipedia.org/wiki/Huffman coding. → page 93[6] INFORMATION TECHNOLOGY DIGITAL COMPRESSION ANDCODING OF CONTINUOUS-TONE STILL IMAGESREQUIREMENTS AND GUIDELINES.https://www.w3.org/Graphics/JPEG/itu-t81.pdf. → page 93[7] JPEG. https://jpeg.org/jpeg/index.html. → page 93[8] nvjpeg. https://developer.nvidia.com/nvjpeg. → pages 94, 96[9] PSNR. https://en.wikipedia.org/wiki/Peak signal-to-noise ratio. → page 94[10] Accelerating AI with smarter memory.https://semiconductor.samsung.com/insights/technology/pim/. → page 97[11] libjpeg-turbo. https://libjpeg-turbo.org. → pages 94, 96106[12] AMD Graphics Core Next (GCN) Architecture.https://www.amd.com/Documents/GCN Architecture whitepaper.pdf,2012. → page 62[13] CL Offline Compiler. https://github.com/HSAFoundation/CLOC, 2017. →page 60[14] 2021. URL https://www.upmem.com/developer/. → pages 77, 94[15] N. Agarwal, D. Nellans, E. Ebrahimi, T. F. Wenisch, J. Danskin, and S. W.Keckler. Selective gpu caches to eliminate cpu-gpu hw cache coherence. InHigh Performance Computer Architecture (HPCA), 2016 IEEEInternational Symposium on, pages 494–506. IEEE, 2016. → page 71[16] S. R. Agrawal, S. Idicula, A. Raghavan, E. Vlachos, V. Govindaraju,V. Varadarajan, C. Balkesen, G. Giannikis, C. Roth, N. Agarwal, and et al.A many-core architecture for in-memory data processing. In Proceedingsof the 50th Annual IEEE/ACM International Symposium onMicroarchitecture, MICRO-50 17, page 245258, New York, NY, USA,2017. Association for Computing Machinery. ISBN 9781450349529.doi:10.1145/3123939.3123985. URLhttps://doi.org/10.1145/3123939.3123985. → pages 73, 74[17] J. Ahn, S. Hong, S. Yoo, O. Mutlu, and K. Choi. A scalableprocessing-in-memory accelerator for parallel graph processing. InProceedings of the 42nd Annual International Symposium on ComputerArchitecture, ISCA 15, page 105117, New York, NY, USA, 2015.Association for Computing Machinery. ISBN 9781450334020.doi:10.1145/2749469.2750386. URLhttps://doi.org/10.1145/2749469.2750386. → pages 73, 74[18] J. Alglave, M. Batty, A. F. Donaldson, G. Gopalakrishnan, J. Ketema,D. Poetzl, T. Sorensen, and J. Wickerson. Gpu concurrency: Weakbehaviours and programming assumptions. ACM SIGARCH ComputerArchitecture News, 43(1):577–591, 2015. → page 47[19] J. Antony, P. P. Janes, and A. P. Rendell. Exploring thread and memoryplacement on numa architectures: Solaris and linux, ultrasparc/fireplaneand opteron/hypertransport. In HiPC, 2006. → pages 8, 37[20] M. Awasthi, D. W. Nellans, K. Sudan, R. Balasubramonian, and A. Davis.Handling the problems and opportunities posed by multiple on-chipmemory controllers. In Proceedings of the 19th International Conference107on Parallel Architectures and Compilation Techniques, PACT ’10, pages319–330, New York, NY, USA, 2010. ACM. ISBN 978-1-4503-0178-7.doi:10.1145/1854273.1854314. URLhttp://doi.acm.org/10.1145/1854273.1854314. → page 38[21] D. H. Bailey, E. Barszcz, J. T. Barton, D. S. Browning, R. L. Carter,L. Dagum, R. A. Fatoohi, P. O. Frederickson, T. A. Lasinski, R. S.Schreiber, H. D. Simon, V. Venkatakrishnan, and S. K. Weeratunga. Thenas parallel benchmarks - summary and preliminary results. InProceedings of the 1991 ACM/IEEE Conference on Supercomputing, SC’91, pages 158–165, New York, NY, USA, 1991. ACM. ISBN0-89791-459-7. doi:10.1145/125826.125925. URLhttp://doi.acm.org/10.1145/125826.125925. → page 23[22] R. Barik, R. Kaleem, D. Majeti, B. T. Lewis, T. Shpeisman, C. Hu, Y. Ni,and A.-R. Adl-Tabatabai. Efficient mapping of irregular c++ applicationsto integrated gpus. In Proceedings of Annual IEEE/ACM InternationalSymposium on Code Generation and Optimization, page 33. ACM, 2014.→ pages 51, 71[23] A. Baumann, P. Barham, P.-E. Dagand, T. Harris, R. Isaacs, S. Peter,T. Roscoe, A. Schu¨pbach, and A. Singhania. The multikernel: A new osarchitecture for scalable multicore systems. In Proceedings of the 22ndACM SIGOPS Symposium on Operating Systems Principles, SOSP ’09,pages 29–44, New York, NY, USA, 2009. ACM. ISBN978-1-60558-752-3. doi:10.1145/1629575.1629579. URLhttp://doi.acm.org/10.1145/1629575.1629579. → page 38[24] J. R. Beveridge, D. Bolme, B. A. Draper, and M. Teixeira. The csu faceidentification evaluation system. Machine Vision and Applications, 16(2):128–138, 2005. doi:10.1007/s00138-004-0144-7. URLhttp://dx.doi.org/10.1007/s00138-004-0144-7. → page 23[25] C. Bienia and K. Li. PARSEC 2.0: A new benchmark suite forchip-multiprocessors. In Proceedings of the 5th Annual Workshop onModeling, Benchmarking and Simulation, June 2009. → page 23[26] S. Blagodurov, S. Zhuravlev, M. Dashti, and A. Fedorova. A case fornuma-aware contention management on multicore systems. In Proceedingsof the 2011 USENIX Annual Technical Conference, ATC ’11, Berkeley,CA, USA, 2011. USENIX Association. → page 38108[27] S. Blagodurov, S. Zhuravlev, M. Dashti, and A. Fedorova. A case fornuma-aware contention management on multicore systems. In Proceedingsof the USENIX Annual Technical Conference, 2011. → page vi[28] W. Bolosky, R. Fitzgerald, and M. Scott. Simple but effective techniquesfor numa memory management. In Proceedings of the 12th ACMSymposium on Operating Systems Principles, SOSP ’89, pages 19–31, NewYork, NY, USA, 1989. ACM. ISBN 0-89791-338-8.doi:10.1145/74850.74854. URL http://doi.acm.org/10.1145/74850.74854.→ pages 8, 36[29] S. Boyd-Wickizer, H. Chen, R. Chen, Y. Mao, F. Kaashoek, R. Morris,A. Pesterev, L. Stein, M. Wu, Y. Dai, Y. Zhang, and Z. Zhang. Corey: Anoperating system for many cores. In Proceedings of the 8th USENIXConference on Operating Systems Design and Implementation, OSDI’08,pages 43–57, Berkeley, CA, USA, 2008. USENIX Association. → pages8, 10[30] S. Boyd-Wickizer, M. F. Kaashoek, R. Morris, and N. Zeldovich. Asoftware approach to unifying multicore caches. Technical ReportMIT-CSAIL-TR-2011-032, MIT, 2011. → pages 38, 39, 40[31] T. Brecht. On the importance of parallel application placement in numamultiprocessors. In USENIX Experiences with Distributed andMultiprocessor Systems, volume 4 of SEDMS ’93, pages 1–18, Berkeley,CA, USA, 1993. USENIX Association. → page 36[32] S. Che, M. Boyer, J. Meng, D. Tarjan, J. W. Sheaffer, S.-H. Lee, andK. Skadron. Rodinia: A benchmark suite for heterogeneous computing. InIEEE International Symposium on Workload Characterization IISWC,2009. → page 49[33] J. Choi, A. Chandramowlishwaran, K. Madduri, and R. Vuduc. A cpu-gpuhybrid implementation and model-driven scheduling of the fast multipolemethod. In Proceedings of Workshop on General Purpose ProcessingUsing GPUs, page 64. ACM, 2014. → pages 51, 71[34] J. Corbet. AutoNUMA: the other approach to NUMA scheduling. LWN,March 2012. → page 22[35] B. Dally. Power, programmability, and granularity: The challenges ofexascale computing. In Proceedings of the 25th International Parallel &Distributed Processing Symposium, IPDPS ’11, pages 878–, Washington,109DC, USA, 2011. IEEE. ISBN 978-0-7695-4385-7.doi:10.1109/IPDPS.2011.420. URLhttp://dx.doi.org/10.1109/IPDPS.2011.420. → page 34[36] M. Dashti and A. Fedorova. Analyzing memory management methods onintegrated cpu-gpu systems. In ACM SIGPLAN International Symposiumon Memory Management (ISMM), pages 59–69. ACM, 2017. → page 41[37] M. Dashti, A. Fedorova, J. Funston, F. Gaud, R. Lachaize, B. Lepers,V. Quema, and M. Roth. Traffic management: a holistic approach tomemory placement on numa systems. In ACM SIGPLAN Notices,volume 48, pages 381–394. ACM, 2013. → page 7[38] M. K. Dashti. Implementation of resource contention management in thelinux kernel for multicore numa systems. Master’s thesis, AppliedSciences: School of Computing Science, 2012. → page vi[39] R. H. Dennard, F. H. Gaensslen, V. L. Rideout, E. Bassous, and A. R.LeBlanc. Design of ion-implanted mosfet’s with very small physicaldimensions. IEEE Journal of Solid-State Circuits, 9(5):256–268, 1974. →page 1[40] F. Devaux. The true processing in memory accelerator. In 2019 IEEE HotChips 31 Symposium (HCS), pages 1–24, Aug 2019.doi:10.1109/HOTCHIPS.2019.8875680. URLhttps://old.hotchips.org/hc31/HC31 1.4 UPMEM.FabriceDevaux.v2 1.pdf.→ page 73[41] A. Devic, S. B. Rai, A. Sivasubramaniam, A. Akel, S. Eilert, and J. Eno. Topim or not for emerging general purpose processing in ddr memorysystems. In Proceedings of the 49th Annual International Symposium onComputer Architecture, pages 231–244, 2022. → page 102[42] A. M. Devices. AMD64 Technology Lightweight Profiling Specification.Sunnyvale, CA, USA, Aug. 2010. → page 35[43] J. Draper, J. Chame, M. Hall, C. Steele, T. Barrett, J. LaCoss, J. Granacki,J. Shin, C. Chen, C. W. Kang, I. Kim, and G. Daglikoca. The architectureof the diva processing-in-memory chip. In Proceedings of the 16thInternational Conference on Supercomputing, ICS 02, page 1425, NewYork, NY, USA, 2002. Association for Computing Machinery. ISBN1581134835. doi:10.1145/514191.514197. URLhttps://doi.org/10.1145/514191.514197. → page 74110[44] P. Drongowski and B. Center. Instruction-based sampling: A newperformance analysis technique for AMD family 10h processors.Sunnyvale, CA, USA, 2007. → page 15[45] H. Esmaeilzadeh, E. Blem, R. St Amant, K. Sankaralingam, and D. Burger.Dark silicon and the end of multicore scaling. In ACM SIGARCH ComputerArchitecture News, volume 39, pages 365–376. ACM, 2011. → page 1[46] M. Ferdman, A. Adileh, O. Kocberber, S. Volos, M. Alisafaee, D. Jevdjic,C. Kaynak, A. D. Popescu, A. Ailamaki, and B. Falsafi. Clearing theclouds: A study of emerging scale-out workloads on modern hardware. InProceedings of the 17th International Conference on Architectural Supportfor Programming Languages and Operating Systems, ASPLOS ’12, pages37–48, New York, NY, USA, 2012. ACM. ISBN 978-1-4503-0759-8.doi:10.1145/2150976.2150982. URLhttp://doi.acm.org/10.1145/2150976.2150982. → pages 39, 40[47] O. Gabbay. AMD HSA kernel driver. https://lwn.net/Articles/605153,2014. → page 44[48] B. Gamsa, O. Krieger, J. Appavoo, and M. Stumm. Tornado: Maximizinglocality and concurrency in a shared memory multiprocessor operatingsystem. In Proceedings of the 3rd Symposium on Operating SystemsDesign and Implementation, OSDI ’99, pages 87–100, Berkeley, CA, USA,1999. USENIX Association. ISBN 1-880446-39-1. → page 38[49] F. Gao, G. Tziantzioulis, and D. Wentzlaff. Computedram: In-memorycompute using off-the-shelf drams. In Proceedings of the 52nd AnnualIEEE/ACM International Symposium on Microarchitecture, MICRO 52,page 100113, New York, NY, USA, 2019. Association for ComputingMachinery. ISBN 9781450369381. doi:10.1145/3352460.3358260. URLhttps://doi.org/10.1145/3352460.3358260. → page 74[50] M. Gao and C. Kozyrakis. Hrl: Efficient and flexible reconfigurable logicfor near-data processing. In 2016 IEEE International Symposium on HighPerformance Computer Architecture (HPCA), pages 126–137, March 2016.doi:10.1109/HPCA.2016.7446059. → page 74[51] S. Ghose, A. Boroumand, J. Kim, J. Gmez-Luna, and O. Mutlu.Processing-in-memory: A workload-driven perspective. IBM Journal ofResearch and Development, 63(6):3:1–3:19, 2019. → page 73111[52] M. Gokhale, B. Holmes, and K. Iobst. Processing in memory: the terasysmassively parallel pim array. Computer, 28(4):23–31, April 1995. ISSN1558-0814. doi:10.1109/2.375174. → page 74[53] J. Gomez-Luna, T. Grass, A. Rico, E. Ayguade, A. J. Pena, et al. Evaluatingthe effect of last-level cache sharing on integrated gpu-cpu systems withheterogeneous applications. In Workload Characterization (IISWC), 2016IEEE International Symposium on, pages 1–10. IEEE, 2016. → page 56[54] J. Go´mez-Luna, I. El Hajjy, L.-W. Changy, V. Garcı´a-Floreszx, S. G.de Gonzaloy, T. B. Jablin, A. J. Pena, and W.-m. Hwu. Chai: collaborativeheterogeneous applications for integrated-architectures. In PerformanceAnalysis of Systems and Software (ISPASS), 2017 IEEE InternationalSymposium on, pages 43–54. IEEE, 2017. → page 1[55] J. Go´mez-Luna, I. E. Hajj, I. Fernandez, C. Giannoula, G. F. Oliveira, andO. Mutlu. Benchmarking a new paradigm: An experimental analysis of areal processing-in-memory architecture. arXiv preprint arXiv:2105.03814,2021. → page 102[56] Google. Snappy - a fast compressor/decompressor. Goole Inc., 2020. URLhttp://google.github.io/snappy/. → page 78[57] J. Hestness, S. W. Keckler, and D. A. Wood. A comparative analysis ofmicroarchitecture effects on cpu and gpu memory system behavior. InWorkload Characterization (IISWC), 2014 IEEE International Symposiumon, pages 150–160. IEEE, 2014. → page 71[58] J. Hestness, S. W. Keckler, and D. A. Wood. Gpu computing pipelineinefficiencies and optimization opportunities in heterogeneous cpu-gpuprocessors. In Workload Characterization (IISWC), 2015 IEEEInternational Symposium on, pages 87–97. IEEE, 2015. → page 71[59] M. D. Hill and M. R. Marty. Amdahl’s law in the multicore era. Computer,41(7), 2008. → page 3[60] A. Isenko, R. Mayer, J. Jedele, and H.-A. Jacobsen. Where is my trainingbottleneck? hidden trade-offs in deep learning preprocessing pipelines.arXiv preprint arXiv:2202.08679, 2022. → pages 87, 103[61] R. Kaleem, R. Barik, T. Shpeisman, B. T. Lewis, C. Hu, and K. Pingali.Adaptive heterogeneous scheduling for integrated gpus. In Proceedings ofthe 23rd international conference on Parallel architectures andcompilation, pages 151–162. ACM, 2014. → pages 51, 71112[62] R. Kaplan, L. Yavits, and R. Ginosar. From processing-in-memory toprocessing-in-storage. In Supercomput. Front. Innov. : Int. J. 4, 2017.doi:10.14529/jsfi170307. URL https://doi.org/10.14529/jsfi170307. →page 74[63] G. Karunaratne, M. Le Gallo, G. Cherubini, L. Benini, A. Rahimi, andA. Sebastian. In-memory hyperdimensional computing. In NatureElectronics, 2020. doi:10.1038/s41928-020-0410-3. URLhttps://doi.org/10.1038/s41928-020-0410-3. → page 78[64] L. Ke, U. Gupta, C.-J. Wu, B. Y. Cho, M. Hempstead, B. Reagen, X. Zhang,D. Brooks, V. Chandra, U. Diril, A. Firoozshahian, K. Hazelwood, B. Jia,H.-H. S. Lee, M. Li, B. Maher, D. Mudigere, M. Naumov, M. Schatz,M. Smelyanskiy, and X. Wang. Recnmp: Accelerating personalizedrecommendation with near-memory processing, 2019. → pages 73, 102[65] L. Ke, X. Zhang, J. So, J.-G. Lee, S.-H. Kang, S. Lee, S. Han, Y. Cho, J. H.Kim, Y. Kwon, et al. Near-memory processing in action: Acceleratingpersonalized recommendation with axdimm. IEEE Micro, 42(01):116–127,2022. → page 102[66] R. Knauerhase, P. Brett, B. Hohlt, T. Li, and S. Hahn. Using osobservations to improve performance in multicore systems. IEEE Micro,28(3):54–66, May 2008. ISSN 0272-1732. doi:10.1109/MM.2008.48. URLhttp://dx.doi.org/10.1109/MM.2008.48. → page 38[67] Y. Kwon, Y. Lee, and M. Rhu. Tensordimm: A practical near-memoryprocessing architecture for embeddings and tensor operations in deeplearning. In Proceedings of the 52nd Annual IEEE/ACM InternationalSymposium on Microarchitecture, MICRO ’52, page 740753, New York,NY, USA, 2019. Association for Computing Machinery. ISBN9781450369381. doi:10.1145/3352460.3358284. URLhttps://doi.org/10.1145/3352460.3358284. → page 73[68] Y.-C. Kwon, S. H. Lee, J. Lee, S.-H. Kwon, J. M. Ryu, J.-P. Son,O. Seongil, H.-S. Yu, H. Lee, S. Y. Kim, Y. Cho, J. G. Kim, J. Choi, H.-S.Shin, J. Kim, B. Phuah, H. Kim, M. J. Song, A. Choi, D. Kim, S. Kim,E.-B. Kim, D. Wang, S. Kang, Y. Ro, S. Seo, J. Song, J. Youn, K. Sohn,and N. S. Kim. 25.4 a 20nm 6gb function-in-memory dram, based on hbm2with a 1.2tflops programmable computing unit using bank-levelparallelism, for machine learning applications. In 2021 IEEE International113Solid- State Circuits Conference (ISSCC), volume 64, pages 350–352,2021. doi:10.1109/ISSCC42613.2021.9365862. → page 74[69] R. Lachaize, B. Lepers, and V. Que´ma. Memprof: A memory profiler fornuma multicore systems. In Proceedings of the 2012 USENIX AnnualTechnical Conference, ATC ’12, pages 53–64, Berkeley, CA, USA, 2012.USENIX Association. → page 19[70] R. P. LaRowe, C. S. Ellis, and M. A. Holliday. Evaluation of numa memorymanagement through modeling and measurements. IEEE Transactions onParallel and Distributed Systems, 3(6):686–701, Nov 1992. ISSN1045-9219. doi:10.1109/71.180624. → page 36[71] D. Lustig, S. Sahasrabuddhe, and O. Giroux. A formal analysis of thenvidia ptx memory consistency model. In Proceedings of theTwenty-Fourth International Conference on Architectural Support forProgramming Languages and Operating Systems, pages 257–270, 2019.→ page 47[72] Z. Majo and T. R. Gross. Memory management in numa multicoresystems: Trapped between cache contention and interconnect overhead. InProceedings of the 2011 International Symposium on MemoryManagement, ISMM ’11, pages 11–20, New York, NY, USA, 2011. ACM.ISBN 978-1-4503-0263-0. doi:10.1145/1993478.1993481. URLhttp://doi.acm.org/10.1145/1993478.1993481. → page 38[73] Z. Majo and T. R. Gross. Memory system performance in a NUMAmulticore multiprocessor. In Proceedings of the 4th Annual InternationalConference on Systems and Storage, SYSTOR ’11, pages 12:1–12:10, NewYork, NY, USA, 2011. ACM. ISBN 978-1-4503-0773-4.doi:10.1145/1987816.1987832. URLhttp://doi.acm.org/10.1145/1987816.1987832. → page 38[74] Y. Mao, R. Morris, and F. Kaashoek. Optimizing mapreduce for multicorearchitectures. Technical report, MIT, 2010. → page 23[75] A. Merkel, J. Stoess, and F. Bellosa. Resource-conscious scheduling forenergy efficiency on multicore processors. In Proceedings of the 5thEuropean Conference on Computer Systems, EuroSys ’10, pages 153–166,New York, NY, USA, 2010. ACM. ISBN 978-1-60558-577-2.doi:10.1145/1755913.1755930. URLhttp://doi.acm.org/10.1145/1755913.1755930. → page 38114[76] D. Merrill, M. Garland, and A. Grimshaw. Scalable gpu graph traversal. InACM SIGPLAN Notices, volume 47, pages 117–128. ACM, 2012. → pages51, 71[77] Z. Metreveli, N. Zeldovich, and M. F. Kaashoek. Cphash: Acache-partitioned hash table. In Proceedings of the 17th ACM SIGPLANSymposium on Principles and Practice of Parallel Programming, PPoPP’12, pages 319–320, New York, NY, USA, 2012. ACM. ISBN978-1-4503-1160-1. doi:10.1145/2145816.2145874. URLhttp://doi.acm.org/10.1145/2145816.2145874. → page 38[78] J. Mohan, A. Phanishayee, A. Raniwala, and V. Chidambaram. Analyzingand mitigating data stalls in dnn training. arXiv preprint arXiv:2007.06775,2020. → pages 4, 87, 88, 89, 90, 102[79] F. Montagna, A. Rahimi, S. Benatti, D. Rossi, and L. Benini. Pulp-hd:Accelerating brain-inspired high-dimensional computing on a parallelultra-low power platform. In 2018 55th ACM/ESDA/IEEE DesignAutomation Conference (DAC), pages 1–6, 2018. → page 78[80] A. Morad, L. Yavits, S. Kvatinsky, and R. Ginosar. Resistive gp-simdprocessing-in-memory. ACM Trans. Archit. Code Optim., 12(4), Jan. 2016.ISSN 1544-3566. doi:10.1145/2845084. URLhttps://doi.org/10.1145/2845084. → page 74[81] S. Mukherjee, Y. Sun, P. Blinzer, A. K. Ziabari, and D. Kaeli. Acomprehensive performance analysis of hsa and opencl 2.0. InPerformance Analysis of Systems and Software (ISPASS), 2016 IEEEInternational Symposium on, pages 183–193. IEEE, 2016. → page 71[82] H. A. D. Nguyen, J. Yu, M. A. Lebdeh, M. Taouil, S. Hamdioui, andF. Catthoor. A classification of memory-centric computing. J. Emerg.Technol. Comput. Syst., 16(2), Jan. 2020. ISSN 1550-4832.doi:10.1145/3365837. URL https://doi.org/10.1145/3365837. → page 74[83] J. Nider, C. Mustard, A. Zoltan, J. Ramsden, L. Liu, J. Grossbard,M. Dashti, R. Jodin, A. Ghiti, J. Chauzi, et al. A case study of{Processing-in-Memory} in {off-the-Shelf} systems. In 2021 USENIXAnnual Technical Conference (USENIX ATC 21), pages 117–130, 2021. →page 73115[84] J. Nider, J. Dagger, N. Gharavi, and D. Ng. Bulk jpeg decoding onin-memory processors. In International Systems and Storage Conference,2022. → page 95[85] I. Pandis, R. Johnson, N. Hardavellas, and A. Ailamaki. Data-orientedtransaction execution. Proc. VLDB Endow., 3(1-2):928–939, Sept. 2010.ISSN 2150-8097. doi:10.14778/1920841.1920959. URLhttp://dx.doi.org/10.14778/1920841.1920959. → pages 38, 39[86] A. Pesterev, N. Zeldovich, and R. T. Morris. Locating cache performancebottlenecks using data profiling. In Proceedings of the 5th EuropeanConference on Computer Systems, EuroSys ’10, pages 335–348, New York,NY, USA, 2010. ACM. ISBN 978-1-60558-577-2.doi:10.1145/1755913.1755947. URLhttp://doi.acm.org/10.1145/1755913.1755947. → page 19[87] G. L. Peterson. Myths about the mutual exclusion problem. InformationProcessing Letters, 12(3):115–116, 1981. → page 57[88] A. Rahimi, S. Benatti, P. Kanerva, L. Benini, and J. M. Rabaey.Hyperdimensional biosignal processing: A case study for emg-based handgesture recognition. In 2016 IEEE International Conference on RebootingComputing (ICRC), pages 1–8, 2016. → page 78[89] O. Russakovsky, J. Deng, H. Su, J. Krause, S. Satheesh, S. Ma, Z. Huang,A. Karpathy, A. Khosla, M. Bernstein, et al. Imagenet large scale visualrecognition challenge. International journal of computer vision, 115(3):211–252, 2015. → page 89[90] T.-I. Salomie, I. E. Subasu, J. Giceva, and G. Alonso. Database engines onmulticores, why parallelize when you can distribute? In Proceedings of the6th European Conference on Computer Systems, EuroSys ’11, pages17–30, New York, NY, USA, 2011. ACM. ISBN 978-1-4503-0634-8.doi:10.1145/1966445.1966448. URLhttp://doi.acm.org/10.1145/1966445.1966448. → pages 38, 39[91] L. T. Schermerhorn. Automatic page migration for linux. In linux.conf.au,2007. → page 37[92] P. Siegl, R. Buchty, and M. Berekovic. Data-centric computing frontiers: Asurvey on processing-in-memory. In Proceedings of the SecondInternational Symposium on Memory Systems, MEMSYS 16, page 295308,New York, NY, USA, 2016. Association for Computing Machinery. ISBN1169781450343053. doi:10.1145/2989081.2989087. URLhttps://doi.org/10.1145/2989081.2989087. → page 74[93] P. Singh, C.-R. M, P. Raghavendra, T. Tye, D. Das, and A. N. Platformcoherency and soc verification challenges. In 13th InternationalSystem-on-Chip (SoC) Conference, Exhibit & Workshops, 2013. → pagesxiv, 1, 61[94] X. Song, H. Chen, R. Chen, Y. Wang, and B. Zang. A case for scalingapplications to many-core with os clustering. In Proceedings of the 6thEuropean Conference on Computer Systems, EuroSys ’11, pages 61–76,New York, NY, USA, 2011. ACM. ISBN 978-1-4503-0634-8.doi:10.1145/1966445.1966452. URLhttp://doi.acm.org/10.1145/1966445.1966452. → page 38[95] H. S. Stone. A logic-in-memory computer. IEEE Transactions onComputers, C-19(1):73–78, Jan 1970. ISSN 2326-3814.doi:10.1109/TC.1970.5008902. → page 74[96] Y. Sun, X. Gong, A. K. Ziabari, L. Yu, X. Li, S. Mukherjee, C. Mccardwell,A. Villegas, and D. Kaeli. Hetero-mark, a benchmark suite for cpu-gpucollaborative computing. In Workload Characterization (IISWC), 2016IEEE International Symposium on, pages 1–10. IEEE, 2016. → pages51, 71[97] D. Tam, R. Azimi, and M. Stumm. Thread clustering: Sharing-awarescheduling on SMP-CMP-SMT multiprocessors. In Proceedings of the 2ndACM SIGOPS European Conference on Computer Systems, EuroSys ’07,pages 47–58, New York, NY, USA, 2007. ACM. ISBN978-1-59593-636-3. doi:10.1145/1272996.1273004. URLhttp://doi.acm.org/10.1145/1272996.1273004. → page 37[98] F. D. (UPMEM). the true processing-in-memory accelerator. → page 94[99] B. Verghese, S. Devine, A. Gupta, and M. Rosenblum. Operating systemsupport for improving data locality on cc-numa compute servers. InProceedings of the 7th International Conference on Architectural Supportfor Programming Languages and Operating Systems, ASPLOS ’96, pages279–289, New York, NY, USA, 1996. ACM. ISBN 0-89791-767-7.doi:10.1145/237090.237205. URLhttp://doi.acm.org/10.1145/237090.237205. → pages 8, 36117[100] J. Vesely, A. Basu, M. Oskin, G. Loh, and A. Bhattacharjee. Observationsand opportunities in architecting shared virtual memory for heterogenoussystems. In IEEE International Symposium on Performance Analysis ofSystems and Software (ISPASS). IEEE, 2016. → pages 64, 65[101] T. Vijayaraghavan, Y. Eckert, G. H. Loh, M. J. Schulte, M. Ignatowski,B. M. Beckmann, W. C. Brantley, J. L. Greathouse, W. Huang,A. Karunanithi, et al. Design and analysis of an apu for exascalecomputing. In HPCA, 2017. → page 44[102] K. Wang, K. Angstadt, C. Bo, N. Brunelle, E. Sadredini, T. Tracy,J. Wadden, M. Stan, and K. Skadron. An overview of micron’s automataprocessor. In 2016 International Conference on Hardware/SoftwareCodesign and System Synthesis (CODES+ISSS), pages 1–3, Oct 2016. →page 74[103] W. Weber and A. Gupta. Exploring the benefits of multiple hardwarecontexts in a multiprocessor architecture: Preliminary results. In The 16thAnnual International Symposium on Computer Architecture, pages273–280, May 1989. doi:10.1109/ISCA.1989.714562. → page 75[104] D. Zhang, N. Jayasena, A. Lyashevsky, J. L. Greathouse, L. Xu, andM. Ignatowski. Top-pim: Throughput-oriented programmable processingin memory. In Proceedings of the 23rd International Symposium onHigh-Performance Parallel and Distributed Computing, HPDC 14, page8598, New York, NY, USA, 2014. Association for Computing Machinery.ISBN 9781450327497. doi:10.1145/2600212.2600213. URLhttps://doi.org/10.1145/2600212.2600213. → page 74[105] M. Zhao, N. Agarwal, A. Basant, B. Gedik, S. Pan, M. Ozdal,R. Komuravelli, J. Pan, T. Bao, H. Lu, et al. Understanding andco-designing the data ingestion pipeline for industry-scale recsys training.arXiv preprint arXiv:2108.09373, 2021. → pages 4, 87, 90, 103[106] J. Zhou and B. Demsky. Memory management for many-core processorswith software configurable locality policies. In Proceedings of the 2012International Symposium on Memory Management, ISMM ’12, pages3–14, New York, NY, USA, 2012. ACM. ISBN 978-1-4503-1350-6.doi:10.1145/2258996.2259000. URLhttp://doi.acm.org/10.1145/2258996.2259000. → page 37[107] S. Zhuravlev, S. Blagodurov, and A. Fedorova. Addressing shared resourcecontention in multicore processors via scheduling. In Proceedings of the11815th International Conference on Architectural Support for ProgrammingLanguages and Operating Systems, ASPLOS ’10, pages 129–142, NewYork, NY, USA, 2010. ACM. ISBN 978-1-60558-839-1.doi:10.1145/1736020.1736036. URLhttp://doi.acm.org/10.1145/1736020.1736036. → page 38119
Thesis/Dissertation
eng
Electrical and Computer Engineering
Vancouver : University of British Columbia Library
Attribution-NonCommercial-NoDerivatives 4.0 International
Graduate
Fedorova, Alexandra
Seltzer, Margo
Towards understanding and mitigating memory-access challenges in computing systems
Text