CUDA学习笔记9——CUDA 共享内存 / Shared Memory

由于共享内存拥有仅次于寄存器的读写速度，比全局内存快得多。因此，能够用共享内存访问替换全局内存访问的场景都可以考虑做对应的优化。

不利用共享内存的矩阵乘法

不利用共享内存的矩阵乘法的直接实现。每个线程读取A的一行和B的一列，并计算C的相应元素，如图。
访问次数 ：
从全局内存中读取A的次数为B.width，读取B的次数为A.height。

在这里插入图片描述

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "error.cuh"
#include <stdlib.h>
#include<math.h>
#include <malloc.h> 
#include <stdlib.h>//利用share memory 和统一内存优化矩阵乘
#define M 80
#define N 2000// 线程块尺寸
#define BLOCK_SIZE 16///-----------没有共享内存的矩阵乘法-----------
// 矩阵以行为主的顺序存储:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct
{int width;int height;float* elements;
} Matrix;// MatMul()调用的矩阵乘法内核
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{// 每个线程通过将结果累积到Cvalue中来计算C的一个元素float Cvalue = 0;int row = blockIdx.y * blockDim.y + threadIdx.y;int col = blockIdx.x * blockDim.x + threadIdx.x;for (int e = 0; e < A.width; ++e)Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];C.elements[row * C.width + col] = Cvalue;
}// 矩阵乘法核的前向声明
//__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
// 矩阵乘法-主机代码
//矩阵维度被假定为BLOCK_SIZE的倍数
void MatMul(const Matrix A, const Matrix B, Matrix C)
{// 将A和B加载到设备内存Matrix d_A;d_A.width = A.width; d_A.height = A.height;size_t size = A.width * A.height * sizeof(float);cudaMalloc(&d_A.elements, size);cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);Matrix d_B;d_B.width = B.width; d_B.height = B.height;size = B.width * B.height * sizeof(float);cudaMalloc(&d_B.elements, size);cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);// 在设备内存中分配CMatrix d_C;d_C.width = C.width; d_C.height = C.height;size = C.width * C.height * sizeof(float);cudaMalloc(&d_C.elements, size);// Invoke kerneldim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);MatMulKernel <<< dimGrid, dimBlock >>>(d_A, d_B, d_C);//从设备内存中读取CcudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);// 释放设备内存cudaFree(d_A.elements);cudaFree(d_B.elements);cudaFree(d_C.elements);
}int main()
{	Matrix matrix_1, matrix_2, matrix_out;int memsize = sizeof(float) * M * N;int memsize_out = sizeof(float) * M * M;matrix_1.width  = matrix_2.height =  M;matrix_2.width  = matrix_1.height = N;matrix_out.width  = matrix_out.height = M;cudaMallocHost((void**)&matrix_1.elements, memsize);cudaMallocHost((void**)&matrix_2.elements, memsize);cudaMallocHost((void**)&matrix_out.elements, memsize_out);for (int y = 0; y < N; y++)for (int x = 0; x < M; x++)matrix_1.elements[y * M + x] = (float)rand() / 1.0E5 ;for (int y = 0; y < M; y++)for (int x = 0; x < N; x++)matrix_2.elements[y * N + x] = (float)rand() / 1.0E5;//for (int y = 0; y < N; y++)//{//	printf("\n matrix_1[%d]:\n", y);//	for (int x = 0; x < M; x++)//	{//		printf("%.2f ", matrix_1.elements[y * M + x]);//	}//}//for (int y = 0; y < M; y++)//{//	printf("\n matrix_2[%d]:\n", y);//	for (int x = 0; x < N; x++)//	{//		printf("%.2f ", matrix_2.elements[y * N + x]);//	}//}cudaEvent_t start, stop_gpu;cudaEventCreate(&start);//创建事件cudaEventCreate(&stop_gpu);//创建事件cudaEventRecord(start, 0);//记录事件MatMul(matrix_1, matrix_2, matrix_out);cudaEventRecord(stop_gpu,0);//记录事件cudaEventSynchronize(stop_gpu);float time_gpu;cudaEventElapsedTime(&time_gpu, start, stop_gpu);//事件计时//printf("\n GPU time: %.4f ms \n", time_gpu);cudaEventDestroy(start);//销毁事件cudaEventDestroy(stop_gpu);for (int y = 0; y < M; y++){printf("\n matrix_out[%d]:\n", y);for (int x = 0; x < M; x++){printf("%.2f ", matrix_out.elements[y * M + x]);}}cudaFreeHost(matrix_1.elements);cudaFreeHost(matrix_2.elements);cudaFreeHost(matrix_out.elements);system("pause");return 0;
}

