一区二区三区三上|欧美在线视频五区|国产午夜无码在线观看视频|亚洲国产裸体网站|无码成年人影视|亚洲AV亚洲AV|成人开心激情五月|欧美性爱内射视频|超碰人人干人人上|一区二区无码三区亚洲人区久久精品

0
  • 聊天消息
  • 系統(tǒng)消息
  • 評論與回復(fù)
登錄后你可以
  • 下載海量資料
  • 學(xué)習(xí)在線課程
  • 觀看技術(shù)視頻
  • 寫文章/發(fā)帖/加入社區(qū)
會員中心
創(chuàng)作中心

完善資料讓更多小伙伴認(rèn)識你,還能領(lǐng)取20積分哦,立即完善>

3天內(nèi)不再提示

CUDA編程共享內(nèi)存

冬至子 ? 來源:指北筆記 ? 作者:張北北 ? 2023-05-19 15:32 ? 次閱讀
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

Shared Memory

共享內(nèi)存是使用__shared__內(nèi)存空間說明符分配的 。

共享內(nèi)存預(yù)期要比全局內(nèi)存快得多 。 它可以用作臨時存儲器(或軟件管理緩存),以最小化來自CUDA block 的全局內(nèi)存訪問 ,如下面的矩陣乘法示例所示。

下面的代碼示例是一個簡單的矩陣乘法實現(xiàn),它不利用共享內(nèi)存。每個線程讀取A的一行和B的一列,并計算C的相應(yīng)元素,如圖1所示。因此, A從全局內(nèi)存中讀取B的width次數(shù),B從全局內(nèi)存中讀取A的height次數(shù) 。

從左到右是x的方向,從上到下是y的方向。 (x,y) x是0-dim,y是1-dim,和正常的 shape 表示是反著的。

圖片

圖1 Matrix Multiplication without Shared Memory

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
    int width;
    int height;
    float* elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    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);

    // Allocate C in device memory
    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<<

下面的代碼示例是一個利用共享內(nèi)存的矩陣乘法的實現(xiàn)。在這個實現(xiàn)中, 每個線程塊負(fù)責(zé)計算C的一個方陣子矩陣Csub,塊中的每個線程負(fù)責(zé)計算Csub中的一個元素 。如圖2所示, Csub等于兩個矩形矩陣的乘積:一個是與Csub具有相同行索引的維數(shù)(A.width, block_size)的子矩陣,另一個是與Csub具有相同列索引的維數(shù)(block_size, A.width)的子矩陣 。為了適應(yīng)設(shè)備的資源,這兩個矩形矩陣根據(jù)需要被分成多個尺寸為block_size的方陣,Csub被計算為這些方陣乘積的和。每一個乘積都是這樣執(zhí)行的:首先將兩個對應(yīng)的方陣從全局內(nèi)存加載到共享內(nèi)存,由一個線程加載每個矩陣的一個元素,然后讓每個線程計算乘積的一個元素。每個線程將每個產(chǎn)品的結(jié)果累積到一個寄存器中,并將結(jié)果寫入全局內(nèi)存。

圖片

圖2 Matrix Multiplication with Shared Memory

通過這種方式阻塞計算,我們利用了快速共享內(nèi)存的優(yōu)勢,并節(jié)省了大量全局內(nèi)存帶寬, 因為A只從全局內(nèi)存讀取(B.width / block_size)次,而B是讀取(a.height / block_size)次 。

前面代碼示例中的Matrix類型使用stride字段進(jìn)行了擴(kuò)充,以便子矩陣可以有效地用相同的類型表示 。__device__函數(shù)用于獲取和設(shè)置元素,并從矩陣中構(gòu)建任何子矩陣。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
    int width;
    int height;
    int stride; 
    float* elements;
} Matrix;

// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
    return A.elements[row * A.stride + col];
}

// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col,
                           float value)
{
    A.elements[row * A.stride + col] = value;
}

// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
 __device__ Matrix GetSubMatrix(Matrix A, int row, int col) 
{
    Matrix Asub;
    Asub.width    = BLOCK_SIZE;
    Asub.height   = BLOCK_SIZE;
    Asub.stride   = A.stride;
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row
                                         + BLOCK_SIZE * col];
    return Asub;
}

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
    // Load A and B to device memory
    Matrix d_A;
    d_A.width = d_A.stride = 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 = d_B.stride = 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);

    // Allocate C in device memory
    Matrix d_C;
    d_C.width = d_C.stride = 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<<
聲明:本文內(nèi)容及配圖由入駐作者撰寫或者入駐合作網(wǎng)站授權(quán)轉(zhuǎn)載。文章觀點僅代表作者本人,不代表電子發(fā)燒友網(wǎng)立場。文章及其配圖僅供工程師學(xué)習(xí)之用,如有內(nèi)容侵權(quán)或者其他違規(guī)問題,請聯(lián)系本站處理。 舉報投訴
  • 寄存器
    +關(guān)注

    關(guān)注

    31

    文章

    5434

    瀏覽量

    124435
  • 存儲器
    +關(guān)注

    關(guān)注

    38

    文章

    7649

    瀏覽量

    167347
  • CUDA
    +關(guān)注

    關(guān)注

    0

    文章

    122

    瀏覽量

    14120
收藏 人收藏
加入交流群
微信小助手二維碼

掃碼添加小助手

