Cache 一致性问题
单核 Cache 中每个 Cache line 有2个标志:dirty 和 valid 标志,它们很好的描述了 Cache 和 Memory 之间的数据关系(数据是否有效,数据是否被修改),而在多核处理器中,多个核会共享一些数据。
只有 Core 0 访问变量 x,它的 Cache line 数据和内存中的数据一致,数据只存在于本 Cache 中。
3个 Core 都访问变量 x,它们对应的 Cache line 数据和内存中的数据一致,数据存在于很多 Cache 中。
Core 0 修改了x的值之后,这个 Cache line 数据被修改了,和内存中的数据不一致,数据只存在于本 Cache 中。其他 Core 对应的 Cache line 数据无效。
内存栅栏
简单来说,内存栅栏(Memory Barrier)就是从本地或工作内存到主存之间的拷贝动作。
仅当写操作线程先跨越内存栅栏而读线程后跨越内存栅栏的情况下,写操作线程所做的变更才对其他线程可见。关键字 synchronized 和 volatile 都强制规定了所有的变更必须全局可见,该特性有助于跨越内存边界动作的发生,无论是有意为之还是无心插柳。
在程序运行过程中,所有的变更会先在寄存器或本地 cache 中完成,然后才会被拷贝到主存以跨越内存栅栏。此种跨越序列或顺序称为 happens-before。
写操作必须要 happens-before 读操作,即写线程需要在所有读线程跨越内存栅栏之前完成自己的跨越动作,其所做的变更才能对其他线程可见。
内存对齐
为何要内存对齐?
1. 平台原因(移植原因):不是所有的硬件平台都能访问任意地址上的任意数据的,某些硬件平台只能在某些地址处取某些特定类型的数据,否则抛出硬件异常。
2. 性能原因:经过内存对齐后,CPU的内存访问速度大大提升。
这是普通程序员心目中的内存印象,由一个个的字节组成,而 CPU 并不是这么看待的。
CPU 把内存当成是一块一块的,块的大小可以是2,4,8,16字节大小,因此 CPU 在读取内存时是一块一块进行读取的。块大小成为 memory accessgranularity(粒度)本人把它翻译为"内存读取粒度"。
假设CPU要读取一个int型4字节大小的数据到寄存器中,分两种情况讨论:
1. 数据从0字节开始
2. 数据从1字节开始
再次假设内存读取粒度为4。
当该数据是从0字节开始时,CPU 只需读取内存一次即可把这4字节的数据完全读取到寄存器中。
当该数据是从1字节开始时,问题变的有些复杂,此时该 int 型数据不是位于内存读取边界上,这就是一类内存未对齐的数据。
此时 CPU 先访问一次内存,读取0—3字节的数据进寄存器,并再次读取4—5字节的数据进寄存器,接着把0字节和6,7,8字节的数据剔除,最后合并1,2,3,4字节的数据进寄存器。对一个内存未对齐的数据进行了这么多额外的操作,大大降低了 CPU 性能。
这还属于乐观情况了,上文提到内存对齐的作用之一为平台的移植原因,因为以上操作只有有部分 CPU 肯干,其他一部分 CPU 遇到未对齐边界就直接*罢*工*了。
相关推荐
5. **同步与内存栅栏**: 为了确保线程间的正确同步,有时需要使用内存栅栏(Memory Barrier)来确保所有线程在继续执行之前完成特定的操作。 在"cuda_使用cuda并行加速实现之LinearAttention"这个压缩包文件中,...
例如,利用线程间的数据共享,避免全局内存的随机访问,以及优化内存对齐等。 8. **CUDA C++的扩展**:CUDA还提供了如动态并行ism(Dynamic Parallelism)、流和事件(Stream and Event)、纹理和表面内存等高级...
详细讨论了CUDA的内存模型,包括内存对齐、内存复制、内存优化策略等。 **第6章:并行计算** 介绍并行计算的高级主题,如线程同步、原子操作、栅栏同步等。 **第7章:设备接口** 这部分描述了如何使用CUDA运行时...
7. **优化技巧**:为了充分发挥GPU的并行计算能力,程序员需要关注数据对齐、内存访问模式、计算效率等因素。例如,通过减少全局内存访问、充分利用共享内存、优化线程配置等方式来提升程序性能。 通过"helloGPGPU-...
5. **内存管理策略**:包括如何分配连续的物理内存供DMA使用,如何高效地回收和重用内存,以及如何处理内存对齐要求。 6. **同步机制(Syncronization)**:比如使用 fences(栅栏)来确保图像处理和传输的顺序,...
书中会讲解如何使用不同的内存类型,并讨论内存对齐和数据传输的最佳实践。 此外,CUDA编程还涉及到线程同步和原子操作。由于GPU执行的是数据并行任务,线程间的同步是确保正确性的关键。书中会通过实例展示如何...
`bandwidthTest` 是一个用于测量内存带宽的工具,它可以测试主机内存与显存之间以及显存内部的数据传输速率,这对于评估硬件性能和优化代码的内存访问模式至关重要。`bicubicTexture` 和 `convolution*` 系列程序...
7. **性能优化**:为了提高性能,这个无锁队列可能还使用了其他技巧,如避免忙等待(busy-waiting)、使用内存栅栏(memory barriers)来确保数据的同步,以及可能的缓存对齐策略来减少缓存冲突。 8. **使用场景**...
该命令要求内存地址对齐,因此选项C是正确的,表示对齐向量存双精度浮点数。 3. CUDA子矩阵数组变量声明: CUDA是NVIDIA开发的一种用于GPU编程的并行计算平台。在优化矩阵乘法的CUDA程序中,为了提升性能,通常会...
- **同步机制**:例如栅栏(barrier)和事件(event),用于协调不同线程组之间的执行顺序。 - **计算效率**:通过优化内存访问、减少计算冗余和有效使用并行度来提高GPU程序的性能。 此外,书中可能还会涵盖: - ...
- **内存对齐**: 保证数据在内存中的位置以优化性能。 5. **并行计算与同步** - **线程块内并行**: 使用`__syncthreads()`同步线程块内的所有线程。 - **流(Streams)**: 多个并行任务可以通过流异步执行,提高...
- 对齐(Alignment)讨论了如何控制对象在内存中的对齐方式。 4. 标准转换(Standard conversions) - 本节讲解了C++中的隐式类型转换,包括lvalue-to-rvalue转换、数组到指针的转换、函数到指针的转换等。 5. ...
- **原因**:某些显示卡为了节省内存空间,采用颜色查找表来表示颜色。 - **原理**:将颜色值映射到一个较小的索引值上,然后通过索引值在颜色查找表中查找实际的颜色值。 **2.12 视频显示器分辨率:** - **x方向...