Cuda编程——使用share memory优化矩阵乘法

在上一篇文章:第一个Cuda程序,矩阵相乘代码,我们设计了一种并行的矩阵乘法程序,效果和使用CPU计算的一样,但时间有了很大的降低,然而,这只是最基本的一种方法,事实上我们完全可以让程序变得更快!

仔细看看,会发现我们使用的是global memory,而share memory的访问速度要远远大于global memory,所以我们将使用share memory优化矩阵乘法,让程序更快!

#include <stdio.h>#include<stdlib.h>#include <cuda.h>#include <cuda_runtime.h>#include <device_launch_parameters.h>#include <device_functions.h>
# define BLOCK_SIZE 8# define M 6# define N 8# define K 6__managed__ float a[M * N];__managed__ float b[N*K];__managed__ float c_gpu[M * K];__managed__ float c_cpu[M * K];__global__ void gpu_matrix(float* a, float* b, float* c, const int m, const int n, const int k){  __shared__ float sub_a[BLOCK_SIZE][BLOCK_SIZE];  __shared__ float sub_b[BLOCK_SIZE][BLOCK_SIZE];  int x = threadIdx.x + blockIdx.x * blockDim.x;  int y = threadIdx.y + blockIdx.y * blockDim.y;  float temp = 0.0;  int step, i;  for (step = 0; step <N / BLOCK_SIZE; step++)  {    if ((step * BLOCK_SIZE + threadIdx.x) >= N || y >= M)    {      sub_a[threadIdx.y][threadIdx.x] = 0.0;    }    else    {      sub_a[threadIdx.y][threadIdx.x] = a[y * N + (step * BLOCK_SIZE + threadIdx.x)];    }
    if ((step * BLOCK_SIZE + threadIdx.y) >= N || x >= K)    {      sub_b[threadIdx.y][threadIdx.x] = 0.0;
    }    else    {      sub_b[threadIdx.y][threadIdx.x] = b[(step * BLOCK_SIZE + threadIdx.y) * K + x];    }    __syncthreads();    for (i = 0; i < BLOCK_SIZE; i++)    {      temp = temp + sub_a[threadIdx.y][i] * sub_b[i][threadIdx.x];    }    __syncthreads();
    if (x < K && y < M)    {      c[y * K + x] = temp;    }  }}
void cpu_matrix(float* a, float* b, float* c, const int m, const int n, const int k){  int y, x, step;  float temp;  for (y = 0; y < M; y++)  {    for (x = 0; x < K; x++)    {      temp = 0;      for (step = 0; step < N; step++)      {        temp += a[y * N + step] * b[step * K + x];      }      c[y * K + x] = temp;    }  }}
int main(){  int x,y;  float item1;  //初始化矩阵  for (y = 0; y < M; y++)  {    for (x = 0; x < N; x++)    {      item1 = x + y;      a[y * N + x] = item1;
    }
  }  for (y = 0; y < N; y++)  {    for (x = 0; x < K; x++)    {      item1 = x + y;      b[y * K + x] = item1;
    }  }  printf("-----------------两个矩阵-#--------------\n");  for (y = 0; y < M; y++)  {
    for (x = 0; x < N; x++)    {      printf("%f ", a[y * N + x]);    }    printf("\n");  }  for (y = 0; y < N; y++)  {    for (x = 0; x < K; x++)    {      printf("%f  ", b[y * K + x]);    }    printf("\n");  }  unsigned int grid_rows = (M + BLOCK_SIZE - 1) / BLOCK_SIZE;  unsigned int grid_cols = (K + BLOCK_SIZE - 1) / BLOCK_SIZE;
  dim3 dimGrid(grid_rows, grid_cols);  dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  gpu_matrix <<<dimGrid, dimBlock>>> (a, b, c_gpu, M, N, K);
  cudaDeviceSynchronize();
  cpu_matrix(a, b, c_cpu, M, N, K);  //  //打印cpu计算结果  printf("---------------cpu计算结果---------#------------------\n");
  for (y = 0; y < M; y++)  {    for (x = 0; x < K; x++)    {      printf("%f  ", c_cpu[y * K + x]);    }    printf("\n");
  }
  printf("------------------gpu计算结果------#------------------\n");  //打印GPU计算结果  for (y = 0; y < M; y++)  {    for (x = 0; x < K; x++)    {      printf("%f  ", c_gpu[y * K + x]);    }    printf("\n");  }  return 0;}
