简介这是一本由Manning出版社出版的《Parallel and High Performance Computing》技术专著PDF面向具备基础编程与计算机体系结构知识的中高级开发者、科研人员及高校师生系统讲解CPU与GPU并行计算的核心原理与工程实践。全书覆盖并行计算动机、性能建模、数据设计、算法模式、向量化优化、OpenMP高效实现及MPI分布式编程等关键模块内容深度适配HPC应用开发、科学计算加速与异构平台调优场景。资源为单文件PDF大小35.83MB排版规范、图文并茂含完整目录、章节习题与作者致谢便于系统研读与重点查阅。已有692人学习下载读者可直接获取原版高清电子书掌握从串行到并行的思维跃迁路径、主流并行框架的实操要点以及真实性能瓶颈的定位与优化方法。1. GPU 并行不是“把循环塞进 kernel 就完事”它本质是一场对内存带宽与计算吞吐的精密协奏你写了一个 CUDA kernel每个线程处理一个数组元素gridDim.x * blockDim.x N看似完美编译通过跑起来也比 CPU 快——但实测加速比只有 2.3×而你的 A100 理论峰值是 CPU 的 20 倍以上。你翻遍文档加__restrict__、改float为float4、甚至手动 unroll 循环结果性能纹丝不动。这不是玄学也不是“显卡不行”而是你正站在 GPU 并行最核心的断层面上GPU 不是更快的 CPU它是为极致数据并行而生的带宽机器它的“快”90% 取决于你有没有让 10240 个线程同时从 L1 缓存里拿到数据而不是在等全局内存吐出下一个字节。这本书《Parallel and High Performance Computing》之所以被某高校 HPC 实验室列为必读材料并非因为它教你怎么写cudaMalloc而是它用整整 Part 39–13 章把 GPU 并行的底层契约拆解成可验证、可测量、可调试的工程事实从 SM 的 warp 调度器如何隐藏访存延迟到 PCI-e 5.0 的 64 GB/s 带宽如何被一次cudaMemcpy吃掉 80%再到__syncthreads()在物理上到底同步了什么层级的 cache line。它不告诉你“该用 OpenACC”而是告诉你当你在#pragma acc parallel loop里漏掉copyin(a[0:n])编译器会默默把整个a数组从主机内存逐页拷贝——而这个动作在 GPU 上可能耗时 37ms远超 kernel 执行的 2ms。搞懂这个你才真正跨过“能跑”和“跑得值”的分水岭。2. GPU 架构的本质不是“更多核”而是“更窄的核 更宽的内存通道”GPU 的并行能力常被简化为“几千个 CUDA 核心”但这完全误导了性能调优的方向。真实情况是现代 GPU如 NVIDIA Ampere的 Streaming MultiprocessorSM里每个 warp32 线程共享一套指令发射单元和寄存器文件但它们的算力爆发依赖于极宽的数据通路——A100 的 L1 缓存带宽高达 2 TB/s是其 L2 缓存带宽2 TB/s的 2 倍更是 CPU L3 带宽~300 GB/s的 6 倍以上。这意味着GPU 的“并行”不是靠单线程快而是靠 32 个线程组成 warp在同一周期内发起 32 次对齐的 32-bit 访存请求由硬件自动合并成一条 1024-bit 的宽总线事务。如果你的数据布局导致这 32 次请求分散在不同 cache line合并失败带宽利用率立刻跌穿 30%。这本书第 9 章开篇就用一张 SM 内部结构图点破本质SM 中的 “warp scheduler” 不调度“计算”它调度的是“访存请求队列”而 “dispatch unit” 发射的不是 ALU 指令是“load/store 向量指令”。理解这点才能看懂为什么float4 a[N]比float a[N]快——不是因为向量化计算而是因为一次ld.global.v4.f32指令能取 16 字节对齐数据完美匹配 warp 的 32 线程并发需求。2.1 Warp 执行模型锁步lockstep不是限制而是带宽优化的强制协议Warp 内 32 个线程必须执行同一条指令SIMT常被初学者视为枷锁。但书中 9.3 节明确指出这是硬件为消除分支发散divergence导致的带宽浪费而设的硬性保障。当 warp 中所有线程访问同一 block 内连续地址如a[tid], a[tid1], ..., a[tid31]硬件可将 32 次 4-byte 请求合并为一次 128-byte 的 cache line 读取一旦出现分支如if (tid % 2 0) a[tid] ...硬件必须序列化执行两组线程带宽利用率直接腰斩。书中给出实测数据在 V100 上无分支的 stencil kernel 达到 85% 的理论内存带宽加入简单条件判断后带宽跌至 41%。这不是 bug是设计——warp lockstep 强制你把“数据访问模式”作为第一设计约束。// ✅ 高效warp 内所有线程访问连续对齐地址 __global__ void good_stencil(float* __restrict__ out, float* __restrict__ in, int n) { int tid blockIdx.x * blockDim.x threadIdx.x; if (tid 1 tid n-1) { // 所有线程读 in[tid-1], in[tid], in[tid1] —— 地址连续 out[tid] 0.25f * (in[tid-1] 2.0f * in[tid] in[tid1]); } } // ❌ 低效tid % 32 导致 warp 内地址跳跃无法合并 __global__ void bad_stencil(float* __restrict__ out, float* __restrict__ in, int n) { int tid blockIdx.x * blockDim.x threadIdx.x; if (tid % 32 0 tid n) { // 仅每 32 个线程中 1 个执行 out[tid] in[tid] * 2.0f; } }提示__restrict__关键字在此处不是可选优化而是向编译器声明out和in指针不重叠——若编译器无法确认这一点它可能插入额外的内存屏障或放弃向量化直接导致带宽下降 15%~20%。2.2 GPU 内存层次L1/L2/Global 不是缓存“大小”问题而是“延迟容忍窗口”问题CPU 工程师谈缓存关注命中率GPU 工程师谈缓存关注延迟掩盖能力。书中 9.3 节用一张表格对比关键参数以 A100 为例内存层级延迟cycle带宽GB/s容量设计目标Register1~10 TB/s256 KB/SM存放 warp 私有变量零等待L1 Cache20–302 TB/s128 KB/SM合并 warp 访存容忍 20 cycle 延迟L2 Cache80–1002 TB/s40 MB服务全 chip容忍 100 cycle 延迟Global Mem800–12002 TB/s40–80 GB用 10240 个 warp 并发隐藏延迟关键洞察在于GPU 不靠“减少延迟”而靠“增加并发”来填满长延迟管道。当一个 warp 因 global memory 访问停顿 1000 cyclesSM 的 warp scheduler 会立即切换到另一个已准备好指令的 warp 执行——只要 SM 中有 ≥32 个活跃 warp就能保证计算单元永不空闲。书中强调cudaOccupancyMaxPotentialBlockSize返回的最优blockSize本质就是让每个 SM 能驻留足够多的 warp通常 32–64 个来覆盖最坏延迟。如果你的 kernel 寄存器用量过高 64/线程SM 能驻留的 warp 数锐减延迟掩盖失效性能雪崩。2.3 PCI-e 总线CPU-GPU 数据搬运不是“拷贝”而是“带宽税”GPU 再快数据进不来也是空谈。书中 9.4 节直指痛点PCI-e 5.0 x16 的理论带宽虽达 64 GB/s但实际cudaMemcpy吞吐常卡在 12–18 GB/s。原因有三协议开销PCI-e 数据包含 20 byte header小数据块传输效率极低CPU 内存控制器瓶颈主机端 DDR5 带宽~80 GB/s虽高但 PCIe root complex 与内存控制器间存在争用同步成本cudaMemcpy是同步操作CPU 线程阻塞等待 DMA 完成期间无法重叠计算。书中给出硬核方案用cudaMemcpyAsync pinned memory页锁定内存将带宽推至 25 GB/s并强调——pinned memory 不是“更快”而是“可预测”它绕过 OS page fault 处理使 DMA 控制器能直接寻址物理页避免 TLB miss 导致的微秒级抖动。# 实测命令用 nvbandwidth 工具验证你的系统真实 PCIe 带宽 $ nvbandwidth -d 0 -t memcpy -s 1G # 测 GPU-GPU 带宽绕过PCIe $ nvbandwidth -d 0 -t h2d -s 1G # 测 Host-Device 带宽走PCIe # 若 h2d 带宽 15 GB/s优先检查是否用了 pinned memory注意cudaMallocHost()分配的 pinned memory 会占用宝贵的物理内存且过多会导致系统 OOM。书中建议只对频繁传输的 buffer如神经网络输入/输出使用中间特征图用cudaMalloc即可。2.4 避坑GPU 并行开发中最易踩的 4 个“带宽陷阱”这些不是理论风险而是某实验室在移植 CFD 代码时血泪记录的真实翻车现场现象kernel 执行时间稳定在 1.2ms但nvprof显示gpu__inst_executed仅 1.5M 条指令远低于 A100 的 100M/ms 理论吞吐。原因kernel 中存在未对齐的float3结构体数组导致每次ld.global.v3.f32无法合并触发 3 次独立 4-byte 访存warp 利用率不足 10%。解决改用float4或alignas(16)强制 16 字节对齐带宽提升 3.2×。现象OpenACC 版本比 CUDA 版本慢 40%nvvp分析显示data movement占总耗时 65%。原因#pragma acc data copyin(a[0:n])中n是运行时变量编译器无法做 size 推导被迫每次拷贝整个a数组实际只需前n个元素。解决改用copyin(a[0:n])为copyin(a[0:n])语法不变但确保n是 const 表达式或显式用acc_copyin()API。现象多 GPU 训练时nvidia-smi显示 GPU 0 利用率 95%GPU 1 仅 30%nsys追踪发现 GPU 1 长时间等待cudaStreamSynchronize()。原因所有 GPU 共享同一个 host 线程调用cudaStreamSynchronize()形成串行等待GPU 1 必须等 GPU 0 的 stream 完成才开始同步。解决为每张 GPU 创建独立 host 线程或改用cudaEventRecord()cudaEventSynchronize()实现异步等待。现象cudaMalloc分配 2GB 显存成功但 kernel 启动报cudaErrorLaunchOutOfResources。原因SM 的寄存器总量有限A100 为 65536/SM若 kernel 每线程用 128 寄存器blockSize1024时需 131072 寄存器 65536编译器静默降频至blockSize512导致 occupancy 不足。解决用nvcc --ptxas-options-v查看寄存器用量或用__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)强制编译器优化。3. 并行算法模式不是“写 for 循环”而是“构造可扩展的计算拓扑”GPU 的强大不在于它能跑多少线程而在于它能让你用极低成本构建大规模、低耦合、高局部性的计算结构。书中第 5 章和第 10 章反复强调成功的 GPU 算法本质是把问题映射成“NDRangeN维网格→ Work-group工作组→ Work-item工作项”三级拓扑且每一级都满足数据局部性约束。例如前缀和prefix sum看似简单但 naive 实现每个线程算完整前缀会导致严重 bank conflict 和重复计算而书中的 work-efficient scan10.6 节将其拆解为先按 work-group 内 32 线程做 local scan利用 shared memory 零延迟再对各 work-group 的 partial result 做 global scan最后 work-group 内广播 offset 并修正——整个过程无原子操作、无全局同步带宽利用率超 92%。这种“分治聚合”的拓扑思维才是 GPU 并行的底层范式。3.1 Spatial Hashing用哈希表打破“网格束缚”实现动态负载均衡传统 stencil 计算假设数据在规则网格上但真实物理模拟如粒子系统中粒子位置随机分布。书中 5.5 节提出的 spatial hashing 是经典解法将三维空间划分为 cell用哈希函数h(x,y,z) (xs) ^ (ys) ^ (zs)将粒子映射到 bucket再用 linked list 或 AoSoA 存储同 bucket 粒子。关键在于——哈希本身不并行但“对每个 bucket 内粒子做 pairwise interaction”可完美并行。书中给出实测在 1M 粒子系统中spatial hashing GPU 并行比 CPU OpenMP 快 47×且扩展性近乎线性16 GPU 时达 72×。其成功根源在于哈希将“全局稀疏问题”转化为“大量独立密集子问题”每个子问题天然适配 GPU 的 massively parallel 模型。3.2 Reduction 模式为什么__syncthreads()不能解决跨 work-group 同步Reduction规约是并行基石但新手常误以为__syncthreads()能同步所有线程。书中 10.4 节一针见血__syncthreads()仅同步同一 block 内的线程而跨 block 的 reduction 必须经 global memory 中转。高效方案是两级 reductionLevel 1每个 block 内用 shared memory __syncthreads()做 local reduction输出 1 个 partial sumLevel 2将所有 partial sum 写入 global memory启动新 kernel或用 atomicAdd做 final reduction。书中强调Level 1 的 shared memory 使用必须规避 bank conflict——若用sdata[tid]存储当tid为 32 的倍数时所有线程访问同一 memory bank带宽归零。正确做法是sdata[tid/32][tid%32]或直接用shfl_down_syncAmpere 支持。// ✅ 无 bank conflict 的 block-level reductionAmpere __device__ float warpReduceSum(float val) { for (int offset 16; offset 0; offset / 2) val shfl_down_sync(0xFFFFFFFF, val, offset); return val; } __global__ void reduction_kernel(float* input, float* output, int n) { extern __shared__ float sdata[]; int tid threadIdx.x; float sum (tid n) ? input[blockIdx.x * blockDim.x tid] : 0.0f; // Block-level reduction into sdata[0] for (int s blockDim.x / 2; s 32; s 1) { if (tid s) sdata[tid] sum sdata[tid s]; __syncthreads(); } if (tid 32) sdata[tid] sum warpReduceSum(sum); // 最后 32 线程用 warp shuffle if (tid 0) output[blockIdx.x] sdata[0]; }提示shfl_down_sync比 shared memory 快 3–5×因它走寄存器级互联无 memory bank 争用。但仅适用于 compute capability ≥ 7.0 的 GPU。3.3 Stencil 计算Ghost Cell 交换为何是 MPIGPU 混合编程的试金石Stencil模板计算如热传导、流体模拟需访问邻居点。在多 GPU 场景下“ghost cell”边界外一层虚拟单元数据必须跨 GPU 传输。书中 8.4 节和 9.5 节联合剖析纯 MPI 方案CPU 间交换 ghost cell导致 GPU 计算与通信串行而高效方案是overlap communication with computationGPU 计算 interior region内部区域同时 CPU 发起MPI_Isend/MPI_Irecv交换 ghost cellGPU 计算完成后用cudaStreamWaitEvent()等待通信完成再计算 boundary region。书中给出关键数据在 4 GPU 的 2D mesh 上重叠后通信隐藏率达 91%整体 time-to-solution 缩短 3.8×。这揭示了混合编程的核心——GPU 负责计算密度高的 interiorCPU/MPI 负责低延迟通信二者通过 streams/events 严格解耦。3.4 避坑算法模式选择中的 3 个致命误判这些是某图像处理团队在移植实时视频增强算法时的惨痛教训现象用thrust::sort对 10M 像素排序GPU 版本比 CPUstd::sort慢 2×。原因thrust::sort是通用算法对小规模数据 1M不如 CPU 的 introsort且像素排序常需 key-value 关联thrust::sort_by_key引入额外内存拷贝。解决改用 radix sort书中 5.5 节推荐或对像素值域有限0–255时用 counting sort prefix sum速度提升 12×。现象cudaMemcpy后立即cudaDeviceSynchronize()nvvp显示 GPU 利用率曲线呈锯齿状高-零-高-零。原因同步操作阻塞 host 线程导致 GPU 空闲等待且cudaMemcpy本身已是同步操作二次Synchronize()纯属冗余。解决删除cudaDeviceSynchronize()改用cudaEventRecord()记录 kernel 结束点用cudaEventSynchronize()按需等待。现象OpenMP offload 版本在 A100 上比 V100 慢 15%nsys显示gpu__inst_executed指令数相同。原因A100 默认启用FP16加速但 OpenMP 编译器未自动启用#pragma omp declare simd导致部分计算仍走 FP32 pipeline。解决添加#pragma omp declare simd(inbranch)并用-foffload-archsm_80显式指定架构。4. 编程模型选择不是“语法糖”而是“对硬件控制权的精确让渡”CUDA、OpenACC、OpenMP、HIP、SYCL……工具链的繁荣常让开发者陷入“选择困难”。但书中 Part 2 和 Part 3 的深层逻辑是每种模型代表一种不同的“控制粒度-开发成本”权衡而非功能替代。CUDA 提供最细粒度控制寄存器分配、warp shuffle、memory fence代价是代码强绑定 NVIDIAOpenACC 用 pragma 声明“我要并行这段”编译器全权负责 kernel 生成与数据移动适合快速原型OpenMP 4.5 则在两者间折中——它允许你用#pragma omp target teams distribute parallel for显式控制 team/workgroup 层级又保留#pragma omp simd指导向量化。关键洞察在于没有“最好”的模型只有“最适合当前约束”的模型。书中 11.2 节用 stream triadc[i] a[i] b[i] * scalar案例证明OpenACC 版本开发耗时 2 小时性能达 CUDA 的 92%而 OpenMP 版本耗时 4 小时性能达 98%因其允许你精细控制map(tofrom:c[0:n])的粒度。选择依据应是项目生命周期短期 demo 选 OpenACC、团队技能熟悉 OpenMP 选 OpenMP、硬件多样性需支持 AMD GPU 选 HIP/SYCL。4.1 OpenACC当“最小改动”成为最高优先级时的终极答案OpenACC 的核心价值不是性能而是可维护性。书中 11.2 节强调一个成熟的 CPU 代码库往往有数十万行 C/Fortran重写为 CUDA 的成本远超收益。OpenACC 的#pragma acc kernels可让编译器自动识别并行区域而#pragma acc data copyin(a[0:n]) copyout(c[0:n])则精确控制数据生命期。书中警告kernelspragma 的危险在于“过度信任编译器”——若循环内有函数调用如log()编译器可能因无法内联而放弃并行。解决方案是#pragma acc routine seq显式标记函数为串行或用#pragma acc kernels loop independent强制并行。// ✅ OpenACC 安全并行化显式声明独立性 #pragma acc kernels loop independent for (int i 0; i n; i) { c[i] a[i] b[i] * s; } // ❌ OpenACC 危险并行化编译器可能因函数调用放弃并行 #pragma acc kernels for (int i 0; i n; i) { c[i] expensive_func(a[i], b[i]); // 若 expensive_func 未标记为 acc routine编译器可能串行化 }4.2 OpenMP OffloadC 开发者的“可控加速”入口OpenMP 4.5 的target指令是 C 生态的利器。书中 11.3 节指出它与 CUDA 的关键差异在于host/device 代码统一编译——无需.cu文件、无需nvccg -fopenmp -foffloadnvptx-none即可生成 GPU 代码。更重要的是#pragma omp target teams distribute parallel for允许你像调优 CPU 一样调优 GPUnum_teams(128)控制 work-group 数量thread_limit(32)绑定 warp 大小schedule(static, 1024)指导数据分块。书中实测对 3D FFTOpenMP 版本比 OpenACC 快 18%因其能更精准地控制teams与threads的乘积匹配 SM 数量。4.3 CUDA/HIP当“最后一公里性能”决定成败时的唯一选择当 OpenACC/OpenMP 无法突破性能天花板CUDA/HIP 是必经之路。书中 12.2 节强调CUDA 的优势不在语法而在对硬件状态的完全可见性——cudaGetLastError()可捕获 warp-level 错误cuda-memcheck能定位 shared memory bank conflictNsight Compute可查看每个 warp 的 IPCinstructions per cycle。而 HIP 的价值在于用几乎相同的 CUDA 语义hipMalloc,hipLaunchKernel编写代码hipcc编译器可一键生成 AMD GPUCDNA或 NVIDIA GPUPTX代码。书中提醒HIP 不是“CUDA 兼容层”而是“语义对齐层”——hipMemcpy的行为与cudaMemcpy完全一致但hipDeviceSynchronize()在 AMD 上可能对应clFinish()需严格测试。4.4 避坑编程模型迁移中的 3 个隐蔽断层这些是某金融风控系统从 OpenACC 迁移至 CUDA 时的深度排查记录现象OpenACC 版本#pragma acc parallel loop正常CUDA 版本cudaMemcpy后 kernel 结果全为 0。原因OpenACC 自动处理struct成员对齐而 CUDA 中__device__ struct Particle { float x,y,z; int id; }的id成员因 4-byte 对齐导致sizeof(Particle)16但 OpenACC 传输时按 12 字节解释数据错位。解决CUDA 中显式用__align__(16)或#pragma pack(4)并用static_assert(sizeof(Particle)12, )验证。现象HIP 版本在 AMD MI250X 上运行正常在 NVIDIA A100 上hipGetLastError()返回hipErrorInvalidValue。原因HIP 代码中调用hipMalloc(d_ptr, size)时size为 0AMD 驱动容忍NVIDIA 驱动严格校验。解决所有hipMalloc前加if (size 0) return;或用hipMallocManaged替代。现象OpenMP offload 版本在#pragma omp target update from(a[0:n])后a[0]值异常。原因update from仅从 device 拷贝数据但若 kernel 中a被#pragma omp target data map(to: a[0:n])映射为只读则update from无效。解决确保map属性为tofrom或alloc或改用#pragma omp target update from(a[0:n])前先#pragma omp target data map(tofrom: a[0:n])。5. 性能剖析不是“看火焰图”而是“用硬件计数器反向推演执行流”性能调优的终点不是“更快”而是“理解为什么快/慢”。书中 Part 3 第 13 章GPU Profiling and Tools彻底颠覆常规认知nvprof的“GPU Time”只是幻觉真正可信的是硬件计数器hardware counter——如sms__sass_thread_inst_executed_op_fadd_pred_on.sum实际执行的浮点加法数、l1tex__t_sectors_pipe_lsu_mem_shared_op_ld.sumshared memory 读取扇区数。这些计数器直接反映 warp 的实际行为不受 host 调度、驱动开销干扰。例如若sms__inst_executed为 10M但sms__sass_thread_inst_executed_op_fadd_pred_on.sum仅 2M说明 80% 指令是分支或访存计算密度极低若l1tex__t_sectors_op_ld.sum远高于l1tex__t_sectors_pipe_lsu_mem_shared_op_ld.sum则 shared memory 使用不当大量请求落到 L1 cache。书中强调真正的 profiling 是“提出假设→用计数器验证→修正代码→再验证”的闭环。它不教你“点击哪里”而是给你一套可复现的验证方法论。5.1 Nsight Compute用ncu命令行直击 warp 级瓶颈GUI 工具易用但命令行ncu才是生产环境真相。书中 13.1 节给出标准流程ncu --set full ./my_app获取全量计数器ncu -k my_kernel -o profile --metrics sms__inst_executed,sms__sass_thread_inst_executed_op_fadd_pred_on,l1tex__t_sectors_op_ld聚焦关键指标ncu -f --set full --replay-mode kernel ./my_app report.txt生成可 grep 的文本报告。书中实测案例某矩阵乘 kernel 的sms__inst_executed为 1.2B但sms__sass_thread_inst_executed_op_fadd_pred_on.sum仅 0.3Bsms__sass_thread_inst_executed_op_fmul_pred_on.sum仅 0.3B而l1tex__t_sectors_op_ld.sum高达 2.1B——结论清晰计算密度不足FLOPs/byte 0.3瓶颈在访存非计算。解决方案不是换算法而是改用mma.sync指令Tensor Core将 FLOPs/byte 提升至 128。5.2 Roofline Model用一张图诊断所有性能问题Roofline 模型书中 9.3 节是 GPU 性能分析的“罗盘”。它以Arithmetic IntensityAIFLOPs/Byte为横轴PerformanceGFLOP/s为纵轴画出两条 roofMemory-Bound Roof斜线斜率 峰值内存带宽如 A100 的 2039 GB/sCompute-Bound Roof水平线高度 峰值算力如 A100 的 19.5 TFLOPS FP16。若 kernel 落在 Memory-Bound Roof 下方优化方向是提升 AI如用 Tensor Core、增大 block size 减少重复加载若落在 Compute-Bound Roof 下方则需提升计算效率如用mma.sync、减少分支。书中强调Roofline 不是理论玩具而是可测量的工程工具——用ncu获取sms__inst_executed和l1tex__t_sectors_op_ld.sum即可算出实际 AI。# 从 ncu 报告中提取关键值示例 $ grep sms__inst_executed report.txt sms__inst_executed 1200000000 $ grep l1tex__t_sectors_op_ld report.txt l1tex__t_sectors_op_ld 2100000000 # 计算 Arithmetic Intensity (AI) # 假设每个 sector 32 bytes, 每条指令平均 2 FLOPs (fma) AI (1200000000 * 2) / (2100000000 * 32) ≈ 0.0357 FLOPs/Byte # A100 的 Memory-Bound Roof 斜率 2039 GB/s 2039 * 10^9 Bytes/s # 对应 Performance AI * 2039e9 0.0357 * 2039e9 ≈ 72.8 GFLOP/s # 实际测得 performance 为 58 GFLOP/s → 确认 memory-bound且有优化空间5.3 从计数器到代码一个 real-world 的优化闭环书中 13.4 节以浅水方程模拟shallow water simulation为例展示完整闭环初始测量ncu -k sw_kernel --metrics sms__inst_executed,sms__sass_thread_inst_executed_op_fadd_pred_on,l1tex__t_sectors_op_ld→ AI 0.021Performance 32 GFLOP/s假设AI 过低因u,v,h三个数组分开存储导致每次计算需 3 次 global memory 访问验证ncu --metrics l1tex__t_sectors_op_ld本文还有配套的精品资源点击获取