最新要闻
- 成都燃气:2022年度净利润约4.92亿元 同比增加0.57%
- 环球速讯:一批“色狼”被官方公布,罚了!
- 6299元起!魅族20 INFINITY无界版入网:双向卫星通信比苹果更强
- 朱一龙、倪妮亮相北影节开幕红毯:新电影《消失的她》即将上映
- 青岛一鸵鸟在车流中狂奔引骚乱:当天就被抓回
- 冰淇淋事件闹大 有车主扣掉车标“宝马”二字 称感觉羞愧
- F-150纯电皮卡充电时起火 三辆新车惨烈烧毁!福特:韩国电池的锅
- 初三英语听力训练技巧_初三英语听力训练 全球滚动
- 联动长三角,走进马鞍山!浦东历史博物馆首次实现展览走出去
- 每一磅碎牛肉使用多少预制炸玉米饼调味料?-天天短讯
- 山西一妈妈为催婚给25岁儿子床头摆稻草人:神回应看醉
- 天天简讯:沃尔沃新车轮胎5个螺丝少4个 女车主崩溃要退车
- 女子吐槽软卧车厢分配3男1女 12306回应:可找工作人员调换
- 目前适配最完美的小折叠!vivo X Flip评测:超好看超能打|世界观焦点
- 环球快资讯丨三星S23 Ultra用户入手小米13 Ultra:拿它当相机来用
- 贵州大方:樱桃产销两旺,果农笑迎丰收_焦点热文
手机
iphone11大小尺寸是多少?苹果iPhone11和iPhone13的区别是什么?
警方通报辅警执法直播中被撞飞:犯罪嫌疑人已投案
- iphone11大小尺寸是多少?苹果iPhone11和iPhone13的区别是什么?
- 警方通报辅警执法直播中被撞飞:犯罪嫌疑人已投案
- 男子被关545天申国赔:获赔18万多 驳回精神抚慰金
- 3天内26名本土感染者,辽宁确诊人数已超安徽
- 广西柳州一男子因纠纷杀害三人后自首
- 洱海坠机4名机组人员被批准为烈士 数千干部群众悼念
家电
Nvidia GPU Virtual Memory Management
(资料图片)
1 常用显存管理
1.1 CUDA Runtime API
在CUDA编程中,编程人员通常会使用以下CUDART同步API进行显存申请和释放,比如调用cudaMalloc并传入所需的显存size即可返回显存的虚拟地址,使用完成后可调用cudaFree进行释放。
__host__ __device__ cudaError_t cudaMalloc(void **devPtr, size_t size);__host__ cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);__host__ cudaError_t cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height);__host__ __device__ cudaError_t cudaFree(void *devPtr);
此外,CUDART也提供了显存申请和释放的异步API供编程人员使用,只需传入CUDA Stream即可调用,返回给编程人员的也是显存的虚拟地址。
__host__ cudaError_t cudaMallocAsync(void **devPtr, size_t size, cudaStream_t hStream);__host__ cudaError_t cudaFreeAsync(void *devPtr, cudaStream_t hStream);
1.2 CUDA Driver API
Nvidia在用户态CUDA Driver提供了一套API用于显存申请和释放,返回的结果与CUDART API没有区别,只是在使用层面与CUDART API有区别,比如调用cuMemAlloc之前需要编程人员手动使用CUDA Driver API进行初始化(cuInit)和创建Context(cuCtxCreate),而对于调用cudaMalloc来说,这些都是隐式完成,对编程人员是透明的。CUresult cuMemAlloc(CUdeviceptr* dptr, size_t bytesize);CUresult cuMemAllocManaged(CUdeviceptr* dptr, size_t bytesize, unsigned int flags);CUresult cuMemAllocPitch(CUdeviceptr* dptr, size_t* pPitch, size_t WidthInBytes, size_t Height, unsigned int ElementSizeBytes);CUresult cuMemFree(CUdeviceptr dptr);
2 Virtual Memory Management
2.1 特性
就常用显存管理API来说,由于编程人员只能获取到显存的虚拟地址,如果有动态调整显存大小的需求(比如GPU上vector扩容),用户必须显示地申请更大的一块显存,并从原始显存中复制数据到新显存,再释放原始显存,然后继续跟踪新分配的显存地址,这样的操作通常会导致应用程序的性能降低和较高的显存带宽峰值利用率。
在CUDA 10.2中引入VMM API为应用程序提供了一种直接管理统一虚拟地址空间的方法,可以将显存的虚拟地址和物理地址解耦,允许编程人员分别处理它们。VMM API允许编程人员在合适的时候将显存的虚拟地址与物理地址进行映射和解映射。借助VMM API可以更好地解决动态调整显存大小的需求,只需要申请额外的物理地址,再与原始虚拟地址扩展的空间进行映射,既不需要更换追踪的显存地址,也不需要将数据从原始显存拷贝到新显存。因此,VMM API能够帮助编程人员构建更高效的动态数据结构,并更好地控制应用程序中的显存使用。参考Introducing Low-Level GPU Virtual Memory Management。
2.2 API
VMM API主要包含显存粒度获取API、虚拟地址管理API、物理地址管理API、映射管理API以及访存管理API。// Calculates either the minimal or recommended granularity.CUresult cuMemGetAllocationGranularity(size_t* granularity, const CUmemAllocationProp* prop, CUmemAllocationGranularity_flags option);// Allocate an address range reservation.CUresult cuMemAddressReserve(CUdeviceptr* ptr, size_t size, size_t alignment, CUdeviceptr addr, unsigned long long flags);// Free an address range reservation.CUresult cuMemAddressFree(CUdeviceptr ptr, size_t size);// Create a CUDA memory handle representing a memory allocation of a given size described by the given properties.CUresult cuMemCreate(CUmemGenericAllocationHandle* handle, size_t size, const CUmemAllocationProp* prop, unsigned long long flags);// Release a memory handle representing a memory allocation which was previously allocated through cuMemCreate.CUresult cuMemRelease(CUmemGenericAllocationHandle handle);// Retrieve the contents of the property structure defining properties for this handle.CUresult cuMemGetAllocationPropertiesFromHandle(CUmemAllocationProp* prop, CUmemGenericAllocationHandle handle);// Maps an allocation handle to a reserved virtual address range.CUresult cuMemMap(CUdeviceptr ptr, size_t size, size_t offset, CUmemGenericAllocationHandle handle, unsigned long long flags);// Unmap the backing memory of a given address range.CUresult cuMemUnmap(CUdeviceptr ptr, size_t size);// Get the access flags set for the given location and ptr.CUresult cuMemGetAccess(unsigned long long* flags, const CUmemLocation* location, CUdeviceptr ptr);// Set the access flags for each location specified in desc for the given virtual address range.CUresult cuMemSetAccess(CUdeviceptr ptr, size_t size, const CUmemAccessDesc* desc, size_t count);
3 使用
参考cuda sample,给出使用VMM API进行显存申请和释放的示例代码。
3.1 显存申请
显存申请主要包括获取显存粒度、申请虚拟地址、申请物理地址、虚拟地址与物理地址映射、释放物理地址handle(注意此处并不会真正释放物理地址)和设置访问权限几个步骤。
cudaError_t vmm_alloc(void **ptr, size_t size) { CUmemAllocationProp prop = {}; memset(prop, 0, sizeof(CUmemAllocationProp)); prop->type = CU_MEM_ALLOCATION_TYPE_PINNED; prop->location.type = CU_MEM_LOCATION_TYPE_DEVICE; prop->location.id = currentDevice; size_t granularity = 0; if (cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } size = ((size - 1) / granularity + 1) * granularity; CUdeviceptr dptr; if (cuMemAddressReserve(&dptr, size, 0, 0, 0) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } CUmemGenericAllocationHandle allocationHandle; if (cuMemCreate(&allocationHandle, size, &prop, 0) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } if (cuMemMap(dptr, size, 0, allocationHandle, 0) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } if (cuMemRelease(allocationHandle) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } CUmemAccessDesc accessDescriptor; accessDescriptor.location.id = prop.location.id; accessDescriptor.location.type = prop.location.type; accessDescriptor.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; if (cuMemSetAccess(dptr, size, &accessDescriptor, 1) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } *ptr = (void *)dptr; return cudaSuccess;}
3.2 显存释放
显存释放主要包括获取显存粒度、虚拟地址与物理地址解映射(注意此处解映射之后物理地址随即释放)和释放虚拟地址几个步骤。
cudaError_t vmm_free(void *ptr, size_t size) { if (!ptr) { return cudaSuccess; } CUmemAllocationProp prop = {}; memset(prop, 0, sizeof(CUmemAllocationProp)); prop->type = CU_MEM_ALLOCATION_TYPE_PINNED; prop->location.type = CU_MEM_LOCATION_TYPE_DEVICE; prop->location.id = currentDevice; size_t granularity = 0; if (cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM) != CUDA_SUCCESS) { return cudaErrorMemoryAllocation; } size = ((size - 1) / granularity + 1) * granularity; if (cuMemUnmap((CUdeviceptr)ptr, size) != CUDA_SUCCESS || cuMemAddressFree((CUdeviceptr)ptr, size) != CUDA_SUCCESS) { return cudaErrorInvalidValue; } return cudaSuccess;}
4 问题
4.1 P2P访问
使用CUDART实现设备的对等访问可以直接调用cudaDeviceEnablePeerAccess API设置,而使用VMM实现设备的对等访问需要调用cuMemSetAccess API设置显存的访问权限。
4.2 带宽
笔者曾经做过一个项目,期间对比测试过VMM和cuMemAlloc申请的显存在H2D、D2H和D2D带宽上的差异(Tesla V100,CUDA 10.2,CUDA Driver 470.80,主机内存为普通内存或pinned memory),发现VMM的带宽略低于cuMemAlloc,尝试过并行优化、异步优化和小包优化,效果都不明显。百思不得其解后向Nvidia反馈,其美研工程师排查后表示是CUDA Driver内部的一个bug,发过来修复版本后测试两者带宽无明显差异。
关键词:
-
003CCE Turbo配置容器网卡动态预热|焦点速看
更新时间:2023-04-13GMT+08:00Reference:& 160;CCETurbo配置容器网卡动态预热_云容器引擎CCE_最佳实践_网
来源: -
Nvidia GPU Virtual Memory Management
1常用显存管理1 1CUDARuntimeAPI在CUDA编程中,编程人员通常会使用以下CUDART同步API进行显存申请和释放,
来源: 003CCE Turbo配置容器网卡动态预热|焦点速看
Nvidia GPU Virtual Memory Management
成都燃气:2022年度净利润约4.92亿元 同比增加0.57%
焦点精选!债市日报:4月21日
环球速讯:一批“色狼”被官方公布,罚了!
6299元起!魅族20 INFINITY无界版入网:双向卫星通信比苹果更强
朱一龙、倪妮亮相北影节开幕红毯:新电影《消失的她》即将上映
青岛一鸵鸟在车流中狂奔引骚乱:当天就被抓回
冰淇淋事件闹大 有车主扣掉车标“宝马”二字 称感觉羞愧
F-150纯电皮卡充电时起火 三辆新车惨烈烧毁!福特:韩国电池的锅
初三英语听力训练技巧_初三英语听力训练 全球滚动
联动长三角,走进马鞍山!浦东历史博物馆首次实现展览走出去
每一磅碎牛肉使用多少预制炸玉米饼调味料?-天天短讯
劝人写码,千刀万剐——“前端已死”难道要成真了?|世界快报
【新华解读】外资3月持有人民币债券规模回升 外资有望享受更多政策红利 播资讯
山西一妈妈为催婚给25岁儿子床头摆稻草人:神回应看醉
天天简讯:沃尔沃新车轮胎5个螺丝少4个 女车主崩溃要退车
女子吐槽软卧车厢分配3男1女 12306回应:可找工作人员调换
目前适配最完美的小折叠!vivo X Flip评测:超好看超能打|世界观焦点
环球快资讯丨三星S23 Ultra用户入手小米13 Ultra:拿它当相机来用
贵州大方:樱桃产销两旺,果农笑迎丰收_焦点热文
看热讯:电脑如何改成光盘启动项-(电脑如何改成光盘启动项设置)
全球即时看!深圳二手房房贷与参考价解绑?深圳多家银行:目前根据评估价放贷
MySQL Execution Plan--DISTINCT语句优化
Python函数与码复用
Appuploader安装指南
全志Uboot fdt修改DTS进行临时调试的方法
热门看点:ReactNative 打包发布 Android 应用
世界实时:BBBA德系“四兄弟”组合解散!宝沃中国申请破产
夏季必备:GLM夏季薄款七分/九分/长裤34.9元新低抄底
鸿星尔克请网友吃冰淇淋 获赞大气!宝马MINI感受下 环球消息
半年内5名机车网红车祸身亡 又见两只脚够不着地女骑:大货身旁疾驰
刘慈欣在联合国被“催更”:当年奥巴马连发两封邮件要新书 环球短讯
它强任它强,驾控随我心,斯巴鲁Crosstrek全新上市!
全球速看:打好软件国产化攻坚战,闪信科技面向人工智能和数字经济进行新一代升级
【快播报】DRF的权限组件(源码分析)
物联网常见协议之Amqp协议及使用场景解析 当前报道
每日报道:本轮系列赛前两战,詹姆斯面对布鲁克斯防守打出14中7
【财经分析】聚焦可持续转型升级 绿色债券领域蓬勃发展
【财经分析】供需趋松压力增大 铁矿石大幅下挫-全球观察
不同档位DLSS对画面帧数的影响究竟有多大?看完秒懂|天天通讯
老人无牌无证驾驶老头乐从河南跑到江苏 还要周游全国
10分钟补能400km 理想首批4C超级充电站上线运营:女生也能拎得动-天天日报
腾讯起诉《王者荣耀》代练赢了:构成不正当竞争 获赔60万
马云被聘为香港大学荣誉教授:聘期三年|天天速看料
全球即时看!项目播报 | 璞华×江苏鼎为云,打造数字化“采云链”体系,赋能新能源机电行业生态化
环球关注:Excel的列数如何用数字表示?
WCF教程_编程入门自学教程_菜鸟教程-免费教程分享 每日速看
当前资讯!启动“盟圈”,开拓聚才引才“新阵地”
热文:推广央行数字货币,牙买加财政部出台激励计划
中国人民解放军陆军第九四七医院与咸阳职业技术学院签订校企合作协议 当前聚焦
每日快报!女子骑摩托车遇沉降路段摔倒遭碾压身亡 竟担主责:官方重申不服可上诉
世界百事通!致敬《流浪地球2》国内首个类ChatGPT模型MOSS正式开源
画面过度拟真被质疑造假:《Unrecord》官方发布引擎录像自证清白_百事通
华为回应孟晚舟要去美国出差传言:纯属造谣
传小鹏P7接亲车队开辅助驾驶连环追尾 小鹏回应:拿普通事故抹黑
贵州:一片茶叶走出的乡村振兴路-速递
环球看点!发现有检查“黑车”司机扔下车跑了 留下一车乘客发呆
【GPT开发】人人都能用ChatGPT4.0做Avatar虚拟人直播
浅谈日出日落的计算方法以及替代工具 - 日出日落 API-精彩看点
这款产品,竟然用了二千多年才爆火?|世界最资讯
当前焦点!比亚迪李云飞:比亚迪今年要卖300万台车
德国玩家竟如此钟爱AMD显卡!销量超过NVIDIA几乎20% 每日速讯
青岛国家深远海绿色养殖试验区养殖鱼类首次成功度夏-天天快讯
环球观点:尚无特效药!国产带状疱疹疫苗全国接种启动:一针能管10年
半价手慢无:卜珂巧克力夹心蛋卷14.9元/3罐 速递
电影情节成真!加拿大机场1.6吨黄金被盗 价值超亿美元_世界快讯
上演“史诗级灾难公关”:宝马蒸发市值可买约5亿个冰淇淋
可以再等等!iQOO 12系列在路上:200W超快闪充-天天新资讯
香菜饼最简单的做法,只需一把香菜,2个鸡蛋!
焦点热门:AIGC的阿克琉斯之踵
超过50多个热门的免费可用 API 分享|滚动
火山引擎 DataTester 3 大功能升级:聚焦敏捷、智能与易用,帮助企业降本增效
天天热讯:三峡艺术高中分数线(三峡艺术高中)
新华指数|钢“财”说:库存延续下行,基本面仍待改善|今日要闻
环球热点!长峰医院火灾系医院内部施工引燃可燃涂料所致
榴莲批发价1斤低至20元:海南国产榴莲也快上市了
2023新款超轻超弹:匹克轻弹风逸跑鞋119元狂促(原价249元)
《惊奇队长2》创纪录!62万踩成漫威最讨人厌的预告
男子被贴膏药秒倒地 官方通报:涉事膏药无迷药成分
驾驶舱可变大床!五菱龙卡本月上市:横竖都能拉
天语e66怎么开机?天语e66手机可以拍视频吗?
华为T8951什么时候上市的?华为T8951手机参数
中兴v987是什么型号?中兴V987手机参数
七月未央作者_七月未央
三星S3850如何恢复出厂设置?三星S3850手机参数
摩托罗拉MT620上市价格是多少?摩托罗拉MT620手机参数
C#写一套最全的SQL server帮助类(包括增删改查)
秦都区吴办吴家堡社区新时代文明实践站开展党建引领 “童”行同乐活动
热点!110万的比亚迪大把人买!仰望U8预售48小时:订单已超1.3万份
NVIDIA大失所望:RTX 4070卖崩了!AMD成市场赢家
天天资讯:用户自驾突发紧急情况!最后用华为P60 Pro获救:用一次续一生
五一假期十大热门城市:看看有你想去的吗?
升空后人工引爆!下次再来 成功三分之一的Starship 新动态
长安C401将被称为长安辛特克斯动力来自马自达1.6 天天微资讯
JS中的进程和线程-动态
linux下查看文件内容工具发布啦!
04 设置工作模式与环境(下)收集信息
天天速讯:全国首单绿色及能源保供双主题ABS在深交所完成发行
环球速讯:园企智推官“我为企业代言”推介大赛开启