CUDA学习(九)想好好解释一下利用shared memory去做matrixMul

2023-11-03

 代码来自CUDA Samples\v9.1\0_Simple\matrixMul

主要注释了matrixMulCUDA()

C=A*B

它的理念是把矩阵分成小块,每个线程利用两层循环,大循环在迭代subMatrix,小循环则是迭代每个小subMatrix的横纵坐标,

可以理解为一个线程计算出的Csub就是对应一个元素,然后你想象所有线程是并行的,所以所有的元素计算就可以理解过来了,但是单纯的一个线程是解决不了这个问题的,因为它小循环迭代的时候是需要所有元素的内容的,这也是使用__syncthreads()的作用,而你在使用global memory的时候,不需要用这个,因为每个线程各做各的,都需要去读global memory,这也是使用shared memory的作用,它可以减少对global的访问。

说的不对的地方请大家给我指出来,这各shard memory的使用是我觉得最难的

 

 

/**
 * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

/**
 * Matrix multiplication: C = A * B.
 * Host code.
 *
 * This sample implements matrix multiplication as described in Chapter 3
 * of the programming guide.
 * It has been written for clarity of exposition to illustrate various CUDA
 * programming principles, not with the goal of providing the most
 * performant generic kernel for matrix multiplication.
 *
 * See also:
 * V. Volkov and J. Demmel, "Benchmarking GPUs to tune dense linear algebra,"
 * in Proc. 2008 ACM/IEEE Conf. on Supercomputing (SC '08),
 * Piscataway, NJ: IEEE Press, 2008, pp. Art. 31:1-11.
 */

// System includes
#include <stdio.h>
#include <assert.h>

// CUDA runtime
#include <cuda_runtime.h>

// Helper functions and utilities to work with CUDA
#include <helper_functions.h>
#include <helper_cuda.h>

/**
 * Matrix multiplication (CUDA Kernel) on the device: C = A * B
 * wA is A's width and wB is B's width
 */
