内存分配的初步解释

关于协调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时,cudaMallocmalloc慢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相关推荐

  1. Unified Memory与unified memory managed详解

    统一寻址(Unified Memory): 可直接访问CPU内存.GPU显存,不需要手动拷贝数据. CUDA 6在现有的内存池结构上增加了一个统一内存系统,程序员可以直接访问任何内存/显存资源,或者在 ...

  2. Spark 内存管理 spark.executor.memory /spark.memory.fraction/spark.memory.offHeap.size【堆外内存/内存管理】 钨丝计划

    spark1.6及之后: 堆内内存: spark.executor.memory 包含 spark.memory.fraction: spark.memory.fraction 包含 spark.me ...

  3. CUDA系列学习(四)Parallel Task类型 与 Memory Allocation

    本文为CUDA系列学习第四讲,首先介绍了Parallel communication patterns的几种形式(map, gather, scatter, stencil, transpose), ...

  4. 杭电ACM-LCY算法进阶培训班-专题训练15

    杭电ACM-LCY算法进阶培训班-专题训练(03-07-11-15) 1012 最短路 #pragma GCC optimize(2) #pragma GCC optimize(3,"Ofa ...

  5. 杭电ACM-LCY算法进阶培训班-专题训练(矩阵快速幂)

    杭电ACM-LCY算法进阶培训班-专题训练(矩阵快速幂)[模板] 传送门 杭电ACM-LCY算法进阶培训班-专题训练(矩阵快速幂)[模板] 矩阵快速幂模板 Count Problem Descript ...

  6. 2022年超全的Android面经(附含面试题|进阶资料)

    前言 前言 面试应该是双向的. 很幸运有接受过很多安卓开发者的面试经历.在整个过程中,我学到了很多东西,也了解了包括腾讯和阿里在内的顶级公司的面试. 在整个过程中,我收集了来自我的朋友.互联网和其他来 ...

  7. 杭电ACM-LCY算法进阶培训班-专题训练(计算几何入门)

    杭电ACM-LCY算法进阶培训班-专题训练(计算几何入门) 传送门 杭电ACM-LCY算法进阶培训班-专题训练(计算几何入门) Shape of HDU Problem Description Inp ...

  8. 杭电ACM-LCY算法进阶培训班-专题训练(KMP)

    杭电ACM-LCY算法进阶培训班-专题训练(KMP) 杭电ACM-LCY算法进阶培训班-专题训练(KMP) 剪花布条 Problem Description Input Output Sample I ...

  9. 2021-06-18杭电ACM-LCY算法进阶培训班-专题训练16

    杭电ACM-LCY算法进阶培训班-专题训练(04-08-12-16) 1009 Intervals #pragma GCC optimize(2) #pragma GCC optimize(3,&qu ...

  10. 杭电ACM-LCY算法进阶培训班-专题训练09

    杭电ACM-LCY算法进阶培训班-专题训练09 1014 剪花布条 #pragma GCC optimize(2) #pragma GCC optimize(3,"Ofast",& ...

最新文章

  1. html5 audio标志改变音量,HTML5之Audio(二)—— processor调节音量
  2. jquery ajax java上传文件_jQuery Ajax方式上传文件的方法
  3. 技术图文:如何在CSDN上写自己的技术Blog?
  4. DataTable转换成IList
  5. ToolBarManager可任意选择你想要的菜单
  6. Delphi调用REST
  7. 针对于高频低频图像的理解
  8. Java中的引用类型(强引用、弱引用)和垃圾回收
  9. android 源码分析notification,# Notification 源码分析
  10. Oracle 执行长SQL
  11. 网站移植到linux上后常犯的错误
  12. 计算机一级excel典型试题,最新excel计算机一级试题合集
  13. 花店管理系统php网站,网上花店管理系统下载_网上花店管理系统官方下载-太平洋下载中心...
  14. 幸运九宫格抽奖系统带后台源码
  15. Android 获取手机充电状态
  16. html中的href属性_href(HTML属性)
  17. 怎样查网站的排名和收录情况
  18. 中国哪里的大米最好吃?这几个地区最出名,网友为此吵翻了
  19. MobileNetv2-SSDlite训练自己的数据集(一)——配置安装caffe-ssd
  20. C语言+EasyX实现数字雨

热门文章

  1. nginx resolver失败
  2. 利用高德地图 API 显示地图信息
  3. 《纽约客》:还原真实的扎克伯格
  4. 前端面试系列-JavaScript作用域和作用域链
  5. ArcGIS Server 发布服务失败
  6. 2022华为机试真题 C++ 实现【数大雁】
  7. 香港服务器要个人信息么,香港个人信息应当遵循服务器23.225合法
  8. 如何将livp文件转换为jpeg图片格式
  9. 移动硬盘无法读取是怎么回事?
  10. python实现决策树 西瓜书_西瓜书学习笔记-决策树