【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.pswp.cn/bicheng/76884.shtml
繁體地址，請注明出處：http://hk.pswp.cn/bicheng/76884.shtml
英文地址，請注明出處：http://en.pswp.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