内存分配的初步解释

关于协调CPU和GPU之间的内存创建和分配以及传输

CPU分配内存

CPU分配内存主要有两种方式：

通过C标准库中的malloc函数完成
调用CUDA中的cudaMallocHost函数

cudaMallocHost函数通过页面锁定，可以提供更高的CPU和GPU传输速率，吞吐量增加是Barracuda10的2.4倍，Barracuda04的2.0倍，Barracuda01的1.5倍。但是缺点是使用cudaMallocHost分配内存比malloc更加慢，
每次调用cudaMallocHost分配1M的内存需要2300微妙左右，当需要分配512M，时间上升到61毫秒，比malloc在分配时慢了3-5个量级

GPU分配内存

GPU的分配方式有：

通过CUDA的cudaMalloc函数
cudaMallocPitch函数
cudaMallocArray函数

cudaMallocPitch 和 cudaMallocArray都没有cudaMalloc来的快。cudaMalloc的分配时间如图1所示，256bytes的调用分配耗时1微妙，2KB到4KB分配耗时有显著提升，接下来在512KB分配大约需要50微妙，大于512K时，耗时显著提升，当分配512M时，需要12.5毫秒左右。总而言之，在小于4M时，cudaMalloc分配速度比malloc慢1.5个量级，大于4MB时，cudaMalloc比malloc慢2-4个量级。

例子来自https://github.com/CoffeeBeforeArch/cuda_programming

https://www.youtube.com/watch?v=LGhduZNudDY&list=PLxNPSjHT5qvu4Q2UElj3HUCh2lpSooQWo&index=4&frags=wn

中的vectorAdd_pinned.cu

vectorAdd_um_prefetch.cu的对比

权威解释来自

《CUDA C++ Programming Guide》

Data Prefetching
Data prefetching means migrating data to a processor’s memory and mapping it in
that processor’s page tables before the processor begins accessing that data. The intent
of data prefetching is to avoid faults while also establishing data locality. This is most
valuable for applications that access data primarily from a single processor at any given
time. As the accessing processor changes during the lifetime of the application, the data
can be prefetched accordingly to follow the execution flow of the application. Since work
is launched in streams in CUDA, it is expected of data prefetching to also be a streamed
operation as shown in the following API:
cudaError_t cudaMemPrefetchAsync(const void *devPtr,
size_t count,
int dstDevice,
cudaStream_t stream);
where the memory region specified by devPtr pointer and count number of bytes, with
ptr rounded down to the nearest page boundary and count rounded up to the nearest
page boundary, is migrated to the dstDevice by enqueueing a migration operation in
stream. Passing in cudaCpuDeviceId for dstDevice will cause data to be migrated to
CPU memory.
Consider a simple code example below:
void foo(cudaStream_t s) {
char *data;
cudaMallocManaged(&data, N);
init_data(data, N); // execute on CPU
cudaMemPrefetchAsync(data, N, myGpuId, s); // prefetch to GPU
mykernel<<<..., s>>>(data, N, 1, compare); // execute on GPU
cudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s); // prefetch to CPU
cudaStreamSynchronize(s);
use_data(data, N);
cudaFree(data);
}
Without performance hints the kernel mykernel will fault on first access to data
which creates additional overhead of the fault processing and generally slows down
the application. By prefetching data in advance it is possible to avoid page faults and
achieve better performance

This API follows stream ordering semantics, i.e. the migration does not begin until all
prior operations in the stream have completed, and any subsequent operation in the
stream does not begin until the migration has completed.

翻译

数据预取：
数据预取意味着将数据迁移到处理器的内存，并在处理器开始访问该数据之前将其映射到该处理器的页表中。数据预取的目的是避免错误，同时建立数据局部性。这对于在任何给定时间主要从单个处理器访问数据的应用程序都是最有价值的。随着访问处理器在应用程序的生命周期内发生变化，可以相应地预取数据以跟踪应用程序的执行流程。由于工作是在CUDA中的流中启动的，因此预计数据预取也是流式操作，如以下API所示：

cudaError_t cudaMemPrefetchAsync(const void *devPtr,size_t count,int dstDevice,cudaStream_t stream);

其中由devPtr指针指定的内存区域和计数字节数（ptr向下舍入到最接近的页边界并向上舍入到最接近的页边界的计数）将通过在流中排列迁移操作而迁移到dstDevice。传递给dstDevice的cudaCpuDeviceId将导致数据迁移到CPU内存。
考虑下面的简单代码示例：

