第四章：Source 细读

第二、三章把Details页读完了。我们已经知道reduce_ill的全局读不合并、共享内存读存在冲突，归并过程中还有大量同步和线程陪跑。接下来切到Source页，把这些线索接回代码，看看问题分别落在哪一行源码、哪一条指令上。

页面初见

先看整页的布局。报告标签栏下面是Source页自己的控件区，最左边View当前选中Source and SASS，于是中间的主体分成左右两块：左边是reduce.cu里的CUDA C++源码，右边是reduce_ill编译后的SASS。SASS是面向具体GPU架构的机器指令的汇编表示，GPU实际执行的代码在这一层查看。源码说明程序想做什么，SASS展示这些操作最后编成了哪些指令。两边属于同一次kernel运行，选中一边的代码行，另一边会高亮与它关联的位置。

左右两块各自都是一张“代码加指标”的表。每块左侧放代码，右侧按列摆放执行次数、采样、寄存器等信息，一行代码与同一行的指标横向对齐。源码的一行可能关联多条SASS，所以左边可以看问题集中在哪行，右边可以继续分辨是哪条指令产生了这些工作。每块底部都有横向滚动条，用来查看屏幕之外的代码或指标列；左侧当前还有蓝色外框，表示它是获得焦点、接受相关操作的一侧，点击右边就能切换过去。

每块最右侧、竖向滚动条旁边还有Heatmap，也就是用颜色表示指标高低的细窄竖带。它把某个指标在整份代码中的分布压缩到一条竖带里，深色对应较低值，最亮的白色对应最高值，让我们滚动之前就能留意到视野之外的热点；在Heatmap上按住鼠标右键，还能查看色标。究竟用哪个指标，由各块上方的Navigate By选择，截图里是Instructions Executed，也就是执行次数。代码右侧的指标格里也有横条，用长度表示数值大小。同一指标列中，数值越大，横条越长，方便我们比较哪几行的值更高。

主体下面是Inline Functions、Source Markers和Statistics三个附加面板，分别补充函数内联调用的信息、代码位置上的提示，以及选中行的统计结果。截图当前打开Inline Functions，底下的一行文字提示我们选择内联函数中的源码行。

至此，整页就能按“顶部选择怎么显示和怎么找，中间读代码与指标，底部查看当前选择的补充信息”来理解。

顶部控件

顶部控件分成两行。第一行从左到右负责视图选择、搜索导航、折叠和数值显示；第二行属于下方对应的代码视图，用来选择文件或函数、指定导航指标、解析源码位置。把两行连着看，才能理解上面的按钮在操作哪份代码、依据哪列数据。

View下拉菜单有六项：

切换View后，性能数据仍然来自已经采集的那次运行，只是按关联信息放到不同代码层次上。编译器会展开循环、重排和合并操作，因此三层代码可能一行对应多行，也可能有多行共用一部分关联指令。序章编译时使用的-lineinfo提供机器指令与源码行之间的关联，采集时的--import-source yes把当时能够找到的源码保存进报告；可用视图取决于程序保留了哪些信息，本报告已经能够显示三种代码。

View右边带P的纸张按钮打开Source View Profiles，用来保存视图配置。挑好要显示的列、调整列宽与顺序、固定常看的列，再选好导航指标，就可以把这套布局存下来，以后打开其他报告时套用。

再往右，截图中这组控件带蓝框，说明它只操作当前获得焦点的一侧。先看放大镜按钮，即打开搜索，也可以按Ctrl+F，搜索针对当前选中的列，没有选中列时默认搜索Source列，查代码可以输入变量名或指令名，查数值可以使用>= 5这样的表达式。

搜索按钮后面有六个箭头，它们读取哪一列，由第二行的Navigate By决定。当前选中Instructions Executed，箭头就围绕执行次数寻找位置；换成其他指标，同一组按钮就帮助我们查找那个指标。六个按钮按截图从左到右的作用如下：

