cuda矩阵相乘_CUDA计算矩阵相乘
1.最简单的 kernel 函数
__global__ void MatrixMulKernel( float* Md, float* Nd, float* Pd, int Width)
{
int tx = threadIdx.x; // cloumn
int ty = threadIdx.y; // row
float Pvalue = 0;
for (int k = 0; k
{
float Mdele = Md[ty*Width + k];
float Ndele = Nd[k*Width + tx];
Pvalue += Mdele * Ndele;
}
Pd[ty*Width + tx] = Pvalue;
}
2.适用更大矩阵
第一节中例子缺点是,假如使用更多的块时,每个块中会计算相同的矩阵。而且矩阵元素不能超过512个线程(块最大线程限制)。
改进方法是,假设每个块维度都是方阵形式,且其维度由变量 TILE_WIDTH 指定。矩阵 Pd 的每一维都划分为部分,每个部分包含 TILE_WIDTH 个元素。
#define TILE_WIDTH 4
#define Width 8
__global__ void MatrixMulKernel( float* Md, float* Nd, float* Pd, int Width)
{
int Col = blockId.x*TILE_WIDTH + threadIdx.x; // cloumn
int Row = blockId.y*TILE_WIDTH + threadIdx.y; // row
float Pvalue = 0;
for (int k = 0; k
{
Pvalue += Md[Row*Width + k] * Nd[k*Width + Col];
}
Pd[ty*Width + tx] = Pvalue;
}
dim3 dimGrid(Width/TILE_WIDTH, Width/TILE_WIDTH);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
CUDA 硬件相关概念
针对 GT200而言
每个 SM 有最多8个块
每个 SM 有最多1024个线程
应当注意每个块中分配的线程数,以便能够充分利用SM
以32个线程为一个 warp,warp 为线程调度单位
3.通过改变储存器提到访问效率
CGMA(Compute to Global Memory Access),尽量提高CGMA比值。
对前面来说,每个for循环内需要两次访问全局内存(Md[Row*Width + k], Nd[k*Width + Col]),两次浮点计算(加法与乘法)。因此 CGMA 为2:2 = 1。
G80 全局存储器带宽为 86.4 GB/s,计算峰值性能 367 Gflops。(每秒取 21.6G 个变量,由于CGMA=1,会进行 21.6G 次浮点计算)
若每个单精度浮点数为 4 字节,那么由于存储器限制,最大浮点操作不会超过 21.6 Gflops。
| 存储器类型 | 变量 | 周期 | 特点 | 访问速度 |
| --- | --- | --- | --- |
| 共享存储器 | 共享变量 | kernel函数 | 每个块中所有线程都可以访问,用于线程间协作高效方式 | 相当快,高度并行访问 |
| 常数存储器 | 常数变量 | 所有网格 | 相当快,并行访问 |
| 寄存器 | 自动变量 | 线程 | 寄存器具有储存容量限制 | 非常块
| 全局存储器 | 全局变量 | | 用于调用不同kernel函数时传递信息 | 慢 |
减少全局存储器流量策略
各个存储器特点:
全局存储器,容量大,访问慢
共享存储器,容量小,访问块
新算法要点:
将共享存储器上数据划分子集,每个子集满足共享存储器容量限制
通过线程协作将 M 和 N 中的元素加载到共享存储器中,每个线程负责块中一个元素赋值(Mds[threadIdx.y][threadIdx.x],Nds[threadIdx.y][threadIdx.x])
通过加载到共享存储器,使得每个块访问全局内存的次数减小为原来的 TILE_WIDTH 分之一(加载到共享内存时读取一次,每个块中使用 TILE_WIDTH 次)
__global__ void kernel_tile(float* M, float* N, float* P){
int i, k;
float Pvalue = 0;
__shared__ float Mds[tile_width][tile_width];
__shared__ float Nds[tile_width][tile_width];
int Row = threadIdx.y + blockIdx.y*tile_width;
int Col = threadIdx.x + blockIdx.x*tile_width;
for( i = 0; i< width/tile_width; i++ ){
Mds[threadIdx.y][threadIdx.x] = M[Row*width + i*tile_width + threadIdx.x];
Nds[threadIdx.y][threadIdx.x] = N[(threadIdx.y + i*tile_width)*width + Col];
__syncthreads();
for (k = 0; k
Pvalue += Mds[threadIdx.y][k]*Nds[k][threadIdx.x];
}
__syncthreads();
}
P[Row*width + Col] = Pvalue;
}
存储器容量限制
G80 硬件中
每个 SM 寄存器大小为 8 KB(8192 B)
每个 SM 共享内存大小为 16 KB
若每个 SM 中容纳线程数为 768,那么每个线程可用寄存器不超过 8 KB/768=10 Bytes(两个单精度变量占用 8 Bytes)
若每个线程占用了多余 10 Bytes,那么就会减少 SM 上线程数,且以块为单位减少
若每个 SM 中容纳 8 个块,那么每个块不能使用超过 2 KB 的存储器
4.数据预取
在 CUDA 中,当某些线程在等待其存储器访问结果时,CUDA 线程模型可以通过允许其他 warp 继续运行,这样就能容许长时间的访问延时。
为了充分利用此特性,需要当使用当前数据元素时预取下一个数据元素,这样就可以正价在储存器访问和已访问的数据使用指令之间的独立指令的数目。
预取技术经常和分块技术结合,解决带宽限制和长时间延迟问题。
例如在第三节中矩阵乘法问题中,使用预取技术后函数流程变为
对应程序为
__global__ void kernel_prefetch(float* M, float* N, float* P){
int i;
float Pvalue = 0;
float Mc, Nc;
int Row = threadIdx.y + tile_width*blockIdx.y;
int Col = threadIdx.x + tile_width*blockIdx.x;
__shared__ float Mds[tile_width][tile_width];
__shared__ float Nds[tile_width][tile_width];
Mc = M[Row*width + threadIdx.x];
Nc = N[Col + threadIdx.y*width];
for (i = 1; i
Mds[threadIdx.y][threadIdx.x] = Mc;
Nds[threadIdx.y][threadIdx.x] = Nc;
__syncthreads();
Mc = M[Row*width + threadIdx.x + i*tile_width];
Nc = N[Col + (threadIdx.y + i*tile_width)*width];
for (int k = 0; k
Pvalue += Mds[threadIdx.y][k]*Nds[k][threadIdx.x];
}
__syncthreads();
}
P[Col + Row*width] = Pvalue;
}
Mc,Nc 为增加的两个储存在寄存器内的变量。
cuda矩阵相乘_CUDA计算矩阵相乘相关推荐
- matlab 解方程组 矩阵,用MATLAB计算矩阵和解线性方程组.ppt
用MATLAB计算矩阵和解线性方程组.ppt 第五章 矩阵与行列式 §5.6 用MATLAB计算 矩阵与行列式 用MATLAB计算矩阵与行列式 行列式的求值 矩阵的基本运算 矩阵的加.减 数与矩阵相乘 ...
- python输出n阶矩阵_python-递归计算矩阵(nxn)的行列式
您确定您的未成年人返回了一个新对象,而不是对原始矩阵对象的引用吗?我使用了您的确切行列式方法,并为您的类实现了一个次要方法,它对我来说很好. 下面是您的矩阵类的快速/脏实现,因为我没有您的实现.为简洁 ...
- 判断矩阵对称以及计算矩阵稀疏度
计算稀疏度 稀疏度与密度,自己之前搞反了,sparsity等于所有零元素的个数除以总元素的个数,density等于所有非零元素的个数除以总元素的个数. import numpy as np impor ...
- 矩阵乘法的计算和来源
矩阵乘法的计算 矩阵,是线性代数中的基本概念之一.一个m×n的矩阵就是m×n个数排成m行n列的一个数阵. 矩阵乘法是一种高效的算法可以把一些一维递推优化到log(n),还可以求路径方案等,所以更是一种 ...
- 用MATLAB计算矩阵和行列式
用MATLAB计算矩阵和行列式 det(X) 函数---计算行列式的值 计算矩阵的乘法 计算矩阵的逆矩阵 计算矩阵的伴随矩阵 计算矩阵的行最简和 秩 det(X) 函数-计算行列式的值 用法: 例: ...
- python如何输入一个矩阵_python怎么输入矩阵
python怎么输入一个数组矩阵 下面是基于python3.4的 这样子: 1 2 3 4 import numpy as np arr = [1,2,3,4,5,6,7,8,9] matrix_a ...
- numpy向量转换为矩阵_Numpy之将矩阵拉成向量的实例
Numpy之将矩阵拉成向量的实例 废话不多说,直接上代码吧! # 矩阵操作 # 将矩阵拉成向量 import numpy as np x = np.arange(10).reshape(2,5) pr ...
- cuda矩阵相乘_CUDA入门实战2:将矩阵乘法速度提升5000倍
本实验采用不同的方法来计算 8192 * 8192 的整型矩阵乘法运算. C语言版 C语言是大家公认的高性能语言,那我们就从C语言开始吧. // 用一位数组表示二维矩阵 mat1 = (int*) m ...
- 两个列向量相乘怎么计算_矩阵:行主序、列主序、行向量、列向量
看龙书的时候发现一个矩阵在传入Shader之前都要转置一下,很好奇为什么要有一步这样的操作. 行主序和列主序 行主序指矩阵在内存中逐行存储,列主序指矩阵在内存中逐列存储. 行主序矩阵内存布局: 列主序 ...
最新文章
- 挨踢项目求生法则-团队建设篇
- java blob压缩_如何从Oracle中用Java压缩的BLOB列中提取XML文档
- NoSQL生态系统——hash分片和范围分片两种分片
- Java虚拟机13:Java类加载机制
- 产品策划七:App界面交互设计规范
- “error LNK2019: 无法解析的外部符号”的几种可能原因
- Unity中Instantiate一个prefab时需要注意的问题
- 小b删列(51Nod-2523)
- Ubuntu 挂载新磁盘
- loadrunner取出字符串的后面几位
- 2017年高考改革地区:浙江、上海
- 【codeup22562】最长回文子串(dp基础题)
- CSMA/CD和拥塞控制AIMD其实是一回事!
- 有道词典 linux 无法运行,ubuntu16.04安装不上有道词典的解决办法
- 彻底解决win10出现的无法访问网络位置*:\XXXXX\XXXXX的问题
- C3 linearization
- 贤鱼的刷题日常【c++动态规划】4978:宠物小精灵之收服
- 安装Adobe软件时显示:您的浏览器或者操作系统不再受支持,您可能需要安装操作系统的最新更新。
- RadioButton设置默认选中后无法取消,可选中多个的问题
- 在已有公众号发布微信小程序
热门文章
- 【概率论与数理统计 第三版 浙江大学 盛骤 谢式千 潘乘毅 编】作业答案
- Android Sensor感应器介绍,获取用户移动方向,指南针原理
- 关于“一个SAPer的网络日志”
- 如何让计算机桌面自动更换,怎么让电脑自动更换桌面背景
- Axure制作微信APP原型(三):消息+个人模块----聊天、浮窗、登录设备、个人主页、更多信息
- 关于Unity绑定手机
- 联想笔记本重装win7系统之后无线不能用
- 【Dash搭建可视化网站】项目5: 利用Dash 实现动态图表
- java没有编译器吗_java有编译器吗
- Springboot自动重启