企业官网建设流程全解析

第一阶段：建立最小闭环（纯 elementwise）

1. Vector Add

#include <cstdio> #include <cstdlib> #include <cmath> #include <cuda_runtime.h> //MXMACA通过 cu-bridge兼容CUDA Runtime API // 错误检查：第一个必须建立的验证习惯 #define CHECK_CUDA(call) \ do { \ cudaError_t err = (call); \ if (err != cudaSuccess) { \ fprintf(stderr,"CUDA error at %s:%d: %s\n", \ __FILE__,__LINE__,cudaGetErrorString(err));\ exit(EXIT_FAILURE); \ } \ }while(0) //========================= // 1. CPU Baseline:只依赖标准C++，不碰任何 GPU API //========================== void vector_add_cpu(int n, const float* a, const float* b, float* c) { for (int i=0; i<n; ++i) { c[i]=a[i] + b[i]; } } //=================== //2.GPU Kernel:每个线程负责一个元素，带边界判断 //=================== __global__ void vector_add_kernel(int n, const float* a, const float*b, float* c) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i<n) { // 向上取整带来的“多余线程”必须在这里退出 c[i]=a[i]+b[i]; } } // ============================================================ // 3. 逐元素误差比对：验证习惯的核心 // 这里用 abs_err <= atol + rtol * |expected| // 对于纯粹的 float 加法，理论上可以要求“逐位完全相等”， // 但用相对/绝对容差是更通用的习惯，后续算子直接复用。 // ============================================================ bool verify(const float* gpu_result, const float* cpu_baseline, int n, float atol = 1e-5f, float rtol = 1e-5f) { int mismatch_count = 0; float max_abs_err = 0.0f; for (int i = 0; i < n; ++i) { float g = gpu_result[i]; float c = cpu_baseline[i]; float abs_err = fabsf(g - c); float tol = atol + rtol * fabsf(c); if (abs_err > tol) { if (mismatch_count < 10) { // 只打印前若干个失配，避免刷屏 fprintf(stderr, "MISMATCH at [%d]: gpu=%.9g cpu=%.9g abs_err=%.3g tol=%.3g\n", i, g, c, abs_err, tol); } ++mismatch_count; } if (abs_err > max_abs_err) max_abs_err = abs_err; } if (mismatch_count > 0) { fprintf(stderr, "VERIFY FAILED: %d / %d elements out of tolerance. max_abs_err=%.3g\n", mismatch_count, n, max_abs_err); return false; } printf("VERIFY PASSED: all %d elements within atol=%.1e, rtol=%.1e (max_abs_err=%.3g)\n", n, atol, rtol, max_abs_err); return true; } // ============================================================ // 4. Host 主流程：分配 → 初始化 → CPU 基准 → GPU 执行 → 回读 → 验证 // ============================================================ int main() { const int N = 1 << 20; // 1M 元素，足够暴露边界和线程配置问题 const size_t bytes = N * sizeof(float); // ---- Host 内存 ---- float* h_a = (float*)malloc(bytes); float* h_b = (float*)malloc(bytes); float* h_c_cpu = (float*)malloc(bytes); // CPU baseline 结果 float* h_c_gpu = (float*)malloc(bytes); // 从 GPU 回读的结果 // 用确定性数据初始化，避免 rand() 依赖种子导致无法复现 for (int i = 0; i < N; ++i) { h_a[i] = static_cast<float>(i % 1024) * 0.001f; h_b[i] = static_cast<float>((i * 7) % 1024) * 0.002f; } // ---- 先跑 CPU baseline，这是后续一切比对的“真值” ---- vector_add_cpu(N, h_a, h_b, h_c_cpu); // ---- Device 内存 ---- float *d_a, *d_b, *d_c; CHECK_CUDA(cudaMalloc(&d_a, bytes)); CHECK_CUDA(cudaMalloc(&d_b, bytes)); CHECK_CUDA(cudaMalloc(&d_c, bytes)); // ---- H2D ---- CHECK_CUDA(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); CHECK_CUDA(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); // ---- Launch ---- const int threads = 256; const int blocks = (N + threads - 1) / threads; // 向上取整，边界判断在 kernel 内兜底 printf("Launching kernel: N=%d, blocks=%d, threads=%d\n", N, blocks, threads); vector_add_kernel<<<blocks, threads>>>(N, d_a, d_b, d_c); CHECK_CUDA(cudaGetLastError()); // 捕获 launch 本身的配置错误 CHECK_CUDA(cudaDeviceSynchronize()); // 等待 kernel 真正完成，再回读 // ---- D2H ---- CHECK_CUDA(cudaMemcpy(h_c_gpu, d_c, bytes, cudaMemcpyDeviceToHost)); // ---- 验证：GPU 结果 vs CPU baseline，逐元素比对 ---- bool ok = verify(h_c_gpu, h_c_cpu, N); // ---- 释放 ---- CHECK_CUDA(cudaFree(d_a)); CHECK_CUDA(cudaFree(d_b)); CHECK_CUDA(cudaFree(d_c)); free(h_a); free(h_b); free(h_c_cpu); free(h_c_gpu); return ok ? 0 : 1; }

Q：容差阈值 ： 允许的误差阈值 = 绝对容差 + 相对容差 × |CPU 值|， 其中绝对容差为人为设定的，一般为1e-5，相对容差为 一个缩放系数，为1e-5。 其用于判断，是否误差在容许范围内

if (abs_err > tol)。 但是为什么要用这样的计算方式呢，为什么不直接用绝对容差呢，或直接用相对容差x |CPU值|呢，而且后面这一坨为什么要这么计算呢？

A：用绝对容差，是避免 在CPU中计算出来的值接近0，导致错判。如 CPU值是0，而绝对误差（|gpu值-CPU值|）为1e-8，这明显是可以容忍的，但判断 abs_err>tol 为true，即程序认为不合理，导致错判。 所以 有绝对容差的意义就体现了。 （为小数兜底）

同时，如果只有后面相对容差那一坨也不行。 如果CPU值计算出来是一个很大数，如 CPU值是1e8，其实如果误差是10的话是容许的，这样就体现它的作用了。（为大数兜底）

为啥要乘|CPU值|呢，因为默认|CPU值|是标准值。

第二阶段：引入规约与线程协作

第三阶段：引入数据重排与分块访存

第四阶段：复杂融合算子与流水线