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);
	// 在设备内存中分配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;
}

在这里插入图片描述

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

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

CVE-2022-0543（Redis 沙盒逃逸漏洞）

CVE-2022-0543（Redis 沙盒逃逸漏洞）

简介 CVE-2022-0543是一个与Redis相关的安全漏洞。在Redis中&#xff0c;用户连接后可以通过eval命令执行Lua脚本&#xff0c;但在沙箱环境中脚本无法执行命令或读取文件。然而&#xff0c;攻击者可以利用Lua沙箱中遗留的变量package的loadlib函数来加载动态链接库liblua5.1.s…
阅读更多...
jenkins 参数构建

jenkins 参数构建

应用保存 [rootjenkins-node1 .ssh]# ssh-keygen Generating public/private rsa key pair. Enter file in which to save the key (/root/.ssh/id_rsa): Enter passphrase (empty for no passphrase): Enter same passphrase again: Your identification has been saved i…
阅读更多...
利用GUI实现渲染二维码效果

利用GUI实现渲染二维码效果

以下是一个简单的 Java 验证码实现示例&#xff1a; import java.awt.Color; import java.awt.Font; import java.awt.Graphics; import java.awt.image.BufferedImage; import java.util.Random;import javax.imageio.ImageIO;public class CaptchaGenerator {public static …
阅读更多...
C++基础(4)——类与对象(默认成员函数)

C++基础(4)——类与对象(默认成员函数)

目录 1.拷贝构造函数&#xff1a; 1.1 为什么要引入拷贝构造&#xff1a; 1.2 拷贝构造函数的定义及特性&#xff1a; 1.3 什么类可以不用编写拷贝构造&#xff1a; 2. 赋值运算符重载&#xff1a; 2.1 为社么要引入运算符重载&#xff1a; 2.2运算符重载的定义以及特性…
阅读更多...
轻松管理文件名：文件批量重命名的技巧与操作

轻松管理文件名：文件批量重命名的技巧与操作

在日常工作中&#xff0c;文件管理是一项至关重要的任务。其中&#xff0c;文件名的管理更是关键。文件名是在查找文件时最直观的线索。一个好的文件名简短而准确地反映文件的内容或用途。然而&#xff0c;随着时间的推移&#xff0c;可能会发现文件名变得冗长、混乱甚至无法反…
阅读更多...
指针变量与指针类型的深入理解

指针变量与指针类型的深入理解

1.知识总结 相关代码展示 #include <stdio.h> int main() {int n 0x11223344;int *pi &n; *pi 0; return 0; } #include <stdio.h> int main() {int n 0x11223344;char *pc (char *)&n;*pc 0;return 0; } #include <stdio.h> int main() {i…
阅读更多...
SSM家具个性定制管理系统开发mysql数据库web结构java编程计算机网页源码eclipse项目

SSM家具个性定制管理系统开发mysql数据库web结构java编程计算机网页源码eclipse项目

一、源码特点 SSM 家具个性定制管理系统是一套完善的信息系统&#xff0c;结合springMVC框架完成本系统&#xff0c;对理解JSP java编程开发语言有帮助系统采用SSM框架&#xff08;MVC模式开发&#xff09;&#xff0c;系统具有完整的源代码和数据库&#xff0c;系统主要采用…
阅读更多...
【jvm】虚拟机之堆

【jvm】虚拟机之堆

目录 一、堆的核心概述二、堆的内存细分&#xff08;按分代收集理论设计&#xff09;2.1 java7及以前2.2 java8及以后 三、堆内存大小3.1 说明3.2 参数设置3.3 默认大小3.4 手动设置3.5 jps3.6 jstat3.7 OutOfMemory举例 四、年轻代与老年代4.1 说明 五、对象分配过程5.1 说明5…
阅读更多...
Jackson无缝替换Fastjson

Jackson无缝替换Fastjson

目录 文章目录 一&#xff0c;Fastjson到Jackson的替换方案方案代码序列化反序列化通过key获取某种类型的值类型替换 二&#xff0c;Springboot工程中序列化的使用场景三&#xff0c;SpringMVC框架中的Http消息转换器1&#xff0c;原理&#xff1a;2&#xff0c;自定义消息转换…
阅读更多...
Visio学习笔记

Visio学习笔记

1. 常用素材 1.1 立方体&#xff1a;张量, tensor 操作路径&#xff1a;更多形状 ⇒ 常规 ⇒ 基本形状 自动配色 在选择【填充】后Visio会自动进行配色&#xff1b;
阅读更多...
【C/C++】排序算法代码实现

【C/C++】排序算法代码实现