在这里插入图片描述

利用共享内存的矩阵乘法

每个线程块负责计算矩阵C的一个方子矩阵Csub，块内的每个线程负责计算Csub的一个元素。

Csub等于两个矩形矩阵的乘积:维度为(A.width,block_size)的A的子矩阵与Csub具有相同的行索引，维度为(A.width,block_size)的B的子矩阵与Csub具有相同的列索引。
为了适应设备的资源，这两个矩形矩阵被分割成尽可能多的尺寸为block_size的方阵，Csub被计算为这些方阵乘积的和。
首先将两个对应的方阵从全局内存加载到共享内存，其中一个线程加载每个矩阵的一个元素，然后让每个线程计算乘积的一个元素。每个线程将这些产品的结果累积到寄存器中，完成后将结果写入全局内存。
访问次数 ：
通过这种方式阻塞计算，我们利用了快速共享内存并节省了大量的全局内存带宽，矩阵A只从全局内存中读取(B.width/block_size)次，矩阵B只从全局内存中读取(A.height/block_size)次。

在这里插入图片描述

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "error.cuh"
#include <stdlib.h>
#include<math.h>
#include <malloc.h> 
#include <stdlib.h>//利用share memory 和统一内存优化矩阵乘#define M 80
#define N 2000// 线程块尺寸
#define BLOCK_SIZE 16///-----------矩阵乘法与共享内存-----------
// 矩阵以行为主的顺序存储:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct
{int width;int height;float* elements;
} Matrix;
// 得到一个矩阵元素
__device__ float GetElement(const Matrix A, int row, int col)
{return A.elements[row * A.width + col];
}
// 设置一个矩阵元素
__device__ void SetElement(Matrix A, int row, int col, float value)
{A.elements[row * A.width + col] = value;
}
// 获取A的BLOCK_SIZExBLOCK_SIZE子矩阵subb，它位于A的左上角的col子矩阵和行子矩阵
__device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{Matrix Asub;Asub.width = BLOCK_SIZE;Asub.height = BLOCK_SIZE;Asub.elements = &A.elements[A.width * BLOCK_SIZE * row + BLOCK_SIZE * col];return Asub;
}// 矩阵乘法核的前向声明
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
// 矩阵乘法-主机代码
// 矩阵维度被假定为BLOCK_SIZE的倍数
void MatMul(const Matrix A, const Matrix B, Matrix C)
{// 将A和B加载到设备内存Matrix d_A;d_A.width = A.width; d_A.height = A.height;size_t size = A.width * A.height * sizeof(float);cudaMalloc(&d_A.elements, size);cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);Matrix d_B;d_B.width = B.width; d_B.height = B.height;size = B.width * B.height * sizeof(float);cudaMalloc(&d_B.elements, size);cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);// 在设备内存中分配CMatrix d_C;d_C.width = C.width; d_C.height = C.height;size = C.width * C.height * sizeof(float);cudaMalloc(&d_C.elements, size);// 调用内核dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);MatMulKernel <<< dimGrid, dimBlock >>>(d_A, d_B, d_C);// 从设备内存中读取CcudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);// 空闲设备内存cudaFree(d_A.elements);cudaFree(d_B.elements);cudaFree(d_C.elements);
}
// MatMul()调用的矩阵乘法内核
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{// 块行和列int blockRow = blockIdx.y;int blockCol = blockIdx.x;// 每个线程块计算C的一个子矩阵CsubMatrix Csub = GetSubMatrix(C, blockRow, blockCol);// 每个线程通过将结果累积到Cvalue中来计算Csub的一个元素float Cvalue = 0;// 线程行和列在Csubint row = threadIdx.y;int col = threadIdx.x;// 遍历计算Csub所需的A和B的所有子矩阵，将每对子矩阵相乘并累加结果for (int i = 0; i < (A.width / BLOCK_SIZE); ++i){// 得到A的子矩阵Matrix Asub = GetSubMatrix(A, blockRow, i);// 得到B的子矩阵BMatrix Bsub = GetSubMatrix(B, i, blockCol);// 用于存储sub和sub的共享内存tively__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];// 将subb和Bsub从设备内存加载到共享内存，每个线程加载每个子矩阵的一个元素As[row][col] = GetElement(Asub, row, col);Bs[row][col] = GetElement(Bsub, row, col);// 同步以确保在开始计算之前加载子矩阵__syncthreads();// 将subb和Bsub相乘for (int j = 0; j < BLOCK_SIZE; ++j)Cvalue += As[row][j] * Bs[j][col];// 同步以确保在下一次迭代中加载两个新的子矩阵A和B之前完成前面的计算__syncthreads();}// 将Csub写入设备内存// 每个线程写入一个元素SetElement(Csub, row, col, Cvalue);
}int main()
{	Matrix matrix_1, matrix_2, matrix_out;int memsize = sizeof(float) * M * N;int memsize_out = sizeof(float) * M * M;matrix_1.width  = matrix_2.height =  M;matrix_2.width  = matrix_1.height = N;matrix_out.width  = matrix_out.height = M;cudaMallocHost((void**)&matrix_1.elements, memsize);cudaMallocHost((void**)&matrix_2.elements, memsize);cudaMallocHost((void**)&matrix_out.elements, memsize_out);for (int y = 0; y < N; y++)for (int x = 0; x < M; x++)matrix_1.elements[y * M + x] = (float)rand() / 1.0E5 ;for (int y = 0; y < M; y++)for (int x = 0; x < N; x++)matrix_2.elements[y * N + x] = (float)rand() / 1.0E5;cudaEvent_t start, stop_gpu;cudaEventCreate(&start);//创建事件cudaEventCreate(&stop_gpu);//创建事件cudaEventRecord(start, 0);//记录事件MatMul(matrix_1, matrix_2, matrix_out);cudaEventRecord(stop_gpu,0);//记录事件cudaEventSynchronize(stop_gpu);float time_gpu;cudaEventElapsedTime(&time_gpu, start, stop_gpu);//事件计时//printf("\n GPU time: %.4f ms \n", time_gpu);cudaEventDestroy(start);//销毁事件cudaEventDestroy(stop_gpu);for (int y = 0; y < M; y++){printf("\n matrix_out[%d]:\n", y);for (int x = 0; x < M; x++){printf("%.2f ", matrix_out.elements[y * M + x]);}}cudaFreeHost(matrix_1.elements);cudaFreeHost(matrix_2.elements);cudaFreeHost(matrix_out.elements);system("pause");return 0;
}

