CUDA：矩阵乘法的实现

最新推荐文章于 2024-04-29 13:34:21 发布
最新推荐文章于 2024-04-29 13:34:21 发布
阅读量530 收藏 11
点赞数 7
文章标签： 矩阵 算法 线性代数 c++
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_43308899/article/details/135504530
版权
本文详细介绍了如何在CPU和GPU上使用共享内存和矩阵分块优化矩阵乘法，通过将大矩阵分解为小块并行计算，展示了CUDA代码实例，以及与CPU版本的性能对比。
摘要由CSDN通过智能技术生成

说在前面：本篇文章给出矩阵乘法在cpu和gpu上的实现，对代码做出详细的解释，并给出完整的可运行的代码。

优化思路

        这里我们结合共享内存和矩阵分块的思路去实现gpu端的矩阵乘法，我们会将矩阵分成很多个小矩阵，每个block对于一个小矩阵的计算，举个例子来说明这个问题，我们将a矩阵和b矩阵平均切分为9个小矩阵，同时也分配9个blcok，1号blcok需要做的是将a和b的1号矩阵相乘，再加上a的2号矩阵乘以b的4号矩阵，在加上a的3号矩阵乘以b的7号矩阵，这样就获得了输出矩阵的1号矩阵，以此类推可以得到最终的整个输出矩阵，需要注意的是，我们在每次进行小矩阵计算的时候需要先将小矩阵的数据读取到共享内存中，因为共享内存是有限的，所以我们对输出的小矩阵的计算需要进行分步。

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

#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;
}

代码解释

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

        这里我们定义了三个宏，M是矩阵a的行数，N是矩阵a的列数（也是矩阵b的行数），K是矩阵b的列数。

__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

        这里我们定义了四个统一变量，分别用于存储a和b矩阵已经矩阵相乘后在cpu和gpu端的结果。BLCOK_SIZE是指每个blcok的大小是16*16。

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;
        }
    }
}

        这是矩阵乘法在cpu端的实现，这里不做过多的解释，比较常规，在主函数中会对cpu和gpu端的结果进行比较。

__global__ void gpu_matrix(int* a, int* b, int* c, int m, int n, int k)

        a和b是输入矩阵，c是输出矩阵，m和n是a的行数和列数，k是b的列数。

 __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;

        在共享内存中定义了两个矩阵用于存储每个小矩阵的数据。x和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();
    }

        tmp是用于存储矩阵计算的结果，idx是当前线程对于的矩阵元素的全局索引。这里的for循环其实就是计算小矩阵的乘法，最终输出的小矩阵的值是通过n/BLOCK_SIZE个步骤实现的。

    if ( x < k && y < m)
    {
        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;
}

        main函数就是实现了输入矩阵的初始化，以及cpu矩阵乘法的调用和gpu核函数的调用，并将两个结果进行比较。

总结

        总体来说，这个核函数利用了 GPU 的并行处理能力和共享内存来高效地执行矩阵乘法。通过将矩阵分割成小块并在共享内存中处理，这种方法大大减少了全局内存访问的次数，从而提高了效率。