加入工程師交流群

    評論

    相關(guān)推薦
    熱點推薦

    內(nèi)存共享原理解析

    內(nèi)存共享是一種在多個進(jìn)程之間共享數(shù)據(jù)的機(jī)制,它允許不同的進(jìn)程直接訪問同一塊內(nèi)存區(qū)域,從而實現(xiàn)數(shù)據(jù)的快速傳遞和通信。
    的頭像 發(fā)表于 02-19 15:11 ?1818次閱讀
    <b class='flag-5'>內(nèi)存</b><b class='flag-5'>共享</b>原理解析

    CUDA編程教程

    Nvidia CUDA 2.0編程教程
    發(fā)表于 03-05 07:30

    cuda程序設(shè)計

      •GPGPU及CUDA介紹   •CUDA編程模型   •多線程及存儲器硬件
    發(fā)表于 11-12 16:12 ?0次下載

    共享內(nèi)存IPC原理,Linux進(jìn)程間如何共享內(nèi)存?

    共享內(nèi)存是在內(nèi)存中單獨開辟的一段內(nèi)存空間,這段內(nèi)存空間有自己特有的數(shù)據(jù)結(jié)構(gòu),包括訪問權(quán)限、大小和最近訪問的時間等。該數(shù)據(jù)結(jié)構(gòu)定義如下
    的頭像 發(fā)表于 07-16 13:43 ?8921次閱讀
    <b class='flag-5'>共享</b><b class='flag-5'>內(nèi)存</b>IPC原理,Linux進(jìn)程間如何<b class='flag-5'>共享</b><b class='flag-5'>內(nèi)存</b>?

    CUDA 6中的統(tǒng)一內(nèi)存模型

    的,并通過PCI-Express總線相連。在CUDA6之前, 這是程序員最需要注意的地方。CPU和GPU之間共享的數(shù)據(jù)必須在兩個內(nèi)存中都分配,并由程序直接地在兩個內(nèi)存之間來回復(fù)制。這給
    的頭像 發(fā)表于 07-02 14:08 ?3140次閱讀

    深入剖析Linux共享內(nèi)存原理

    不同進(jìn)程之間進(jìn)行通信,需要讓不同進(jìn)程共享相同的物理內(nèi)存,Linux通過? 共享內(nèi)存 ?來實現(xiàn)這個功能。下面先來介紹一下Linux系統(tǒng)的共享
    的頭像 發(fā)表于 10-30 09:52 ?2533次閱讀
    深入剖析Linux<b class='flag-5'>共享</b><b class='flag-5'>內(nèi)存</b>原理

    通過使用CUDA GPU共享內(nèi)存

    共享內(nèi)存是編寫優(yōu)化良好的 CUDA 代碼的一個強(qiáng)大功能。共享內(nèi)存的訪問比全局內(nèi)存訪問快得多,因為
    的頭像 發(fā)表于 04-11 10:03 ?7900次閱讀

    CUDA簡介: CUDA編程模型概述

    CUDA 編程模型中,線程是進(jìn)行計算或內(nèi)存操作的最低抽象級別。 從基于 NVIDIA Ampere GPU 架構(gòu)的設(shè)備開始,CUDA 編程
    的頭像 發(fā)表于 04-20 17:16 ?3361次閱讀
    <b class='flag-5'>CUDA</b>簡介: <b class='flag-5'>CUDA</b><b class='flag-5'>編程</b>模型概述

    CUDA編程模型的統(tǒng)一內(nèi)存

      內(nèi)存空間的統(tǒng)一意味著主機(jī)和設(shè)備之間不再需要顯式內(nèi)存傳輸。在托管內(nèi)存空間中創(chuàng)建的任何分配都會自動遷移到需要的位置。
    的頭像 發(fā)表于 05-07 14:47 ?1645次閱讀

    Linux系統(tǒng)的共享內(nèi)存的使用

    但有時候為了讓不同進(jìn)程之間進(jìn)行通信,需要讓不同進(jìn)程共享相同的物理內(nèi)存,Linux通過 共享內(nèi)存 來實現(xiàn)這個功能。下面先來介紹一下Linux系統(tǒng)的共享
    的頭像 發(fā)表于 11-14 11:55 ?1586次閱讀

    使用CUDA進(jìn)行編程的要求有哪些

    CUDA是NVIDIA的一種用于GPU編程的技術(shù),CUDA核心是GPU上的一組小型計算單元,它們可以同時執(zhí)行大量的計算任務(wù)。
    的頭像 發(fā)表于 01-08 09:20 ?2862次閱讀

    介紹CUDA編程模型及CUDA線程體系

    CUDA 編程模型主要有三個關(guān)鍵抽象:層級的線程組,共享內(nèi)存和柵同步(barrier synchronization)。
    的頭像 發(fā)表于 05-19 11:32 ?2483次閱讀
    介紹<b class='flag-5'>CUDA</b><b class='flag-5'>編程</b>模型及<b class='flag-5'>CUDA</b>線程體系

    CUDA編程分布式共享內(nèi)存

    計算能力9.0中引入的線程塊集群為線程塊集群中的線程提供了訪問集群中所有參與線程塊的共享內(nèi)存的能力。
    的頭像 發(fā)表于 05-19 15:35 ?1641次閱讀

    Linux進(jìn)程間如何實現(xiàn)共享內(nèi)存通信

    在上面的例程中,我們首先使用ftok()函數(shù)生成一個key值作為共享內(nèi)存的標(biāo)識符。然后使用shmget()函數(shù)創(chuàng)建共享內(nèi)存區(qū)域,shmaddr指向
    發(fā)表于 06-19 09:55 ?795次閱讀

    CUDA核心是什么?CUDA核心的工作原理

    CUDA核心(Compute Unified Device Architecture Core)是NVIDIA圖形處理器(GPU)上的計算單元,用于執(zhí)行并行計算任務(wù)。每個CUDA核心可以執(zhí)行單個線程的指令,包括算術(shù)運算、邏輯操作和內(nèi)存
    發(fā)表于 09-27 09:38 ?1w次閱讀
    <b class='flag-5'>CUDA</b>核心是什么?<b class='flag-5'>CUDA</b>核心的工作原理