cpp文件

#include "stdafx.h"
#include <stdio.h>
#include <stdlib.h> //为rand()及srand()提供函数声明
#include <time.h>

extern "C" int mulWithCuda(float *c, const float *a, const float *b, int size);

int _tmain(int argc, _TCHAR* argv[])
{
   int i = 0, j = 0, k = 0;
   float sum = 0;
   int size = 8;
   srand(time(NULL));
   float * matrix_a = (float *)malloc(size * size * sizeof(float));     //创建一维数组
   float * matrix_b = (float *)malloc(size * size * sizeof(float));     //创建一维数组
   float * matrix_c = (float *)malloc(size * size * sizeof(float));     //创建一维数组
   float * matrix_d = (float *)malloc(size * size * sizeof(float));     //创建一维数组

for (i = 0; i < size; i++)
   {
       for (j = 0; j < size; j++)
       {
           //生成随机数
           *(matrix_a + i * size + j) = (float)rand() / (RAND_MAX / 10);
           *(matrix_b + i * size + j) = (float)rand() / (RAND_MAX / 10);
       }
   }
   for (i = 0; i < size; i++)
   {
       for (j = 0; j < size; j++)
       {
           printf("%f ", *(matrix_a + i * size + j));
       }
       printf("\n");
   }
   printf("\n");

for (i = 0; i < size; i++)
   {
       for (j = 0; j < size; j++)
       {
           printf("%f ", *(matrix_b + i * size + j));
       }
       printf("\n");
   }
   printf("\n");
   clock_t start = clock();
   for (i = 0; i < size; i++)
   {
       for (j = 0; j < size; j++)
       {
           sum = 0;
           for (k = 0; k < size; k++)
           {
               sum = sum + *(matrix_a + i * size + k) * (*(matrix_b + k * size + j));
           }
           *(matrix_d + i * size + j) = sum;
           printf("%f ", sum);
       }
       printf("\n");
   }
   clock_t end = clock();
   double interval = double(end - start) / CLK_TCK;
   printf("CPU运行时间为:%lf\n", interval);

// Add vectors in parallel.
   clock_t start1 = clock();
   int cudaStatus = mulWithCuda(matrix_c, matrix_a, matrix_b, size);
   clock_t end1 = clock();
   double interval1 = double(end1 - start1) / CLK_TCK;
   printf("GPU运行时间为:%lf\n", interval1);
   //printf("加速比为:%lf\n", interval / interval1);
   for (i = 0; i < size; i++)
   {
   for (j = 0; j < size; j++)
   {
   //*(matrix_c + i * size + j) = *(matrix_a + i * size + j) + *(matrix_b + i * size + j);
   printf("%f ", *(matrix_c + i * size + j));
   }
   printf("\n");
   }
   // cudaDeviceReset must be called before exiting in order for profiling and
   // tracing tools such as Nsight and Visual Profiler to show complete traces.

return 0;
}

kernel.cu

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#define dimen 16

__global__ void mulKernel1(float* c, float* a, float* b, int size)
{
   float sum = 0;
    //创建共享存储器数组
    __shared__ float A[dimen][dimen];
    __shared__ float B[dimen][dimen];

//计算线程的不同维度索引
    int bx=blockIdx.x;
    int by=blockIdx.y;
    int tx=threadIdx.x;
    int ty=threadIdx.y;
    //计算得到结果在线程块中的索引
    int row = by*dimen+ty;
    int col = bx*dimen+tx;

for (int k=0;k<(size-1)/dimen+1;k++)
    {
        //加载数据到共享存储器
       if(k * dimen + tx < size && row < size)
       {
           A[ty][tx]=a[row*size+k*dimen+tx];
       }
       else
       {
           A[ty][tx] = 0;
       }
       if(k * dimen + ty< size && col < size)
       {
           B[ty][tx]=b[(k*dimen+ty)*size+col];
       }
       else
       {
           B[ty][tx] = 0;
       }
        __syncthreads(); //同步块内其他线程，使A和B共享存储器中数据填满

for(int m=0;m<dimen;m++)
            sum +=A[ty][m]*B[m][tx];
        __syncthreads(); //同步矩阵上移动的共享块，使得sum的值完整
    }
    //将sum值赋值到线程对应的矩阵位置
   if(row < size && col < size)
   {
        c[row*size+col]=sum;
   }

}
__global__ void mulKernel(float *c,float *a, float *b, int size)
{
   int threadx=threadIdx.x;
   int thready=threadIdx.y;
   int blockx=blockIdx.x;
   int blocky=blockIdx.y;
   int row = blockx*blockDim.x+threadx;
   int col = blocky*blockDim.x+thready;
    float sum=0;
   for(int i=0;i<size;i++)
       sum+=a[row*size+i]*b[i*size+col];
   if(row < size && col < size)           //增加判断条件，避免由于矩阵分块不匹配导致的计算错误
   {
       c[row*size+col]=sum;
   }
}