pqin
关注
  • 7
    点赞
  • 踩
  • 11
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
strassen-cuda:Strassen矩阵乘法算法的CUDA实现
05-26
结合"strassen-cuda"项目，我们可以理解这是一个利用CUDA实现Strassen矩阵乘法算法的程序。该项目可能包含以下关键知识点： 1. **CUDA编程模型**：了解CUDA编程的基础，包括设备和主机的概念、内存层次（全局内存、...
CUDA学习(十一):原子操作实现向量内积
hjxu2016的博客
11-20 1312
文章目录一、原子操作的概念二、一次原子操作替换两次归约三、Block内归约，block间原子操作 本文章为 《 GPU编程与优化 大众高性能计算》的读书笔记，例子也都取自书中教程。 向量内积运算中，需要用到的原子操作函数是原子加法函数，GPU的向量内积运算包含的两个累加部分都可以采用原子操作替换，分别是block内thread间归约和block间归约，当然也可以只用一次原子操作。 原子操作又可以分为全局存储原子操作和共享存储原子操作，抢号能对应向量内积中的两个累加部分。 一、原子操作的概念 摘自 CUDA
基于CUDA的矩阵相乘
身是菩提树、心是明镜台、时时勤拂拭、莫使惹尘埃
03-09 9408
这几天研究了一下CUDA，发现其并行的思想和普通的CPU多线程思想不太一致，但还是挺不错。主要是将任务划分成一个个block，然后每个block里面再划分成细的线程。然后每个线程做自己做的 事情。这种并行思想很适用于像矩阵运算这些元素与元素之间的运算并不耦合得很厉害，但整体数据很大的情况，这只是我对CUDA的初步感觉。 矩阵相乘的CPU程序如下： //C = A*B void
CUDA-矩阵乘法实现（含提速对比，适合上手）
最新发布
翟天保的博客
04-29 732
CUDA-矩阵乘法实现（含提速对比，适合上手）
CUDA之通用矩阵乘法：从入门到熟练！
CV_Autobot的博客
03-24 317
作者|紫气东来 编辑|汽车人原文链接：https://zhuanlan.zhihu.com/p/657632577点击下方卡片，关注“自动驾驶之心”公众号戳我-> 领取自动驾驶近15个方向学习路线>>点击进入→自动驾驶之心『模型部署』技术交流群本文只做学术分享，如有侵权，联系删文通用矩阵乘法 (General Matrix Multiplication，GEMM) 是各...
CUDA~矩阵乘运算
whaosoft143ai的博客
11-01 2808
心爱的cuda文章终于又找到一个 赶紧搬啊本文主要介绍用CUDA实现矩阵乘法运算（C = A x B）的几个基本方法，帮助大家理解矩阵在GPU上面的运算与CPU上的有何异同，通过实践上手CUDA的优化计算，相比基础方法，能提速10倍以上。本文内容涉及到CUDA矩阵1D运算、2D运算、共享内存、CUBLAS的使用。代码:https://github.com/CalvinXKY/BasicCUDA/tree/master/matrix_multiply。​ whaosoft aiothttp://143a
CUDA:矩阵乘法的实现(Share Memory)
m0_46197553的博客
07-07 4041
记录了使用Share Memory优化的CUDA矩阵乘法
CUDA编程之矩阵乘法
qq_43398345的博客
05-04 2436
对于平铺矩阵算法，对于每一个Block，首先将其存储在共享内存中, 花费 * Block_Size ^ 2，对BLock_SIze中每一个元素进行读取需要花费: 2 * k / Block_Size 个单位时间，一共有Block_Size ^ 2个元素， 一共有 (m * n) / (block_size * block_size)个Block，所以总时间就是这些值进行相乘。n])的矩阵乘法运算，每个线程都要读取A的一整行和B的一整列。• 共享内存属于线程块，可以被一个线程块内的所有线程访问。
CUDA编程实现矩阵乘法（代码解析）
ing100的博客
10-13 1175
算法思路+代码，CUDA编程之矩阵乘法全面解析，涉及到共享内存与全局内存的使用以及矩阵分块问题等
CUDA实现矩阵相乘
wulele2的博客
07-22 2016
#include <stdio.h> #define BLOCK_NUM 8 #define THREAD_NUM 32 #define R_SIZE BLOCK_NUM * THREAD_NUM #define M_SIZE R_SIZE*R_SIZE void __global__ matmul1(int *da, int *db, int *dres); void __global__ matmul1(int *da, int *db, int *dres) { // 获
基于CUDA架构矩阵乘法的研究
10-19
实验结果显示，基于CUDA的矩阵乘法相对于CPU实现，可以获得显著的加速比，最高可达1079.64，这意味着GPU的计算效率远超CPU。此外，GPU的浮点运算能力得到了充分利用，峰值比最高达到30.85%，这充分展示了CUDA架构在...
并行程序分块矩阵乘法
05-26
并行mpi 分块矩阵乘法 windows7下平台 vc6.0
cuda.zip_cuda矩阵_矩阵乘法
09-23
学习CUDA矩阵乘法不仅可以提升计算性能，还能帮助理解并行计算的基本原理，这对于进行大规模数值模拟和数据分析至关重要。在实际应用中，结合适当的优化技术，如张量核心利用和混合精度计算，可以进一步提升计算效率...
CUDA实现稀疏大矩阵乘法
02-21
稀疏矩阵的DIA/ELLPACK/COO/CSR/HYB表示形式，以及各表示形式下的稀疏矩阵乘法（稀疏大矩阵*矢量）的CUDA实现。对于矩阵中每一行稀疏元素个数较统一的情况，ELLPACK表示最佳，其次是HYB（ELL+COO）。关于稀疏矩阵的...
mul.zip_cuda程序_矩阵乘法
09-20
这个项目对于初学者来说是一个很好的起点，因为它提供了简单的矩阵乘法实现，同时内含注释，有助于理解CUDA编程的基本概念。对于有经验的开发者，它也是一个基础示例，可以在此基础上进行性能优化，例如引入共享内存...
cuda矩阵乘法(简单理解)
Super5311的博客
07-18 1004
cuda矩阵乘法
CUDA编程（九）并行矩阵乘法
热门推荐
MingChao_Sun
04-09 4万+
CUDA编程（九）矩阵乘法 在之前我们一直围绕着一个非常简单的求立方和的小程序学习CUDA，不过这个立方和的小程序没有什么实际意义，这篇博客我们用CUDA并行了矩阵乘法，问题也比较简单，基于上一个立方和程序的经验，完成这个程序也不算太难，但是运行结果存在很多问题，毕竟我们还没有开始优化，不过除了速度问题，GPU浮点数运算的精度也成了一个大问题，不过这些之后都会慢慢解决。
矩阵乘法——基于GPU的并行编程模型CUDA程序设计
千色千寻的博客
08-21 3607
CUDA并行编程——矩阵乘法 目录CUDA并行编程——矩阵乘法1 题目描述2 设计思路实验环境3 源码3.1 串行程序3.2 并行程序3.3 性能对比与分析 1 题目描述 题目1：编写一个矩阵乘法的GPU并行程序，并且与对应规模的串行程序进行运行时间的比对（n=500，1000，1500，2000，3000，5000），画出规模和时间对比图。 矩阵A（n，n） 矩阵B（n，n） C = A x B 要求：   1、完成程序的开发并验证其正确性，完成一个实验报告（程序源代码、变量和语句的详细说明；   2、
cuda计算矩阵的乘法
qq_46463876的博客
12-07 372
【代码】cuda计算矩阵的乘法。
cuda cudatoolkit
08-18
CUDA（Compute Unified Device Architecture）是由 NVIDIA 开发的一种并行计算平台和编程模型，用于利用 NVIDIA GPU 进行通用目的的并行计算。 CUDA Toolkit 是一个用于开发和优化 CUDA 应用程序的软件开发工具包。它包含了 GPU 加速计算所需的各种库、编译器、工具和文档。CUDA Toolkit 提供了针对 NVIDIA GPU 的并行计算功能，使开发者能够以高效的方式利用 GPU 进行加速计算。 开发者可以使用 CUDA C/C++、CUDA Fortran、CUDA Python 等编程语言来编写 CUDA 应用程序。CUDA Toolkit 提供了相应的编译器和运行时库，能够将 CUDA 代码转化为 GPU 可执行的指令。 如果你想在使用 CUDA 的项目中进行 GPU 加速计算，需要先安装 CUDA Toolkit，并确保你的系统支持 NVIDIA 的 GPU。你可以从 NVIDIA 官方网站下载适合你系统的 CUDA Toolkit 版本，并按照官方文档进行安装和配置。 希望这个解答对你有所帮助！如果你还有其他问题，请随时提问。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值