接在导航按钮右边的是折叠与数值显示控件。这里有一个链条，用来切换这一组控件同时作用于两侧，还是只作用于当前获得焦点的一侧；后面的方框减号／加号折叠或展开文件、函数内容。折叠后保留标题和汇总指标，适合在多个文件或函数之间先看整体分布。

最后三个按钮改变数值的显示方式。%在适用的指标上切换相对百分比与绝对值；+−在绝对值显示下切换便于阅读的缩写与完整精度；饼图形状的按钮切换百分比和横条采用的范围，是整个kernel结果，还是当前文件、函数。折叠状态使用整个结果范围。我们先按默认的整个结果范围读，核对具体次数时，再切到绝对值和完整精度。默认情况下，适用的计数指标显示为“当前位置的值 ÷ 本次kernel的该指标总量 × 100 %”，单元格里横条的长度则参照这个指标的最大值。某条指令即使只占总量12.50 %，只要它已经是该列最大值，也能画出满格横条。数字回答它占总量多少，横条帮助比较各个位置。Live Registers这样的数量和Avg. Predicated-On Threads Executed这样的平均值，仍然按各自含义读。

再看第二行。Source下拉框选择当前文件或函数，图中左侧是reduce.cu，右侧是reduce_ill，切换条目后会跳到对应内容的开头。然后是Navigate By，决定导航箭头使用的指标和Heatmap展示的指标：找最常执行的位置选择Instructions Executed，追共享内存的多余工作可以选择L1 Wavefronts Shared Excessive，检查采样分布则选择Warp Stall Sampling (All Samples)。Navigate By前面的链条默认联动两侧的指标选择，一边换指标，另一边有同名指标时也跟着换；断开后可以分别选择。

Navigate By下拉菜单里还有More…，点开就是Column Chooser列选择窗口，右键表头也能打开。窗口按当前视图和报告已有数据列出可用列，左侧复选框选中的是当前显示的列，三角形用于展开分组，顶部Check All和Uncheck All分别全选和取消全选，点击OK应用。我们可以围绕当前问题挑出几列，再配合导航查找相关代码。Source和PTX视图的选项分组相同，SASS视图还提供机器指令地址和寄存器相关选项，图中底部列出了这些额外选项。

Warp Stalls Aggregated汇总warp状态的采样结果，可以看全部样本，也可以只看调度器没有发射指令时的样本。Instructions Executed这一组统计执行工作量，既有warp级的指令执行次数，也有线程级的次数和平均参与线程数，指令可以带执行条件，因此线程级统计还会区分全部活跃线程与其中满足条件、实际参与操作的线程。Memory Access描述访存操作本身，访问哪种内存空间、读还是写、每个线程一次访问多宽。L1 Cache这一组提供请求数量、共享内存冲突和wavefront等信息，帮助我们检查L1TEX处理访存请求时做了多少工作。L2 Cache这一组则提供全局访存的理论sector需求、理想需求与多余量，以及显式缓存策略等选项。Warp Stall和Warp Stalls (Not Issued)两组，把不同warp状态拆成独立的列，例如可以只看同步等待的采样，前一组使用全部采样，后一组使用调度器未发射指令时的采样。

SASS视图还提供Register Dependencies和Register Access两组。Register Dependencies用连线展示指令之间的数据依赖：一条指令需要读取某个寄存器里的值，工具会寻找这个值可能由哪些指令写入，再把相应位置连起来。Register Access则按单个寄存器提供访问选项，用来查看某个寄存器的访问位置。

下面逐项说明截图中已经勾选的项，这些也是默认的勾选项：

要把共享访存的实际工作和额外工作并排核对，可以同时勾选L1 Wavefronts Shared、Ideal和Excessive；追全局加载则选择对应的L2 Theoretical Sectors Global几列。选好以后，右键表头还能通过Freeze和Unfreeze固定或取消固定列，Source列默认固定在左边。配合Navigate By和Heatmap，就能先找热点，再把相关代码与数字留在视野里。

