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

这篇具有很好参考价值的文章主要介绍了【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核心。它使用了两个共享内存数组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;
    }
}

变量存储类别 关键字总结

用于指定变量的存储类别,这些关键字决定了变量的存储位置以及如何在不同线程和线程块之间共享:文章来源地址https://www.toymoban.com/news/detail-799127.html

关键字 描述 作用域 生命周期
device 用于在GPU的全局内存中声明变量。 所有线程 应用程序执行期间
global 用于定义在主机上调用但在设备上执行的函数(即CUDA核心函数)。 - -
host 用于定义在主机上调用并执行的函数。 - -
shared 用于声明位于共享内存中的变量。 同一个线程块内的线程 线程块的执行期间
constant 用于声明位于常量内存中的变量。 所有线程 应用程序执行期间
managed 用于声明在主机和设备之间共享的统一内存变量。 所有线程和主机 应用程序执行期间
  • __device__:这些变量存储在设备的全局内存中,可以被所有线程访问,但访问延迟较高。
  • __global__:定义的是CUDA核心函数,这种函数可以从主机(CPU)调用并在设备(GPU)上异步执行。
  • __host__:定义的是常规的C++函数,仅在主机上执行。
  • __shared__:声明的变量位于共享内存中,这是一种较快的内存类型,但仅在同一个线程块内的线程之间共享
  • __constant__:用于声明常量内存中的变量,这种内存对于所有线程来说是只读的,访问速度快,但空间有限。
  • __managed__:Unified Memory(统一内存)中的变量,可以被GPU和CPU共同访问,CUDA运行时负责管理内存的迁移

到了这里,关于【cuda】四、基础概念:Cache Tiled 缓存分块技术的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • Spring Boot 3.0系列【25】数据篇之Spring Cache缓存技术使用详解

    有道无术,术尚可求,有术无道,止于术。 本系列Spring Boot版本3.0.5 源码地址:https://gitee.com/pearl-organization/study-spring-boot3

    2023年04月14日
    浏览(36)
  • CUDA基础(三)CPU架构,指令,GPU架构

    CPU中央处理器,负责执行用户和操作系统下发的指令。CPU只能接受01二进制语言,0和1用来控制高低电位。比如,一个加法运算,在x86处理器上的的二进制代码为: 01001000 00000001 11000011 这样一行代码被称为机器码,它执行了加法操作。除了这样的加法,CPU的电路还要实现很多

    2024年02月08日
    浏览(38)
  • Conda、Git、pip设置代理教程 解决Torch not compiled with CUDA enabled问题 pip缓存坑 No module named “Crypto“

    总结 在使用Conda时,如果您需要通过代理访问网络资源,可以按照以下步骤配置代理: 打开终端并运行以下命令以设置HTTP代理: 请将“代理服务器”和“端口号”替换为您的代理服务器和端口号。例如,如果您使用的代理服务器是“proxy.example.com”,端口号是“8080”,则命

    2024年02月09日
    浏览(51)
  • GPU加速02:超详细Python Cuda零基础入门教程,没有显卡也能学!

    Python是当前最流行的编程语言,被广泛应用在深度学习、金融建模、科学和工程计算上。作为一门解释型语言,它运行速度慢也常常被用户诟病。著名Python发行商Anaconda公司开发的Numba库为程序员提供了Python版CPU和GPU编程工具,速度比原生Python快数十倍甚至更多。使用Numba进行

    2024年02月02日
    浏览(60)
  • 构建Docker基础镜像(ubuntu20.04+python3.9.10+pytorch-gpu-cuda11.8)

    内容如下 访问官网下载页 https://www.python.org/downloads/release/python-3910/ 下拉选择 Gzipped 包 ps:创建镜像名为 ub2004py3910pytorchgpucuda118 标签为 latest 的镜像,从当前路径下的 DockerFile 文件打包

    2024年02月05日
    浏览(88)
  • 【多版本cuda自由切换】在ubuntu上安装多个版本的CUDA,并且可以随时切换cuda-11.3//cuda-11.8//cuda-11.6//cuda-11.2

    问题描述         项目开发中,不同的项目可能对不同的cuda版本有所要求,常见的是这几种cuda-11.3//cuda-11.8//cuda-11.6,按照之前的认知,一个主机只能安装一个版本的cuda,否则会引起环境混乱,知道cuda底层逻辑的人都知道这有多么扯蛋,对吧。         也正是因为受到这个

    2024年02月03日
    浏览(54)
  • 【CUDA】判断电脑是否安装CUDA

    打开命令行窗口(例如 PowerShell 或命令提示符),然后输入以下命令: nvcc --version 如果安装了 CUDA,它会显示 CUDA 编译工具的版本信息,类似于你上面提供的输出。 一般是 如果你安装了 NVIDIA 显卡驱动程序,你可以在系统托盘中找到 NVIDIA 控制面板。在其中的信息部分可能会

    2024年02月08日
    浏览(39)
  • 了解NVIDAI显卡驱动(包括:CUDA、CUDA Driver、CUDA Toolkit、CUDNN、NCVV)

    转载 一篇 背景   开发过程中需要用到GPU时,通常在安装配置GPU的环境过程中遇到问题;CUDA Toolkit和CUDNN版本的对应关系;CUDA和电脑显卡驱动的版本的对应关系;CUDA Toolkit、CUDNN、NCVV是什么呢? 举个例子 安装TensorFlow2.1过程中,想要使用到电脑的显卡来进行开发,但是发现默

    2024年02月13日
    浏览(57)
  • cuda 安装(windows)图解+指定版本cuda

    除了在linux上用,部分初学者也会在windows上用 需要在linux上安装的参考这篇:cuda linux安装 有的教程讲的很复杂,起始很简单,nvidia官方已经给了全家桶了,直接一个包装完,不用考虑其中的N个组件 1.下载Cuda Toolkit(cuda全家桶) cuda-toolkit官方页面 找到需要的版本 2.安装 双击

    2024年02月22日
    浏览(44)
  • 问题7:虚拟机+ubuntu+安装cuda(傻瓜式操作)+cuda path配置+查看cuda的版本

    目录 1.cuda的安装 2.cuda path的配置 3.检查cuda的版本号(为安装pytorch做准备) ... 建议看看下面的几条ps,可以避免踩坑! ps:本文所用ubuntu系统版本为v-22.04(如果打算安装可以参考博主的另一篇文章),ubuntu-v-22.04最高支持cuda-v-11.7.0 ps:此文为2023.2.4所写,此时pytorch支持的c

    2024年02月02日
    浏览(39)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包