CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)

这篇具有很好参考价值的文章主要介绍了CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)
本系列教程将介绍具体的CUDA编程代码的细节

CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)文章来源地址https://www.toymoban.com/news/detail-509272.html

#include <stdio.h>
#include <math.h>

// a[][] * b[][] = c[][]
// 
//                         b00 b01 b02 b03
//                         b10 b11 b12 b13
//                         b20 b21 b22 b23
//                         b30 b31 b32 b33
//
// a00 a01 a02 a03         c00 c01 c02 c03
// a10 a11 a12 a13         c10 c11 c12 c13     block(1, 0) -> shared memory
// a20 a21 a22 a23         c20 c21 c22 c23     c20 c21
// a30 a31 a32 a33         c30 c31 c32 c33     c30 c31
//
//                              b00 b01->  sub_b_step_0
//                              b10 b11
//
//                              b20 b21->  sub_b_step_1
//                              b30 b31
// sub_a_step_0 sub_a_step_1    sub_c
// a20 a21      a22 a23         c20 c21
// a30 a31      a32 a33         c30 c31
//
// sub_c = sub_a_step_0 * sub_b_step_0 + sub_a_step_1 * sub_b_step_1;
//
// for(int step =0; step < N/block_size; step++ )
//      load sub_a_step to shared memory;
//      load sub_b_step to shared memory;
//      tmp += sub_a_step_on_sharedmemory * sub_b_step_on_sharedmemory;
// sub_c = tmp;
//
// cudaMalloc -> global memory
// data global memory -> shared memory
// threads shared memory -> register
// shared memory SM(stream multi-processor) same block same shared memory
//
// c21 = a20 * b01 + a21 * b11 + a22 * b21 + a23 * b31
// a00 a01 a02 a03 a10 a11 a12 a13 a20 a21 a22 a23 a30 a31 a32 a33
// 0   1   2   3   4   5   6   7   8   9   10  11  12  13  14  15
// b00 b01 b02 b03 b10 b11 b12 b13 b20 b21 b22 b23 b30 b31 b32 b33

#define M 1000
#define N 500
#define K 1000

__managed__ int a[M*N];
__managed__ int b[N*K];
__managed__ int c_gpu[M*K];
__managed__ int c_cpu[M*K];

#define BLOCK_SIZE 16

__global__ void gpu_matrix(int* a, int* b, int* c, int m, int n, int k)
{
    __shared__ int sub_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int sub_b[BLOCK_SIZE][BLOCK_SIZE];

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    int tmp =0;
    int idx;
    for(int step=0; step <= n/BLOCK_SIZE; step++)
    {
        int step_x = step * BLOCK_SIZE + threadIdx.x;
        int step_y = y;
        idx = step_y * n + step_x;
        if(step_x >= n || step_y >= m)
        {
            sub_a[threadIdx.y][threadIdx.x] =0;
        }
        else
        {
            sub_a[threadIdx.y][threadIdx.x] = a[idx];
        }

        step_x = x;
        step_y = step * BLOCK_SIZE + threadIdx.y;
        idx = step_y * k +step_x;
        if(step_x >= k || step_y >= n)
        {
            sub_b[threadIdx.y][threadIdx.x] = 0;
        }
        else
        {
            sub_b[threadIdx.y][threadIdx.x] = b[idx];
        }

        __syncthreads();

        for(int i = 0; i < BLOCK_SIZE; i++)
        {
            tmp +=sub_a[threadIdx.y][i] * sub_b[i][threadIdx.x];
        }
        __syncthreads();
    }

    if ( x < k && y < m)
    {
        c[y*k + x] = tmp; 
    }
}

void cpu_matrix(int* a, int* b, int* c, int m, int n, int k)
{
    for( int y = 0; y < m; y++)
    {
        for(int x = 0; x < k; x++)
        {
            int tmp = 0;
            for(int step =0; step < n; step++)
            {
                tmp += a[y*n + step] * b[step*k + x];
            }
            c[y * k + x] = tmp;
        }
    }
}

int main()
{
    for(int y=0; y<M; ++y)
    {
        for(int x=0; x<N; ++x)
        {
            a[y * N + x] = rand()%1024;
        }
    }

    for(int y=0; y<N; ++y)
    {
        for(int x=0; x<K; ++x)
        {
            b[y*K + x] = rand()%1024;
        }
    }

    unsigned int grid_x = (K + BLOCK_SIZE -1)/BLOCK_SIZE;
    unsigned int grid_y = (M + BLOCK_SIZE -1)/BLOCK_SIZE;

    dim3 dimGrid(grid_x, grid_y);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);

    gpu_matrix<<<dimGrid, dimBlock>>>(a, b, c_gpu, M, N, K);

    cpu_matrix(a, b, c_cpu, M, N, K);

    bool errors = false;

    for(int y=0; y<M; y++)
    {
        for(int x=0; x<K; x++)
        {
            if(fabs(c_cpu[y*K + x] - c_gpu[y*K+x]) > (1.0e-10))
            {
                errors = true;
                printf("c_cpu: %d. c_gpu: %d", c_cpu[y*K + x], c_gpu[y*K+x]);
            }
        }
    }

    printf("Result: %s\n", errors?"Error":"Pass");

    return 0;
}



























