", launched the fundamentals of CUDA programming by showing how to jot down a easy program that allotted two arrays of numbers in memory accessible to the GPU and then added them collectively on the GPU. To do this, I launched you to Unified Memory, which makes it very easy to allocate and entry knowledge that may be utilized by code running on any processor within the system, CPU or GPU. I finished that put up with a few easy "exercises", one among which encouraged you to run on a recent Pascal-based GPU to see what occurs. I was hoping that readers would strive it and comment on the results, and a few of you probably did! I recommended this for 2 reasons. First, because Pascal GPUs such because the NVIDIA Titan X and the NVIDIA Tesla P100 are the primary GPUs to include the Web page Migration Engine, which is hardware help for Unified Memory web page faulting and migration.

The second reason is that it supplies a fantastic opportunity to study more about Unified Memory. Fast GPU, Fast Memory… Right! However let’s see. First, I’ll reprint the results of running on two NVIDIA Kepler GPUs (one in my laptop and one in a server). Now let’s attempt working on a extremely quick Tesla P100 accelerator, based mostly on the Pascal GP100 GPU. Hmmmm, that’s below 6 GB/s: slower than running on my laptop’s Kepler-primarily based GeForce GPU. Don’t be discouraged, though; we can fix this. To grasp how, I’ll must let you know a bit more about Unified Memory. What is Unified Memory? Unified Memory is a single memory address house accessible from any processor in a system (see Figure 1). This hardware/software program know-how allows applications to allocate knowledge that can be read or written from code working on both CPUs or GPUs. Allocating Unified Memory is so simple as changing calls to malloc() or new with calls to cudaMallocManaged(), an allocation function that returns a pointer accessible from any processor (ptr in the following).

When code working on a CPU or GPU accesses information allocated this way (typically known as CUDA managed information), the CUDA system software and/or the hardware takes care of migrating memory pages to the memory of the accessing processor. The vital level here is that the Pascal GPU architecture is the primary with hardware support for virtual memory page faulting and migration, by way of its Web page Migration Engine. Older GPUs primarily based on the Kepler and Maxwell architectures additionally assist a more limited form of Unified Memory. What Occurs on Kepler When i call cudaMallocManaged()? On methods with pre-Pascal GPUs like the Tesla K80, calling cudaMallocManaged() allocates size bytes of managed memory on the GPU device that's energetic when the decision is made1. Internally, the driver additionally units up web page desk entries for all pages covered by the allocation, so that the system is aware of that the pages are resident on that GPU. So, in our instance, operating on a Tesla K80 GPU (Kepler architecture), x and y are each initially absolutely resident in GPU memory.

Then in the loop beginning on line 6, the CPU steps by way of both arrays, initializing their elements to 1.0f and 2.0f, respectively. Because the pages are initially resident in system memory, a page fault occurs on the CPU for every array web page to which it writes, and the GPU driver migrates the web page from gadget memory to CPU memory. After the loop, all pages of the 2 arrays are resident in CPU memory. After initializing the information on the CPU, the program launches the add() kernel so as to add the elements of x to the weather of y. On pre-Pascal GPUs, upon launching a kernel, the CUDA runtime must migrate all pages beforehand migrated to host memory or to another GPU back to the gadget memory of the gadget operating the kernel2. Since these older GPUs can’t page fault, all information should be resident on the GPU simply in case the kernel accesses it (even when it won’t).