【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的一个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核心。它使用了两个共享内存数组MdsNds来存储两个输入矩阵的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运行时负责管理内存的迁移

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/635964.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

2024第十二届中国(上海)国际流体机械展览会(CFME2024)

2024上海流体机械展|泵展|风机展|压缩机展 2024第十二届中国(上海)国际流体机械展览会(CFME2024) 泵阀门风机压缩机减变速机真空干燥设备展 时间&#xff1a;2024年11月25日-27日 地点&#xff1a;国家会展中心&#xff08;上海虹桥&#xff09; ◆展会背景&#xff1a; &q…

从 Context 看 Go 设计模式:接口、封装和并发控制

文章目录 Context 的基本结构Context 的实现和传递机制为什么 Context 不直接传递指针案例&#xff1a;DataStore结论 在 Go 语言中&#xff0c; context 包是并发编程的核心&#xff0c;用于传递取消信号和请求范围的值。但其传值机制&#xff0c;特别是为什么不通过指针传递…

服务器与普通电脑的区别,普通电脑可以当作服务器用吗?

服务器在我们日常应用中非常常见&#xff0c;手机APP、手机游戏、PC游戏、小程序、网站等等都需要部署在服务器上&#xff0c;为我们提供各种计算、应用服务。服务器也是计算机的一种&#xff0c;虽然内部结构相差不大&#xff0c;但是服务器的运行速度更快、负载更高、成本更高…

智慧工地解决方案及案例:PPT全文26页,附下载

关键词&#xff1a;智慧工地解决方案&#xff0c;智慧工地建设&#xff0c;智慧工地整体架构&#xff0c;数字化工地&#xff0c;智慧工程 一、智慧工地建设对传统建筑业的影响 1、提高了施工效率和质量&#xff1a;智慧工地建设采用了先进的信息技术和管理方法&#xff0c;可…

力扣343. 整数拆分(动态规划)

Problem: 343. 整数拆分 文章目录 题目描述思路解题方法复杂度Code 题目描述 思路 该题目可以抽象成动态规划中的爬楼梯模型&#xff0c;将整数的拆分类比为上台阶&#xff1a; 1.每个阶段可以从整数中划分出1、2、…k的一个整数 2.int dp[n 1] dp[i]表示为i的整数划分的最大…

Next-GPT: Any-to-Any Multimodal LLM

Next-GPT: Any-to-Any Multimodal LLM 最近在调研一些多模态大模型相关的论文&#xff0c;发现Arxiv上出的论文根本看不过来&#xff0c;遂决定开辟一个新坑《一页PPT说清一篇论文》。自己在读论文的过程中会用一页PPT梳理其脉络和重点信息&#xff0c;旨在帮助自己和读者快速了…

2024年华为OD机考高分攻略-完整题库-两周350分

华为OD是个不错的机会,很适合非软件行业到软件行业的转身。 但是很多同学之前没有软件基础,不知道该如何高效的准备OD机考。 我是一名软件培训老师,我的学生有上百人顺利通过了华为OD机考,并取得了高分,我将经验分享给大家,华为OD机试2周350分,高效复习策略: 1、牛客…

C语言辨析——int a=5;为什么++a=1能编译通过而a++=1不行呢?

1. 问题 有人问&#xff1a;int a5; 为什么a1能编译通过a&#xff1d;1编译错误呢&#xff1f; 解释&#xff1a;不管是a1还是a1在C编译环境下都无法编译通过&#xff0c;但在C中&#xff0c;a1可以编译通过&#xff0c;而a1无法编译通过&#xff0c;这也是C和C的一个不同。因…

线程池的简单介绍及使用

线程池 线程池的参数介绍拒绝策略 线程池的任务处理流程使用Executors创建常见的线程池 线程池的参数介绍 corePoolSize: (核心线程数)这是线程池中始终存在的线程数&#xff0c;即使这些线程处于空闲状态。maximumPoolSize:(最大线程数) 是线程池允许的最大线程数。keepAliveT…

免费200万Tokens 用科大讯飞API调用星火大模型服务