这里&#xff0c;汇总了常见的排序算法具体代码实现。使用C语言编写。 排序算法实现 插入排序冒泡排序选择排序快速排序希尔排序归并排序 插入排序 #include <stdio.h> #include <stdlib.h>void InsertSort(int arr[],int n){int i,j,temp;for(i 1;i < n;i){ …
阅读更多...
初识Java 18-3 泛型

初识Java 18-3 泛型

目录 边界 通配符 编译器的能力范畴 逆变性 无界通配符 捕获转换 本笔记参考自&#xff1a; 《On Java 中文版》 边界 在泛型中&#xff0c;边界的作用是&#xff1a;在参数类型上增加限制。这么做可以强制执行应用泛型的类型规则&#xff0c;但还有一个更重要的潜在效果…
阅读更多...
vue el-table （固定列+滚动列）【横向滚动条】确定滚动条是在列头还是列尾

vue el-table （固定列+滚动列）【横向滚动条】确定滚动条是在列头还是列尾

效果图&#xff1a; 代码实现&#xff1a; html&#xff1a; <script src"//unpkg.com/vue2/dist/vue.js"></script> <script src"//unpkg.com/element-ui2.15.14/lib/index.js"></script> <div id"app" style&quo…
阅读更多...
实战JVM高CPU、内存问题分析定位

实战JVM高CPU、内存问题分析定位

背景&#xff1a; 业务中台组件MOSC开展压测工作&#xff0c;并发场景下发现CPU使用率达到100%&#xff0c;虽然程序没有报错&#xff0c;但是这种情况显然已经达到性能瓶颈&#xff0c;对服务带来了验证的效能影响&#xff0c;所以针对该CPU问题必须进行详细的根因分析处理。…
阅读更多...
浅谈Python中的鸭子类型和猴子补丁

浅谈Python中的鸭子类型和猴子补丁

文章目录 前言一、鸭子类型二、猴子补丁关于Python技术储备一、Python所有方向的学习路线二、Python基础学习视频三、精品Python学习书籍四、Python工具包项目源码合集①Python工具包②Python实战案例③Python小游戏源码五、面试资料六、Python兼职渠道 前言 Python 开发者可能…
阅读更多...
AdaBoost提升分类器性能

AdaBoost提升分类器性能

目录 AdaBoost算法原理 AdaBoost工作详情 初始权重分配 第一轮 第二轮 后续轮次 最终模型 AdaBoost的API解释 AdaBoost 对房价进行预测 AdaBoost 与决策树模型的比较 结论 AdaBoost算法原理 在数据挖掘中&#xff0c;分类算法可以说是核心算法&#xff0c;其中 Ada…
阅读更多...
如何应用ChatGPT撰写、修改论文及工作报告，提供写作能力及优化工作？？

如何应用ChatGPT撰写、修改论文及工作报告，提供写作能力及优化工作？？

如果我想让gpt从pdf文档中提取相关关键词的内容&#xff0c;可以怎么做呢&#xff1f;&#xff1f;我们评论区讨论 ChatGPT 在论文写作与编程方面也具备强大的能力。无论是进行代码生成、错误调试还是解决编程难题&#xff0c;ChatGPT都能为您提供实用且高质量的建议和指导&am…
阅读更多...
flink和机器学习模型的常用组合方式

flink和机器学习模型的常用组合方式

背景 flink是一个低延迟高吞吐的系统&#xff0c;每秒处理的数据量高达数百万&#xff0c;而机器模型一般比较笨重&#xff0c;虽然功能强大&#xff0c;但是qps一般都比较低&#xff0c;日常工作中&#xff0c;我们一般是如何把flink和机器学习模型组合起来一起使用呢? fli…
阅读更多...
9.Docker的虚悬镜像-Dangling Image

9.Docker的虚悬镜像-Dangling Image

1.虚悬镜像的概念 虚悬镜像 (Dangling Image) 指的是仓库名 (镜像名) 和标签 TAG 都是 的镜像。 2.构建本地虚悬镜像 这里我以unbuntu为例来说明。 2.1 编写Dockerfile文件 FROM ubuntu:22.042.2 根据Dockerfile文件构建虚悬镜像 docker build .上面这段命令&#xff0c…
阅读更多...
C#开发的OpenRA游戏之属性RenderSprites（8）

C#开发的OpenRA游戏之属性RenderSprites（8）

C#开发的OpenRA游戏之属性RenderSprites（8） 本文开始学习RenderSprites属性，这个属性是跟渲染有关的，因此它就摄及颜色相关的内容，所以我们先来学习一下调色板，这是旧游戏的图片文件保存的格式，如果放在现代来看，不会再采用这种方法，毕竟现在存储空间变大，便宜了，并…
阅读更多...
最新文章

Copyright @ 2022~2023