顶部控件最后还剩下一个Redo Resolve按钮用于重新解析源码文件位置。报告带到另一台机器后，如果源码没有导入报告，原路径又不存在，页面可能提示File Not Found，此时可通过Resolve指定本机上的源码，Redo Resolve允许重新做这一步。选择的文件应与当初编译时相符，文件属性不匹配会提示File Mismatch。本报告已经能显示源码，可以直接继续读。

Source：从源码找到问题所在

把View切成Source，左右两块合成一个完整的源码视图，留给指标的宽度也增加了。

左侧是reduce.cu的行号和代码，右侧汇总与该行关联的机器指令信息。我们以第90行从input读数、再写入smem的语句为例，沿右侧各列从左到右读。行号旁的黄色三角形是附在这行代码上的分析提示。

先看Live Registers，为23，说明走到这处代码时需要保留23个寄存器的值。图中第92行同步处的值是5，对比可见，输入装载位置需要同时保留更多中间值，跟踪寄存器需求时可以先检查这段代码。第三章已经算出，本kernel每线程使用36个寄存器，按分配粒度取整后仍允许100 %理论occupancy，说明当前寄存器容量足以支持这份配置满驻，因此应先处理已经出现的访存和同步问题。

紧接着的Warp Stall Sampling (All Samples)为23.50 %。采样器会定期选取一个活跃warp，记录它当时的指令位置和状态，一条这样的记录就是一个样本。ncu再根据机器指令与源码的对应关系，统计每行源码对应的样本数量。本次All Samples采样共记录17,547个样本，其中4,123个样本的指令位置对应源码第90行，4,123 ÷ 17,547 × 100 % ≈ 23.50 %。把鼠标停在单元格的不同色块上，可以查看这4,123个样本分别记录了哪些warp状态，判断warp当时在做什么、等什么。其中2,714个样本记录为MIO Throttle，即等待MIO指令队列入口腾出空位，超过第90行样本的一半。这把第三章的MIO队列等待定位到了输入搬运阶段；第90行同时包含全局读取和共享存储，分析时还要检查共享存储给MIO通路带来的工作。

Instructions Executed为15.67 %。本次kernel运行中，所有warp执行第90行关联机器指令的次数合计为3,211,264，执行整个kernel全部机器指令的次数合计为20,488,192，3,211,264 ÷ 20,488,192 × 100 % ≈ 15.67 %。这里一个warp执行一条机器指令记一次，第90行的合计次数包括读写数据、计算地址等关联指令的执行。因此，优化这段输入搬运时，可以同时检查地址计算和访存指令能否精简，再把修改后的执行次数与3,211,264次对比。再往右，Avg. Predicated-On Threads Executed为32，表示平均每次warp级指令执行时，都有32个活跃线程满足条件并参与操作。这表明装载阶段的线程参与是充分的，继续判断访问效率时，重点应转向这32个线程访问的地址是否连续。

接下来的三列可以合起来读。Address Space显示Global(8)、Shared(8)，Access Operation显示Load(8)、Store(8)，Access Size显示32(16)。第90行关联八条全局加载和八条共享存储，每条指令让每个参与线程访问32 bit，也就是一个float的4字节。循环中八轮“从input读取，再写入smem”的操作，在编译后展开成了这16条访存指令，括号统计的就是编译后的访存指令条数。前面的3,211,264次则是第90行全部关联机器指令在运行中的执行次数，里面还包括地址计算等操作；切到SASS选中某一条指令，才能单独查看它的执行次数。这些信息把第90行定位为全局内存到共享内存的搬运阶段；检查源码时，全局读取看globalIdx的计算，共享写入看smem_idx的布局。