template <int BLOCK_SIZE> __global__ void
matrixMulCUDA(float *C, float *A, float *B, int wA, int wB)
{
    // Block index
    int bx = blockIdx.x;
    int by = blockIdx.y;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    // Index of the first sub-matrix of A processed by the block
    //这是A的第一个细分矩阵的索引,利用block处理(by=blockIdx.y)
    //也是从这里让aBegin根据不同block进行处理
    //不同block拥有不同的aBegin,即不同的block线程下的循环初始值的不同
    int aBegin = wA * BLOCK_SIZE * by;

    // Index of the last sub-matrix of A processed by the block
    //这是A的最后一个细分矩阵的索引,利用block处理
    //注意wA是A的width
    int aEnd   = aBegin + wA - 1;

    // Step size used to iterate through the sub-matrices of A
    //A的细分矩阵的步距
    int aStep  = BLOCK_SIZE;

    // Index of the first sub-matrix of B processed by the block
    //这是B的第一个细分矩阵的索引,也是利用block处理(bx=blockIdx.x)
    int bBegin = BLOCK_SIZE * bx;

    // Step size used to iterate through the sub-matrices of B
    //B的细分矩阵的迭代用步距
    int bStep  = BLOCK_SIZE * wB;

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

CUDA学习(九)想好好解释一下利用shared memory去做matrixMul 的相关文章

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

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

    CUDA 版本 10 1 帕斯卡 GPU 所有命令都发送到默认流 void ptr cudaMalloc ptr launch kernel lt lt lt gt gt gt ptr cudaDeviceSynchronize Is th
  • CUDA、NPP 滤波器

    CUDA NPP 库支持使用 nppiFilter 8u C1R 命令过滤图像 但不断出现错误 我可以毫无问题地启动并运行 boxFilterNPP 示例代码 eStatusNPP nppiFilterBox 8u C1R oDeviceS
  • CUDA - 将 CPU 变量传输到 GPU __constant__ 变量

    与 CUDA 的任何事情一样 最基本的事情有时也是最难的 所以 我只想将变量从 CPU 复制到 GPUconstant变量 我很难过 这就是我所拥有的 constant int contadorlinhasx d int main int
  • 内联 PTX 汇编代码强大吗?

    我看到一些代码示例 人们在 C 代码中使用内联 PTX 汇编代码 CUDA工具包中的文档提到PTX很强大 为什么会这样呢 如果我们在 C 代码中使用这样的代码 我们会得到什么好处 内联 PTX 使您可以访问未通过 CUDA 内在函数公开的指
  • 通过 cuFFT 进行逆 FFT 缩放

    每当我使用 cuFFT 绘制程序获得的值并将结果与 Matlab 的结果进行比较时 我都会得到相同形状的图形 并且最大值和最小值位于相同的点 然而 cuFFT 得到的值比 Matlab 得到的值大得多 Matlab代码是 fs 1000 s
  • “gld/st_throughput”和“dram_read/write_throughput”指标之间有什么区别?

    在 CUDA 可视化分析器版本 5 中 我知道 gld st requested throughput 是应用程序请求的内存吞吐量 然而 当我试图找到硬件的实际吞吐量时 我很困惑 因为有两对似乎合格的指标 它们是 gld st throug
  • cudaDeviceScheduleBlockingSync 和 cudaDeviceScheduleYield 之间有什么区别?

    正如这里所说 如何减少 CUDA 同步延迟 延迟 https stackoverflow com questions 11953722 how to reduce cuda synchronize latency delay 等待设备结果有
  • 云或烟雾的粒子系统

    我正在尝试使用 OpenGL 和 CUDA 制作一个简单的用于云和烟雾模拟的粒子系统 如何使粒子系统中的粒子表现得像真正的云或烟雾在低湍流风中的表现 我现在遇到的一些问题是 颗粒聚集成一个大球 粒子扩散到无限远 粒子突然弹射离开 我已经完成
  • 设备内存刷新cuda

    我正在运行一个 C 程序 其中调用了两次 cuda 主机函数 我想清理这两个调用之间的设备内存 有没有办法可以刷新 GPU 设备内存 我使用的是计算能力为2 0的Tesla M2050 如果你只想将内存归零 那么cudaMemset可能是最
  • OpenCV 2.4.3rc 和 CUDA 4.2:“OpenCV 错误:没有 GPU 支持”

    我在这张专辑中上传了几张截图 https i stack imgur com TELST jpg https i stack imgur com TELST jpg 我正在尝试在 Visual Studio 2008 中的 OpenCV 中
  • 如何在CUDA应用程序中正确应用线程同步?

    一般来说 我在应用程序中偶尔会使用线程同步 因为我并不经常需要此功能 我并不是真正的高级 C C 程序员 但我也不是初学者 我开始学习 CUDA C 对当今 GPU 与 CPU 的能力相比感到兴奋 我意识到 CUDA 编程主要是关于并行线程
  • C 中的 CUDA:如何使用 cudaMemcpyAsync 修复错误 11

    我目前正在尝试使用 CUDA 运行一个简单的多 GPU 程序 它的基本作用是将一个包含一些虚拟数据的大型数组复制到 GPU GPU 进行一些数学计算 然后将结果数组复制回来 我在 VS2017 的输出中没有收到任何错误 但我设置的一些错误消
  • Cuda:最小二乘求解,速度较差

    最近 我使用Cuda编写了一个名为 正交匹配追踪 的算法 在我丑陋的 Cuda 代码中 整个迭代需要 60 秒 而 Eigen lib 只需 3 秒 在我的代码中 矩阵 A 是 640 1024 y 是 640 1 在每一步中 我从 A 中
  • 在 CUDA 中的设备内存上分配 2D 数组

    如何在 Cuda 中的设备内存中分配和传输 往返于主机 2D 数组 我找到了解决这个问题的方法 我不必展平阵列 内置的cudaMallocPitch 函数完成了这项工作 我可以使用以下命令将阵列传输到设备或从设备传输阵列cudaMemcpy
  • 将内核链接到 PTX 函数

    我可以使用 PTX 文件中包含的 PTX 函数作为外部设备函数 将其链接到另一个应调用该函数的 cu 文件吗 这是另一个问题CUDA 将内核链接在一起 https stackoverflow com questions 20636800 c
  • 无法编译cuda_ndarray.cu:libcublas.so.7.5:无法打开共享对象文件

    我正在尝试在 aws 实例中导入 theano 库以使用 GPU 我已经使用 boto 编写了一个 python 脚本来自动执行 aws 设置 该脚本本质上会从我的本地计算机对实例执行 ssh 然后启动一个 bash 脚本 其中我执行 py
  • 一维纹理内存访问比一维全局内存访问更快吗?

    我正在测量标准纹理和 1Dtexture 内存访问之间的差异 为此 我创建了两个内核 global void texture1D float doarray int size int index calculate each thread
  • 了解流式多处理器 (SM) 和流式处理器 (SP)

    我正在尝试了解 GPU 的基本架构 我已经阅读了很多材料 包括这个非常好的答案 https stackoverflow com a 2213744 2386113 但我仍然很困惑 无法得到一个好的图片 我的理解 GPU 包含两个或多个流式多
  • 将 cuda 数组传递给 Thrust::inclusive_scan

    我可以对 cpu 上的数组使用包容性扫描 但是否可以对 gpu 上的数组执行此操作 注释是我知道有效但我不需要的方式 或者 是否有其他简单的方法可以对设备内存中的数组执行包含扫描 Code include

随机推荐

  • 华为od机试题-2023-最新真题-完整题库-两周350分

    华为OD机试题库每半年刷新一次 目前已经整理了90道原题 并提供了java python C 三种解法 这就是2023年Q2的完整题库了 华为OD机试2周350分 高效复习策略 1 牛客网刷基础算法题 每个算法都了解一下 用一周时间 牛客网
  • 基于JDBC的图书管理系统

    文章目录 1 图书管理系统类图 2 程序设计 2 1Student类 2 2Teacher类 2 3Sever类 2 4BookManage类 2 5Manage类 2 6Mbin类 2 7Menu类 3 功能 3 1管理员功能 3 2学生
  • 应急响应LINUX&Windows

    应急响应LINUX Windows linux 文件名 说明 etc passwd 用户信息文件 etc crontab 定时任务文件 etc anacrontab 异步定时任务文件 etc rc d rc local 开机启动项 var
  • angular2 通过 ControlValueAccessor NG_VALUE_ACCESSOR实现自定义表单控件

    一 演示自定义控件的双向绑定 二 如何实现 cform component ts 自定义表单控件CformComponent 的实现 注意下里面的注释段落 import Component OnInit Input forwardRef f
  • QT自定义Widget控件及其使用

    今天来给大家分享一下QT自定义widget控件及其使用 当ui设计器提供的界面不满足实际需求时 可以从QWidget继承自定义的界面组件 有两种方法一种是提升法 另一种是ui设计器自定义界面组件widget组件 我们本次先说提升法 通过这个
  • 预告:Intel、Hulu、阿里、京东、携程等大数据实战直播

    前言 由CSDN主办的SDCC 2017之大数据技术实战线上峰会将在CSDN学院举行 作为SD系列技术峰会的一部分 本次线上峰会秉承干货实料 案例 的内容原则 将邀请圈内顶尖的布道师 技术专家和技术引领者 共话大数据平台构建 优化提升大数据
  • 学生选课系统---数据库课程设计SQL Server

    可以直接从我的GitHub中获取文档 学生选课系统GitHub 一 题目 学生选课系统 二 需求分析 1 根据学生专业学年学期等信息 录入课程完成课程计划 2 根据课程计划 录入任课教师信息 3 学生可以根据学年学期等信息 选择课程完成选课
  • 2022年第三届“大湾区杯”粤港澳金融数学数学建模竞赛

    2022年第三届 大湾区杯 粤港澳金融数学数学建模竞赛 文章目录 2022年第三届 大湾区杯 粤港澳金融数学数学建模竞赛 前言 一 分析B题 券商研报对公司股票走势的影响和投资策略 1 解读题目 2 第一问解析 3 第二问解析 4 第三问解
  • python读excel文件解析

    目录 一 题目描述 二 示例 三 分析 1 获取工作簿对象 2 获取表单的名字 3 通过表单名获取指定的表单对象 4 获取单元格中的值 一 题目描述 例如在当前文件夹下有一个名为 阿里巴巴2020年股票数据 xls 的Excel文件 读取并
  • Android Studio 自动更新

    一丶最近项目开发在用Android Studio 不过无奈自动更新检查时总是提示 Connection failed Please check your network connection and try again 怀疑可能是因为 不可
  • 关于vue3的使用插件FileSaver并且excel导出的实现方法

    参考文章 1vue elementui表格前端导出excel以及自定义导出样式 前端自定义报表模板 elementui 小郑有点困了的博客 CSDN博客 2Vue3 XLSX FileSaver导出导入功能 菜鸟学习JAVA开发的博客 CS
  • ESP8266 下载报错please check partition type 6 addr:3fd000 len:3000

    问题现象 串口打印 ESP8266 12F 的错误如下 ets Jan 8 2013 rst cause 2 boot mode 3 7 load 0x40100000 len 23948 room 16 tail 12 chksum 0x
  • 《信号与系统》解读 前言:为什么要解读《信号与系统》?

    什么需要解读 信号与系统 为什么需要开辟这个专题呢 1 基础理论的基础 2 工程技术的基础 3 建立系统思维 1 基础理论的基础 信号与系统 是 数字信号处理DSP 通信原理 数字通信原理 等其他基础课程的基础 传统的信号与系统偏硬件 现代
  • 写五个以科技为主题的html页面,要求用css实现,请写出实现代码

    在此基础上 我们可以创建五个以科技为主题的 HTML 页面 并使用 CSS 来实现设计 科技新闻页面 HTML 代码
  • Fibonacci数列(斐波那契数列)

    斐波那契数列 Fibonacci sequence 又称黄金分割数列 因数学家列昂纳多 斐波那契 Leonardoda Fibonacci 以兔子繁殖为例子而引入 故又称为 兔子数列 指的是这样一个数列 1 1 2 3 5 8 13 21
  • TCP的超时时间间隔的简单理解

    网上查了许多资料 都比较抽象 自我理解记录篇 本文适用于查了许多资料 但觉得还是比较抽象 有基础相关知识的人阅读 样本RTT SampleRTT EstimatedRTT 1 EstimatedRTT aSampleRTT 参考值 0 12
  • JS--变量

    变量在使用时分为两步 声明变量 赋值 1 声明变量 var age 声明一个名称为age的变量 var是一个JS关键词 用来声明变量 使用该关键词声明变量后 计算机会 自动为变量分配内存空间 不需要程序员管 age是程序员定义的变量名 我们
  • 《C++ primer plus》精炼(OOP部分)——对象和类(4)

    学习是人类进步的阶梯 也是个人成功的基石 罗伯特 肯尼迪 文章目录 友元函数 利用友元函数重载 lt lt 运算符 重载部分示例 矢量类 友元函数 先看看在上一章中我们作为例子的代码 class Student string name in
  • QComboBox 设置代理组件

    背景 QComboBox 是Qt中比较常用的一个输入控件 用于实现一个文本下拉列表 在简单的应用场景中 QComboBox 完全可以满足要求 但是项目实践过程中会遇到以下问题 需要实现比较复杂的下拉列表 比如同时显示图标 文字 按钮等 QC
  • CUDA学习(九)想好好解释一下利用shared memory去做matrixMul

    代码来自CUDA Samples v9 1 0 Simple matrixMul 主要注释了matrixMulCUDA C A B 它的理念是把矩阵分成小块 每个线程利用两层循环 大循环在迭代subMatrix 小循环则是迭代每个小subM