In this sgemm case, the performance impact of using managed memory is negligible. If your CUDA Fortran program is bogged down with two copies of variables, like variable_name for the host, and variable_name_d for the device, then managed memory might be the answer for you. The data movement still happens, but rather than being explicit, now it is controlled by the unified memory management system behind the scenes, much like the way an OS manages virtual memory. Real, allocatable, dimension(:,:) :: goldĬall sgemmNN_16x16>(A, B, C, m, N, k, alpha, beta)Īs you can see, we no longer need two copies of the arrays A, B, and C. Real, managed, allocatable, dimension(:,:) :: A, B, C Now, here's the same code using managed data: Istat = cudaEventElapsedTime(time, start, stop) Real, allocatable, device, dimension(:,:) :: dA, dB, dCĬall sgemm_cpu(A, B, gold, m, N, k, alpha, beta)Ĭall sgemmNN_16x16>(dA, dB, dC, m, N, k, alpha, beta) Real, allocatable, dimension(:,:) :: A, B, C, gold Here's the code before CUDA managed data: Here's an example of how we've used it to simplify the host code in our sgemm test, a CUDA Fortran example we've used frequently here at PGI: In CUDA Fortran, we've added the managed keyword to the language, which can be used in host code similarly to the device keyword. Specifically, in CUDA C/C++, managed memory usage consists of replacing calls to cudaMalloc() with cudaMallocManaged(), and removing explicit uses of cudaMemcpy() to transfer data between host and device. The system (driver and OS) controls the physical page movement. Unlike zero-copy memory, managed memory is not pinned and static, but migrates between the host and device on access. Unified Memory provides a way to obtain a single pointer to user allocated data that can be used in both host and device code. Using managed memory simplifies many coding tasks, makes source code cleaner, and enables a unified view of complicated data structures across host and device memories.Ī good explanation of Unified Memory can be found on the NVIDIA Parallel Forall blog, and Appendix J of the CUDA C Programming Guide gives a detail explanation.
0 Comments
Leave a Reply. |
AuthorWrite something about yourself. No need to be fancy, just an overview. ArchivesCategories |