第90行的L1 Wavefronts Shared Excessive没有显示数值，接着看L2 Theoretical Sectors Global Excessive，为100.00 %。整个kernel的全局访存一共多出14,680,064个理论sector，第90行关联访存指令的多余量也是14,680,064个。这一行的多余量除以整个kernel的多余总量，结果就是100 %，说明本次全部多余全局理论sector都来自第90行对应的访存操作。源码中globalIdx = gridDim.x * smem_idx + blockIdx.x，同一warp相邻线程的smem_idx相差1，globalIdx就相差8,192，读取地址相隔8,192 × 4 = 32,768字节。32个线程读取的数据分散到32个sector，非合并访问就来自这种地址排布。要减少这些多余sector，应从input[globalIdx]的读取方式着手，让同一warp的线程读取相邻数据。复测时应看多余理论sector的绝对数是否下降：只要其余位置都没有多余量，这一行即使改善，分布比例仍可能保持100 %。

再把同样的读法用到第99行。Address Space、Access Operation和Access Size分别是Shared(2)、Load(2)、128(2)，对应两条128位共享内存加载。每条加载让每个线程取四个float，两条正好取齐求和需要的八个数。L1 Wavefronts Shared Excessive为100.00 %，整个kernel多出的共享wavefront合计为524,288个，第99行对应的两条加载就产生了这524,288个，所以该行占全部多余共享wavefront的100 %。一行源码里逐个读取、累加的八个数，编译后已经成组加载，PTX视图可以把这种展开形式直接摆出来。编译器已经生成两条宽加载，因此这里应重点检查两组共享内存地址的bank映射，并用524,288这个多余wavefront总数验证布局调整的效果。

最后看看归并和写回的位置。第110行条件判断的Instructions Executed为48.62 %，对应9,961,472次warp级指令执行；第111行归并操作为12.79 %，对应2,621,440次。这两个百分比的分母都是整个kernel的20,488,192次warp级指令执行，条件判断这一行的关联指令执行次数，是归并操作所在行的3.8倍。若目标是减少归并阶段的指令执行次数，应先检查反复执行的取模条件和相关计算。再看第119行写回，Avg. Predicated-On Threads Executed为1，说明平均每次执行写回相关指令时只有一个线程参与，与第90行的32形成对照。这与每个block只由thread 0输出一个部分和的代码吻合，写回阶段的线程参与少是工作划分的结果；判断写回效率时，还要看实际写回流量和sector利用情况。若要单独分析归并加法，就沿第111行的关联找到FADD；第三章算出的12.45 %，表示这条FADD的全部活跃线程执行次数中，满足条件并真正执行加法的线程次数所占的比例。归并越往后，满足加法条件的线程越少，block同步却仍在每轮执行，因此要把线程参与方式和同步次数一起检查。

这张Source图已经给出三个可以动手检查的位置：第90行的全局读取地址排布、第99行的共享内存读取布局，以及归并阶段反复执行的判断和同步。后续沿PTX和SASS查看这些位置，就是要弄清机器实际执行了哪些操作，为修改访问方式和归并组织提供依据。

PTX：看源码中的操作怎样展开

把View切成PTX，表格左侧换成中间指令，右侧数值来自与这条PTX指令关联的SASS机器指令。这样既能看编译过程怎样组织操作，也能把实际执行产生的数据对到PTX的位置。

截图从装载数据的末段走到共享内存求和的开始。第150、165行是全局内存加载，第154、169行把读回的数写进共享内存，第171行同步，之后第175和180行各加载四个float，再由加法指令求和。顺着“读数据、写数据、同步、求和”来看，就能和源码接上。先把其中一组加载、存储放到一起：

ld.global.f32 %f59, [%rd17];
st.shared.f32 [%r25], %f59;