// Helper function for using CUDA to add vectors in parallel.
extern "C" int mulWithCuda(float *c, float *a, float *b, int size)
{
    float *dev_a = 0;
    float *dev_b = 0;
    float *dev_c = 0;
    cudaError_t cudaStatus;

// Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

// Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * size * sizeof(float));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

cudaStatus = cudaMalloc((void**)&dev_a, size * size * sizeof(float));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

cudaStatus = cudaMalloc((void**)&dev_b, size * size * sizeof(float));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

// Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * size * sizeof(float), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

cudaStatus = cudaMemcpy(dev_b, b, size * size * sizeof(float), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

dim3 grid((size-1)/dimen+1,(size-1)/dimen+1,1);
dim3 block(dimen,dimen,1);

// Launch a kernel on the GPU with one thread for each element.
mulKernel1<<<grid, block>>>(dev_c, dev_a, dev_b, size);

// Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

// Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}

CUDA 分块矩阵乘法相关推荐

分块矩阵乘法+乒乓操作
本文用system verilog实现了分块矩阵乘法中计算输出矩阵的某一块,并且进行了pingpang操作,以掩盖数据传输时间. 这是顶层模块的代码: `timescale 1ns / 1ps // ...
cache 在X86和ARM的性能比较 - 矩阵累加和分块矩阵乘法
有一段时间在x86和arm服务器下面做开发,需要平台之间的移植,然后经常发现同一段代码在不同平台下面的表现不一样,有一大部分原因是不同平台对cache处理方法不一样. 大部分参考资料上说,cache有 ...
MPI编程——分块矩阵乘法(cannon算法)
要求: 分析本题难点在于不同process之间的通信,算法主要利用了cannon算法,cannon算法描述如下: 以上算法主要分为两个过程:分配初始位置.进行乘-加运算.循环单步移位.为了方便,下面 ...
c++实现矩阵乘法和分块矩阵乘法
矩阵A大小 : m * p,矩阵B大小 : p * m,结果矩阵C大小 : m * n,分块的大小为k * k. 废话少说,原理也不提,直接上代码 #include "iostream&qu ...
编译器O2优化下，分块矩阵乘法的TLB分析猜想
直接将写在实验报告里的那段放进去就算了,好累. 3.3(2分)对最优分块大小的分析实验表明,分块大小为 32 时性能最好.这个结果和你的预期一致吗? 不一致 .如果不一致,其原因在于使用perf工 ...
C语言分块矩阵乘法,c语言矩阵相乘
该楼层疑似违规已被系统折叠隐藏此楼查看此楼程序清单 #include&nbsp int&nbspmain(void) { &nbsp&nbsp&nbsp&a ...
CUDA加速计算矩阵乘法进阶玩法(共享内存)
CUDA加速计算矩阵乘法&进阶玩法~共享内存一.基础版矩阵乘法二.为什么可以利用共享内存加速矩阵乘法 1.CUDA内存读写速度比较 2.申请共享内存三.改进版矩阵乘法(利用共享内存) 一 ...
CUDA: 矩阵乘法优化
矩阵乘法是有实用价值的程序,我们会使用浮点数. 虽然矩阵乘法有点老套,不过因为它相当简单,而且也可以用来介绍一些有关 CUDA 的有趣性质. 矩阵乘法为了单纯起见,我们这里以方形的矩阵为例子.基本上 ...
矩阵乘法——CUDA 优化记录
CUDA 是 NVIDIA 的 GPGPU 模型,它使用 C 语言为基础,可以直接以大多数人熟悉的 C 语言,写出在显示芯片上执行的程序,而不需要去学习特定的显示芯片的指令或是特殊的结构." ...

CUDA 分块矩阵乘法

cpp文件

kernel.cu

CUDA 分块矩阵乘法相关推荐

最新文章

热门文章