在这里插入图片描述

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

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

相关文章

『Linux升级路』基础开发工具——gcc/g++篇

『Linux升级路』基础开发工具——gcc/g++篇

&#x1f525;博客主页&#xff1a;小王又困了 &#x1f4da;系列专栏&#xff1a;Linux &#x1f31f;人之为学&#xff0c;不日近则日退 ❤️感谢大家点赞&#x1f44d;收藏⭐评论✍️ 目录 一、快速认识gcc/g 二、预处理 &#x1f4d2;1.1头文件展开 &#x1f4d2;1…
阅读更多...
java字符串的常见用法

java字符串的常见用法

java字符串的常见用法 Java中的字符串是一个非常常用的对象&#xff0c;它属于Java的内置类String类的实例。字符串在Java中是不可变的&#xff0c;即一旦创建了一个字符串对象&#xff0c;就不能修改它的值。 下面是一些关于Java字符串的详细用法&#xff1a; 1&#xff09;创…
阅读更多...
从零开始，用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合

从零开始，用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合

&#x1f38f;&#xff1a;你只管努力&#xff0c;剩下的交给时间 &#x1f3e0; &#xff1a;小破站 "从零开始&#xff0c;用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合 前言准备工作编写docker-compose.yml文件为什么使用本机ip为什么skywa…
阅读更多...
代码随想录-刷题第六天