第一条的ld表示加载，.global指出全局内存，.f32表示32位浮点数：从%rd17提供的地址读出一个值，放进%f59。第二条的st表示存储，.shared指出共享内存，把刚才的%f59写到%r25提供的地址。一行“从input读取，再赋给smem”的源码，在这里拆成了读取和写入两种指令。带百分号的名字是PTX虚拟寄存器，后续编译还会分配机器寄存器，因此这些编号不能直接用来数Registers Per Thread。

完成装载后，bar.sync 0;执行block内的屏障同步，在图中对应装载后的__syncthreads()。它让block里的线程先完成共享内存写入，再继续求和。往下看第175行，出现了更宽的加载：

ld.shared.v4.f32 {%f25, %f26, %f27, %f28}, [%r36];

.v4表示四个元素，每个元素是.f32，所以每个参与线程一次读取4 × 32 = 128 bit，即16字节，分别放入四个虚拟寄存器。第180行的[%r36+16]把地址向后移16字节，再读下一组四个float，两条加载正好取到八个数，后面的add.f32把它们相加。Source第99行Access Size出现128，现在有了对应的编译结果。这说明编译器已经完成了成组读取，接下来应检查这两条宽加载的地址布局是否仍造成额外工作。

再看两条加载右侧的L1 Wavefronts Shared Excessive，各为50.00 %。分母是整个kernel多出的524,288个共享wavefront，分子是每条加载各自多出的262,144个，262,144 ÷ 524,288 × 100 % = 50 %。两条加载的多余量相加为524,288个，按源码行汇总时都对应第99行，所以Source第99行显示100 %。第二章定位到的两条LDS.128，正是这两组PTX加载最终生成的机器指令；PTX在这里帮助我们看清了八个数怎样分组读取。两条加载各占一半，排查共享内存布局时就要同时检查两组地址；只改善其中一组，另一组的多余工作仍会保留。

SASS：读GPU实际执行的指令

把View切成SASS，左侧增加了机器指令地址，右侧还能看到寄存器依赖连线。

我们以图中选中的第24行为例。最左边#里的24是SASS列表行号，Address列给出这条机器指令的地址，Source列里的内容是：

@!P2 LDG.E R5, [R10.64]

这条指令从全局内存加载数据，读出的值放进R5，[R10.64]使用以R10起始的64位地址操作数提供读取地址。最前面的@!P2是执行条件：P2是保存真假的谓词寄存器，!表示取反，P2为假的线程执行这次加载。

然后沿右侧指标从左往右读。Live Registers为19，这处代码需要保留19个寄存器的值。它比Source第90行显示的23小，说明同一行源码关联的不同机器指令位置，也会有不同的存活寄存器需求。沿关联指令比较这一列，可以继续定位哪些位置需要同时保留更多中间值。

Warp Stall Sampling (All Samples)为4.14 %，因为本次17,547个All Samples样本中，有727个记录的指令位置就是这条LDG.E，727 ÷ 17,547 × 100 % ≈ 4.14 %。Source第90行的4,123个样本，是这一行关联的全部机器指令位置上的样本总和，这条LDG.E位置的727个样本是其中一部分。这727个样本中有620个为MIO Throttle，说明在这条指令位置采样时，warp大多在等待MIO指令队列入口空位。需要结合这段输入搬运前后的共享访问等MIO指令，检查队列压力来自哪些操作。

Instructions Executed为0.32 %，可以用启动配置直接算一遍：8,192个block，每个block有256 ÷ 32 = 8个warp，共65,536个warp。这条展开后的加载指令由每个warp执行一次，合计65,536次。整个kernel的warp级指令执行总次数是20,488,192，所以65,536 ÷ 20,488,192 × 100 % ≈ 0.32 %。再看Avg. Predicated-On Threads Executed，为32，表示warp每次执行这条LDG.E时，平均32个活跃线程都满足P2为假的条件并参与加载。这条加载的执行次数和线程参与都与当前启动配置吻合，排查重点就落到了每次访问这些数据时产生的内存工作上。

