【cuda学习日记】5.2.1 共享内存额外篇

共享内存(Shared Memory)
1.是一种低延迟、高带宽的片上内存
2.由同一个Block内的所有线程共享
3.生命周期与Block相同
4.访问速度比全局内存快约100倍

Block(线程块)
1.GPU执行的基本单位，包含一组线程
2.多个Block组成Grid(网格)
3.Block内的线程可以通过共享内存通信
4.Block之间是独立执行的

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <cuda_runtime.h>#define CHECK(call) \{\const cudaError_t error = call; \if (error != cudaSuccess)\{\printf("Error: %s: %d\n", __FILE__, __LINE__);\printf("code :%d reason :%s\n", error , cudaGetErrorString(error));\exit(1);\}\
}void initialInt( int * ip, int size)
{for (int i =0; i < size; i ++){ip[i] = i;}
}void printMatrix(int *C, const int nx, const int ny)
{int *ic = C;printf("\n matrix : (%d, %d)\n", nx, ny);for (int iy = 0; iy < ny; iy++){for (int ix =0; ix < nx; ix++){printf("%3d",ic[ix]);}ic += nx;printf("\n");}printf("\n");
}__global__ void printThreadIndex(int *A, const int nx, const int ny)
{int bx = blockIdx.x;int by = blockIdx.y;int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int tx  = threadIdx.x;int ty = threadIdx.y;unsigned int idx = iy*nx + ix;const int BM = 2; const int BN = 4;__shared__ float smem[BM][BN];smem[ty][tx] = float(A[idx]);printf("threadidx: (%d ,%d) blockidx:(%d ,%d) coordinate: (%d ,%d) global index: (%2d ival %2d), smem val (%f) \n", threadIdx.x, threadIdx.y,blockIdx.x, blockIdx.y,ix, iy,idx, A[idx],//smem[ty][tx]smem[0][0]);}int main(int argc , char **argv)
{printf("%s starting\n", argv[0]);int dev = 0;cudaDeviceProp deviceprop;CHECK(cudaGetDeviceProperties(&deviceprop,dev));printf("Using Device %d : %s\n", dev, deviceprop.name);CHECK(cudaSetDevice(dev));// set matrix int nx = 8;int ny = 6;int nxy = nx * ny;int nBytes = nxy * sizeof(float);// malloc host memoryint * h_A;h_A = (int *) malloc(nBytes);//initial intinitialInt(h_A, nxy);printMatrix(h_A, nx, ny);// deviceint *d_MatA;cudaMalloc((void **)&d_MatA, nBytes);cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);dim3 block(4,2);dim3 grid ((nx + block.x - 1)/block.x, (ny + block.y - 1)/ block.y);printf("execution config grid (%d, %d), block (%d, %d)\n", grid.x, grid.y, block.x, block.y);printThreadIndex<<<grid, block>>>(d_MatA, nx, ny);cudaDeviceSynchronize();cudaFree(d_MatA);free(h_A);cudaDeviceReset();return 0;
}

示例代码中， block的大小是（4，2）, 所以在核函数中声明4x2大小的 SMEM, 只需要一次load操作，则8个线程会将数据load进 4x2大小的SMEM里。
在做printf的时候，因为SMEM对block 可见，所以访问SMEM[0][0] 打印出来的都是block里第一个线程load进去的数据。

输出如下：

matrix : (8, 6)0  1  2  3  4  5  6  78  9 10 11 12 13 14 1516 17 18 19 20 21 22 2324 25 26 27 28 29 30 3132 33 34 35 36 37 38 3940 41 42 43 44 45 46 47execution config grid (2, 3), block (4, 2)
threadidx: (0 ,0) blockidx:(0 ,1) coordinate: (0 ,2) global index: (16 ival 16), smem val (16.000000) 
threadidx: (1 ,0) blockidx:(0 ,1) coordinate: (1 ,2) global index: (17 ival 17), smem val (16.000000)
threadidx: (2 ,0) blockidx:(0 ,1) coordinate: (2 ,2) global index: (18 ival 18), smem val (16.000000)
threadidx: (3 ,0) blockidx:(0 ,1) coordinate: (3 ,2) global index: (19 ival 19), smem val (16.000000)
threadidx: (0 ,1) blockidx:(0 ,1) coordinate: (0 ,3) global index: (24 ival 24), smem val (16.000000)
threadidx: (1 ,1) blockidx:(0 ,1) coordinate: (1 ,3) global index: (25 ival 25), smem val (16.000000)
threadidx: (2 ,1) blockidx:(0 ,1) coordinate: (2 ,3) global index: (26 ival 26), smem val (16.000000)
threadidx: (3 ,1) blockidx:(0 ,1) coordinate: (3 ,3) global index: (27 ival 27), smem val (16.000000)
threadidx: (0 ,0) blockidx:(1 ,1) coordinate: (4 ,2) global index: (20 ival 20), smem val (20.000000)
threadidx: (1 ,0) blockidx:(1 ,1) coordinate: (5 ,2) global index: (21 ival 21), smem val (20.000000)
threadidx: (2 ,0) blockidx:(1 ,1) coordinate: (6 ,2) global index: (22 ival 22), smem val (20.000000)
threadidx: (3 ,0) blockidx:(1 ,1) coordinate: (7 ,2) global index: (23 ival 23), smem val (20.000000)
threadidx: (0 ,1) blockidx:(1 ,1) coordinate: (4 ,3) global index: (28 ival 28), smem val (20.000000)
threadidx: (1 ,1) blockidx:(1 ,1) coordinate: (5 ,3) global index: (29 ival 29), smem val (20.000000)
threadidx: (2 ,1) blockidx:(1 ,1) coordinate: (6 ,3) global index: (30 ival 30), smem val (20.000000)
threadidx: (3 ,1) blockidx:(1 ,1) coordinate: (7 ,3) global index: (31 ival 31), smem val (20.000000)
threadidx: (0 ,0) blockidx:(1 ,0) coordinate: (4 ,0) global index: ( 4 ival  4), smem val (4.000000)
threadidx: (1 ,0) blockidx:(1 ,0) coordinate: (5 ,0) global index: ( 5 ival  5), smem val (4.000000)
threadidx: (2 ,0) blockidx:(1 ,0) coordinate: (6 ,0) global index: ( 6 ival  6), smem val (4.000000)
threadidx: (3 ,0) blockidx:(1 ,0) coordinate: (7 ,0) global index: ( 7 ival  7), smem val (4.000000)
threadidx: (0 ,1) blockidx:(1 ,0) coordinate: (4 ,1) global index: (12 ival 12), smem val (4.000000)
threadidx: (1 ,1) blockidx:(1 ,0) coordinate: (5 ,1) global index: (13 ival 13), smem val (4.000000)
threadidx: (2 ,1) blockidx:(1 ,0) coordinate: (6 ,1) global index: (14 ival 14), smem val (4.000000)
threadidx: (3 ,1) blockidx:(1 ,0) coordinate: (7 ,1) global index: (15 ival 15), smem val (4.000000)
threadidx: (0 ,0) blockidx:(0 ,2) coordinate: (0 ,4) global index: (32 ival 32), smem val (32.000000)
threadidx: (1 ,0) blockidx:(0 ,2) coordinate: (1 ,4) global index: (33 ival 33), smem val (32.000000)
threadidx: (1 ,0) blockidx:(0 ,2) coordinate: (1 ,4) global index: (33 ival 33), smem val (32.000000)
threadidx: (2 ,0) blockidx:(0 ,2) coordinate: (2 ,4) global index: (34 ival 34), smem val (32.000000)
threadidx: (3 ,0) blockidx:(0 ,2) coordinate: (3 ,4) global index: (35 ival 35), smem val (32.000000)
threadidx: (0 ,1) blockidx:(0 ,2) coordinate: (0 ,5) global index: (40 ival 40), smem val (32.000000)
threadidx: (1 ,1) blockidx:(0 ,2) coordinate: (1 ,5) global index: (41 ival 41), smem val (32.000000)
threadidx: (0 ,1) blockidx:(0 ,2) coordinate: (0 ,5) global index: (40 ival 40), smem val (32.000000)
threadidx: (1 ,1) blockidx:(0 ,2) coordinate: (1 ,5) global index: (41 ival 41), smem val (32.000000)
threadidx: (2 ,1) blockidx:(0 ,2) coordinate: (2 ,5) global index: (42 ival 42), smem val (32.000000)
threadidx: (3 ,1) blockidx:(0 ,2) coordinate: (3 ,5) global index: (43 ival 43), smem val (32.000000)
threadidx: (3 ,1) blockidx:(0 ,2) coordinate: (3 ,5) global index: (43 ival 43), smem val (32.000000)
threadidx: (0 ,0) blockidx:(0 ,0) coordinate: (0 ,0) global index: ( 0 ival  0), smem val (0.000000)
threadidx: (0 ,0) blockidx:(0 ,0) coordinate: (0 ,0) global index: ( 0 ival  0), smem val (0.000000)
threadidx: (1 ,0) blockidx:(0 ,0) coordinate: (1 ,0) global index: ( 1 ival  1), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (0 ,1) blockidx:(0 ,0) coordinate: (0 ,1) global index: ( 8 ival  8), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (2 ,0) blockidx:(0 ,0) coordinate: (2 ,0) global index: ( 2 ival  2), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (3 ,0) blockidx:(0 ,0) coordinate: (3 ,0) global index: ( 3 ival  3), smem val (0.000000)
threadidx: (0 ,1) blockidx:(0 ,0) coordinate: (0 ,1) global index: ( 8 ival  8), smem val (0.000000)
threadidx: (1 ,1) blockidx:(0 ,0) coordinate: (1 ,1) global index: ( 9 ival  9), smem val (0.000000)
threadidx: (2 ,1) blockidx:(0 ,0) coordinate: (2 ,1) global index: (10 ival 10), smem val (0.000000)
threadidx: (3 ,1) blockidx:(0 ,0) coordinate: (3 ,1) global index: (11 ival 11), smem val (0.000000)
threadidx: (0 ,0) blockidx:(1 ,2) coordinate: (4 ,4) global index: (36 ival 36), smem val (36.000000)
threadidx: (1 ,0) blockidx:(1 ,2) coordinate: (5 ,4) global index: (37 ival 37), smem val (36.000000)
threadidx: (2 ,0) blockidx:(1 ,2) coordinate: (6 ,4) global index: (38 ival 38), smem val (36.000000)
threadidx: (3 ,0) blockidx:(1 ,2) coordinate: (7 ,4) global index: (39 ival 39), smem val (36.000000)
threadidx: (0 ,1) blockidx:(1 ,2) coordinate: (4 ,5) global index: (44 ival 44), smem val (36.000000)
threadidx: (1 ,1) blockidx:(1 ,2) coordinate: (5 ,5) global index: (45 ival 45), smem val (36.000000)
threadidx: (2 ,1) blockidx:(1 ,2) coordinate: (6 ,5) global index: (46 ival 46), smem val (36.000000)
threadidx: (3 ,1) blockidx:(1 ,2) coordinate: (7 ,5) global index: (47 ival 47), smem val (36.000000)

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

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

相关文章

[250411] Meta 发布 Llama 4 系列 AI 模型 | Rust 1.86 引入重大语言特性

[250411] Meta 发布 Llama 4 系列 AI 模型 | Rust 1.86 引入重大语言特性

目录 Llama 4 家族登场&#xff1a;开启原生多模态 AI 创新新纪元Rust 1.86.0 版本发布亮点主要新特性与改进其他重要信息 Llama 4 家族登场&#xff1a;开启原生多模态 AI 创新新纪元 Meta AI 近日发布了其最新、最先进的 Llama 4 系列人工智能模型&#xff0c;标志着 AI 技术…
阅读更多...
ArrayList 和 数组 的区别

ArrayList 和 数组 的区别

定义与本质 数组&#xff1a;是 Java 语言内置的数据结构&#xff0c;是存储相同类型元素的连续内存空间。它是一个基本的语言特性&#xff0c;在内存中是一块连续的区域。ArrayList&#xff1a;是 Java 集合框架中的一个类&#xff0c;属于动态数组。它是基于数组实现的&#…
阅读更多...
​‌FireCrawl‌爬虫工具​, Craw4ai

​‌FireCrawl‌爬虫工具​, Craw4ai

‌FireCrawl‌是一款开源的AI爬虫工具&#xff0c;专门用于Web数据提取&#xff0c;并将其转换为Markdown格式或其他结构化数据。FireCrawl特别适合处理使用JavaScript动态生成的网站&#xff0c;能够自动抓取网站及其所有可访问的子页面内容&#xff0c;并将其转换为适合大语言…
阅读更多...
通信原理-非线性调制

通信原理-非线性调制

今天给大家带来的是关于通信原理中非线性调制的内容,一起来看看吧&#xff01;&#xff01;&#xff01; 1.角度调制 2.FM与PM的区别 3.单音调制FM 4.窄带调频 5.宽带调频 5.1FM信号的频谱 5.2FM信号的带宽 5.3FM信号的功率分配 6.FM信号的产生与解调 6.1FM信号的产生 6.2FM…
阅读更多...
文心一言开发指南03——千帆大模型平台产品优势

文心一言开发指南03——千帆大模型平台产品优势

版权声明 本文原创作者&#xff1a;谷哥的小弟作者博客地址&#xff1a;http://blog.csdn.net/lfdfhl 千帆大模型平台作为百度智能云推出的企业级大模型一站式平台&#xff0c;具有显著的产品优势。千帆大模型平台以其基础强大、流程完善、运行稳定和安全可靠的产品优势成为企…
阅读更多...
mysql DQL

mysql DQL

一.基本查询 1.查询多个字段 2.查看所有字段 3.设置别名 4.去除重复记录 二.条件查询 1.大于小于等于 2.查询 身份证为空的 没有所以没有记录 3.在15到20这个区间范围内 4.or/in 或者 4.like 匹配 &#xff08;_匹配单个字符 %匹配多个字符&#xff09; 查询员工信…
阅读更多...
Logisim——1位比较器，2位比较器，4位无符号比较器，16位无符号比较器（头歌平台实测通过）

Logisim——1位比较器，2位比较器，4位无符号比较器，16位无符号比较器（头歌平台实测通过）

1位比较器 2位比较器 4位无符号比较器 16位无符号比较器
阅读更多...
关于 软件开发模型 的分类、核心特点及详细对比分析，涵盖传统模型、迭代模型、敏捷模型等主流类型

关于 软件开发模型 的分类、核心特点及详细对比分析，涵盖传统模型、迭代模型、敏捷模型等主流类型

以下是关于 软件开发模型 的分类、核心特点及详细对比分析&#xff0c;涵盖传统模型、迭代模型、敏捷模型等主流类型&#xff1a; 一、软件开发模型分类及核心特点 1. 瀑布模型&#xff08;Waterfall Model&#xff09; 核心特点&#xff1a; 线性阶段划分&#xff1a;需求分…
阅读更多...
2025年第十六届蓝桥杯省赛C++ A组真题

2025年第十六届蓝桥杯省赛C++ A组真题

2025年第十六届蓝桥杯省赛C A组真题 1.说明2.题目A&#xff1a;寻找质数&#xff08;5分&#xff09;3.题目B&#xff1a;黑白棋&#xff08;5分&#xff09;4. 题目C&#xff1a;抽奖&#xff08;10分&#xff09;5. 题目D&#xff1a;红黑树&#xff08;10分&#xff09;6. 题…
阅读更多...
JVM初探——走进类加载机制|三大特性 | 打破双亲委派SPI机制详解

JVM初探——走进类加载机制|三大特性 | 打破双亲委派SPI机制详解

目录 JVM是什么&#xff1f; 类加载机制 Class装载到JVM的过程 装载&#xff08;load&#xff09;——查找和导入class文件 链接&#xff08;link&#xff09;——验证、准备、解析 验证&#xff08;verify&#xff09;——保证加载类的正确性 准备&#xff08;Prepare&…
阅读更多...
分布式微服务系统架构第106集：jt808，补充类加载器

分布式微服务系统架构第106集：jt808，补充类加载器

加群联系作者vx&#xff1a;xiaoda0423 仓库地址&#xff1a;https://webvueblog.github.io/JavaPlusDoc/ https://1024bat.cn/ 类加载器 类与类加载器 判断类是否“相等” 任意一个类&#xff0c;都由加载它的类加载器和这个类本身一同确立其在 Java 虚拟机中的唯一性&#xf…
阅读更多...
利用 pyecharts 实现地图的数据可视化——第七次人口普查数据的2d、3d展示（关键词：2d 、3d 、map、 geo、涟漪点）

利用 pyecharts 实现地图的数据可视化——第七次人口普查数据的2d、3d展示（关键词：2d 、3d 、map、 geo、涟漪点）

参考文档&#xff1a;链接: link_pyecharts 官方文档 1、map() 传入省份全称&#xff0c;date_pair 是列表套列表 [ [ ],[ ] … ] 2、geo() 传入省份简称&#xff0c;date_pair 是列表套元组 [ ( ),( ) … ] 1、准备数据 population_data&#xff1a;简称经纬度 population_da…
阅读更多...
Enovia许可释放

Enovia许可释放

随着企业规模的扩大和业务的不断增长&#xff0c;Enovia许可证的管理变得至关重要。在许多情况下&#xff0c;企业可能面临许可证资源浪费或不足的问题。为了解决这一问题&#xff0c;Enovia提供了许可释放功能&#xff0c;帮助企业更加灵活地管理和使用许可证资源。本文将介绍…
阅读更多...
每日一道leetcode（回来了！！！）

每日一道leetcode（回来了！！！）

236. 二叉树的最近公共祖先 - 力扣&#xff08;LeetCode&#xff09; 题目 给定一个二叉树, 找到该树中两个指定节点的最近公共祖先。 百度百科中最近公共祖先的定义为&#xff1a;“对于有根树 T 的两个节点 p、q&#xff0c;最近公共祖先表示为一个节点 x&#xff0c;满足…
阅读更多...
【Redis】布隆过滤器应对缓存穿透的go调用实现

【Redis】布隆过滤器应对缓存穿透的go调用实现

布隆过滤器 https://pkg.go.dev/github.com/bits-and-blooms/bloom/v3 作用&#xff1a; 判断一个元素是不是在集合中 工作原理&#xff1a; 一个位数组&#xff08;bit array&#xff09;&#xff0c;初始全为0。多个哈希函数&#xff0c;运算输入&#xff0c;从而映射到位数…
阅读更多...
【ROS2】行为树 BehaviorTree（四）：组合使用子树

【ROS2】行为树 BehaviorTree（四）：组合使用子树

1、大树调用子树 如下图，左边为大树主干： 1）如果门没有关，直接通过； 2）如果门关闭了，执行开门动作，然后通过 右边为子树，主要任务是开门 1）尝试直接开门； 2）尝试开锁开门，最多尝试5次； 3）最后尝试砸门！ XML如何描述大树主干调佣子树：使用关键字 SubTree 来…
阅读更多...
【口腔粘膜鳞状细胞癌】文献阅读

【口腔粘膜鳞状细胞癌】文献阅读

写在前面 看看文章&#xff0c;看看有没有思路 文献 The regulatory role of cancer stem cell marker gene CXCR4 in the growth and metastasis of gastric cancer IF:6.8 中科院分区:1区 医学WOS分区: Q1 目的&#xff1a;通过 scRNA-seq 结合大量 RNA-seq 揭示癌症干细胞…
阅读更多...
【ComfyUI】蓝耘元生代 | ComfyUI深度解析：高性能AI绘画工作流实践

【ComfyUI】蓝耘元生代 | ComfyUI深度解析：高性能AI绘画工作流实践

【作者主页】Francek Chen 【专栏介绍】 ⌈ ⌈ ⌈人工智能与大模型应用 ⌋ ⌋ ⌋ 人工智能&#xff08;AI&#xff09;通过算法模拟人类智能&#xff0c;利用机器学习、深度学习等技术驱动医疗、金融等领域的智能化。大模型是千亿参数的深度神经网络&#xff08;如ChatGPT&…
阅读更多...
深入理解Java中的队列：核心操作、实现与应用

深入理解Java中的队列：核心操作、实现与应用

队列&#xff08;Queue&#xff09;是计算机科学中最基础且重要的数据结构之一&#xff0c;遵循 先进先出&#xff08;FIFO&#xff09; 的规则。Java通过java.util.Queue接口及其丰富的实现类为开发者提供了强大的队列工具。本文将详细解析Java队列的核心操作、常见实现类及其…
阅读更多...
idea里面不能运行 node 命令 cmd 里面可以运行咋回事啊

idea里面不能运行 node 命令 cmd 里面可以运行咋回事啊

idea里面不能运行 node 命令 cmd 里面可以运行咋回事啊 在 IntelliJ IDEA&#xff08;或其他 JetBrains 系列 IDE&#xff09;中无法运行某些命令&#xff0c;但在系统的命令提示符&#xff08;CMD&#xff09;中可以正常运行&#xff0c;这种情况通常是由于以下原因之一导致的…
阅读更多...
最新文章

Copyright @ 2022~2023