代码随想录-刷题第六天

242. 有效的字母异位词 题目链接&#xff1a;242. 有效的字母异位词 思路&#xff1a;哈希法。利用数组来记录出现的字母个数&#xff0c;然后判断是否为字母异位词。 时间复杂度&#xff1a;O(n) class Solution {public boolean isAnagram(String s, String t) {int[] co…
阅读更多...
【云备份】第三方库的认识与使用

【云备份】第三方库的认识与使用

文章目录 json库粗略认识详细认识writer 类reader类jsoncpp序列化实现jsoncpp反序列化实现 bundle文件压缩库简单认识bundle库实现文件压缩bundle库实现文件解压缩 httplib库Request类Response类Server类Client类 json库 粗略认识 json是一种数据交换格式&#xff0c;采用完全…
阅读更多...
激光切割设备中模组的作用有哪些？

激光切割设备中模组的作用有哪些？

激光切割设备是一种高精度的自动化加工设备&#xff0c;用于对金属、非金属等材料进行精确切割。直线模组作为激光切割设备的重要组成部分&#xff0c;在激光切割设备中起着重要的作用&#xff0c;为设备的运动系统提供了高精度、高稳定性和高效率的运动控制。 1、高精度的位置…
阅读更多...
excel单元格加背景颜色不生效？

excel单元格加背景颜色不生效？

如果在 Excel 中设置单元格背景颜色而发现不生效&#xff0c;可能有几个原因。以下是一些常见的解决方法&#xff1a; 1. **单元格锁定&#xff1a;** 检查所在单元格是否被锁定。如果单元格被锁定&#xff0c;并且工作表被保护&#xff0c;你可能无法更改其背景颜色。在工作表…
阅读更多...
mysql 优化器的AST树是啥

mysql 优化器的AST树是啥

from ChatGPT: MySQL中的优化器&#xff08;optimizer&#xff09;使用AST&#xff08;Abstract Syntax Tree&#xff0c;抽象语法树&#xff09;来表示查询的语法结构。AST是一种树状结构&#xff0c;它反映了查询语句的语法层次&#xff0c;是一个抽象表示&#xff0c;用于更…
阅读更多...
Linux - 文件系统 - 理解目录 - 理解 软/硬链接

Linux - 文件系统 - 理解目录 - 理解 软/硬链接

前言 在上篇博客当中&#xff0c;我们对 文件系统 和 inode 做了初步了解&#xff0c;本博客将在上篇博客的基础之上&#xff0c;对于 文件系统当中的目录进行进步一阐述。 Linux - 进一步理解 文件系统 - inode - 机械硬盘-CSDN博客 目录 一个文件有一个 inode&#xff0c;…
阅读更多...
Redis打包事务,分批提交

Redis打包事务,分批提交

一、需求背景 接手一个老项目&#xff0c;在项目启动的时候&#xff0c;需要将xxx省整个省的所有区域数据数据、以及系统字典配置逐条保存在Redis缓存里面&#xff0c;这样查询的时候会更快; 区域数据字典数据一共大概20000多条,&#xff0c;前同事直接使用 list.forEach…
阅读更多...
Windows安装MongoDB

Windows安装MongoDB

