内存分配的初步解释

关于协调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. python函数返回元组平均数_关于python:使用函数中的单个项返回元组
  2. 日本“AI 鱼脸识别”项目,每分钟识别 100 条
  3. 消息队列的四大典型使用场景
  4. 一个Eclipse代码显示主题
  5. RabbitMQ,RabbitMQ 的工作模式,Spring 整合 RabbitMQ,Springboot 整合RabbitMQ
  6. 麒麟linux创建用户组,麒麟Linux系统用户和组管理指南(21页)-原创力文档
  7. python数据结构的应用场景不包括,Python 数据结构学习
  8. jenkins执行bat失败_关于批处理文件:即使在BAT脚本中成功执行了ROBOCOPY命令,JENKINS作业也会失败...
  9. oracle立即关闭数据库,Oracle数据库的起步和关闭
  10. android订阅管理,RXJAVA取消订阅封装-kotlin-Android
  11. Spring JDBC事务支持类jdbcTemplate(了解)
  12. [控件] 将字符串转换成贝塞尔曲线并执行动画
  13. Python 100 例 练习实例1
  14. html实现pdf导出excel表格,纯前端文件导出-Excel/PDF
  15. 跳转系统设置token
  16. vue-admin websocket接收消息推送+语音提示(详细代码)
  17. 邮箱个性签名html模板,邮件个性签名大全_邮件的经典个性签名模板
  18. R语言主成分回归(PCR)、 多元线性回归特征降维分析光谱数据和汽车油耗、性能数据...
  19. AIX报错(errpt)处理小记
  20. 音频之耳机(Exynos7872)

热门文章

  1. 计算机系统xp和w7,对比分析老电脑装xp还是win7纯净版好
  2. Linux #! /bin/sh的意思
  3. SEO是什么?前端如何进行SEO优化
  4. 解决“远程会话已断开连接,因为访问被拒绝导致许可证存储的创建失败,请使用提升的权限运行远程桌面客户端”问题...
  5. 深度学习手记(八)之PTB实现LSTM模型
  6. 使用Arduino和TTP223电容式触摸传感器制作触摸检测器
  7. 触摸式开关简单入门教程
  8. 在手机与计算机之间进行文件传输的方式,电脑与手机快速传输文件的方法
  9. Priest John's Busiest Day (2-sat)
  10. 推荐免费下载华软源码430套大型企业管理源码,下载地址:http://www.hur.cn/tg/linkin.asp?linkid=205389 源码语言:PB/Delphi/VB/Java/.Ne