CUDA编程 -- 3矩阵乘法

版权声明:林微原创,转载请注明出处: https://blog.csdn.net/Canhui_WANG/article/details/82913855

我们已经知道了threads/blocks在CUDA端的组织方式,接下来我们学学多维度空间下的多线程模型,下面以矩阵乘法为例。


1. 行优先

存储方式

二维矩阵在内存中的存储方式受到编程语言的影响,主要可以分为两种:行优先和列优先。对于编程语言诸如C/C++/CUDA而言,数据在内存中的组织方式是行优先。举例说明行优先的存储方式,如下,

给定一个3×\times 3大小的矩阵AAA3×3=[a1,1a1,2a1,3a2,1a2,2a2,3a3,1a3,2a3,3]A_{3\times 3}=\begin{bmatrix} a_{1,1} & a_{1,2} & a_{1,3}\\ a_{2,1} & a_{2,2} & a_{2,3}\\ a_{3,1} & a_{3,2} & a_{3,3} \end{bmatrix},在行优先的内存模型下面,二维矩阵将逐行地顺序存储,比方说,将第一行矩阵的数据依次存储完毕之后,接着存储第二行矩阵的数据;在第二行矩阵数据存储完毕之后,接着存储第三行矩阵的数据,以此类推,直到将整个矩阵存储完毕为止。直观上,矩阵AA在内存中的存储位置如下,

A3×3=(a1,1a1,2a1,3a2,1a2,2a2,3a3,1a3,2a3,3)A_{3\times 3}=\begin{pmatrix} a_{1,1} & a_{1,2} & a_{1,3} & a_{2,1} & a_{2,2} & a_{2,3} & a_{3,1} & a_{3,2} & a_{3,3} \end{pmatrix}

位置索引

基于给定的n×mn \times m大小的矩阵AA,则ai,ja_{i,j}在内存中的位置索引为i×m+ji\times m + j,其中,0in10 \leq i \leq n-10jm10 \leq j \leq m-1.



2. CPU实现和分析

CPU核心算法

根据行优先的位置索引,我们可以容易地计算矩阵的乘法如下,

for(int row = 0; row < matrix_width; row++)
{
    for(int col = 0; col < matrix_width; col++)
    {
        int single_element = 0;
        for(int k = 0; k < matrix_width; k++)
        {
            single_element += matrix_a_host[row * matrix_width + k] * matrix_b_host[matrix_width  * k + col];
        }
        matrix_c_host[row * matrix_width + col] = single_element;
    }
}

注:矩阵乘法需要保证第一个矩阵的列数等于第二个矩阵的行数。

算法复杂度

基于CPU实现的矩阵乘法的算法复杂度为O(n3)O\left ( n^{3} \right )



3. GPU实现和分析

Threads/Blocks组织方式

出于编程上的方便,我们采用一维的blocks组织方式,二维的threads组织方式,如下,

dim3 dimGrid(1, 1, 1);
dim3 dimBlock(3, 3, 1);

注:默认情况下threads/blocks均采用一维的组织方式。

GPU核心算法
__global__ void matrixs_1D_multiplication(int *matrix_a_dev, int *matrix_b_dev, int *matrix_c_dev, int matrix_width)
{
    int row = threadIdx.y;
    int col = threadIdx.x;

    if(row < matrix_width && col < matrix_width)
    {
        for(int k = 0; k < matrix_width; k++)
        {
            matrix_c_dev[row * matrix_width + col] += matrix_a_dev[row * matrix_width + k] * matrix_b_dev[k * matrix_width + col]; 
        }
    }
}
算法复杂度

基于GPU实现的矩阵乘法的算法复杂度为O(n)O\left ( n \right )



4. 编译调试

CPU版本

源代码:matrix_mul_cpu.cu

编译

nvcc matrix_mul_cpu.cu -o main

运行

./main
GPU版本

源代码:matrix_mul_gpu.cu

编译

nvcc matrix_mul_gpu.cu -o main

运行

./main
展开阅读全文

没有更多推荐了,返回首页