个人网站与企业网站区别,广昌建设局官方网站,公司注册流程及资料,房地产 网站模板缓存分块是一种内存优化技术#xff0c;主要用于提高数据的局部性#xff08;Locality#xff09;#xff0c;以减少缓存未命中#xff08;Cache Miss#xff09;的次数。在现代计算机体系结构中#xff0c;处理器#xff08;CPU#xff09;的速度通常比内存快得多。因…缓存分块是一种内存优化技术主要用于提高数据的局部性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][TILE_WIDTH]; // 定义共享内存用于存储Md的一个Tile__shared__ float Nds[TILE_WIDTH][TILE_WIDTH]; // 定义共享内存用于存储Nd的一个Tileint 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内的乘积并累加到Pvaluefor (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() 出现在两个关键位置:
加载数据到共享内存之后这里的 __syncthreads() 确保了所有线程都完成了对共享内存的写入操作。即使这个写入操作是在 for 循环中完成的我们也需要确保每个线程都完成了当前迭代的加载操作才能安全地开始使用这些共享内存中的数据进行计算。计算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运行时负责管理内存的迁移。