接下来的Address Space、Access Operation和Access Size分别是Global、Load、32，合起来就是：每个参与线程从全局内存读取32 bit，即一个float的4字节。指令中的.64描述地址操作数的宽度，这里的32描述读出数据的宽度，顺着具体指令就能把两处数字对应清楚。32个线程各读4字节，一个warp实际需要的数据是128字节；这些数据会落到多少个sector，取决于线程地址的排布。

这一行的L1 Wavefronts Shared Excessive留空，L2 Theoretical Sectors Global Excessive为12.50 %。分子是这条LDG.E多出的1,835,008个理论sector，分母是整个kernel全部全局访存多出的14,680,064个理论sector，1,835,008 ÷ 14,680,064 × 100 % = 12.50 %。八条LDG.E的多余量相加正好是14,680,064个，它们都关联到Source第90行，因此该源码行显示100 %；PTX第150、165行的两条全局加载也各显示12.50 %。查找非合并全局访问时，就可以沿这个多余sector指标定位到具体加载指令。这条指令只占全部warp级指令执行次数的0.32 %，却产生全部多余全局理论sector的12.50 %，说明执行次数占比低的指令也可能带来较多额外访存工作。优化时应调整同一warp的读取地址排列，再比较该加载的多余sector绝对数和kernel运行时间。

最后看右侧的寄存器依赖图。Register Dependencies展示普通寄存器的读写关联，Predicate Dependencies展示谓词寄存器的读写关联。图中的三角形与左侧SASS指令逐行对齐：朝左的实心三角形表示该指令写入寄存器，朝右的空心三角形表示读取寄存器。连线把同一寄存器可能的写入位置和读取位置连起来，同一行同时读写这个寄存器时，两种三角形都会出现。颜色渐变帮助辨认连线跨越的代码范围。

以第24行的@!P2 LDG.E R5, [R10.64]为例，先在Register Dependencies中找到它读取地址寄存器的空心右箭头。沿连线向上，可以找到第21行IMAD.WIDE写入地址寄存器的实心左箭头：第21行算出以R10起始的64位地址，第24行再用这个地址加载数据。第24行还把加载结果写入R5，对应这一行的实心左箭头；从这个写入位置沿线寻找后面的读取箭头，就能查找哪些指令使用了加载结果。再看Predicate Dependencies，第24行读取P2作为执行条件，沿读取箭头追到写入箭头，就能查找这个条件由哪条指令生成。

这样，我们就能把一条加载的地址来源、结果去向和执行条件串起来。本例已经发现非合并全局访问，接下来可以沿R10的计算过程向前追，再与源码中的globalIdx对照，解释同一warp的读取地址为什么相隔很远。沿R5向后追，则能查看加载结果怎样交给后续的共享存储。要进一步分析等待，就查看相关指令所在行的Warp Stall Sampling色块，把采到的warp状态与这些读写关系一起分析。暂时只查执行次数或访存问题时，可以用Column Chooser隐藏连线列，把宽度留给当前要读的指标。

并排视图：把代码与指令连起来

现在把View切回Source and SASS，就回到了开头的左右布局。左边看源码，右边查具体指令，选中一边，另一边高亮关联位置。我们用共享内存求和这一处走一遍，把前面分开读过的代码和数字放到同一屏上。

左侧选中第99行threadSum += smem[8 * threadIdx.x + i];，右侧滚到了SASS第55到90行附近。高亮的指令中有第71、72行的两条LDS.128，以及后面的FADD，它们分别完成共享内存读取和浮点求和。计算读取地址的机器指令也可能通过编译行号信息对应到第99行，因此选中源码行之后，还要逐条查看高亮的SASS，分清加载、加法和地址计算各自对应哪些指令。

