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

news2024/9/21 22:54:30

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

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

不利用共享内存的矩阵乘法的直接实现。每个线程读取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);
	// 在设备内存中分配C
	Matrix 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 kernel
	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);
	//从设备内存中读取C
	cudaMemcpy(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);
	// 在设备内存中分配C
	Matrix 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);
	// 从设备内存中读取C
	cudaMemcpy(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的一个子矩阵Csub
	Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
	// 每个线程通过将结果累积到Cvalue中来计算Csub的一个元素
	float Cvalue = 0;
	// 线程行和列在Csub
	int 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的子矩阵B
		Matrix 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.coloradmin.cn/o/1245175.html

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

相关文章

Vector - CANoe - Vector Hardware Manager以太网

Vector - CANoe - Vector Hardware Manager以太网

前面的文章中有介绍过基于Network based mode和channel base mode的环境配置&#xff0c;不过我们都是使用比较旧的办法&#xff0c;在我使用了一段时间Vector Hardware Manager配置之后发现这个更加好用结合之前的配置方法&#xff0c;使用起来也更加的灵活&#xff0c;今天就…
阅读更多...
Kafka 常用功能总结（不断更新中....）

Kafka 常用功能总结（不断更新中....）

kafka 用途 业务中我们经常用来两个方面 1.发送消息 2.发送日志记录 kafka 结构组成 broker&#xff1a;可以理解成一个单独的服务器&#xff0c;所有的东西都归属到broker中 partation&#xff1a;为了增加并发度而做的拆分&#xff0c;相当于把broker拆分成不同的小块&…
阅读更多...
98年阿里P6测试猿晒出工资单，看完扎心了。。。

98年阿里P6测试猿晒出工资单，看完扎心了。。。

最近一哥们跟我聊天装逼&#xff0c;说他最近从阿里跳槽了&#xff0c;我问他跳出来拿了多少&#xff1f;哥们表示很得意&#xff0c;说跳槽到新公司一个月后发了工资&#xff0c;月入5万多&#xff0c;表示很满足&#xff01;这样的高薪资着实让人羡慕&#xff0c;我猜这是税后…
阅读更多...
使用Git bash切换Gitee、GitHub多个Git账号

使用Git bash切换Gitee、GitHub多个Git账号

Git是分布式代码管理工具&#xff0c;使用命令行的方式提交commit、revert回滚代码。这里介绍使用Git bash软件来切换Gitee、GitHub账号。     假设在gitee.com上的邮箱是alicefoxmail.com 、用户名为alice&#xff1b;在github上的邮箱是bobfoxmail.com、用户名为bob。 账号…
阅读更多...
【Java 进阶篇】Redis持久化之RDB：数据的安全守护者

【Java 进阶篇】Redis持久化之RDB：数据的安全守护者

Redis&#xff0c;作为一款高性能的键值存储系统&#xff0c;支持多种持久化方式&#xff0c;其中RDB&#xff08;Redis DataBase&#xff09;是其最常用的一种。RDB可以将当前时刻的数据快照保存到磁盘&#xff0c;以便在Redis重启时快速恢复数据。本文将深入探讨RDB的原理、配…
阅读更多...
一文讲明Mybatis 的使用 超详细 【爆肝两万字教程】

一文讲明Mybatis 的使用 超详细 【爆肝两万字教程】

我 | 在这里 &#x1f575;️ 读书 | 长沙 ⭐软件工程 ⭐ 本科 &#x1f3e0; 工作 | 广州 ⭐ Java 全栈开发&#xff08;软件工程师&#xff09; &#x1f383; 爱好 | 研究技术、旅游、阅读、运动、喜欢流行歌曲 &#x1f3f7;️ 标签 | 男 自律狂人 目标明确 责任心强 ✈️公…
阅读更多...
Linux进程通信——信号（一）

Linux进程通信——信号（一）

原理 对于 Linux来说&#xff0c;实际信号是软中断&#xff0c;许多重要的程序都需要处理信号。 信号&#xff0c;为 Linux 提供了一种处理异步事件的方法。比如&#xff0c;终端用户输入了ctrlc来中断程序&#xff0c;会通过信号机制停止一个程序。 概述 信号的名字和编号 …
阅读更多...
【CCF-PTA】第03届Scratch第04题 -- 数字加密

【CCF-PTA】第03届Scratch第04题 -- 数字加密

数字加密 【题目描述】 "狼群战术"是第二次世界大战中德军对大西洋上盟军商船所使用的潜艇战术&#xff0c;一度遏制住英国的海上贸易。直到艾伦图灵成功破译了德国的英格尔码密码&#xff0c;成为二战的一个重要转折点。时至今日&#xff0c;图灵仍然是计算机的一…
阅读更多...
基于JavaWeb+SSM+Vue微信阅读小程序的设计和实现

基于JavaWeb+SSM+Vue微信阅读小程序的设计和实现

基于JavaWebSSMVue微信阅读小程序的设计和实现 源码获取入口Lun文目录前言主要技术系统设计功能截图订阅经典源码专栏[Java 源码获取 源码获取入口 Lun文目录 第1章 绪论 1 1.1 课题背景 1 1.2 课题意义 1 1.3 研究内容 1 第2章 开发环境与技术 3 2.1 MYSQL数据库 3 2.2 JSP技…
阅读更多...
Proteus下仿真AT89C51报“串行口通信失败，请检查电平适配是否正确。”解决办法

Proteus下仿真AT89C51报“串行口通信失败，请检查电平适配是否正确。”解决办法

在Proteus下进行AT89C51串行口仿真时&#xff0c;如果遇到“串行口通信失败&#xff0c;请检查电平适配是否正确”的错误提示&#xff0c;以下是一些解决办法&#xff1a; 1. 了解AT89C51和外部设备的电平要求&#xff1a; 首先&#xff0c;了解AT89C51和外部设备之间的电平…
阅读更多...
【Pytorch】Visualization of Feature Maps（3）

【Pytorch】Visualization of Feature Maps（3）

学习参考来自&#xff1a; Image Style Transform–关于图像风格迁移的介绍github&#xff1a;https://github.com/wmn7/ML_Practice/tree/master/2019_06_03 文章目录 风格迁移 风格迁移 风格迁移出处&#xff1a; 《A Neural Algorithm of Artistic Style》&#xff08;ar…
阅读更多...
Peter算法小课堂—前缀和数组的应用

Peter算法小课堂—前缀和数组的应用

桶 相当于计数排序&#xff0c;看一个视频 桶排序 太戈编程1620题 算法解析 #include <bits/stdc.h> using namespace std; const int R11; int cnt[R];//cnt[t]代表第t天新增几人 int s[R];//s[]数组是cnt[]数组的前缀和数组 int n,t; int main(){cin>>n;for(…
阅读更多...
rabbitMQ发布确认-交换机不存在或者无法抵达队列的缓存处理

rabbitMQ发布确认-交换机不存在或者无法抵达队列的缓存处理

rabbitMQ在发送消息时&#xff0c;会出现交换机不存在&#xff08;交换机名字写错等消息&#xff09;&#xff0c;这种情况如何会退给生产者重新处理&#xff1f;【交换机层】 生产者发送消息时&#xff0c;消息未送达到指定的队列&#xff0c;如何消息回退&#xff1f; 核心&…
阅读更多...
企业数字化转型的作用是什么？_光点科技

企业数字化转型的作用是什么？_光点科技

在当今快速变化的商业环境中&#xff0c;数字化转型已成为企业发展的重要策略。企业数字化转型指的是利用数字技术改造传统业务模式和管理方式&#xff0c;以提升效率、增强竞争力和创造新的增长机会。 提升运营效率&#xff1a;数字化转型通过引入自动化工具和智能系统&#x…
阅读更多...
数据分析基础之《matplotlib（2）—折线图》

数据分析基础之《matplotlib（2）—折线图》

一、折线图绘制与保存图片 1、matplotlib.pyplot模块 matplotlib.pyplot包含了一系列类似于matlab的画图函数。它的函数作用于当前图形&#xff08;figure&#xff09;的当前坐标系&#xff08;axes&#xff09; import matplotlib.pyplot as plt 2、折线图绘制与显示 展示城…
阅读更多...
基于springboot实现校园在线拍卖系统项目【项目源码】

基于springboot实现校园在线拍卖系统项目【项目源码】

基于springboot实现校园在线拍卖系统演示 Javar技术 JavaScript是一种网络脚本语言&#xff0c;广泛运用于web应用开发&#xff0c;可以用来添加网页的格式动态效果&#xff0c;该语言不用进行预编译就直接运行&#xff0c;可以直接嵌入HTML语言中&#xff0c;写成js语言&…
阅读更多...
ubuntu22.04安装wvp-gb28181-pro 2023-11-23最新版本（一键安装）

ubuntu22.04安装wvp-gb28181-pro 2023-11-23最新版本（一键安装）

下载程序 输入下面命令&#xff0c;输入普通用户密码&#xff0c;切换到 root用户 sudo su git clone -b ubuntu_wvp_online_install_2023_0425 https://gitcode.net/zenglg/ubuntu_wvp_online_install.git 等待下载完成 安装 进入到克隆下来的路径中 cd /home/tuners/ub…
阅读更多...
HTTPS攻击怎么防御？

HTTPS攻击怎么防御？

HTTPS 简介 超文本传输安全协议&#xff08; HTTPS &#xff09;是一种通过计算机网络进行安全通信的传输协议。HTTPS 经由 HTTP 进行通信&#xff0c;但利用 SSL/TLS 来加密数据包。 HTTPS 开发的主要目的&#xff0c;是提供对网站服务器的身份认证&#xff0c;保护交换数据的…
阅读更多...
搭配：基于OpenCV的边缘检测实战

搭配：基于OpenCV的边缘检测实战

引言 计算机中的目标检测与人类识别物体的方式相似。作为人类&#xff0c;我们可以分辨出狗的形象&#xff0c;因为狗的特征是独特的。尾巴、形状、鼻子、舌头等特征综合在一起&#xff0c;帮助我们把狗和牛区分开来。 同样&#xff0c;计算机能够通过检测与估计物体的结构和性…
阅读更多...
掌握未来技术趋势，成为领先者——深度解析2023年技术热点

掌握未来技术趋势，成为领先者——深度解析2023年技术热点

掌握未来技术趋势&#xff0c;成为领先者——深度解析2023年技术热点 摘要&#xff1a;本文探讨当前最热门的技术趋势。我们将介绍人工智能、大数据、区块链、5G等前沿技术&#xff0c;并阐述它们如何改变我们的生活。最后&#xff0c;我们将总结如何利用这些技术趋势&#xf…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023