0
  • 聊天消息
  • 系统消息
  • 评论与回复
登录后你可以
  • 下载海量资料
  • 学习在线课程
  • 观看技术视频
  • 写文章/发帖/加入社区
会员中心
创作中心

完善资料让更多小伙伴认识你,还能领取20积分哦,立即完善>

3天内不再提示

CUDA编程共享内存

冬至子 来源:指北笔记 作者:张北北 2023-05-19 15:32 次阅读

Shared Memory

共享内存是使用__shared__内存空间说明符分配的

共享内存预期要比全局内存快得多它可以用作临时存储器(或软件管理缓存),以最小化来自CUDA block 的全局内存访问 ,如下面的矩阵乘法示例所示。

下面的代码示例是一个简单的矩阵乘法实现,它不利用共享内存。每个线程读取A的一行和B的一列,并计算C的相应元素,如图1所示。因此, A从全局内存中读取B的width次数,B从全局内存中读取A的height次数

从左到右是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<<

下面的代码示例是一个利用共享内存的矩阵乘法的实现。在这个实现中, 每个线程块负责计算C的一个方阵子矩阵Csub,块中的每个线程负责计算Csub中的一个元素 。如图2所示, Csub等于两个矩形矩阵的乘积:一个是与Csub具有相同行索引的维数(A.width, block_size)的子矩阵,另一个是与Csub具有相同列索引的维数(block_size, A.width)的子矩阵 。为了适应设备的资源,这两个矩形矩阵根据需要被分成多个尺寸为block_size的方阵,Csub被计算为这些方阵乘积的和。每一个乘积都是这样执行的:首先将两个对应的方阵从全局内存加载到共享内存,由一个线程加载每个矩阵的一个元素,然后让每个线程计算乘积的一个元素。每个线程将每个产品的结果累积到一个寄存器中,并将结果写入全局内存。

图片

图2 Matrix Multiplication with Shared Memory

通过这种方式阻塞计算,我们利用了快速共享内存的优势,并节省了大量全局内存带宽, 因为A只从全局内存读取(B.width / block_size)次,而B是读取(a.height / block_size)次

前面代码示例中的Matrix类型使用stride字段进行了扩充,以便子矩阵可以有效地用相同的类型表示__device__函数用于获取和设置元素,并从矩阵中构建任何子矩阵。

// 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<<
声明:本文内容及配图由入驻作者撰写或者入驻合作网站授权转载。文章观点仅代表作者本人,不代表电子发烧友网立场。文章及其配图仅供工程师学习之用,如有内容侵权或者其他违规问题,请联系本站处理。 举报投诉
  • 寄存器
    +关注

    关注

    31

    文章

    5336

    浏览量

    120235
  • 存储器
    +关注

    关注

    38

    文章

    7484

    浏览量

    163767
  • CUDA
    +关注

    关注

    0

    文章

    121

    浏览量

    13620
收藏 人收藏

    评论

    相关推荐

    内存共享原理解析

    内存共享是一种在多个进程之间共享数据的机制,它允许不同的进程直接访问同一块内存区域,从而实现数据的快速传递和通信。
    的头像 发表于 02-19 15:11 1289次阅读
    <b class='flag-5'>内存</b><b class='flag-5'>共享</b>原理解析

    CUDA编程教程

    Nvidia CUDA 2.0编程教程
    发表于 03-05 07:30

    cuda程序设计

      •GPGPU及CUDA介绍   •CUDA编程模型   •多线程及存储器硬件
    发表于 11-12 16:12 0次下载

    共享内存IPC原理,Linux进程间如何共享内存

    共享内存是在内存中单独开辟的一段内存空间,这段内存空间有自己特有的数据结构,包括访问权限、大小和最近访问的时间等。该数据结构定义如下
    的头像 发表于 07-16 13:43 8628次阅读
    <b class='flag-5'>共享</b><b class='flag-5'>内存</b>IPC原理,Linux进程间如何<b class='flag-5'>共享</b><b class='flag-5'>内存</b>?

    CUDA 6中的统一内存模型

    的,并通过PCI-Express总线相连。在CUDA6之前, 这是程序员最需要注意的地方。CPU和GPU之间共享的数据必须在两个内存中都分配,并由程序直接地在两个内存之间来回复制。这给
    的头像 发表于 07-02 14:08 2810次阅读

    深入剖析Linux共享内存原理

    不同进程之间进行通信,需要让不同进程共享相同的物理内存,Linux通过  共享内存  来实现这个功能。下面先来介绍一下Linux系统的共享
    的头像 发表于 10-30 09:52 2280次阅读
    深入剖析Linux<b class='flag-5'>共享</b><b class='flag-5'>内存</b>原理

    通过使用CUDA GPU共享内存

    共享内存是编写优化良好的 CUDA 代码的一个强大功能。共享内存的访问比全局内存访问快得多,因为
    的头像 发表于 04-11 10:03 7394次阅读

    CUDA简介: CUDA编程模型概述

    CUDA 编程模型中,线程是进行计算或内存操作的最低抽象级别。 从基于 NVIDIA Ampere GPU 架构的设备开始,CUDA 编程
    的头像 发表于 04-20 17:16 2998次阅读
    <b class='flag-5'>CUDA</b>简介: <b class='flag-5'>CUDA</b><b class='flag-5'>编程</b>模型概述

    CUDA编程模型的统一内存

      内存空间的统一意味着主机和设备之间不再需要显式内存传输。在托管内存空间中创建的任何分配都会自动迁移到需要的位置。
    的头像 发表于 05-07 14:47 1287次阅读

    Linux系统的共享内存的使用

    但有时候为了让不同进程之间进行通信,需要让不同进程共享相同的物理内存,Linux通过 共享内存 来实现这个功能。下面先来介绍一下Linux系统的共享
    的头像 发表于 11-14 11:55 1315次阅读

    使用CUDA进行编程的要求有哪些

    CUDA是NVIDIA的一种用于GPU编程的技术,CUDA核心是GPU上的一组小型计算单元,它们可以同时执行大量的计算任务。
    的头像 发表于 01-08 09:20 2530次阅读

    介绍CUDA编程模型及CUDA线程体系

    CUDA 编程模型主要有三个关键抽象:层级的线程组,共享内存和栅同步(barrier synchronization)。
    的头像 发表于 05-19 11:32 1884次阅读
    介绍<b class='flag-5'>CUDA</b><b class='flag-5'>编程</b>模型及<b class='flag-5'>CUDA</b>线程体系

    CUDA编程分布式共享内存

    计算能力9.0中引入的线程块集群为线程块集群中的线程提供了访问集群中所有参与线程块的共享内存的能力。
    的头像 发表于 05-19 15:35 1245次阅读

    Linux进程间如何实现共享内存通信

    在上面的例程中,我们首先使用ftok()函数生成一个key值作为共享内存的标识符。然后使用shmget()函数创建共享内存区域,shmaddr指向
    发表于 06-19 09:55 636次阅读

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

    CUDA核心(Compute Unified Device Architecture Core)是NVIDIA图形处理器(GPU)上的计算单元,用于执行并行计算任务。每个CUDA核心可以执行单个线程的指令,包括算术运算、逻辑操作和内存
    发表于 09-27 09:38 8674次阅读
    <b class='flag-5'>CUDA</b>核心是什么?<b class='flag-5'>CUDA</b>核心的工作原理