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);// 在設備內存中分配CMatrix 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 kerneldim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);MatMulKernel <<< dimGrid, dimBlock >>>(d_A, d_B, d_C);//從設備內存中讀取CcudaMemcpy(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);// 在設備內存中分配CMatrix 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);// 從設備內存中讀取CcudaMemcpy(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的一個子矩陣CsubMatrix Csub = GetSubMatrix(C, blockRow, blockCol);// 每個線程通過將結果累積到Cvalue中來計算Csub的一個元素float Cvalue = 0;// 線程行和列在Csubint 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的子矩陣BMatrix 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.pswp.cn/news/165703.shtml
繁體地址，請注明出處：http://hk.pswp.cn/news/165703.shtml
英文地址，請注明出處：http://en.pswp.cn/news/165703.shtml

如若內容造成侵權/違法違規/事實不符，請聯系多彩編程網進行投訴反饋email:809451989@qq.com，一經查實，立即刪除！

相關文章

『Linux升級路』基礎開發工具——gcc/g++篇

『Linux升級路』基礎開發工具——gcc/g++篇

&#x1f525;博客主頁&#xff1a;小王又困了 &#x1f4da;系列專欄&#xff1a;Linux &#x1f31f;人之為學&#xff0c;不日近則日退 ??感謝大家點贊&#x1f44d;收藏?評論?? 目錄 一、快速認識gcc/g 二、預處理 &#x1f4d2;1.1頭文件展開 &#x1f4d2;1…
閱讀更多...
java字符串的常見用法

java字符串的常見用法

java字符串的常見用法 Java中的字符串是一個非常常用的對象&#xff0c;它屬于Java的內置類String類的實例。字符串在Java中是不可變的&#xff0c;即一旦創建了一個字符串對象&#xff0c;就不能修改它的值。 下面是一些關于Java字符串的詳細用法&#xff1a; 1&#xff09;創…
閱讀更多...
從零開始，用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合

從零開始，用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合

&#x1f38f;&#xff1a;你只管努力&#xff0c;剩下的交給時間 &#x1f3e0; &#xff1a;小破站 "從零開始&#xff0c;用Docker-compose打造SkyWalking、Elasticsearch和Spring Cloud的完美融合 前言準備工作編寫docker-compose.yml文件為什么使用本機ip為什么skywa…
閱讀更多...
代碼隨想錄-刷題第六天

代碼隨想錄-刷題第六天

242. 有效的字母異位詞 題目鏈接&#xff1a;242. 有效的字母異位詞 思路&#xff1a;哈希法。利用數組來記錄出現的字母個數&#xff0c;然后判斷是否為字母異位詞。 時間復雜度&#xff1a;O(n) class Solution {public boolean isAnagram(String s, String t) {int[] co…
閱讀更多...
【云備份】第三方庫的認識與使用

【云備份】第三方庫的認識與使用

文章目錄 json庫粗略認識詳細認識writer 類reader類jsoncpp序列化實現jsoncpp反序列化實現 bundle文件壓縮庫簡單認識bundle庫實現文件壓縮bundle庫實現文件解壓縮 httplib庫Request類Response類Server類Client類 json庫 粗略認識 json是一種數據交換格式&#xff0c;采用完全…
閱讀更多...
激光切割設備中模組的作用有哪些？

激光切割設備中模組的作用有哪些？

激光切割設備是一種高精度的自動化加工設備&#xff0c;用于對金屬、非金屬等材料進行精確切割。直線模組作為激光切割設備的重要組成部分&#xff0c;在激光切割設備中起著重要的作用&#xff0c;為設備的運動系統提供了高精度、高穩定性和高效率的運動控制。 1、高精度的位置…
閱讀更多...
excel單元格加背景顏色不生效？

excel單元格加背景顏色不生效？

如果在 Excel 中設置單元格背景顏色而發現不生效&#xff0c;可能有幾個原因。以下是一些常見的解決方法&#xff1a; 1. **單元格鎖定&#xff1a;** 檢查所在單元格是否被鎖定。如果單元格被鎖定&#xff0c;并且工作表被保護&#xff0c;你可能無法更改其背景顏色。在工作表…
閱讀更多...
mysql 優化器的AST樹是啥

mysql 優化器的AST樹是啥

from ChatGPT: MySQL中的優化器&#xff08;optimizer&#xff09;使用AST&#xff08;Abstract Syntax Tree&#xff0c;抽象語法樹&#xff09;來表示查詢的語法結構。AST是一種樹狀結構&#xff0c;它反映了查詢語句的語法層次&#xff0c;是一個抽象表示&#xff0c;用于更…
閱讀更多...
Linux - 文件系統 - 理解目錄 - 理解 軟/硬鏈接

Linux - 文件系統 - 理解目錄 - 理解 軟/硬鏈接

