Unix/Linux服务器技术分享机器学习与数据挖掘CUDA

GPU编程(三): CPU与GPU的矩阵乘法对比

2019-01-21  本文已影响25人  sean_depp

目录


前言

在上一篇的最后, 我提到了一个矩阵乘法, 这次与CPU进行对比, 从中可以很明显GPU在并行计算上的优势.


计时函数

在贴出代码之前, 来看下我常用的计时函数, 可以精确到微秒级. 首先头文件是#include<sys/time.h>. 结构体为:

struct timeval{
    long tv_sec; /*秒*/
    long tv_usec; /*微秒*/
};

来看下使用的小栗子:

struct timeval start, end;
double timeuse;
int sum = 0;

gettimeofday (&start, NULL);
for (int i = 0; i < 10000; i++){
    sum += i;
}
gettimeofday (&end, NULL);

timeuse = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec)/1000000.0;
printf("Use Time:%f\n",timeuse);

代码

其实CPU部分的代码就是for循环. 你可能会考虑到用多线程, 但是我实测效果不太好, 这篇有代码, 可以去看看. 所以用的基础for循环.

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>    
#include <unistd.h>

#define w 1000

struct Matrix
{
    int width;
    int height;
    float *elements;
};

void matMul(float * M, float * N, float * P, int width){
    for (int i = 0; i < width; i++){
        for (int j = 0; j < width; j++){
            float sum = 0;
            for (int k = 0; k < width; k++){
                float a = M[i * width + k];
                float b = N[k * width + j];
                sum += a * b;
            }
            P[i * width + j] = sum;
        }
    }
}

int main(){
    int width = w;
    int height = w; 
    
    float * m = (float *)malloc (width * height * sizeof (float));
    float * n = (float *)malloc (width * height * sizeof (float));
    float * p = (float *)malloc (width * height * sizeof (float));

    for (int i = 0; i < width * height; i++){
        m[i] = 1.0;
        n[i] = 2.0;
    }

    struct timeval t1,t2;
    gettimeofday(&t1,NULL);
    double timeuse;

    matMul(m, n, p, w);

    gettimeofday(&t2,NULL);
    timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0;
    printf("Use Time:%f\n",timeuse);

    return 0;
}

cuda部分的代码直接贴出来, 解析可以看之前的文章.

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#include <unistd.h>

#define w 1000

struct Matrix
{
    int width;
    int height;
    float *elements;
};

__device__ float getElement(Matrix *A, int row, int col)
{
        return A->elements[row * A->width + col];
}

__device__ void setElement(Matrix *A, int row, int col, float value)
{
        A->elements[row * A->width + col] = value;
}

__global__ void matMulKernel(Matrix *A, Matrix *B, Matrix *C)
{
        float Cvalue = 0.0;
        int row = threadIdx.y + blockIdx.y * blockDim.y;
        int col = threadIdx.x + blockIdx.x * blockDim.x;
        
        for (int i = 0; i < A->width; ++i)
        {
                Cvalue += getElement(A, row, i) * getElement(B, i, col);
        }
        setElement(C, row, col, Cvalue);
}

int main()
{
    int width = w;
    int height = w;

    Matrix *A, *B, *C;

    cudaMallocManaged((void**)&A, sizeof(Matrix));
    cudaMallocManaged((void**)&B, sizeof(Matrix));
    cudaMallocManaged((void**)&C, sizeof(Matrix));

    int nBytes = width * height * sizeof(float);

    cudaMallocManaged((void**)&A->elements, nBytes);
    cudaMallocManaged((void**)&B->elements, nBytes);
    cudaMallocManaged((void**)&C->elements, nBytes);

    A->height = height;
    A->width = width;
    B->height = height;
    B->width = width;
    C->height = height;
    C->width = width;

    for (int i = 0; i < width * height; ++i)
    {
        A->elements[i] = 1.0;
        B->elements[i] = 2.0;
    }

    dim3 blockSize(32, 32);
    dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
        (height + blockSize.y - 1) / blockSize.y);

    struct timeval t1,t2;
    gettimeofday(&t1,NULL);
    double timeuse;

    matMulKernel << < gridSize, blockSize >> >(A, B, C);

    cudaDeviceSynchronize();

    gettimeofday(&t2,NULL);
    timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0;
    printf("Use Time:%f\n", timeuse);

    return 0;
}

来看下结果图:

结果图

gpu是gt750m, cpu是i7-4700mq. 其实cpu是比gpu好很多的, 但是并行计算上gpu的优势依旧明显.


最后

喜欢记得点赞哦, 有意见或者建议评论区见~


上一篇下一篇

猜你喜欢

热点阅读