CUDA编程练习（二） 向量相加

没什么好说的，对应位置相加，除了零代码复制、控制好最大线程在运作外，应该没什么访存优化的地方（对全局内存一定是读\*2，写\*1）

可能以后我会觉得上面的话有问题，有问题能暴露出来是好事

%%cuda

#include <iostream>
#include <assert.h>
#include <cstdlib>
#include <cmath>

#define VECTOR_LENGTH 1000000
#define MAX_ERR 1e-4

__global__ void dummy_vectorAdd(float *out, float *a, float *b, int n) 
{
    for(int i = 0; i < n; i++)
    {
        out[i] = a[i] + b[i];
    }
}

__global__ void vectorAdd(float *out, float *a, float *b, int n) 
{
    int leap=gridDim.x*blockDim.x;
    int tid = blockIdx.x*blockDim.x+threadIdx.x;
    for(int i = tid; i < n; i+=leap)
    {
        out[i] = a[i] + b[i];
    }
}

double cpuGpu_gap(float *a, float *b, int n) {
    double tmp=0;
    for (int i=0; i<n; i++) {
        tmp += abs(a[i]-b[i]);
    }
    return tmp;
}

int main() {
    // 声明变量
    float *a, *b, *cpuout, *out, *out1;
    float *d_a, *d_b, *d_out, *d_out1; 
    a = new float[VECTOR_LENGTH];
    b = new float[VECTOR_LENGTH];
    cpuout = new float[VECTOR_LENGTH];
    out = new float[VECTOR_LENGTH];
    out1 = new float[VECTOR_LENGTH];
    
    // 初始化
    for(int i = 0; i < VECTOR_LENGTH; i++)
    {
        a[i] = std::rand();
        b[i] = std::rand();
    }
    // cpu端计算
    for (int i=0; i<VECTOR_LENGTH; i++) {
        cpuout[i] = a[i]+b[i];
    }
    
    
    // 创建gpu端空间，复制到gpu端
    cudaMalloc(&d_a, sizeof(float) * VECTOR_LENGTH);
    cudaMalloc(&d_b, sizeof(float) * VECTOR_LENGTH);
    cudaMalloc(&d_out, sizeof(float) * VECTOR_LENGTH);  // 线性空间申请
    cudaMalloc(&d_out1, sizeof(float) * VECTOR_LENGTH);
    cudaMemcpy(d_a, a, sizeof(float) * VECTOR_LENGTH, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(float) * VECTOR_LENGTH, cudaMemcpyHostToDevice);
    
    //记录时间
    float time_gpu, time_gpu1;
    cudaEvent_t start, stop;
    
    //创建Event
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    //记录当前时间
    cudaEventRecord(start);
    cudaEventQuery(start);
    
    dummy_vectorAdd<<<1,1>>>(d_out, d_a, d_b, VECTOR_LENGTH);
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);    //等待事件完成。记录之前的任务
    cudaEventElapsedTime(&time_gpu, start, stop);    //计算时间差
    printf("The time for dummy_vectorAdd:\t%f(ms)\n", time_gpu);
    // 转移回CPU计算差值
    cudaMemcpy(out, d_out, sizeof(float) * VECTOR_LENGTH, cudaMemcpyDeviceToHost);
    printf("The gap for dummy_vectorAdd:\t%f\n", cpuGpu_gap(cpuout, out, VECTOR_LENGTH));
    cudaEventDestroy(start);    //消除Event
    cudaEventDestroy(stop);
    
    //  ------------------------------------------
    //创建Event
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    //记录当前时间
    cudaEventRecord(start);
    cudaEventQuery(start);
    
    vectorAdd<<<384,512>>>(d_out1, d_a, d_b, VECTOR_LENGTH);  // SM 72 线程最大1024
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);    //等待事件完成。记录之前的任务
    cudaEventElapsedTime(&time_gpu, start, stop);    //计算时间差
    printf("The time for vectorAdd:\t%f(ms)\n", time_gpu);
    // 转移回CPU计算差值
    cudaMemcpy(out1, d_out1, sizeof(float) * VECTOR_LENGTH, cudaMemcpyDeviceToHost);
    printf("The gap for vectorAdd:\t%f\n", cpuGpu_gap(cpuout, out1, VECTOR_LENGTH));
    cudaEventDestroy(start);    //消除Event
    cudaEventDestroy(stop);
    
    
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_out);
    cudaFree(d_out1);
    
    std::cout<<"以上就是并行与不并行的速度差距"<<std::endl;
}

结果如下

The time for dummy_vectorAdd:	70.540031(ms)
The gap for dummy_vectorAdd:	0.000000
The time for vectorAdd:	0.037600(ms)
The gap for vectorAdd:	0.000000
以上就是并行与不并行的速度差距

乍一看，仅仅是双元素累加，因为顺序一定是固定的，不会出现精度损失；

注意了，这里面是有问题的，统计时间没有考虑传输时间。不过这里只是为了展示并行的必要性以及尝试并行方式，就先不纠结了。