缓存分块是一种内存优化技术,主要用于提高数据的局部性(Locality),以减少缓存未命中(Cache Miss)的次数。在现代计算机体系结构中,处理器(CPU)的速度通常比内存快得多。因此,如果CPU在处理数据时需要频繁地等待数据从内存中加载,就会大大降低程序的执行效率。Cache Tiled技术通过将数据分割成较小的块(Tiles),并确保这些小块能够完全装入CPU的高速缓存(Cache),来减少这种等待时间。
CUDA编程中,用于优化内存访问模式,以减少全局内存(DRAM)访问次数并提高内存带宽的利用率。它的核心思想是将数据分成小块(称为“tiles”或“blocks”),这样每个块可以完全加载到共享内存中。共享内存是一种CUDA核心内的高速缓存内存,其访问速度比全局内存快得多。
基本原理
见啥使用DRAM,也就是全局内存。转而多用L1 Cache。缓存分块是有的时候数据太多了,每次只能加载一部分。
- 减少内存延迟:通过将数据加载到共享内存中,可以减少对全局内存的访问次数,从而减少延迟。
- 提高内存带宽利用率:将数据划分为小块后,可以更有效地利用内存带宽。
- 协同工作:多个线程可以协作加载一个Tile,然后从共享内存中高效读取数据。
实现步骤
- 定义Tile的大小:确定目标内存以及GPU的共享内存大小。计算index用于加载到共享内存。
- 加载数据到共享内存:在CUDA核心中,多个线程协作将全局内存中的数据加载到共享内存。
- 同步线程:确保所有数据都加载到共享内存后,再进行处理。
- 处理数据:从共享内存读取数据,进行计算。
- 将结果写回全局内存:如果需要,将处理后的数据写回到全局内存。
Coding
TILE_WIDTH
是一个预定义的常量,它定义了Tile的大小。
__syncthreads()
是一个同步原语,用于确保一个线程块内的所有线程都达到这一点后才能继续执行。这在使用共享内存时尤其重要,因为它确保在所有线程开始读取共享内存中的数据之前,所有的写入操作都已完成。
#define TILE_WIDTH 16*16*4 // b c bit 定义每个Tile的宽度
// CUDA核心函数,用于矩阵乘法
__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) {
__shared__ float Mds[TILE_WIDTH