void foo(cudaStream_t s) {char *data;cudaMallocManaged(&data, N);init_data(data, N); // execute on CPUcudaMemPrefetchAsync(data, N, myGpuId, s); // prefetch to GPUmykernel << <..., s >> >(data, N, 1, compare); // execute on GPUcudaMemPrefetchAsync(data, N, cudaCpuDeviceId, s); // prefetch to CPUcudaStreamSynchronize(s);use_data(data, N);cudaFree(data);
}

如果没有性能提示，内核mykernel会在首次访问数据时出错，这会造成额外的故障处理开销，并且通常会降低应用程序的速度。通过预先预取数据，可以避免页面错误并实现更好的性能。
此API遵循流排序语义，即直到流中的所有先前操作都完成后才开始迁移，并且在迁移完成之前，流中的任何后续操作都不会开始。
数据使用提示：
当多个处理器需要同时访问相同的数据时，数据预取本身是不够的。在这种情况下，应用程序提供有关如何实际使用数据的提示很有用。以下咨询API可用于指定数据使用情况：

参考链接：

https://devblogs.nvidia.com/maximizing-unified-memory-performance-cuda/

https://zhuanlan.zhihu.com/p/82651065

https://satisfie.github.io/2016/06/06/Caffe%E8%A7%A3%E8%AF%BB1-Pinned-Memory-Vs-Non-Pinned-Memory/

CUDA进阶资料专题（一）pinned memory 和 unified memory相关推荐

Unified Memory与unified memory managed详解
统一寻址(Unified Memory): 可直接访问CPU内存.GPU显存,不需要手动拷贝数据. CUDA 6在现有的内存池结构上增加了一个统一内存系统,程序员可以直接访问任何内存/显存资源,或者在 ...
Spark 内存管理 spark.executor.memory /spark.memory.fraction/spark.memory.offHeap.size【堆外内存/内存管理】钨丝计划
spark1.6及之后: 堆内内存: spark.executor.memory 包含 spark.memory.fraction: spark.memory.fraction 包含 spark.me ...
CUDA系列学习（四）Parallel Task类型与 Memory Allocation
本文为CUDA系列学习第四讲,首先介绍了Parallel communication patterns的几种形式(map, gather, scatter, stencil, transpose), ...
杭电ACM-LCY算法进阶培训班-专题训练15
杭电ACM-LCY算法进阶培训班-专题训练(03-07-11-15) 1012 最短路 #pragma GCC optimize(2) #pragma GCC optimize(3,"Ofa ...
杭电ACM-LCY算法进阶培训班-专题训练（矩阵快速幂）
杭电ACM-LCY算法进阶培训班-专题训练(矩阵快速幂)[模板] 传送门杭电ACM-LCY算法进阶培训班-专题训练(矩阵快速幂)[模板] 矩阵快速幂模板 Count Problem Descript ...
2022年超全的Android面经（附含面试题|进阶资料）
前言前言面试应该是双向的. 很幸运有接受过很多安卓开发者的面试经历.在整个过程中,我学到了很多东西,也了解了包括腾讯和阿里在内的顶级公司的面试. 在整个过程中,我收集了来自我的朋友.互联网和其他来 ...
杭电ACM-LCY算法进阶培训班-专题训练（计算几何入门）
杭电ACM-LCY算法进阶培训班-专题训练(计算几何入门) 传送门杭电ACM-LCY算法进阶培训班-专题训练(计算几何入门) Shape of HDU Problem Description Inp ...
杭电ACM-LCY算法进阶培训班-专题训练（KMP）
杭电ACM-LCY算法进阶培训班-专题训练(KMP) 杭电ACM-LCY算法进阶培训班-专题训练(KMP) 剪花布条 Problem Description Input Output Sample I ...
2021-06-18杭电ACM-LCY算法进阶培训班-专题训练16
杭电ACM-LCY算法进阶培训班-专题训练(04-08-12-16) 1009 Intervals #pragma GCC optimize(2) #pragma GCC optimize(3,&qu ...
杭电ACM-LCY算法进阶培训班-专题训练09
杭电ACM-LCY算法进阶培训班-专题训练09 1014 剪花布条 #pragma GCC optimize(2) #pragma GCC optimize(3,"Ofast",& ...

CUDA进阶资料专题（一）pinned memory 和 unified memory

内存分配的初步解释

CPU分配内存

GPU分配内存

CUDA进阶资料专题（一）pinned memory 和 unified memory相关推荐

最新文章

热门文章