前言 在上篇博客當中&#xff0c;我們對 文件系統 和 inode 做了初步了解&#xff0c;本博客將在上篇博客的基礎之上&#xff0c;對于 文件系統當中的目錄進行進步一闡述。 Linux - 進一步理解 文件系統 - inode - 機械硬盤-CSDN博客 目錄 一個文件有一個 inode&#xff0c;…
閱讀更多...
Redis打包事務,分批提交

Redis打包事務,分批提交

一、需求背景 接手一個老項目&#xff0c;在項目啟動的時候&#xff0c;需要將xxx省整個省的所有區域數據數據、以及系統字典配置逐條保存在Redis緩存里面&#xff0c;這樣查詢的時候會更快; 區域數據字典數據一共大概20000多條,&#xff0c;前同事直接使用 list.forEach…
閱讀更多...
Windows安裝MongoDB

Windows安裝MongoDB

1、下載MongoDB的zip&#xff0c;解壓 2、創建目錄 mkdir D:\JavaSoftware\Database\MongoDB\mongodb-win32-x86_64-windows-5.0.8\data\db mkdir D:\JavaSoftware\Database\MongoDB\mongodb-win32-x86_64-windows-5.0.8\data\log 3、創建一個配置文件mongod.cfg&#xff0c…
閱讀更多...
使用一個接口的結果作為第二個接口的參數并將兩者的數據放置成下拉框的格式

使用一個接口的結果作為第二個接口的參數并將兩者的數據放置成下拉框的格式

背景 我使用下拉框實現選擇id 但是只有兩個接口 一個是所有的id 另一個是id對應的具體信息 我想把id傳入另一個接口并且獲取其name然后寫成類似這樣的數組 [ { value: 1, label: ‘名稱1’ }&#xff0c; { value: 2, label: ‘名稱2’ } { value: 3, label: ‘名稱3’ } ] 然…
閱讀更多...
【PPspliT】ppt轉pdf-保留過渡動畫

【PPspliT】ppt轉pdf-保留過渡動畫

網址 http://www.maxonthenet.altervista.org/ppsplit.php 下載安裝 使用 再次打開ppt&#xff0c;就能在上方的選項欄里頭看到了&#xff1a;
閱讀更多...
GEE生物量碳儲量——利用紅和近紅外波段和OTSU大津法提取純凈森林面積

GEE生物量碳儲量——利用紅和近紅外波段和OTSU大津法提取純凈森林面積

簡介： 如何利用紅和近紅外波段和OTSU大津法提取純凈森林面積？本文的主要邏輯是利用特定時期的遙感影像的波段，提取指定范圍的內的DN值，然后分別統計發生閾值變化的峰值區域，從而作為篩選森林的臨界點，如果研究區較大的話則需要先進行影像分割，分割成為相同大小的區域，…
閱讀更多...
前端開發Java學習

前端開發Java學習

注釋&#xff1a; 單行注釋 // 多行注釋 /* */ 文件注釋 /** */ 1 關鍵字 &#xff08;關鍵字一定是小寫&#xff09; 2 常量 字符串常量"HelloWord","你好世界"整數常量12&#xff0c;-33小數常量3.14&#xff0c;22.1字符常量A&#xff0c;a &…
閱讀更多...
不是默認進入Linux|總是自動進入windows系統

不是默認進入Linux|總是自動進入windows系統

問題描述 不是默認進入Linux系統無法主動出現boot引導自動進入windows系統 嘗試無效 修復引導無效重裝Grub無效重裝系統無效 環境 Ubuntu 22.04 LST微星主板 解決方案 修改引導順序&#xff1a; 開機狂按Del鍵&#xff0c;進入BIOS系統&#xff0c;左側Settings 設置&…
閱讀更多...
src和href的區別

src和href的區別

前言 持續學習總結輸出中&#xff0c;src和href都是HTML中特定元素的屬性&#xff0c;都可以用來引入外部的資源。兩者區別如下&#xff1a; 1、作用 href 用于在當前文檔和引用資源之間確立聯系。 src 用于替換當前內容。 2、范圍用途 src&#xff08;source&#xff09…
閱讀更多...
項目整個管理論文之6

項目整個管理論文之6

閱讀更多...
RabbitMQ基礎教程

RabbitMQ基礎教程

1.什么是消息隊列 消息隊列&#xff08;Message Queue&#xff09;&#xff0c;我們一般簡稱為MQ。消息隊列中間件是分布式系統中重要的組件&#xff0c;具有異步性、松耦合、分布式、可靠性等特點。用于實現高性能、高可用、可伸縮和最終一致性架構。是大型分布式系統不可缺少…
閱讀更多...
Angular11 MSAL B2C登錄實例 （一）

Angular11 MSAL B2C登錄實例 （一）

前言 因為項目需求&#xff0c;需要把Angular 11項目中登錄方式改成B2C登錄&#xff0c;所以在參考了一系列文檔后&#xff0c;成功通過MSAL將項目的登錄方式改成B2C登錄。下面介紹了詳細步驟及一些注意事項。 步驟&#xff1a; 1. 安裝MSAL 在項目中安裝msal npm i azure/…
閱讀更多...
最新文章

Copyright @ 2022~2023