到了这里,关于CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • multiprocessing.shared_memory --- 从进程直接访问的共享内存(1)

    模块提供了一个类 SharedMemory ,用于分配和管理共享内存,供多核或对称多处理器 (SMP) 机器上的一个或多个进程访问。 为了协助共享内存的生命周期管理,尤其是跨不同进程的共享内存, multiprocessing.managers 模块中还提供了一个 BaseManager 子类 SharedMemoryManager 。 在这个模块中

    2024年04月22日
    浏览(34)
  • PostgreSQL数据库动态共享内存管理器——Dynamic shared memory areas

    dsm.c提供的功能允许创建后端进程间共享的共享内存段。DSA利用多个DSM段提供共享内存heap;DSA可以利用已经存在的共享内存(DSM段)也可以创建额外的DSM段。和系统heap使用指针不同的是,DSA提供伪指针,可以转换为backend-local指针,但是该伪指针可以在后端进程之间共享,可

    2024年02月15日
    浏览(43)
  • CUDA编程模型系列三(矩阵乘)

    本系列教程将介绍具体的CUDA编程代码的细节 CUDA编程模型系列三(矩阵乘)

    2024年02月12日
    浏览(78)
  • 重磅!苹果官方发布大模型框架:一个可以充分利用苹果统一内存的新的大模型框架MLX,你的MacBook可以一键运行LLaMA了

    本文来自DataLearnerAI官方网站: 重磅!苹果官方发布大模型框架:一个可以充分利用苹果统一内存的新的大模型框架MLX,你的MacBook可以一键运行LLaMA了 | 数据学习者官方网站(Datalearner) https://www.datalearner.com/blog/1051701871117729 苹果刚刚发布了一个全新的机器学习矿机MLX,这是一个

    2024年02月04日
    浏览(38)
  • Stable Diffusion WebUI内存不够爆CUDA Out of memory怎么办?

    在我们运行SD的时候,我们经常会爆CUDA Out of memory。 我们应该怎么办呢? 这是因为我们的显存或者内存不够了。 如果你是用cpu来跑图的则表示内存不够,这个时候就需要换个大点的内存了。 如果你是用gpu来跑图的就说明你显存不够用咯,这时候咋办呢? 下面我将一一述说解

    2024年02月08日
    浏览(130)
  • Semantic Kernel 入门系列:?Memory内存

    了解的运作原理之后,就可以开始使用Semantic Kernel来制作应用了。 Semantic Kernel将embedding的功能封装到了Memory中,用来存储上下文信息,就好像电脑的内存一样,而LLM就像是CPU一样,我们所需要做的就是从内存中取出相关的信息交给CPU处理就好了。 使用Memory需要注册 embedding

    2023年04月13日
    浏览(42)
  • Linux 内存模型(Memory: the flat, the discontiguous, and the sparse)

    计算机系统中的物理内存是一种宝贵的资源,因此人们付出了大量的努力来有效地管理它。由于现代系统上存储器体系结构的复杂性,这项任务变得更加困难。有几个抽象层处理如何布局物理内存的细节;其中一个简单地称为“内存模型”。内核中支持三个模型,但其中一个

    2024年02月14日
    浏览(41)
  • 【javaEE面试题(五)在JMM(Java Memory Model (Java 内存模型))下谈volatile的作用】【保证内存可见 和 指令有序】

    volatile 能保证内存可见性 volatile 修饰的变量, 能够保证 “内存可见性”. 代码在写入 volatile 修饰的变量的时候 改变线程工作内存中volatile变量副本的值 将改变后的副本的值从工作内存 刷新到主内存 代码在读取 volatile 修饰的变量的时候 从主内存中读取volatile变量的最新值到

    2024年02月16日
    浏览(41)
  • 【javaEE面试题(五)在JMM(Java Memory Model (Java 内存模型))下谈volatile的作用】

    volatile 能保证内存可见性 volatile 修饰的变量, 能够保证 “内存可见性”. 代码在写入 volatile 修饰的变量的时候 改变线程工作内存中volatile变量副本的值 将改变后的副本的值从工作内存 刷新到主内存 代码在读取 volatile 修饰的变量的时候 从主内存中读取volatile变量的最新值到

    2024年02月13日
    浏览(47)
  • CUDA编程入门系列(二) GPU硬件架构综述

    一、Fermi GPU         Fermi GPU如下图所示,由16个SM(stream multiprocessor)组成,不同的SM之间通过L2 Cache和全局内存进行相连。整个架构大致分为两个层次,①总体架构由多个SM组成 ②每个SM由多个SP core(stream processor)组成。SP之间通过互连的网络和L1 Cache和Warp Scheduler等结构进行

    2024年02月07日
    浏览(51)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包