简介 自ChatGPT火了之后&#xff0c;国内的大模型发展如雨后春笋。其中的佼佼者之一就是科大讯飞研发的星火大模型,现在大模型已经更新到V3 版本&#xff0c;而且对开发者也是相当友好&#xff0c;注册就送200万tokens,讯飞1tokens 约等于 1.5 个中文汉字 或者 0.8 个英文单词…

安卓开发-day

一、安卓项目结构 1、manifests文件夹 Android系统配置文件夹&#xff0c;包含一个AndroidManifest.xml文件&#xff1b; AndroidMainifest.xml文件是每个android项目必须要包含的文件&#xff08;项目唯一&#xff09;&#xff0c;创建项目时默认就会生成这个文件&#xff0…

【Python学习】Python学习21- 正则表达式(1)

目录 【Python学习】Python学习21- 正则表达式&#xff08;1&#xff09; 前言re.match函数实例 re.search方法re.match与re.search的区别参考 文章所属专区 Python学习 前言 本章节主要说明Python的正则表达式。 正则表达式是一个特殊的字符序列&#xff0c;它能帮助你方便的…

java常见面试题:如何使用Java进行JMS(Java Message Service)编程?

Java Message Service (JMS) 是一个用于创建、发送、接收和读取消息的 API。它允许应用程序通过消息队列或主题进行通信。以下是使用 Java 进行 JMS 编程的详细步骤&#xff1a; 添加 JMS 依赖 如果你使用 Maven&#xff0c;你可以添加以下依赖到你的 pom.xml 文件&#xff1…

2401vim,vim标号

标号简介 提供高亮,快速告诉用户有用信息.如,调试器在左侧列中有个表示断点的图标. 另一例可能是表示(PC)程序计数器的箭头.标号功能允许在窗口左侧放置标号或图标,并定义应用行的高亮. 此外,调试器还支持8到10种不同的标号和高亮颜色,见|NetBeans|. 使用标号有两个步骤: 1…

动态规划基础(一)引入

T a l k Talk Talk i s is is c h e a p cheap cheap , , , s h o w show show m e me me t h e the the c o d e code code 数字三角形 题目大意 给定数字金字塔&#xff0c;每个单位有自己的权值&#xff0c;问从顶端出发&#xff0c;到底端任意一点的所有路径中&a…

算法训练营Day37(贪心6)

738.单调递增的数字 力扣&#xff08;LeetCode&#xff09;官网 - 全球极客挚爱的技术成长平台 注意特例1000 class Solution:def monotoneIncreasingDigits(self, n: int) -> int:# 将整数转换为字符串strNum str(n)# flag用来标记赋值9从哪里开始# 设置为字符串长度&am…

HarmonyOS 页面跳转控制整个界面的转场动画

好 本文 我们来说 页面间的转场动画 就是 第一个界面到另一个界面 第一个界面的退场和第二个界面的进场效果 首先 我这里 创建了两个页面文件 Index.ets和AppView.ets index组件 编写代码如下 import router from "ohos.router" Entry Component struct Index {b…

Day39 动态规划part02 62. 不同路径 63. 不同路径 II

动态规划part02 62. 不同路径 63. 不同路径 II 62. 不同路径 class Solution { public:int uniquePaths(int m, int n) {vector<vector<int>> dp(m, vector(n,0)); //第一步&#xff0c;确定dp[i][j]就是移动到[i,j]位置的方法数量for(int i 0; i<dp.size();…

解决你所有ArcGIS标注问题!二分式、三分式、条件分类标注、上下标、牵引线、文字格式化…

解决你所有ArcGIS标注问题&#xff01;二分式、三分式、条件分类标注、上下标、牵引线、文字格式化… 标注是将描述性文本放置在地图中的要素上或要素旁的过程。 本文整理了ArcGIS中的各种标注方法、可能遇到的问题和细节&#xff0c;内容比较杂&#xff0c;想到哪写到哪。 …

tui.calender日历创建、删除、编辑事件、自定义样式

全是坑&#x1f573;&#xff01;全是坑&#x1f573;&#xff01;全是坑&#x1f573;&#xff01;能不用就不用&#xff01; 官方文档&#xff1a;https://github.com/nhn/tui.calendar/blob/main/docs/en/apis/calendar.md 实例的一些方法&#xff0c;比如创建、删除、修改、…