程序员经验分享-hwhale

cuda 矩阵乘法，从最容易理解到算得最快（第二版源码-tile机制+共享内存）

2023-11-19

下面我们仅仅引入tiling方法，在共享内存中进行分块矩阵的乘法运算。先分析一下能够减少多少次对全局存储区的访问。

当M=N=K=4096时，用第一版的代码，忽略cache的缓存时，需要从全局存储区读取2*(4096^3)个float变量。

为了让思路简单一点，我们假设A、B和C三个矩阵都是4096*4096的方阵。按照64*64大小的tile来分块，于是ABC都可以变成64x64的以64*64的tile为元素的分块矩阵。按照分块矩阵的乘法规则等同于普通矩阵的乘法规则。C的每个分块，需要64个分块矩阵的乘法操作，并将结果累加。

C的一个分块需要多少次全局内存的访问呢？重复64次，每次先分别取A和B的64*64个元素进共享内存。每个分块需要读取64*（64*64+64*64）个全局内存中的float变量；C总共有64x64个分块元素。所以总共需要读取从A、B读取     (64x64)*64*(64*64+64*64)=4096*64*2*4096=128*4906^2次个float变量。

[2*(4096^3) ] 除以 [128*4906^2]，结果为64倍。也就是说，使用了tile机制，将A、B的数据读取到共享内存再计算矩阵乘法，对全局内存的访问量可以减少为原来的1/64。这将会减少程序的运行时间。让我们把这个思路实现出来，并测试对比一下运行时间。

gridDim.x=流处理器的个数，每个block一次负责一个64*64的C的分块元素的计算；从左向右，从上往下，依次迭代负责。为思考的简单起见，假设C的M和N都是64的倍数。不为倍数的时候，差不多，细节后面再补充。

本文内容由网友自发贡献，版权归原作者所有，本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容，请联系:hwhale#tublm.com(使用前将#替换为@)

CUDA

线性代数

机器学习

算法