本图的L1 Wavefronts Shared Excessive用缩写绝对值显示：左侧源码第99行为524K，即这行对应的两条加载合计多出524,288个wavefront；右侧两条LDS.128各为262K，即各多出262,144个。切回百分比显示时，分母仍是整个kernel的524,288个多余共享wavefront，所以源码第99行为100 %，每条LDS.128各为50 %，与前面PTX两条共享加载的50 %一致。

这条关联把分析线索串起来了：源码要求每个线程取八个float求和，PTX把加载组织成两组四元素读取，SASS确认最终使用两条128位共享加载，性能指标指出多余wavefront集中在这两条指令上。后续调整共享内存布局时，就可以回到这两条加载，比较修改前后的多余wavefront绝对数，确认两组读取的额外工作都减少了。

同样，选中第90行可以检查八条全局加载和八条共享存储，选中归并第111行可以找到对应的加载、FADD和存储，再单独核对加法的线程参与。Source and PTX、PTX and SASS也采用同样的关联方式，前者方便看源码怎样展开，后者方便看中间操作最后编成什么。这里两侧属于同一个kernel结果；不同报告之间的性能比较由标签栏右侧的Compare另行选择对象。

底部附加面板

页面最下面的三个标签围绕上方代码补充信息。Inline Functions区分函数内联后的调用位置，Source Markers集中列出代码提示，Statistics统计选中的一段代码。

Inline Functions

假设kernel两次调用同一个小函数，编译器把函数体直接展开到两个调用位置，这就叫内联，最终可能有两组SASS对应同一份函数源码。在Source表格中，函数内某行的执行次数，会把两个调用位置对应机器指令的执行次数加在一起；Inline Functions则按调用位置分别列出SASS及其指标，方便我们比较两次调用各执行了多少工作。

本例的Inline Functions面板保持为页面初见图里的选择提示，当前没有可展示的内联调用位置。源码、指令和性能指标仍在主体表格中，分析这份reduce_ill报告时，沿前面的关联视图继续即可。

Source Markers

前面代码行号旁边的黄色三角形，在Source Markers中有完整的提示列表。分析规则把警告、错误或通知附到相应代码位置，面板负责集中展示，点击Line/Address里的行号可以跳回对应行。

截图顶部保留All，列表分成reduce.cu和reduce_ill两组，分别对应源码和SASS。组名后面的4和12表示各组有多少条标记，同一行可以有多条提示，例如源码第99行同时有Warp Stall和Uncoalesced Shared Accesses。File这一列列出分组和提示类别，Line/Address给出对应位置，Marker写出具体说明，Source保留那一行代码或指令。

先看与访存有关的两类提示。源码第90行的Uncoalesced Global Accesses写着87.50 %，第99行的Uncoalesced Shared Accesses写着50.00 %，SASS组里对应的访存指令也有这些标记。这里的分子是这一行多出的sector或wavefront，分母是这一行实际需要的sector或wavefront，计算如下：

第90行按理想情况需要2,097,152个sector，实际需要16,777,216个，多出的14,680,064个占实际需求的87.50 %。第99行理想需要524,288个wavefront，实际需要1,048,576个，多出的524,288个占实际需求的50.00 %。回到Source表格，第90行的全局理论sector多余量占整个kernel全部多余全局理论sector的100 %；第99行的共享wavefront多余量，占整个kernel全部多余共享wavefront的100 %。标记中的比例衡量这一行自身有多少额外工作，Excessive列的比例则衡量整个kernel的多余工作分布在哪些代码行。按实际需求除以理想需求，第90行的sector需求达到理想的8倍，第99行的wavefront需求达到理想的2倍。这两处都有明确的额外访存工作：全局读应检查warp内地址连续性，共享读应检查bank映射，再分别核对修改后的sector和wavefront绝对量。

