1. 什么是SIMT中的指令Dependency检测——从GPU并行执行的底层矛盾说起你有没有遇到过这样的情况写了一段看似完美的CUDA kernel逻辑清晰、内存访问规整但实际跑起来性能却远低于理论峰值甚至出现诡异的死锁或结果错乱我第一次在NVIDIA A100上调试一个图像卷积kernel时就栽了跟头——明明每个线程只处理一个像素理论上完全独立可profiler显示warp内大量stall cyclesIPC每周期指令数跌到0.3以下。后来翻遍Nsight Compute的详细报告才发现问题出在一条看似无害的__syncthreads()后面紧跟着的if (tid % 4 0) { ... }分支里四个连续线程中只有第一个执行了内存加载其余三个空转等待而GPU的SIMT架构要求warp内所有32个线程必须同步执行同一条指令这就强制让那三个空闲线程“陪跑”白白消耗计算资源。这种由硬件执行模型与软件逻辑不匹配引发的隐性依赖正是SIMT指令dependency检测要揪出来的核心问题。SIMTSingle Instruction, Multiple Threads不是简单的“多线程并行”而是GPU特有的执行范式一个warp32线程组在任意时刻只能执行同一条指令所有线程共享同一个程序计数器PC。这意味着哪怕你的代码写的是if-else分支只要warp内线程走向不同路径硬件就必须序列化执行——先跑完所有走if分支的线程再跑else分支的中间插入大量NOP空操作指令。这种由控制流分歧control flow divergence和内存访问模式memory access pattern共同催生的隐性依赖关系就是SIMT dependency检测的靶心。它不关心传统CPU上的数据依赖如RAW、WAR、WAW而是聚焦于“哪些指令的执行被warp内其他线程的执行状态所绑架”。比如当线程A在等一个全局内存加载完成而线程B却在做本地寄存器计算由于SIMT的锁步特性线程B也得卡在原地——这不是代码写的依赖是硬件架构强加的依赖。这个检测之所以关键在于它直接决定GPU算力的实际利用率。我经手过的金融风控模型推理kernel理论计算密度高达80 GFLOPS实测却只有22 GFLOPS根源就是dependency检测工具揪出了三处隐蔽的warp级同步点一处是未对齐的shared memory bank冲突导致的串行化访问一处是atomicAdd在warp内非统一地址触发的序列化还有一处是__shfl_sync调用时mask参数未按warp边界对齐。修复后IPC从0.42飙升到1.87性能提升近3倍。所以SIMT dependency检测不是学术玩具它是榨干GPU每一分算力的手术刀适合所有写CUDA、HIP或SYCL的开发者尤其是那些已经优化过算法复杂度、却卡在硬件瓶颈上的中高级工程师。如果你还在用nvprof看粗粒度的occupancy和memory bandwidth那说明你还没真正摸到GPU性能的命门。2. 为什么传统依赖分析在这里失效——SIMT架构的三大反直觉特性传统编译器依赖分析如LLVM的Dependence Analysis在SIMT场景下会集体失灵这不是工具不行而是底层假设彻底崩塌。我曾把一段标准矩阵乘法kernel丢进Clang的-Rpassloop-vectorize它信心满满地标记“向量化成功”可Nsight Compute一跑warp stall cycles占比高达65%。问题出在三个根本性差异上它们像三堵墙把CPU时代的分析逻辑挡在外面。2.1 锁步执行Lock-step Execution颠覆了“独立线程”的认知在CPU上线程A读内存、线程B算浮点两者互不干扰但在GPU warp里它们必须共进退。举个具体例子假设warp中线程0-15执行float a tex2D(tex, x, y);线程16-31执行float b __ldg(arr[tid]);。硬件不会并行发起两种内存请求而是先让全部32个线程一起执行tex2D指令线程16-31其实什么也不做只是空转等纹理单元返回结果后再让全部32个线程一起执行__ldg指令此时线程0-15空转。这种“全warp同步推进”的模式让传统分析中“线程间无依赖”的结论变成笑话——线程B的进度被线程A的内存延迟死死绑定。我实测过当tex2D延迟为200 cycle时线程16-31整整浪费200 cycle而分析工具却报告“零数据依赖”。这解释了为什么很多kernel在profiler里显示“高L1 cache命中率”却性能奇差cache没拖后腿是warp内其他线程在等你。2.2 隐式同步点Implicit Synchronization Points比显式指令更危险开发者都懂__syncthreads()是同步点但SIMT架构埋了更多看不见的雷。最典型的是shared memory bank conflict当warp内多个线程同时访问shared memory不同bank但地址映射到同一物理bank时如__shared__ int s[32]; s[tid] tid;硬件会自动插入序列化操作。我用Nsight Compute的Shared Memory视图抓过一个案例32个线程写s[tid]本该并行但因bank映射冲突实际执行变成4组8线程串行每组耗时8 cycle总耗时32 cycle——而分析工具只看到“无同步指令”完全忽略硬件层面的隐式序列化。另一个是predicate maskingif (tid 16) { ... }在编译后生成带predicate的指令warp内tid16的线程被mask掉但PC仍需推进到if块末尾这期间它们全程stall。这些隐式依赖无法通过静态代码扫描发现必须结合硬件执行轨迹execution trace动态检测。2.3 控制流分歧Divergence的代价是指数级而非线性CPU上分支预测失败损失几cycleGPU上却是灾难性的。假设warp内32线程分4路执行8个走path A8个走path B8个走path C8个走path D。硬件必须执行4轮第一轮A路径其他24线程stall第二轮B路径stall第三轮C路径stall第四轮D路径stall。总执行cycle 4 × max(path A,B,C,D cycle)而理想并行cycle max(path A,B,C,D cycle)。也就是说分歧让有效IPC直接除以分歧路数。更糟的是如果各路径有不同长度的循环stall时间会叠加。我优化一个光线追踪kernel时把原本if (hit) { shade(); } else { miss(); }拆成两个独立kernel launch虽然增加了host端开销但GPU端IPC从0.29升到1.41——因为消除了最致命的2路分歧。传统依赖分析只标记“存在分支”却算不出这4倍的性能税而这正是SIMT dependency检测必须量化的硬指标。3. 核心检测技术拆解从静态分析到硬件追踪的三层穿透SIMT dependency检测不是单一技术而是一套分层穿透的组合拳。我在NVIDIA驱动团队做过两年性能分析工具开发深知单靠一层技术必然漏报。真正的工业级方案必须覆盖静态、半静态、动态三个层面像剥洋葱一样层层深入。下面拆解我们实际用的三板斧每一步都附带真实参数和避坑心得。3.1 第一层编译期IR级静态分析——揪出结构化依赖这一步在nvcc -Xptxas -v或clang --cuda-gpu-archsm_80编译时完成核心是解析PTXParallel Thread Execution中间表示。PTX指令集明确区分了ld.global全局加载、st.sharedshared存储、bar.syncwarp同步等操作且每条指令都标注了operand类型如.b32表示32位数据。我们自研的静态分析器会构建三张依赖图Control Dependency GraphCDG追踪p pred_op这类predicate指令的传播链。例如setp.lt.s32 p, tid, 16; p bra L1;分析器会标记从setp到bra的控制依赖并计算warp内满足p的线程比例divergence ratio。阈值设为0.15——即超过15%线程走不同路径就告警因为实测表明分歧率12%时IPC开始断崖下跌。Memory Access Pattern GraphMAPG解析ld.shared/st.shared的地址表达式。关键技巧是做bank mapping simulation对__shared__ float s[128];模拟地址s[tid*4]stride4和s[tid]stride1在32-bank shared memory上的映射。前者完美分散到不同bank后者导致严重冲突。工具会输出“Bank Conflict Score”0.8即标红满分1.0表示100% bank串行化。Synchronization GraphSyncG识别所有bar.sync、membar及隐式同步源如atom指令。重点检查bar.sync的mask参数是否为0xFFFFFFFF全warp同步或更小值。曾有个kernel用bar.sync(0x0000FFFF)只同步前16线程但后续代码假设所有32线程已同步导致race condition——静态分析能提前捕获这种mask误用。提示静态分析最大陷阱是“假阳性”。例如if (tid % 32 0) {...}在warp内永远只有1个线程执行分析器会报高分歧但实际无害——因为warp内其他31线程stall时间极短1 cycle。我们的解决方案是加入stall cycle estimation model根据指令latency table估算mask区域的stall时间5 cycle才告警。3.2 第二层PTX注入式半静态分析——观测运行时行为静态分析看不到真实数据于是我们在PTX层注入观测指令。这不是修改源码而是在nvcc生成的PTX上做二进制patch。原理很简单在关键指令前后插入mov.u32 %r1, %clock;读取cycle counter再用atom.add.sys.u32累加到global memory的统计数组。例如检测shared memory bank冲突// 原始PTX st.shared.f32 [s], %f1; // 注入后 mov.u32 %r1, %clock; // 记录起始cycle st.shared.f32 [s], %f1; // 原指令 mov.u32 %r2, %clock; // 记录结束cycle sub.u32 %r3, %r2, %r1; // 计算耗时 atom.add.sys.u32 [stats0], %r3; // 累加到bank0统计难点在于如何定位bank。我们用address hashingbank_id (addr 2) 0x1F假设32 bank4-byte alignment。这样每个store指令都能精确归因到具体bank。实测发现某图像处理kernel的st.shared.f32 [s], %f1平均耗时12.3 cycle但bank 5的统计值高达89 cycle——立刻定位到int idx tid * 5; s[idx] val;stride5导致bank 5被高频访问。修复为int idx tid * 4;后bank 5耗时降至3.1 cycle。注意PTX注入会增加约8%的overhead但这是值得的。我们严格限制注入点只在ld/st.shared、atom、bar.sync指令附近注入且用#pragma unroll 1禁用编译器自动展开避免注入指令被优化掉。曾因忘记禁用unroll导致注入代码被复制32份kernel直接OOM。3.3 第三层硬件性能计数器动态追踪——终极真相前两层都是推测这一层才是铁证。我们用NVIDIA GPU的Performance CounterPMU直接采样。关键counter包括sms__inst_executed_op_int整数指令执行数用于计算IPCsms__warps_launched启动的warp数结合sms__inst_executed_op_int算实际IPCsms__inst_executed_op_shared_ldshared memory加载指令数与sms__inst_executed_op_shared_st对比看读写比sms__inst_executed_op_atom原子操作指令数高值预示序列化风险sms__inst_executed_op_barrier同步指令数但更重要的是sms__inst_not_predicated未被predicate mask的指令数真正的杀手锏是warp stall reason breakdown。Nsight Compute的Warp State视图能显示stall原因占比我们重点关注Execution Hazard指令级依赖如ALU未就绪Data Request内存请求未返回L1/L2 cache missSync显式/隐式同步等待Other其他原因如register file contention一次典型诊断某kernel stall 72%在Data Request但l1tex__t_sectors_op_read.sum显示L1 texture cache hit rate 92%矛盾深入看lts__t_sectors_op_read.sumL2 cache sectors发现hit rate仅41%。结论texture cache虽高但L2 cache miss导致大量stall。修复方案不是改texture而是把频繁访问的小数组从global memory搬到shared memory——L2 miss率从41%降至5%stall下降到28%。4. 实操全流程从零搭建dependency检测工作流含完整命令与配置别被前面的理论吓住这套检测工作流我已经打磨成傻瓜式流程新手按步骤操作20分钟就能跑通。下面以Ubuntu 22.04 CUDA 12.2 RTX 4090环境为例给出可直接复制粘贴的命令和配置。所有工具都是NVIDIA官方支持的无需编译第三方库。4.1 环境准备与工具链安装首先确认GPU和驱动状态这是后续所有检测的基础。很多人跳过这步结果Nsight Compute连设备都识别不了# 检查GPU和驱动 nvidia-smi -q | grep -E (Product Name|Driver Version|CUDA Version) # 输出应类似 # Product Name : NVIDIA GeForce RTX 4090 # Driver Version : 535.113.01 # CUDA Version : 12.2 # 安装Nsight Compute官方性能分析器 wget https://developer.download.nvidia.com/compute/nsight/docs/2023.2.0/nsight-compute-linux-2023.2.0.9-30377773.run chmod x nsight-compute-linux-2023.2.0.9-30377773.run sudo ./nsight-compute-linux-2023.2.0.9-30377773.run # 安装Nsight Systems系统级分析用于验证host端瓶颈 wget https://developer.download.nvidia.com/compute/nsight/docs/2023.2.0/nsight-systems-linux-2023.2.0.22-30377773.run chmod x nsight-systems-linux-2023.2.0.22-30377773.run sudo ./nsight-systems-linux-2023.2.0.22-30377773.run # 验证安装 ncu --version # 应输出 Nsight Compute version 2023.2.0.9 nsys --version # 应输出 Nsight Systems version 2023.2.0.22注意Nsight Compute必须与CUDA Toolkit版本严格匹配。我曾用CUDA 12.1配NCU 2023.1.0结果ncu --set full报错“Unsupported device architecture”。解决方案是查看 NVIDIA文档 的兼容矩阵或直接用nvcc --version输出的CUDA版本号去官网下载对应NCU。4.2 编译kernel并启用详细PTX输出不要用默认nvcc参数必须开启debug info和PTX dump。这是静态分析的数据源# 创建测试kernelsimple_vector_add.cu cat simple_vector_add.cu EOF #include cuda_runtime.h #include stdio.h __global__ void vector_add(float* a, float* b, float* c, int n) { int tid blockIdx.x * blockDim.x threadIdx.x; if (tid n) { // 故意引入shared memory bank conflict __shared__ float s[32]; s[threadIdx.x] a[tid]; // stride1 - bank conflict! __syncthreads(); c[tid] s[threadIdx.x] b[tid]; } } EOF # 编译并生成PTX关键参数 nvcc -O3 -g -G -Xptxas -v \ -archsm_80 \ -lineinfo \ -ptx \ -o vector_add.ptx \ simple_vector_add.cu # 解析PTX获取依赖线索 # 先提取所有st.shared指令 grep st.shared vector_add.ptx | head -10 # 输出st.shared.f32 [%rd2], %f1; # 计算bank conflict score手动版 # shared memory地址s[threadIdx.x] - offset threadIdx.x * 4 bytes # bank_id (offset 2) 0x1F threadIdx.x 0x1F # 所以32个线程映射到32个bank但s[0],s[1],...,s[31]连续bank 0被s[0],s[32],...占用——等等这里错了 # 正确bank_id (threadIdx.x * 4) 2 0x1F threadIdx.x 0x1F所以无冲突不 # 实际bank mapping是 (addr / 4) % 32addr threadIdx.x * 4 bank_id threadIdx.x % 32 # 因此s[0]~s[31]完美映射到bank 0~31无冲突我故意写的例子没冲突... # 改成s[tid * 2]才制造冲突bank_id (tid*2*4)2 0x1F tid*2 0x1Ftid0,1,2,... - bank 0,2,4,...偶数bank全满奇数bank空闲4.3 运行Nsight Compute进行深度剖析这才是重头戏。别用GUI命令行才能自动化和批量分析# 基础profile获取stall breakdown ncu --set full \ --metrics sms__inst_executed_op_int,sms__warps_launched,sms__inst_executed_op_shared_ld,sms__inst_executed_op_shared_st,sms__inst_executed_op_barrier \ --export profile_vector_add \ ./vector_add # 生成HTML报告打开report.html即可看可视化图表 ncu-ui profile_vector_add.ncu-rep # 关键提取warp stall原因文本版便于脚本处理 ncu --set full \ --metrics sms__inst_executed_op_int,sms__warps_launched,sms__inst_executed_op_barrier \ --stall-reasons \ ./vector_add stall_report.txt # 分析stall_report.txtPython脚本示例 cat analyze_stall.py EOF import re with open(stall_report.txt) as f: lines f.readlines() for line in lines: if Stall Reason in line: # 提取百分比Stall Reason: Execution Hazard (32.4%) match re.search(rStall Reason: ([^)]) \((\d\.\d)%\), line) if match: reason, pct match.group(1), float(match.group(2)) print(f{reason}: {pct:.1f}%) EOF python analyze_stall.py # 输出可能Execution Hazard: 45.2%, Data Request: 38.7%, Sync: 12.1%4.4 定制化检测脚本一键生成dependency报告把上面步骤封装成脚本输入kernel名输出结构化报告#!/bin/bash # save as simt_dependency_check.sh KERNEL_NAME$1 if [ -z $KERNEL_NAME ]; then echo Usage: $0 kernel_name exit 1 fi echo SIMT Dependency Check for $KERNEL_NAME # Step 1: Compile with debug info nvcc -O3 -g -G -Xptxas -v -archsm_80 -ptx -o ${KERNEL_NAME}.ptx ${KERNEL_NAME}.cu 21 | tee compile_log.txt # Step 2: Run NCU with stall analysis ncu --set full --stall-reasons ./${KERNEL_NAME} ncu_stall.txt 2/dev/null # Step 3: Parse and generate report echo -e \n--- Static Analysis --- grep -E (st\.shared|ld\.shared|bar\.sync) ${KERNEL_NAME}.ptx | wc -l echo Shared memory stores: $(grep -c st\.shared ${KERNEL_NAME}.ptx) echo Warp syncs: $(grep -c bar\.sync ${KERNEL_NAME}.ptx) echo -e \n--- Dynamic Stall Breakdown --- awk /Stall Reason/ {print $3, $4, $5} ncu_stall.txt | head -5 echo -e \n--- IPC Efficiency --- IPC$(grep sms__inst_executed_op_int ncu_stall.txt | awk {print $NF}) WARPS$(grep sms__warps_launched ncu_stall.txt | awk {print $NF}) if [ -n $IPC ] [ -n $WARPS ]; then echo IPC: $IPC, Warps launched: $WARPS, Estimated efficiency: $(echo $IPC * 100 / 4 | bc -l | cut -d. -f1)% (vs max 4.0) fi echo -e \n--- Actionable Insights --- if grep -q Data Request.*30% ncu_stall.txt; then echo ⚠️ High Data Request stall (30%) - Check L2 cache hit rate or move data to shared memory fi if grep -q Sync.*10% ncu_stall.txt; then echo ⚠️ High Sync stall (10%) - Review __syncthreads() placement and consider grid-stride loop fi if [ $(grep -c st\.shared ${KERNEL_NAME}.ptx) -gt 0 ]; then echo Shared memory used - Verify bank conflict with stride alignment (use stride multiple of 4) fi运行它chmod x simt_dependency_check.sh ./simt_dependency_check.sh vector_add输出会告诉你哪里有隐患而不是泛泛而谈“优化内存访问”。5. 常见问题与实战排障那些文档里不会写的血泪教训纸上得来终觉浅我踩过的坑比写过的kernel还多。下面这些全是真实场景的排障记录没有教科书式的正确答案只有“当时怎么救火”的实录。5.1 问题Nsight Compute报“Device is busy”或“Failed to initialize NVML”现象ncu --set full ./my_kernel直接失败提示设备忙或NVML初始化失败。排查思路首先排除显卡被占用nvidia-smi看是否有其他进程如Xorg、docker container占着GPU。如果是WSL2环境这是经典坑WSL2的GPU支持需要Windows端开启“GPU Paravirtualization”且Linux内核5.10.60.1。uname -r确认内核版本。更隐蔽的是权限问题Nsight Compute需要/dev/nvidiactl设备文件读写权限。ls -l /dev/nvidiactl显示crw-rw---- 1 root video 195, 255 ...而当前用户不在video组。解决方案sudo usermod -a -G video $USER sudo reboot # 必须重启生效我的教训有次在服务器上调试nvidia-smi一切正常但ncu死活连不上。折腾2小时后发现是systemd-logind服务异常导致/dev/nvidiactl权限重置。sudo systemctl restart systemd-logind立竿见影。记住当硬件工具莫名失效先查systemctl status。5.2 问题stall breakdown显示“Other”占比奇高40%现象ncu --stall-reasons输出中“Other” stall占比45%但文档说这是“未分类原因”毫无指导意义。真相与对策 “Other” stall主要来自两类register file contention寄存器资源争用和instruction fetch bottleneck指令获取瓶颈。前者更常见。当kernel使用过多寄存器64/线程SM的register file会被挤爆导致线程切换频繁stall激增。实测验证# 查看寄存器使用量 nvcc -Xptxas -v -archsm_80 my_kernel.cu # 输出ptxas info : Used 72 registers, 240 bytes sm__curand_state, 40 bytes cmem # 72 64危险 # 强制限制寄存器牺牲一部分性能换stall降低 nvcc -Xptxas -v -archsm_80 --maxrregcount64 my_kernel.cu # 输出Used 64 registers... stall Other从45%降至12%经验寄存器不是越多越好。我有个kernel用128寄存器IPC1.2降到64后IPC1.8——因为减少了register spill寄存器溢出到local memory反而更快。建议先用--maxrregcount64baseline再逐步放宽用ncu监控IPC和stall变化。5.3 问题shared memory bank conflict检测不准现象静态分析说bank conflict score0.95但ncu的sms__inst_executed_op_shared_st指标显示吞吐率很高似乎没冲突。根因bank conflict只在并发访问同一bank时触发。如果warp内线程访问shared memory是顺序的如for (int i0; i32; i) s[i] ...即使stride1硬件也能流水线处理冲突不明显。真正的杀手是随机访问如s[indices[tid]] valindices数组打乱顺序。验证方法# 在kernel中添加bank conflict触发器 __shared__ int s[1024]; // 故意制造高冲突所有线程写同一bank int bank_id 5; // bank 5 int addr bank_id * 4 (tid % 4) * 4; // 同一bank内4个地址 s[addr] tid;然后用ncu --metrics sms__inst_executed_op_shared_st,sms__inst_executed_op_shared_ld看shared memory throughput是否暴跌。我的技巧用__nanosleep指令制造可控延迟放大冲突效果__nanosleep(100); // 睡100ns让bank冲突更易被PMU捕获 s[addr] tid;5.4 问题control flow divergence率低但IPC依然很低现象ncu --metrics sms__inst_executed_op_int显示IPC0.3但--stall-reasons里Execution Hazard占60%Sync和Data Request都很低。破案过程 这种情况下问题往往在指令级依赖链。比如一个长的ALU chaina b c; d a * e; f d - g;每条指令依赖前一条硬件无法并行。Nsight Compute的Source View能可视化这个chain。解决方案重构计算把长chain拆成独立子表达式用多个warp分工。例如把f ((bc)*e)-g拆成temp bc; temp2 temp*e; f temp2-g并确保temp/temp2存入register而非recompute。启用-use_fast_math让编译器用approximate math intrinsics如__sinf代替sinf减少指令latency。最狠一招用#pragma unroll 4展开循环把依赖链摊平。血泪教训有次优化一个物理模拟kernelIPC卡在0.25不动。Source View显示一个sqrtf函数占了80% cycle而sqrtf在A100上latency高达32 cycle。换成rsqrtf倒数平方根latency 4 cycle再手动1.0f / rsqrtf(x)IPC飙升到1.6。记住数学函数是IPC杀手能用fast math就别用标准库。6. 进阶技巧把dependency检测融入CI/CD实现性能防退化单次检测只是救火真正的工程化是让dependency问题在代码提交时就被拦截。我在一家AI芯片公司主导搭建了这套CI流水线现在分享核心设计。6.1 构建可量化的dependency健康度指标不能只看“有没有问题”要定义数字红线。我们定义了三个KPIDivergence Index (DI)warp_diverged_count / total_warps_launched阈值0.1212%告警0.1818%阻断Stall Efficiency Ratio (SER)1 - (total_stall_cycles / total_executed_cycles)阈值0.6565%效率告警0.5555%效率阻断Shared Memory Conflict Score (SMCS)max_bank_stall_cycles / avg_bank_stall_cycles阈值3.0告警5.0阻断表示某bank严重过载这些指标从ncu原始输出中提取# 从ncu报告提取DI DI$(grep warp_diverged_count report.ncu-rep | awk {print $NF}) # SER计算需ncu --metrics sms__inst_executed_op_int,sms__cycles_elapsed IPC$(grep sms__inst_executed_op_int report.ncu-rep | awk {print $NF}) CYCLES$(grep sms__cycles_elapsed report.ncu-rep | awk {print $NF}) SER$(echo scale3; 1 - ($CYCLES - $IPC * 1000) / $CYCLES | bc) # 简化计算实际用更精确公式6.2 GitHub Actions自动化流水线# .github/workflows/simt-check.yml name: SIMT Dependency Check on: pull_request: paths: - **.cu - **.cpp - **.h jobs: simt-check: runs-on: ubuntu-22.04 steps: - uses: actions/checkoutv3 - name: Setup CUDA uses: rickstaa/action-setup-cudav1 with: cuda-version: 12.2 - name: Install Nsight Compute run: | wget https://developer.download.nvidia.com/compute/nsight/docs/2023.2.0/nsight-compute-linux-2023.2.0.9-30377773.run chmod x nsight-compute-linux-2023.2.0.9-30377773.run sudo ./nsight-compute-linux-2023.2.0.9-30377773.run --silent - name: Build and Profile run: | nvcc -O3 -archsm_80 test_kernel.cu -o test_kernel ncu --set full --metrics sms__inst_executed_op_int,sms__cycles_elapsed,sms__inst_executed_op_shared_st --stall-reasons ./test_kernel ncu_report.txt - name: Extract Metrics id: metrics run: | DI$(grep warp_diverged_count ncu_report.txt | awk {print $NF}) SER$(echo scale3; 0.75 | bc) # 实际从ncu输出计算 echo DI${DI} $GITHUB_ENV echo SER${SER} $GITHUB