cuda 矩阵乘法，从最容易理解到算得最快（第二版源码-tile机制+共享内存） 的相关文章

  • 优化三角矩阵计算的 CUDA 内核的执行

    我正在开发我的第一个 Cuda 应用程序 并且我的内核 吞吐量低于预期 这似乎是目前最大的瓶颈 内核的任务是计算一个 N N 大小的矩阵 DD 包含数据矩阵上所有元素之间的平方距离 数据矩阵 Y 的大小为 N D 以支持多维数据 并存储为行

  • 为什么GK110有192个核心和4个扭曲？

    我想感受一下开普勒的架构 但这对我来说没有意义 如果一个 warp 有 32 个线程 其中 4 个被调度 执行 则意味着 128 个核心正在使用 64 个核心处于空闲状态 白皮书中提到了独立指令 那么64核是为这些指令保留的吗 如果是这样

  • Golang调用CUDA库

    我正在尝试从 Go 代码中调用 CUDA 函数 我有以下三个文件 test h int test add void test cu global void add int a int b int c c a b int test add v

  • CUDA：如何在设备上填充动态大小的向量并将其内容返回到另一个设备函数？

    我想知道哪种技术可以填充设备上的动态大小数组 int row 在下面的代码中 然后返回其内容 以供另一个设备函数使用 为了将问题置于上下文中 下面的代码尝试使用在 GPU 上运行的高斯 勒让德求积来跨越勒让德多项式基组中的任意函数 incl

  • 如何将CUDA时钟周期转换为毫秒？

    我想用一些代码来测量时间within我的内核需要 我已经关注了这个问题 https stackoverflow com questions 11209228 timing different sections in cuda kernel连

  • 为什么 gcc 和 NVCC (g++) 会看到两种不同的结构大小？

    我正在尝试将 CUDA 添加到 90 年代末编写的现有单线程 C 程序中 为此 我需要混合两种语言 C 和 C nvcc 是 c 编译器 问题在于 C 编译器将结构视为特定大小 而 C 编译器将相同的结构视为略有不同的大小 那很糟 我对此感

  • CUDA程序导致nvidia驱动程序崩溃

    当我超过大约 500 次试验和 256 个完整块时 我的 monte carlo pi 计算 CUDA 程序导致我的 nvidia 驱动程序崩溃 这似乎发生在 monteCarlo 内核函数中 任何帮助都会受到赞赏 include

  • CUDA 估计 2D 网格数据的每块线程数和块数

    首先我要说的是 我已经仔细阅读了所有类似的问题 确定每个块的线程和每个网格的块 https stackoverflow com questions 4391162 cuda determining threads per block blo

  • Nvcc 的版本与 CUDA 不同

    我安装了 cuda 7 但是当我点击 nvcc version 时 它打印出 6 5 我想在 GTX 960 卡上安装 Theano 库 但它需要 nvcc 7 0 我尝试重新安装cuda 但它没有更新nvcc 当我运行 apt get i

  • 无法在 CUDA 中执行设备内核

    我正在尝试在全局内核中调用设备内核 我的全局内核是矩阵乘法 我的设备内核正在查找乘积矩阵每列中的最大值和索引 以下是代码 device void MaxFunction float Pd float max int x threadIdx

  • 设置最大 CUDA 资源

    我想知道是否可以设置 CUDA 应用程序的最大 GPU 资源 例如 如果我有一个 4GB GPU 但希望给定的应用程序只能访问 2GB 如果它尝试分配更多 就会失败 理想情况下 这可以在进程级别或 CUDA 上下文级别上设置 不 目前没有允

  • 加速Cuda程序

    要更改哪一部分来加速此代码 代码到底在做什么 global void mat Matrix a Matrix b int tempData new int 2 tempData 0 threadIdx x tempData 1 blockI

  • 在 __device/global__ CUDA 内核中动态分配内存

    根据CUDA 编程指南 http developer download nvidia com compute cuda 3 2 prod toolkit docs CUDA C Programming Guide pdf 第 122 页 可

  • 使用 CUDA 进行逐元素向量乘法

    我已经在 CUDA 中构建了一个基本内核来执行逐元素两个复向量的向量 向量乘法 内核代码插入如下 multiplyElementwise 它工作正常 但由于我注意到其他看似简单的操作 如缩放向量 在 CUBLAS 或 CULA 等库中进行了

  • TensorRT 多线程

    我正在尝试使用 python API 来使用 TensorRt 我试图在多个线程中使用它 其中 Cuda 上下文与所有线程一起使用 在单个线程中一切正常 我使用 docker 和 tensorrt 20 06 py3 图像 onnx 模型和

  • __device__ __constant__ 常量

    有什么区别吗 在 CUDA 程序中定义设备常量的最佳方法是什么 在 C 主机 设备程序中 如果我想将常量定义在设备常量内存中 我可以这样做 device constant float a 5 constant float a 5 问题 1

  • CUDA - 将 CPU 变量传输到 GPU __constant__ 变量

    与 CUDA 的任何事情一样 最基本的事情有时也是最难的 所以 我只想将变量从 CPU 复制到 GPUconstant变量 我很难过 这就是我所拥有的 constant int contadorlinhasx d int main int

  • 最小化 MC 模拟期间存储的 cuRAND 状态数量

    我目前正在 CUDA 中编写蒙特卡罗模拟 因此 我需要生成lots使用随机数cuRAND图书馆 每个线程处理一个巨大的元素floatarray 示例中省略 并在每次内核调用时生成 1 或 2 个随机数 通常的方法 参见下面的示例 似乎是为每

  • 通过 cuFFT 进行逆 FFT 缩放

    每当我使用 cuFFT 绘制程序获得的值并将结果与 Matlab 的结果进行比较时 我都会得到相同形状的图形 并且最大值和最小值位于相同的点 然而 cuFFT 得到的值比 Matlab 得到的值大得多 Matlab代码是 fs 1000 s

  • 大型跨平台软件项目的技巧/资源

    我将开始一个大型软件项目 涉及跨平台 GUI 和大量的数字运算 我计划用 C 和 CUDA 编写大部分应用程序后端 并用 Qt4 编写 GUI 我计划使用 Make 作为我的构建系统 这将是一个只有两名开发人员的项目 一旦我相对深入地了解它

