#include <stdio.h>
#include <math.h>
#define N 100
const double EPS = 1E-6;

void __global__ add(const double *x, const double *y, double *z, int n) {
    // 获取全局索引
    const int index = blockDim.x * blockIdx.x + threadIdx.x;
    
    // 步长
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride) {
        z[i] = x[i] + y[i];
    }
}

// 误差检测
void check(const double *z, const int n) {
    bool error = false;
    double maxError = 0;
    
    for (int i = 0; i < n; i++) {
        maxError = fmax(maxError, fabs(z[i]-70));
        if (fabs(z[i] - 70) > EPS) {
            error = true;
        }
    }
    
    printf("%s\n", error ? "Errors" : "Pass");
    printf("最大误差: %lf\n", maxError);
}

int main() {
    const int arraySize = sizeof(double) * N;

    // 申请host锁定内存
    double *h_x, *h_y, *h_z;
    cudaMallocHost(&h_x, arraySize);
    cudaMallocHost(&h_y, arraySize);
    cudaMallocHost(&h_z, arraySize);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_x[i] = 50;
        h_y[i] = 20;
    }

    // 申请device显存
    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, arraySize);
    cudaMalloc((void **)&d_y, arraySize);
    cudaMalloc((void **)&d_z, arraySize);
    
    // host数据传输到device
    cudaMemcpy(d_x, h_x, arraySize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, arraySize, cudaMemcpyHostToDevice);

    // 核函数执行配置
    dim3 blockSize(128);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    
    // 执行核函数
    add<<<gridSize, blockSize>>>(d_x, d_y, d_z, N);

    // 将device得到的结果传输到host
    cudaMemcpy(h_z, d_z, arraySize, cudaMemcpyDeviceToHost);
    
    // 检查执行结果
    check(h_z, N);

    // 释放host锁定内存
    cudaFreeHost(h_x);
    cudaFreeHost(h_y);
    cudaFreeHost(h_z);
    
    // 释放device显存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    
    return 0;
}