-----------------两个矩阵-#--------------0.000000 1.000000 2.000000 3.000000 4.000000 5.000000 6.000000 7.0000001.000000 2.000000 3.000000 4.000000 5.000000 6.000000 7.000000 8.0000002.000000 3.000000 4.000000 5.000000 6.000000 7.000000 8.000000 9.0000003.000000 4.000000 5.000000 6.000000 7.000000 8.000000 9.000000 10.0000004.000000 5.000000 6.000000 7.000000 8.000000 9.000000 10.000000 11.0000005.000000 6.000000 7.000000 8.000000 9.000000 10.000000 11.000000 12.0000000.000000  1.000000  2.000000  3.000000  4.000000  5.0000001.000000  2.000000  3.000000  4.000000  5.000000  6.0000002.000000  3.000000  4.000000  5.000000  6.000000  7.0000003.000000  4.000000  5.000000  6.000000  7.000000  8.0000004.000000  5.000000  6.000000  7.000000  8.000000  9.0000005.000000  6.000000  7.000000  8.000000  9.000000  10.0000006.000000  7.000000  8.000000  9.000000  10.000000  11.0000007.000000  8.000000  9.000000  10.000000  11.000000  12.000000---------------cpu计算结果---------#------------------140.000000  168.000000  196.000000  224.000000  252.000000  280.000000168.000000  204.000000  240.000000  276.000000  312.000000  348.000000196.000000  240.000000  284.000000  328.000000  372.000000  416.000000224.000000  276.000000  328.000000  380.000000  432.000000  484.000000252.000000  312.000000  372.000000  432.000000  492.000000  552.000000280.000000  348.000000  416.000000  484.000000  552.000000  620.000000------------------gpu计算结果------#------------------140.000000  168.000000  196.000000  224.000000  252.000000  280.000000168.000000  204.000000  240.000000  276.000000  312.000000  348.000000196.000000  240.000000  284.000000  328.000000  372.000000  416.000000224.000000  276.000000  328.000000  380.000000  432.000000  484.000000252.000000  312.000000  372.000000  432.000000  492.000000  552.000000280.000000  348.000000  416.000000  484.000000  552.000000  620.000000

到了这里,我们能够使得矩阵乘法变得相当快(与仅使用CPU计算相比),这在实际应用中非常重要,尤其是数据计算量非常大的情况。

也许到了这里,这两个程序你并没有完全了解,但,不要担心,先把这些代码运行一下,体会使用GPU计算的魅力,为以后的学习打下基础。

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

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

相关文章

服务器网络安全防护措施有哪些?

由于服务器发挥着至关重要的作用&#xff0c;因此存储在服务器上的机密数据和信息非常具有价值。如今有一种流行的说法&#xff0c;“数据就是新的石油”。 如果不确定如何保护服务器安全&#xff0c;或者不确定是否已涵盖所有基础知识&#xff0c;那么可以了解下面提供一些可…

Pandas实战100例 | 案例 10: 应用函数 - 使用 `apply`

案例 10: 应用函数 - 使用 apply 知识点讲解 Pandas 的 apply 函数是一个非常强大的工具&#xff0c;允许你对 DataFrame 中的行或列应用一个函数。这对于复杂的数据转换和计算非常有用。你可以使用 apply 来执行任意的函数&#xff0c;这些函数可以是自定义的&#xff0c;也…

Unity游戏图形学 Shader结构

shader结构 shader语言 openGL&#xff1a;SLG跨平台 >GLSL&#xff1a;openGL shaderlauguge DX&#xff1a;微软开发&#xff0c;性能很好&#xff0c;但是不能跨平台 >HLSL&#xff1a;high level shader language CG&#xff1a;微软和Nvidia公司联合开发&#xff…

open3d相关操作总结

open3d其实有很多交互式命令&#xff0c;在运行程序打开了open3d渲染的窗口后&#xff0c;鼠标点击窗口&#xff0c;按H就会弹出&#xff0c;交互命令的帮助&#xff0c;如下图所示&#xff1a; 其中比较常用的有&#xff1a; Q &#xff1a;退出当前窗口 H&#xff1a;打印帮…

5 - 异常处理

目录 1. 总览 1.1 Exception 与 Error 1.2 checked unchecked 异常 1&#xff09;使用 try-catch 进行捕获 2&#xff09;使用 throws 关键字抛出 1.3 throw 与 throws 1&#xff09;throw 2&#xff09;throws 3&#xff09;区别 1.4 try-catch-finally 2. try wit…

Airflow大揭秘:如何让大数据任务调度变得简单高效?

介绍&#xff1a;Airflow是一个开源的、用于创建、调度和监控数据管道的工作流平台。这个平台使用Python编写&#xff0c;并通过有向无环图&#xff08;Directed Acyclic Graph, DAG&#xff09;来管理任务流程&#xff0c;使得用户不需要知道业务数据的具体内容&#xff0c;只…

Python爬虫学习笔记(一)---Python入门

一、pycharm的安装及使用二、python的基础使用1、字符串连接2、单双引号转义3、换行4、三引号跨行字符串5、命名规则6、注释7、 优先级not>and>or8、列表&#xff08;list&#xff09;9、字典&#xff08;dictionary&#xff09;10、元组&#xff08;tuple&#xff09;11…

SDRAM小项目——写模块

写模块跟着视频看了一个多星期&#xff0c;一开始始终有点弄不清楚&#xff0c;现在记录一下理解的过程。 阅读文档信息&#xff1a; 首先阅读文档信息&#xff0c;了解SDRAM写过程的状态转换和时序图 SDRAM整体状态流程如图所示&#xff1a; 在SDRAM整体系统中&#xff0c…