1、下载MongoDB的zip&#xff0c;解压 2、创建目录 mkdir D:\JavaSoftware\Database\MongoDB\mongodb-win32-x86_64-windows-5.0.8\data\db mkdir D:\JavaSoftware\Database\MongoDB\mongodb-win32-x86_64-windows-5.0.8\data\log 3、创建一个配置文件mongod.cfg&#xff0c…
阅读更多...
使用一个接口的结果作为第二个接口的参数并将两者的数据放置成下拉框的格式

使用一个接口的结果作为第二个接口的参数并将两者的数据放置成下拉框的格式

背景 我使用下拉框实现选择id 但是只有两个接口 一个是所有的id 另一个是id对应的具体信息 我想把id传入另一个接口并且获取其name然后写成类似这样的数组 [ { value: 1, label: ‘名称1’ }&#xff0c; { value: 2, label: ‘名称2’ } { value: 3, label: ‘名称3’ } ] 然…
阅读更多...
【PPspliT】ppt转pdf-保留过渡动画

【PPspliT】ppt转pdf-保留过渡动画

网址 http://www.maxonthenet.altervista.org/ppsplit.php 下载安装 使用 再次打开ppt&#xff0c;就能在上方的选项栏里头看到了&#xff1a;
阅读更多...
GEE生物量碳储量——利用红和近红外波段和OTSU大津法提取纯净森林面积

GEE生物量碳储量——利用红和近红外波段和OTSU大津法提取纯净森林面积

简介： 如何利用红和近红外波段和OTSU大津法提取纯净森林面积？本文的主要逻辑是利用特定时期的遥感影像的波段，提取指定范围的内的DN值，然后分别统计发生阈值变化的峰值区域，从而作为筛选森林的临界点，如果研究区较大的话则需要先进行影像分割，分割成为相同大小的区域，…
阅读更多...
前端开发Java学习

前端开发Java学习

注释&#xff1a; 单行注释 // 多行注释 /* */ 文件注释 /** */ 1 关键字 &#xff08;关键字一定是小写&#xff09; 2 常量 字符串常量"HelloWord","你好世界"整数常量12&#xff0c;-33小数常量3.14&#xff0c;22.1字符常量A&#xff0c;a &…
阅读更多...
不是默认进入Linux|总是自动进入windows系统

不是默认进入Linux|总是自动进入windows系统

问题描述 不是默认进入Linux系统无法主动出现boot引导自动进入windows系统 尝试无效 修复引导无效重装Grub无效重装系统无效 环境 Ubuntu 22.04 LST微星主板 解决方案 修改引导顺序&#xff1a; 开机狂按Del键&#xff0c;进入BIOS系统&#xff0c;左侧Settings 设置&…
阅读更多...
src和href的区别

src和href的区别

前言 持续学习总结输出中&#xff0c;src和href都是HTML中特定元素的属性&#xff0c;都可以用来引入外部的资源。两者区别如下&#xff1a; 1、作用 href 用于在当前文档和引用资源之间确立联系。 src 用于替换当前内容。 2、范围用途 src&#xff08;source&#xff09…
阅读更多...
项目整个管理论文之6

项目整个管理论文之6

阅读更多...
RabbitMQ基础教程

RabbitMQ基础教程

1.什么是消息队列 消息队列&#xff08;Message Queue&#xff09;&#xff0c;我们一般简称为MQ。消息队列中间件是分布式系统中重要的组件&#xff0c;具有异步性、松耦合、分布式、可靠性等特点。用于实现高性能、高可用、可伸缩和最终一致性架构。是大型分布式系统不可缺少…
阅读更多...
Angular11 MSAL B2C登录实例 （一）

Angular11 MSAL B2C登录实例 （一）

前言 因为项目需求&#xff0c;需要把Angular 11项目中登录方式改成B2C登录&#xff0c;所以在参考了一系列文档后&#xff0c;成功通过MSAL将项目的登录方式改成B2C登录。下面介绍了详细步骤及一些注意事项。 步骤&#xff1a; 1. 安装MSAL 在项目中安装msal npm i azure/…
阅读更多...
最新文章

Copyright @ 2022~2023