ggJ &SSKrSSKrSSKJrJrJr \"SS\R SS5\"SS\R SS5\"S5\"S S \R SS5/rS rS r S r Sr Sr Sr SrSrSrSrSrSrSrSrSrSrSrSrSrg)N) Importance MetricRequestRequestedMetricsParserlaunch__block_size block_sizeFlaunch__grid_size grid_size&device__attribute_multiprocessor_count launch__waves_per_multiprocessor num_waves zSee the @url:Hardware Model:https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-hw-model@ description for more details on launch configurations.cg)NLaunchConfigurationr4nsight-compute-2025.1.1/sections/LaunchStatistics.pyget_identifierr,s rcg)NzLaunch Configurationrrrrget_namer0s !rcg)Nz$Kernel launch configuration analysisrrrrget_descriptionr4s 1rcg)N LaunchStatsrrrrget_section_identifierr8s rcS/$)N SOLBottleneckrrrrget_parent_rules_identifiersr<s  rcXSn/nUR5S:XaUR5[-S:gnX4$[UR55HVnURU5nU[-S:wdM#SnUR 5RU5nUR XT45 MX X4$)NFrT) num_instancesvalue WARP_SIZErangecorrelation_idsappend)block_size_metrichas_block_size_issueproblematic_block_sizes instance_idr launch_ids rget_problematic_block_sizesr*@s  &&(A-06689DI 88!!2!@!@!BCK*00=JI%*'+$-==?EEkR '.. /FG D 88rc[R"U[- 5nU[-nUS:XdUS:XaSnOSU- SU[- - -S-n[RR U4$Nrd)mathceilr!NvRules IFrontendSpeedupType_GLOBAL)r num_warpsnum_threads_last_warpimprovement_percents r get_estimated_speedup_block_sizer7Rsq *y01I&2!Y!^]q#89#DD E K     / /1D DDrcSnSnSnUHupE[U5upgXq:dMUnUnUnM! [RRXU4$Nr)r7r1r2SpeedupType_LOCAL) block_sizes max_speeduplaunch_id_at_max_speedupblock_size_at_max_speedupr)r_ speedup_values r$get_max_estimated_speedup_block_sizerA`s`K  !!, ;JG  &'K'0 $(2 % "-    . . Wp pprc SnUSnUcg[U5upVU(dgUR5[RR:Xa8[ UR 55nSRU5n[U5upO.[U5uppSR[UUS:XaSOSU 5nS R[U[[5n UR[RRX5n URXU 5 UR!XR#5U[RR$S R[55 g) a2Execute the block size rule and generate a rule output if necessary. Check whether the block size of (each) kernel launch is a multiple of the warp size, and generate a rule message and a speedup estimate otherwise. In case of range results, find and report the launch with the largest potential speedup. z Block SizerNzCThis kernel launch is configured to execute {:d} threads per block.zSome kernel launches of this workload are configured to use a number of threads per block that is not a multiple of {} (e.g., {} {} for launch ID {}).r-threadthreadsaThreads are executed in groups of {} threads called warps. {} Consequently, some threads in a warp are masked off and those hardware resources are unused. Try changing the number of threads per block to be a multiple of {} threads. Between 128 and 256 threads per block is a good initial range for experimentation. Use smaller thread blocks rather than one large thread block per multiprocessor if latency affects performance. This is particularly beneficial to kernels that frequently call __syncthreads(). {}z>Arrange the number of threads per block to be a multiple of {})r* workload_typer1IActionWorkloadType_KERNELintr formatr7rAr!HARDWARE_MODEL_REFmessager2MsgType_MSG_OPTIMIZATIONspeedup focus_metricnameSeverity_SEVERITY_LOW)feactionmetrics rule_namer%r&r'rworkload_specific_part speedup_typer@r) rule_messagemsg_ids rapply_block_size_rulerYqs_I -   $$562  !D!DD*0023 !f!m!mnx!y&Fz&R# m 11H I ; Y //5v&!O 0   f  "    ZZ))BBL \FJJv]3OOF224j'BSBSBiBilllslst}l~rcNX- U- S-n[RRU4$)Nr.)r1r2r3)r num_smsr6s rget_estimated_speedup_grid_sizer\s-".'9C?    / /1D DDrcSnSnSn[UR55HKnURU5n[Xa5upxX:dM(UnUR 5RU5nUnMM [ R RX#U4$r9)r"rr r\r#r1r2r:) grid_size_metricr[r<r=grid_size_at_max_speedupr(r r?r@s r#get_max_estimated_speedup_grid_sizer`sK  -;;=> $**;7 :9N  &'K'7'G'G'I'O'OP['\ $'0 $ ?    . . Wo oorc (^SnUSn[USR55mUcgUR5S:XaUR5/nO7[UR55Vs/sHodRU5PM nn[ U4SjU55(Ga*UR 5[ RR:Xa7[US5nSRXwS:XaS OS T5n[UT5upO*[UT5uppS RTXwS:XaS OS U 5nS RU[5n UR[ RRX5n UR!XU 5 UR#U UR%5U[ RR&S RT55 g[ U4SjU55(Ga(UR 5[ RR:Xa#[US5nSRUT- 5nOP[)U5nUR+U5nUR-5RU5n SRUT- U 5nSRU5n UR[ RRX5n UR#U UR%5U[ RR.SRST-55 ggs snf)aExecute the grid size rule and generate a rule output if necessary. Check whether the grid size of (each) kernel launch is at least the number of SMs, and generate a rule message and a speedup estimate otherwise. In case of range results, find and report the launch with the largest potential speedup. Also check if the grid size is less than twice the number of SMs for any launch, and suggest to use at least two blocks per SM when __syncthreads() is used. z Small Gridr r Nrc3,># UH oT:v M g7f)Nr.0sizer[s r 'apply_grid_size_rule..s 1jd'>jszrThe grid for this launch is configured to execute only {:d} {}, which is less than the GPU's {:d} multiprocessors.r-blockblockszAt least one kernel launch of this workload has a grid which is configured to execute fewer blocks than the GPU's {:d} multiprocessors (e.g., {:d} {} for launch ID {:d}).a#{} This can underutilize some multiprocessors. If you do not intend to execute this kernel concurrently with other workloads, consider reducing the block size to have at least one block per multiprocessor or increase the size of the grid to fully utilize the available hardware resources. {}zCIncrease the grid size towards the number of multiprocessors ({:d})c32># UH oST-:v M g7f)Nrrcs rrfrg s 7JDAK Jsz2(compared to the currently executed {:.1f} blocks)zE(compared to the currently executed {:.1f} blocks for launch ID {:d})zIf you execute __syncthreads() to synchronize the threads of a block, it is recommended to have at least two blocks per multiprocessor {} This way, blocks that aren't waiting for __syncthreads() can keep the hardware busy.zIIncrease the grid size towards twice the number of multiprocessors ({:d})rk)rHr rr"anyrEr1rFrGrIr\r`rJrKr2rLrMrNrOSeverity_SEVERITY_HIGHminindexr#rP)rQrRrSrTr^ grid_sizesir rUrVr@r)rWrXr(r[s @rapply_grid_size_rulerrsI{+'BCIIKLG %%'1,&,,./ 9>?O?]?]?_9`a9`A,,Q/9` a 1j 111    !W__%H%H HJqM*IFFLfq.whG # +J)U\*] 'L-44DgN >L77=vY>xQZ8 #  &/1CD G--FF ` 67    ! ! #     4 4**0&/  7J 7 7 7    !W__%H%H HJqM*IDKK' # JI$**95K(88:@@MI''-vi'.A9'M #  +,262H+I G--FF `    ! ! #     3 3**0&W*= 3 8]bs*Lc/nUR5S:XaXUR5n[R"U5up4US:=(a US:nU(aSnUR XcU45 XQ4$[ UR55HYnURU5n[R"U5up4US:=(a US:nU(dMFUR XcU45 M[ [ U5S:U4$)Nrr-)rr r/modfr$r"len)num_waves_metricproblematic_launchesr partial_waves whole_waveshas_tail_effectr(s rget_problematic_num_wavesr{.s%%'1,$**, %)YYy%9" '!+@ q0@ K ' '[(Q R44-;;=> $**;7 %)YYy%9" '!+@ q0@ ? ' '[(Q R ? # $q (*> >>rcUS:Xa[RRS4$SSU-- S-n[RRU4$r,)r1r2r:r3)rxryr6s r!get_estimated_speedup_tail_effectr}EsM  22A55q;/#5    / /1D DDrcSnSnSnSnUH"upVn[Xg5upX:dMU nUnUnUnM$ [RRXX44$r9)r}r1r2r:) rwr<r=partial_waves_at_max_speedupwhole_waves_at_max_speedupr)rxryr?r@s r%get_max_estimated_speedup_tail_effectrMsuK #$ !"1E- +<]X  &'K'0 $+8 ()4 & 2F    . . Ws PPrcSnUSnUSnUcgSn[U5upxU(GaUR5[RR:XalUSupn UR 5n UR 5n [ X5up[R"XU - -5nSR[U 5U5nO[U5upnpUR5R U5nUR U5n UR U5n [R"XU - -5nSRU[U 5U5nX:agS RUU[5nUR[RR UU5nUR#UX5 UR%UUR'5U [RR(S 5 gg) zExecute the tail effect rule and generate a rule output if necessary. Check whether any kernel launch has a partial wave of thread blocks. In case of range results, report the launch with the largest (local) speedup. z Tail Effectr r NrzWThis kernel launch results in {:d} full waves and a partial wave of {:d} thread blocks.zAt least one kernel launch of this workload results in a partial wave (e.g., for launch ID {:d}, {:d} full waves and a partial wave of {:d} thread blocks are executed).a#A wave of thread blocks is defined as the maximum number of blocks that can be executed in parallel on the target GPU. The number of blocks in a wave depends on the number of multiprocessors and the theoretical occupancy of the kernel. {} Under the assumption of a uniform execution duration of all thread blocks, this partial wave may account for up to {:.1f}% of the total runtime of this kernel. Try launching a grid with no partial wave. The overall impact of this tail effect also lessens with the number of full waves executed for a grid. {}zQDecrease the number of partial waves (the fractional part of the number of waves))r{rEr1rFrGr r}r/r0rIrHrr#rJrKr2rLrMrNrOSeverity_SEVERITY_DEFAULT)rQrRrSrTr^rvspeedup_thresholdrzrwr?rxryr r rVr@partial_wave_blocksrUr(r)rWrXs rapply_tail_effect_ruler`s I{+{+ ,EFV,W)O    !W__%H%H H,@,C )Ak(..0I(..0I*KM*g 'L"&))I9R,S"T //5v $&90 #66JK QLm(88:@@MI(..{;I(..{;I"&))I9R,S"T 55;Vs;/1D6 #  ,  &&" "G--FF V_` 6<7    ! ! #     7 7@ erc[R"U5nURS5RS5nUR 5n[ X5R [5n[X2U5 [X2U5 [X2U5 gr9) r1 get_context range_by_idx action_by_idxfrontendrparserequested_metricsrYrrr)handlectxrRrQrSs rapplyrsl   f %C   a . .q 1F B$V4::;LMG"g.W-2w/r)r/r1RequestedMetricsrrrOPTIONALrr!rJrrrrrr*r7rArYr\r`rrr{r}rrrrrrrs2 NN& j6I6I4QVW%{J4G4GuU:;4k:CVCVX\^cd   > !"29$ Eq"7tEp$g T?.EP&N b 0r