CUDA实战2

1.clock()函数是C/C++中的计时函数，相关的数据类型是clock_t，使用clock函数可以计算运行某一段程序所需的时间，如下所示程序计算从10000000逐渐减一直到0所需的时间。

注：每次运行所需时间可能会不一样

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>
int main()
{
    //测试clock_t的使用
    clock_t start, end;
    long n = 10000000L;
    double duration;
    printf("使 %ld 循环减一变为 0 所需的时间是：",n);
    //开始时间
    start = clock();
    //循环减一
    while(n--);
    //结束时间
    end = clock();
    //计算整个过程的时间结束时间减开始时间）, 
    //CLOCKS_PER_SEC是"time.h"文件中定义的常量，
    //表示一秒钟包含多少时钟计时单元（即毫秒）。
    duration = (double)(end-start) / CLOCKS_PER_SEC;
    printf(" %f 秒
",duration);
    return 0;
}

 

 

 2.

CLOCKS_PER_SEC，它用来表示一秒钟会有多少个时钟计时单元，其定义如下：

#define CLOCKS_PER_SEC ((clock_t)1000)

可以看到每过千分之一秒（1毫秒），调用clock()函数返回的值就加1。

可以使用公式clock()/CLOCKS_PER_SEC来计算一个进程自身的运行时间。

 

矢量求和运算

假设我们有两组数据，我们需要将这两组数据中对应的元素两两相加，并将结果保存在第三个数组中。

//CUDA的头文件
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
//C语言的头文件
#include <stdio.h>
#include <time.h>

#define N 6000
#define thread_num 1024

//GPU函数声明
__global__ void add(int* a, int* b, int* c);
//CPU函数声明
void add_CPU(int *a, int *b,int *c);

int main()
{
    //GPU方法计时声明
    float time_CPU, time_GPU;
    cudaEvent_t start_GPU, stop_GPU, start_CPU, stop_CPU;
    //CPU方法计时声明
    float time_cpu, time_gpu;
    clock_t start_cpu, stop_cpu, start_gpu, stop_gpu;
    int a[N], b[N], c[N], c_CPU[N];
    int *dev_a, *dev_b, *dev_c;
    
    int block_num;
    block_num = (N + thread_num - 1)/thread_num;
    
    //在GPU上分配内存
    cudaMalloc((void**)&dev_a, N*sizeof(int));
    cudaMalloc((void**)&dev_b, N*sizeof(int));
    cudaMalloc((void**)&dev_c, N*sizeof(int));
    
    //在CPU上进行赋值
    for(int i = 0; i < N; i++)
    {
        a[i] = -i;
        b[i] = i*i;
    }
    
 
    //记录当前时间
    start_cpu = clock();
    
    add_CPU(a, b, c_CPU);
    
    stop_cpu = clock();
    //记录当前时间
    printf("Tne time for CPU:	%f(ms)
", (float)(stop_cpu - start_cpu) / CLOCKS_PER_SEC);
    
    
    //输出CPU结果
    printf("
Result from CPU:
");
    for(int i = 0; i<N; i++)
    {
        printf("CPU:	%d+%d=%d
",a[i],b[i],c_CPU[i]);
    }
    
    //GPU计算
    cudaMemcpy(dev_a,a,N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b,b,N*sizeof(int), cudaMemcpyHostToDevice);
    
    //创建Event
    cudaEventCreate(&start_GPU);
    cudaEventCreate(&stop_GPU);
    
    //记录当时时间
    cudaEventRecord(start_GPU,0);
    start_gpu = clock();
    //调用核函数
    add<<<block_num,thread_num>>>(dev_a,dev_b,dev_c);
    
    stop_gpu = clock();
    //记录当时时间
    cudaEventRecord(stop_GPU,0);
    cudaEventSynchronize(start_GPU);
    cudaEventSynchronize(stop_GPU);
    cudaEventElapsedTime(&time_GPU, start_GPU, stop_GPU);
    printf("
The time from GPU :	%f(ms)
",time_GPU);
    
    //将device复制到host
    cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);
    //将GPU中的结果拷贝出来
    cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);
    
    //输出
    printf("
Result from GPU:
");
    for(int i = 0; i<N; i++)
    {
        printf("GPU:	%d+%d=%d
",a[i],b[i],c[i]);
    }
    cudaEventDestroy(start_GPU);
    cudaEventDestroy(stop_GPU);
    
    //释放内存
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
    printf("
The time for CPU by event:	%f(ms)
", time_CPU);
    printf("The time for GPU by event:	%f(ms)
", time_GPU);
    
    time_cpu = (float)(stop_cpu - start_cpu) / CLOCKS_PER_SEC;
    time_gpu = (float)(stop_gpu - start_gpu) / CLOCKS_PER_SEC;
    printf("
The time for CPU by host:	%f(ms)
", time_cpu);
    printf("The time for GPU by host:	%f(ms)
", time_gpu);
    
    
    return 0;
}
//GPU函数
__global__ void add(int *a, int *b, int *c)
{
    int tid = blockIdx.x*blockDim.x+threadIdx.x;//计算该索引处的数据
    if (tid < N)
    {
        c[tid] = a[tid] + b[tid];
    }
}

//CPU函数
void add_CPU(int *a, int *b, int *c)
{
    for (int i = 0; i < N; i++)
    {
        c[i] = a[i] + b[i];
    }
}

 3.CUDA实现矩阵乘

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define M 1024
#define K 1024
#define N 1024

void initial(double* list,int row,int col)
{
    double *num = list;
//    srand((unsigned)time(NULL));
    for (int i=0; i<row*col; i++)
    {
        num[i] = rand()%10;
    }
}

void CpuMatrix(double *A,double *B,double *C)
{
       int i,j,k;
       
       for( i=0; i<M; i++)
       {
           for(j=0; j<N; j++)
           {
               double sum = 0;
               for(int k=0; k<K; k++)
               {
                   sum += A[i*K + k] * B[k * N + j];
               }
               C[i * N + j] = sum;
           }
       }
}

__global__ void GpuMatrix(double *dev_A,double *dev_B,double *dev_C)
{
    int ix = threadIdx.x + blockDim.x * blockIdx.x;
    int iy = threadIdx.y + blockDim.y * blockIdx.y;
    
    if(ix<K && iy<M)
    {
    double sum = 0;
    for( int k = 0; k < K;k++)
    {
        sum += dev_A[iy*K + k] * dev_B[k*N + ix];
    }
    dev_C[iy * N + ix] = sum;
   }
}

void printMatrix(double *list,int row,int col)
{
    double *p = list;
    for(int i=0; i<row; i++)
    {
        for(int j=0; j<col; j++)
        {
            printf("%10lf",p[j]);
        }
        p = p + col;
        printf("
");
    }
}

int main(int argc,char **argv)
{
    clock_t start_cpu,stop_cpu,start_gpu,stop_gpu;
    double time_cpu,time_gpu;
    
    float time_CPU, time_GPU;
    cudaEvent_t start_GPU, stop_GPU, start_CPU, stop_CPU;
    
    //printf("Amatrix:(%d*%d)
",M,K);
    int Axy = M*K;
    int Abytes = Axy * sizeof(double);
    
//    printf("Bmatrix:(%d*%d)
",K,N);
    int Bxy = K*N;
    int Bbytes = Bxy * sizeof(double);
    
    int nxy = M*N;
    int nbytes = nxy * sizeof(double);
    
    double *host_A, *host_B, *host_C, *c_CPU;
    host_A = (double*)malloc(Abytes);
    host_B = (double*)malloc(Bbytes);
    host_C = (double*)malloc(nbytes);
    c_CPU = (double*)malloc(nbytes);
    
    //初始化
    initial(host_A,M,K);
    //输出
    printf("A:(%d,%d):
",M,K);
//    printMatrix(host_A,M,K);
    
    initial(host_B,K,N);
    //输出
    printf("B:(%d,%d):
",K,N);
//    printMatrix(host_B,K,N);
   
    start_cpu = clock();
    CpuMatrix(host_A,host_B,host_C);
    stop_cpu = clock();
    
    printf("The time from CPU is 	%f(ms)
",(float) (stop_cpu-start_cpu) / CLOCKS_PER_SEC);
    //输出
    printf("Host_C:(%d,%d):
",M,N);
//    printMatrix(host_C,M,N);
    
    //GPU计算
    double *dev_A,*dev_B,*dev_C;
    cudaMalloc((void**)&dev_A,Axy*sizeof(double));
    cudaMalloc((void**)&dev_B,Bxy*sizeof(double));
    cudaMalloc((void**)&dev_C,nxy*sizeof(double));
    dim3 block(1024,1024);
    dim3 grid(32,32); 
    cudaMemcpy(dev_A,host_A,Abytes,cudaMemcpyHostToDevice);
    cudaMemcpy(dev_B,host_B,Bbytes,cudaMemcpyHostToDevice);
    //创建Event 
    cudaEventCreate(&start_GPU);
    cudaEventCreate(&stop_GPU);
    cudaEventRecord(start_GPU,0);
    start_gpu = clock();
    
    GpuMatrix<<<grid,block>>>(dev_A,dev_B,dev_C);
    
    stop_gpu = clock();
    cudaEventRecord(stop_GPU,0);
    cudaEventSynchronize(start_GPU);
    cudaEventSynchronize(stop_GPU);
    //计算时间差
    cudaEventElapsedTime(&time_GPU,start_GPU,stop_GPU);
    printf("The time from GPU is 	%f(ms)
",time_GPU);
    //消除Event
    cudaEventDestroy(start_GPU);
    cudaEventDestroy(stop_GPU);
    
    cudaMemcpy(c_CPU,dev_C,nbytes,cudaMemcpyDeviceToHost);
    //输出
    printf("device_C:(%d,%d):
",M,N);
//    printMatrix(c_CPU,M,N);
    
    //释放内存
     cudaFree(dev_A);
     cudaFree(dev_B);
     cudaFree(dev_C);
     free(host_A);
     free(host_B);
     free(host_C);
     free(c_CPU);
     
     time_cpu = (float) (stop_cpu-start_cpu) / CLOCKS_PER_SEC;
     time_gpu = (float) (stop_gpu-start_gpu) / CLOCKS_PER_SEC;
     printf("
The time for CPU by host:	%f(ms)
", time_cpu);
     printf("The time for GPU by host:	%f(ms)
", time_gpu);
     return 0;

}