【cuda】四、基础概念：Cache Tiled 缓存分块技术

缓存分块是一种内存优化技术，主要用于提高数据的局部性（Locality），以减少缓存未命中（Cache Miss）的次数。在现代计算机体系结构中，处理器（CPU）的速度通常比内存快得多。因此，如果CPU在处理数据时需要频繁地等待数据从内存中加载，就会大大降低程序的执行效率。Cache Tiled技术通过将数据分割成较小的块（Tiles），并确保这些小块能够完全装入CPU的高速缓存（Cache），来减少这种等待时间。

CUDA编程中，用于优化内存访问模式，以减少全局内存(DRAM)访问次数并提高内存带宽的利用率。它的核心思想是将数据分成小块（称为“tiles”或“blocks”），这样每个块可以完全加载到共享内存中。共享内存是一种CUDA核心内的高速缓存内存，其访问速度比全局内存快得多。

基本原理

见啥使用DRAM，也就是全局内存。转而多用L1 Cache。缓存分块是有的时候数据太多了，每次只能加载一部分。

• 减少内存延迟：通过将数据加载到共享内存中，可以减少对全局内存的访问次数，从而减少延迟。
• 提高内存带宽利用率：将数据划分为小块后，可以更有效地利用内存带宽。
• 协同工作：多个线程可以协作加载一个Tile，然后从共享内存中高效读取数据。

实现步骤

1. 定义Tile的大小：确定目标内存以及GPU的共享内存大小。计算index用于加载到共享内存。
2. 加载数据到共享内存：在CUDA核心中，多个线程协作将全局内存中的数据加载到共享内存。
3. 同步线程：确保所有数据都加载到共享内存后，再进行处理。
4. 处理数据：从共享内存读取数据，进行计算。
5. 将结果写回全局内存：如果需要，将处理后的数据写回到全局内存。

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][TILE_WIDTH]; // 定义共享内存，用于存储Md的一个Tile
    __shared__ float Nds[TILE_WIDTH][TILE_WIDTH]; // 定义共享内存，用于存储Nd的一个Tile

    int bx = blockIdx.x;  // 获取当前块的x坐标
    int by = blockIdx.y;  // 获取当前块的y坐标
    int tx = threadIdx.x; // 获取当前线程在块中的x坐标
    int ty = threadIdx.y; // 获取当前线程在块中的y坐标

    // 计算Pd矩阵中的行号和列号
    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;

    float Pvalue = 0; // 初始化计算值

    // 遍历Md和Nd矩阵的Tile，计算Pd矩阵的元素
    for (int m = 0; m < Width/TILE_WIDTH; ++m) {
        // 协作加载Md和Nd的Tile到共享内存
        Mds[ty][tx] = Md[Row*Width + (m*TILE_WIDTH + tx)];
        Nds[ty][tx] = Nd[Col + (m*TILE_WIDTH + ty)*Width];
        __syncthreads(); // 确保所有线程都加载完毕

        // 计算Tile内的乘积并累加到Pvalue
        for (int k = 0; k < TILE_WIDTH; ++k) {
            Pvalue += Mds[ty][k] * Nds[k][tx];
        }
        __syncthreads(); // 确保所有线程都计算完毕
    }
    // 将计算结果写入Pd矩阵
    Pd[Row*Width + Col] = Pvalue;
}

在这个示例中，MatrixMulKernel 是用于矩阵乘法的CUDA核心。它使用了两个共享内存数组Mds和Nds来存储两个输入矩阵的Tile。每个线程块处理输出矩阵Pd的一个Tile。线程块中的每个线程共同工作，加载输入矩阵的相应部分到共享内存，然后使用这些数据来计算输出矩阵的一个元素。
__syncthreads() 出现在两个关键位置:

1. 加载数据到共享内存之后：这里的 __syncthreads() 确保了所有线程都完成了对共享内存的写入操作。即使这个写入操作是在 for 循环中完成的，我们也需要确保每个线程都完成了当前迭代的加载操作，才能安全地开始使用这些共享内存中的数据进行计算。
2. 计算Tile内的乘积并累加到Pvalue之后：第二个 __syncthreads() 确保了所有线程都完成了当前Tile的计算。在开始处理下一个Tile之前，这是必要的，因为下一个Tile的计算可能依赖于共享内存中的新数据。

在这两种情况下，__syncthreads() 的作用是确保所有线程在继续执行之前都达到同一点。

对比原始矩阵乘法的代码：

__global__ void MatrixMulSimple(float* A, float* B, float* C, int Width) {
    int Row = blockIdx.y * blockDim.y + threadIdx.y;
    int Col = blockIdx.x * blockDim.x + threadIdx.x;

    if (Row < Width && Col < Width) {
        float Pvalue = 0;
        for (int k = 0; k < Width; ++k) {
            Pvalue += A[Row * Width + k] * B[k * Width + Col];
        }
        C[Row * Width + Col] = Pvalue;
    }
}

变量存储类别 关键字总结

用于指定变量的存储类别，这些关键字决定了变量的存储位置以及如何在不同线程和线程块之间共享：

关键字	描述	作用域	生命周期
device	用于在GPU的全局内存中声明变量。	所有线程	应用程序执行期间
global	用于定义在主机上调用但在设备上执行的函数（即CUDA核心函数）。	-	-
host	用于定义在主机上调用并执行的函数。	-	-
shared	用于声明位于共享内存中的变量。	同一个线程块内的线程	线程块的执行期间
constant	用于声明位于常量内存中的变量。	所有线程	应用程序执行期间
managed	用于声明在主机和设备之间共享的统一内存变量。	所有线程和主机	应用程序执行期间

• __device__：这些变量存储在设备的全局内存中，可以被所有线程访问，但访问延迟较高。
• __global__：定义的是CUDA核心函数，这种函数可以从主机（CPU）调用并在设备（GPU）上异步执行。
• __host__：定义的是常规的C++函数，仅在主机上执行。
• __shared__：声明的变量位于共享内存中，这是一种较快的内存类型，但仅在同一个线程块内的线程之间共享。
• __constant__：用于声明常量内存中的变量，这种内存对于所有线程来说是只读的，访问速度快，但空间有限。
• __managed__：Unified Memory（统一内存）中的变量，可以被GPU和CPU共同访问，CUDA运行时负责管理内存的迁移。