【并行计算-CUDA开发】有关CUDA当中global memory如何实现合并访问跟内存对齐相关的问题
2023-09-27 14:20:27 时间
ps:这是英伟达二面面的一道相关CUDA的题目。《NVIDIA CUDA编程指南》第57页开始
在合并访问这里,不要跟shared memory的bank conflict搞混淆了,这里很重要。
global memory没有被缓存(面试答错了!),因此,使用正确的存取模式来获得最大的内存带宽,更为重要,尤其是如何存取昂贵的设备内存device
memory。
首先,设备device有能力,在一个单一指令下,从global memory中读取32-bit, 64-bit, 128-bit字进入寄存器register。
分配如下:
__device__ type device[32];
type data = device[tid];
编译一个单一加载指令,type必须是sizeof(type) = 4、8 或者 16 这样的,
而且要求内置类型像float2 或 float4 一样的自动完成的。
对于结构、大小和队列要求可以通过编译器强制使用队列指定的__align__(8) 或 __align__(16),
例如
struct __align(8)__{
float a;
float b;
};
//或者是
struct __align(16)__{ float a; float b;
float c; };
对于结构 > 16字节的,编译器生成几个加载指令,来保证它生成最低数量的指令,这样的结构应该用__align__(16)定义
例如:
struct __align(16)__{ float a; float b;
float c
float d
float e;
};
//!被编译成两个128-bit的加载指令,而不是5个32-bit加载指令
其次,全局内存地址同时被每线程的一个half-warp访问(执行读和写指令)时,应该排列好,以便内存的存取可以结合进入一个接近单一的,排列好的内存存取。
它意味着,在每一个half-warp中,在half-warp中的第N个线程应该访问该地址。
HalfWarpBaseAddress + N
这里,HalfWarpBaseAddress 是类型 type* ,而 type 应该符合之前讨论过的大小和队列要求。
HalfWarpBaseAddress应该排列成16 * sizeof ( type ) 字节,例如16 * sizeof ( type ) 的倍数。
任何一个驻留在全局内存BaseAddress的变量的地址,或者一个来自 D.5 或 E.8 部分内存分配规则返回的地址,问题被排列成至少256个字节,以此来满足内存队列的约束,
HalfWarpBaseAddress - BaseAddress 应该是 16 * sizeof ( type ) 的倍数
注意:如果一个half-warp满足了上面的所有需求,那么每线程的内存访问被联合了,即使half-warp的一些线程实际上没有访问内存。
建议:对于整个warp满足这个需求,而不是分开的,半个半个的。
因为未来的设备将默认为必要的要求。
示例:联合的 64-bit 访问会比联合的 32-bit 访问内存带宽低一点,
联合的128-bit访问会比联合的32-bit访问内存带宽要低很多。
一个公共的全局内存访问样式是,当每个带有线程ID tid 的线程访问位于一个数组的一个元素时,元素的地址位于类型 type* 的 BaseAddress,使用以下地址
BaseAddress + tid
为了获得内存的联合,type 必须符合之前讨论过的大小和队列的要求。
如果 type 的结构 > 16字节,它应该被分成几个满足要求的结构,并且数据应该在内存中被划分成关于这些结构的几个数组,而不是一个类型 type* 的单一数组。
另一个公共的全局内存访问样式是,
当带有索引 (tx, ty) 的每条线程访问地址位于类型 type* 的BaseAddress 和宽度 width 的2D数组的一个元素使用以下地址:
BaseAddress + width * ty + tx
在这样的情况下,获得 half-warp 的所有块线程的内存结合,只有当
a. 块线程的宽度是一半warp大小的倍数
b. width 是16 的倍数
特别是,这意味着,宽度不是16的倍数的数组将被更高效地访问,如果它实际上分配是宽度被传入到最接近16的倍数而且,它的列因此被填充了。
cudaMallocPitch() 和 cuMemAllocPitch() 函数和在D.5和E.8 部分描述的相关内存拷贝的函数,使开发人员能够编写非硬件独立的代码,来分配遵循这些约束的数组。
相关文章
- 【优化】Java开发中注意内存问题,影响JVM
- 用实例带你深入理解Java内存模型
- 【VS开发】Visual C++内存泄露检测—VLD工具使用说明
- 【CUDA开发】CUDA面内存拷贝用法总结
- 【ARM-Linux开发】Linux内存管理:ARM Memory Layout以及mmu配置
- 【VS开发】【C/C++开发】传递双重指针申请内存,典型用法
- 【操作系统】【C/C++开发】内存管理
- 【Qt开发】将内存图像数据封装成QImage V2
- 【VS开发】内存泄漏相关问题
- 【VS开发】内存映射文件进程间共享内存
- 【VS开发】malloc申请内存错误分析
- 【C++】内存管理、底层
- 哪个更快:Java堆还是本地内存
- Stegoloader--可隐藏在PNG图片直接进驻内存的木马
- AE开发使用内存图层
- OpenSSH内存耗尽漏洞CVE-2016-8858 将会消耗服务器37.5G内存
- Android内存泄漏的检测流程、捕捉以及分析
- iOS开发篇-内存管理(上)
- JVM之对象创建流程及对象内存布局
- 统计本机内存情况和使用状况的脚本
- 《挑战30天C++入门极限》新手入门:C++中堆内存(heap)的概念和操作方法
- Oracle 内存参数调优设置
- 【Java 虚拟机原理】垃圾回收算法 ( Java 虚拟机内存分区 | 垃圾回收机制 | 引用计数器算法 | 引用计数循环引用弊端 )
- 对开发中常见的内存泄露,GDI泄露进行检测
- ODA开发/CAD二次开发/C#开发-- “System.AccessViolationException尝试读取或写入受保护的内存,这通常指示其他内存已损坏”异常快速排除思路与建议(踩坑血泪史)
- Linux 段页式内存管理
- 【汇编语言/底层开发】3、通过寄存器进行内存访问