【算法小课堂】动态规划

动态规划 动态规划相信大家都知道&#xff0c;动态规划算法也是新手在刚接触算法设计时很苦恼的问题&#xff0c;有时候觉得难以理解&#xff0c;但是真正理解之后&#xff0c;就会觉得动态规划其实并没有想象中那么难。网上也有很多关于讲解动态规划的文章&#xff0c;大多都…

Java--业务场景:在Spring项目启动时加载Java枚举类到Redis中(补充)

文章目录 前言步骤测试结果 前言 通过Java–业务场景&#xff1a;在Spring项目启动时加载Java枚举类到Redis中,我们成功将Java项目里的枚举类加载到Redis中了&#xff0c;接下来我们只需要写接口获取需要的枚举值数据就可以了&#xff0c;下面一起来编写这个接口吧。 步骤 在…

mysql-bin日志清理,并设置expire_logs_days时间,mysql占用空间过大问题

mysql-bin日志清理&#xff0c;并设置expire_logs_days时间&#xff0c;mysql占用空间过大问题 文章目录 问题查看mysql配置参数解决全局修改参数清理日志规则手动清理my.cnf 外传 问题 最近发现生产环境的服务器磁盘空间吃紧&#xff0c;查下到底是哪里占用的空间比较大&…

leetcode238:除自身以外数组的乘积

文章目录 1.使用除法&#xff08;违背题意&#xff09;2.左右乘积列表3.空间复杂度为O(1)的方法 在leetcode上刷到了这一题&#xff0c;一开始并没有想到好的解题思路&#xff0c;写篇博客再来梳理一下吧。 题目要求&#xff1a; 不使用除法在O(n)时间复杂度内 1.使用除法&am…

Tomcat Notes: URL Mapping

This is a personal study notes of Apache Tomcat. Below are main reference material. - YouTube Apache Tomcat Full Tutorial&#xff0c;owed by Alpha Brains Courses. https://www.youtube.com/watch?vrElJIPRw5iM&t801s 1、URL Mapping To Resources1.1、What w…

新一代数字原住民:市场痛点与“繁”思维应对之道

随着科技的迅速发展&#xff0c;尤其是互联网的普及&#xff0c;新一代数字原住民经营者已经逐渐成为市场的主力军。不同于传统的消费者&#xff0c;有着独特的消费习惯和心理需求。企业要在这激烈的市场竞争中获得优势&#xff0c;深入了解这一群体的特征和心理、行为&#xf…

有趣的事,讲给有趣的人听

哈哈哈&#xff0c;今天不写技术了&#xff0c;今天分享一下生活&#xff0c;技术我们什么时候都可以学&#xff0c;但是生活更值得我们现在就去更好的体验&#xff01; 两年多的涤生大数据&#xff0c;认识了形形色色的小伙伴&#xff0c;陆续沟通下来6000多人&#xff0c;彼时…

数据库锁表原因、排查、解决

一.场景 场景1场景2二.原因三.排查四.解决方案 一.场景 场景1 锁表通常发生在DML&#xff08; insert 、update 、delete &#xff09; A操作进行全量数据同步&#xff0c;对整个表的粒度进行上锁&#xff0c;导致B操作只能等待A操作完成才能进入插入数据。此时就出现了锁表…

Pandas实战100例 | 案例 14: 数据透视表 - 使用 `pivot_table`

案例 14: 数据透视表 - 使用 pivot_table 知识点讲解 数据透视表是一种常见的数据汇总工具&#xff0c;用于按照一个或多个键对数据进行分类汇总。Pandas 的 pivot_table 函数提供了一种快速创建数据透视表的方法。你可以指定行索引、列索引&#xff0c;以及用于聚合的数据和…

Elasticsearch windows开箱即用【记录】

一、准备工作 安装ES之前要在本机安装好JDK&#xff0c;对应的兼容性见官网链接&#xff1a;https://www.elastic.co/cn/support/matrix ES官网链接&#xff1a;https://www.elastic.co/cn/, 我本机安装的是JDK8&#xff0c;测试使用的是7.3.0版本的ES和Kibana。 1、首先去…

Windows平台程序和Android平台程序的差异

Windows平台程序和Android平台程序的差异 1 Windows平台环境和Android平台JVM虚拟机的差异&#xff1a; 1&#xff09;由于JVM虚拟机上的数据是大端处理的&#xff0c;而Windows平台上的数据是小端的&#xff0c;所以在一些数据的处理上需要进行转换&#xff1b; 2&#xf…

vmware创建嵌套虚拟机

嵌套虚拟机的搭建 在vmware 虚拟机设置中&#xff0c;打开处理器的虚拟化Intel VT-x/EPT 或AMD-V/RVI&#xff08;v&#xff09;配置虚拟机yum 源&#xff0c;安装qemu、qemu-kvm、libvirt从阿里镜像源下载centos iso 阿里源 centos-7-x86准备虚拟机创建所需xml&#xff0c;ce…