第二章按整个kernel计算时，全局理论sector总数是16,785,408，比第90行多8,192。这部分来自第119行的写回：8,192个block各由thread 0写一个float，每次写回需要一个sector，共8,192个。该写回的理想需求也是8,192个，所以它增加了全kernel的sector总数，却没有增加Excessive多余量。由此，全kernel的多余比例为14,680,064 ÷ 16,785,408 × 100 % ≈ 87.46 %，第90行单独计算则是87.50 %。写回每个sector只用4字节，仍有前面病态四的利用率问题；Excessive衡量的是相对当前参与线程与访问宽度的理想sector需求，还要结合实际使用的字节数判断访问效率。

列表里还有Warp Stall提示。源码第99、109行的提示让我们继续检查对应的SASS位置。第71行LDS.128位置采到1,757个All Samples样本，占整个kernel的17,547个样本约10.0 %；这1,757个样本中有1,583个记录为barrier，占1,583 ÷ 1,757 × 100 % ≈ 90.1 %。第115行位置有4,105个样本，占17,547个样本约23.4 %；其中4,069个为barrier，占4,069 ÷ 4,105 × 100 % ≈ 99.1 %。每条提示的第一个百分比说明样本分布在哪个指令位置，第二个百分比说明这个位置采到的样本中有多少在等待同步。

这些标记要和附近代码一起读。第71行是一条共享内存加载LDS.128，该位置的样本却主要记录为barrier，图中前一条第70行正是同步指令BAR.SYNC。指令名告诉我们该位置准备做什么操作，采样状态告诉我们采到warp时它处于什么状态。看到这些barrier样本，就应向前检查同步点，并结合归并循环与第三章的Barrier统计，分析warp在等待哪些线程到达同步点。第115行分支前的第113行也是BAR.SYNC，对应归并循环中的块级同步。结合该分支位置几乎全是barrier样本，以及归并加法参与线程逐轮减少的现象，优化归并时应把线程分工与每轮块级同步一起检查，在保持数据依赖正确的前提下评估能否减少同步。

Statistics

Statistics根据上面选中的代码行计算统计结果，按住Shift可以选择多行代码。前面的SASS图只选中一条LDG.E，第99行关联图也只选中一行源码，所以各自面板里的汇总、最大值、最小值和平均值会出现相同数值。现在改为选择一段代码，就能观察区间内部的差别。

图中保留源码、Instructions Executed和对应统计列，便于把选区与计算结果对着读。下面四个百分比都以整个kernel的warp级指令执行总次数20,488,192为分母；分子分别取选区内的执行次数总和、单行最大值、单行最小值和各行平均值：

将面板结果对应回报告，参与这组执行次数统计的源码行是109、110、111、113、118、119和121，共7行，其中第121行对应函数收尾指令。这7行都有非零的warp级指令执行次数，所以Average先把这7行的次数相加，再除以7；选区中的注释、空行不参与这次平均。第110行关联指令执行9,961,472次，第111行为2,621,440次，前者是后者的3.8倍，与前面Source视图中看到的条件判断和归并操作的执行次数对比一致。

这些统计规则要逐列看。Instructions Executed的Aggregate把选区内的warp级指令执行次数相加；Live Registers的Aggregate则取选区内最大的存活寄存器数，表示这段代码的峰值寄存器需求。检查寄存器时，可以用这个峰值找出需要继续细分的代码区间，再结合Launch Statistics的实际分配量和Occupancy判断是否影响驻留。Address Space这样的文字属性会在汇总结果中列出选区出现过的地址空间，它的Maximum、Minimum和Average留空。报告未采集的指标，或当前指标不支持的统计项，也会留空。

这一章到这里就结束了。我们先认清Source页的布局，再从源码的性能分布出发，沿PTX和SASS看操作怎样展开，用并排关联把同一处问题串起来。Source Markers帮助找到需要检查的位置，Statistics帮助比较一段代码里的工作分布。面对Details中发现的异常，就能把硬件侧的数字接回具体指令，再解释回源码的访问和执行方式。下一章沿标签栏继续往后，读Context、Comments、Raw和Session，把这份报告的分析过程收束起来。