随机推荐

  • 微信小程序报错-errCode: -1

    1 报错截图 2 报错原因 Collection remove 需要小程序端2 9 4版本或之后的版本才支持 看了眼我的调试基础库 我的版本是2 8 1 所以 3 解决方法 点击右上角的详情 本地设置 调试基础库选择2 14 0

  • js数组相加相减函数

    数组相减 reduceArray arr1 arr2 for var i arr1 length 1 i gt 0 i var a arr1 i for var j arr2 length 1 j gt 0 j var b arr2 j i

  • Qt打开包含头文件以及打开函数声明和定义的方法

    当第一次拿到被人的程序的时候不知道如何向VS或者keil这样的编译器可以直接右键打开头文件 Qt可以通过 1 鼠标放在你需要打开的头文件或者函数声明的地方 如下图我鼠标放在MainWindow上 2 Ctrl 鼠标左键单击 或者使用快捷键F

  • python-算法时间复杂度和空间复杂度

    大O表示法 O 名称 举例 1 常量时间 一次赋值 logn 对数时间 折半查找 n 线性时间 线性查找 nlogn 对数线性时间 快速排序 n 2 平方 两重循环 n 3 立方 三重循环 2 n 指数 递归求斐波那契数列 n 阶乘 旅行商

  • html文本元素

    文章目录 h p span pre code 实体字符 strong i em del s h h head 标题 一共有六级标题 hKaTeX parse error Expected got EOF at end of input 6

  • 【编译原理】 CS143 斯坦福大学公开课 第一周：简介

    youtube 1 1 Introduction to Compilers and interpreters 1 1 Introduction to Compilers and interpreters 编译器解释器介绍 两种主要的实现编程

  • three.js中聚光灯及其属性介绍

    一 聚光灯及其属性介绍 Three js中的聚光灯 SpotLight 是一种用于在场景中创建聚焦光照的光源类型 它有以下属性 color 聚光灯的颜色 intensity 聚光灯的强度 distance 聚光灯的有效距离 angle 聚光

  • [毕业设计]2023-2024年最新电子科学与技术专业毕设选题题目推荐汇总

    文章目录 1前言 2 如何选题 3 选题方向 3 1 嵌入式开发方向 3 2 物联网方向 3 3 人工智能方向 3 4 算法研究方向 3 5 学长作品展示 4 最后 1前言 近期不少学弟学妹询问学长关于电子科学与技术专业相关的毕设选题 学长

  • java如何检测连接池连接情况,如何检查是否使用了连接池

    I use HSQLDB EclipseLink Gemini on OSGI framework Felix In spite that I ve set pool in persistence xml I have serious su

  • 全网最全谷粒商城记录_01、简介-项目介绍（2022-07-06更新完成）

    声明 本教程不收取任何费用 欢迎转载 尊重作者劳动成果 不得用于商业用途 侵权必究 目录 分布式基础 全栈开发篇 分布式高级 微服务架构篇 高可用集群 架构师提升篇 希望大家 微服务架构图简单介绍 项目简介 1 项目背景 1 电商模式 1

  • JavaWeb学习笔记-part1

    互联网通信 什么是互联网通信 两台计算机通过网络实现文件共享行为 就是互联网通信 互联网通信中的角色划分 客户端 用于发送请求的计算机 服务端 用于接受请求 并满足请求的计算机 互联网通信模型 C S通信模型 client software

  • Handler机制与原理

    为什么会出现内存泄漏问题呢 分析 Handler使用是用来进行线程间通信的 所以新开启的线程是会持有Handler引用的 如果在Activity等中创建Handler 并且是非静态内部类的形式 就有可能造成内存泄漏 非静态内部类是会隐式持有

  • uniapp 开发微信小程序之新版隐私协议

    自从微信小程序官方更新隐私协议 用户必须同意之后 才能获取个人信息 这就导致在获取用户信息之前 需要有个隐私协议弹窗 大致如下图 微信小程序官方提供的API和 uniapp 开发的稍微有点区别 这里只记录 uniapp 开发的 如果需要微信

  • 高中学历的程序员，以包装的方式进入现在的公司，想跳槽咋办？

    网友自述 我在现在广州这家公司工作了两年 技术上有一定提升 但这两年我过得一直不是很快乐 因为我学历包装 所以我不敢跟同事交往太深 一直孤身一人 非常难受 可能这就是代价吧 现在我想换一个公司 我不想再用假身份了 但不知道用高中学历是否能够

  • Java对点、线、面生成栅格瓦片jpg，并渲染呈现

    Java对点 线 面生成栅格瓦片jpg 并渲染呈现 1 效果图 2 原理 2 1 面瓦片的生成 2 2 线瓦片的生成 2 3 多点瓦片的生成 3 源码 参考 这篇博客将介绍从前端HTML页面到后端预生成栅格瓦片jpg 并提供查询接口供前端h

  • Python文件操作

    1 with open E 信息 docx rb as f 2 read data f read 3 f closed rb 以二进制形式读取指定路径的文件 再以二进制形式写入指定路径 wb 1 with open E 信息 2 docx

  • Go Web编程实战（6）----反射

    目录 反射 反射的3大原则 接口类型变量 转换为 反射类型对象 反射类型对象 转换为 接口类型变量 反射类型对象 修改 值必 可写的 反射 与其他语言一样 Go语言的反射同样是指 计算机程序在运行时 可以访问 检测和修改它本身状态或行为的一

  • MAC 怎么终端怎么退出和进入Anaconda环境

    mac安装完anaconda 后 命令行窗口默认使用conda的 取消默认 用以下一行代码在命令行运行即可 重启终端 conda config set auto activate base false 将false改为true设置默认环境为

  • Codeforces 1475C. Ball in Berland(二元容斥)

    题目传送门 题意 一个班级有a个男生和b个女生 现在这个班级有k对男女愿意一起出席毕业典礼 这里注意k对男女中可能会有某个男生或女生出现在多个pair中 你从这k对中找出两对 使得这两对中的男生不相同 女生不相同 即一个男生或女生不可能在一

  • cuda 矩阵乘法，从最容易理解到算得最快（第二版源码-tile机制+共享内存）

    下面我们仅仅引入tiling方法 在共享内存中进行分块矩阵的乘法运算 先分析一下能够减少多少次对全局存储区的访问 当M N K 4096时 用第一版的代码 忽略cache的缓存时 需要从全局存储区读取2 4096 3 个float变量 为了

热门标签