N j(I 8 % S SK Jr S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SK r S SK
r
S SK r S SK r S SK
r
S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SKrS SKrS SKrS SKrS SKJrJrJ r J r J r J r J r S SK!J!r! S SKJ"r" S SK J#r# S SK$J%r%J&r&J'r'J(r(J)r)J*r*J+r+J,r,J-r-J.r.J/r/J0r0J1r1 S SK2J3r3J4r4J5r5 S S KJ6r6 S SK7r7S SK8r8S SK9J:s J;r< S S
K=J>r> S S K?J@r@ S S KAJBrB S S
KCJDrD S SKEJFrF S SK9JGrGJHrH SS/rIS SKJJKrKJLrLJMrMJNrN \-( ah S SKJOrOJPrPJQrQ S SKRJSrS S SK8JTrTJUrUJVrV S SKWJXrX S SKYJZrZ S SK[J\r\ S SK]J^r^ SSK_J`r` SS KaJbrb SS KcJdrd SS KeJfrf SS KgJhrhJiriJjrjJkrkJlrlJmrm SS KnJoro SS!KpJqrqJrrr / S"Qrs\0" S#5 rt\R GSGS$ j5 rvS S%KwJxrx S S&KyJzrz S S'K{J|r| S S(K}J~r~ S S)KJr S S*KJr S S+KJrJrJrJrJr S S,KJrJr S S-KJrJr SS.KJr SS/KJr \R S0:H r\ GR* " \5 r\0" S15 r\\7GR4 \7GR4 4 r\+\1\8GR8 \\8R 4 r\( a S2O\GR> " S3S45 rS5S6S7\ 3S8.rS9rS9rS9rS:r\F" \8GRL \8GRN \8GRP \8GRR \8GRT \8GRV \8GRX \8GRZ \8GR\ \8GR^ \8GR` \8GRb \8GRd \8GRf \8GRh /5 rS;\S<' S=r\\S-
- S :X a \S>: d S?5 eGSHS@ jrGSISA jr " SB SC\7GRt 5 r\GRx " SDSE9 " SF SG5 5 rGSJGSKSH jjr GSL GSMSJ jjr GSL GSMSK jjr\R GSNSL j5 rGSOSM jrGSPSN jrGSQSO jrGSRSP jr GSSSQ jrGSTSR jr GSUSS jrGSVST jrGSWSU jr GSXSV jrGSYSW jrSX 4 GSZSY jjr GS[S[ jrGS\GS]S\ jjr GS^ GS_S] jjr GS` GSaS^ jjrGSbS_ jrGScS` jrGSdSa jrGSeSb jrGSfSc jr\4" Sd5 r\0" SeSDSf9r\\'\%\4 \4 r " Sg Sh\,\(\\4 5 rGSgSi jr GSgSj jr GShSk jr GSiSl jr GSjSm jr GSkSn jr GSl GSmSo jjr GSnSp jrGSoSq jrGSpSr jrGSqSs jrGSrSt jrGSsSu jrGStSv jrGSuSw jrGSvSx jrGSwSy jr\" / SzQ5 r GSxS{ jrGSyS| jrGSzS} jrS SKrGS{S~ jr/ rSZ\S' GS|S jrGS{S jr\GR GS}S j5 r\GR GS~ GSS jj5 r\r\r\rSIS.GSS jjrSIS. GSS jjr\RD " S>5 GSS j5 r " S S\*5 r\GRx " S S5 5 r " S S5 Gr " S SG\ 5 Gr\GR GSS j5 Gr " S S5 Gr " S SG\5 Gr\R GSGSS jj5 Gr\RD GSS j5 Gr\RD GSNS j5 GrGSS jGr GSl GSS jjGr GSS jGr
GSS jGr GSS jGr SISISDS. GSS jjGr
SSIS. GSS jjGrSIS. GSS jjGrSIS. GSS jjGr GSS jGr\RD " SS9GSNS j5 Gr\RD " SS9GSNS j5 Gr\RD " SS9GSNS j5 Gr GSS jGrGSS jGr GS GSS jjGrGSS jGr\1\\7GR4 4 GrS\S' \R GS GSS jj5 Gr\R GSS j5 Gr\R GSS j5 Gr \R GSS j5 Gr \R GSS j5 Gr GSS jGr GSS jGr GSS jGr!GSS jGr"GSS jGr# GSS jGr$ GS GSS jjGr%GSNS jGr& " S S5 Gr' GSS jGr( GSS jGr)GSS jGr*GSS jGr+GSS jGr, GSS jGr- GSS jGr.\GR GSS j5 Gr/ GSl GSS jjGr0GSS jGr1GSS jGr2GSS jGr3GSS jGr4GSS jGr5GSS jGr6\GR GSS j5 Gr7GSS jGr8\R GSS j5 Gr9\R GSS j5 Gr:\R GSS j5 Gr;GSS jGrGSS jGr?GSNS jGr@GSNS jGrAGSS jGrBGSwS jGrC " S S\GR 5 GrE GSS jGrFGSS jGrG GSS jGrH GSl GSS jjGrIGSS jGrJ GSl GSS jjGrKGSS jGrL GSS jGrM GSS jGrNS 4 GSS jjGrOS 4 GSS jjGrPGSS jGrQGSS jGrR\GRx " S S5 5 GrS\GR GSS j5 GrTGSS jGrUGSS jGrVGSNS jGrWGSS jGrXGSS jGrY GSS jGrZGSS jGr[GSS jGr\GSS jGr]GSS jGr^ GSS jGr_GSS jGr` GSS jGraGSS jGrb GSl GSS jjGrc GSS jGrdGSS jGre GSS jGrfGSNGS jGrgGSGS jGrhGSGSGSGSGSGSGSGS.GriG\iGR 5 V Vs0 s H u pX_M snn Grk\GR " GS 5 GrmGSGS
jGrnGSGS jGroGSGS jGrpGSGS
jGrq\R GSGS j5 Grr\GRx " GS GS5 5 Grs0 GrtGS\GS' GSGS jGru\F" 5 GrvGS\GS' GSGS jGrwGSlGSGS jjGrxGSGS jGry\0" GS5 Grz\0" GS5 Gr{ " GS GS \ G\zG\{4 5 Gr|\3" SDGS 9GSlSDSE.GSGS jjj5 Gr}GSGS jGr~ " GS GS!\GR 5 Gr\R GSGS" j5 GrGSNGS# jGrGSGS$ jGrGSGS% jGrGSGS& jGrGSGS' jGrGSGGS( jGrGSGS) jGrGSNGS* jGrGSGS+ jGrGS,GrGSGS- jGrGSGS. jGrGSGS/ jGr GS GSGS0 jjGrGSGS1 jGrGSGS2 jGrGSNGS3 jGrGSGS4 jGrGSGS5 jGr\GRx " SDSE9 " GS6 GS75 5 Gr\GS8\%4 Gr\G\G\/G\4 Gr " GS9 GS:5 GrG\" 5 GrGSGS; jGrGSGS< jGrGSGS= jGrGSGS> jGrGSGS? jGr\F" / GS@Q5 GrGSGSA jGr\"GSGSB j5 GrGSGSC jGr GSGSD jGr GSGSE jGr GS GSGSF jjGrgs snn f ( ) annotationsN)Callable
Collection GeneratorIteratorMappingMutableMapping
MutableSet)datetime) lru_cache)StringIO)
Anycast ConcatenateGenericLiteral
NamedTupleOptionalProtocol
TYPE_CHECKING TypeAlias TypeGuardTypeVarUnion)dataclass_transform ParamSpecSelf)mock)datasheet_tops)DeviceProperties)_needs_inductor_compile) dtype_abbrs)
OrderedSet) tree_flatten
tree_map_only!activation_quantization_aten_pass inductor_autotune_lookup_table) free_symbolsfree_unbacked_symbols IterateExprsShapeEnv)IterableSequence
ValuesView)Path)SymBoolSymFloatSymInt) ELEMENTWISE_TYPE_PROMOTION_KIND) GraphModule)Node) ScalingType ) WorkspaceArgPythonWrapperCodegen)Dep
GraphLowering)Buffer ExternKernelIRNodeLayout OperationReinterpretViewCompiledFxGraph)BaseSchedulerNodeSchedulerBuffer)cudampsxpumtiaTc [ V s/ s H* n [ [ U 5 R 5 ( d M( U PM, nn [ U5 S:: d e[ U5 S:X a SnU$ UR 5 nU$ s sn f )Nr7 r rH ) GPU_TYPESgetattrtorch is_availablelenpop)x
avail_gpusgpu_types f/root/GenerationalWealth/GenerationalWealth/venv/lib/python3.13/site-packages/torch/_inductor/utils.py get_gpu_typerX j sh &KY'%*;*H*H*J!YJK z?a Z A-vH O 4>>>3CH O Ls
'A2A2)get_interface_for_device)detect_fake_mode)
DeviceType) EventList)GraphTransformObserver) ShapeProp)CeilDivCleanDivFloorDivIdentityModularIndexing) make_symbolSymT) bound_sympy ValueRangesconfig)ceildivwin32_Tspv TORCHINDUCTOR_XPU_KERNEL_FORMATzebinz.cubinz.hsaco.)rH hiprJ zOrderedSet[torch.dtype]_TMA_SUPPORTED_DTYPES@ zmust be power of 2c * U [ - S-
[ * - $ )z/Round up to the nearest multiple of ALIGN_BYTESr7 ) ALIGN_BYTES)nbytess rW _alignrz s [ 1 $ 44 c [ U [ R [ R 45 ( a# [ [ [ U R 5 5 $ [ U [ 5 =( d" [ R " U [ 5 [ :H $ )z:v can be statically proven to be a multiple of ALIGN_BYTES)
isinstancesympyAddMaxallmap _is_alignedargsaligngcdrx )vs rW r r sT !eii+,,3{AFF+,, a K599Q #< #KKr{ c 4 \ rS rSrSrSrSr\SS j5 rSr g) r z 5 c MG UPMK sn5 nU( a# U[@ R& " S U 5 5 S- - n[* R- SU5 U$ s snf s snf ! , ( d f GN= fs sn
n f ! , ( d f GNK= fs snf ):
Returns benchmark results by examining torch profiler events.
This could be more accurate as it doesn't count CPU side overhead.
However, this also requires manually excluding irrelevant event, e.g.
vectorized_elementwise_kernel which is used to fill L2 cache,
various CUDA events, etc, so could also be fragile.
ArH dtypedeviceT
enable_timing r7
activities
RunCudaModuleN
raw eventsself_device_time_totalsort_by row_limitzfused_abs_max_\dc 3 8 # U H oR v M g 7fr device_time_total.0events rW fp8_bench..( s Q 3 3 @@profiling results: %s ms)!rP rH synchronizeemptyr float16Eventrecordrangezero_ elapsed_timemaxprofilerprofileProfilerActivityCUDAnvtxtensorzipmeanitemlogdebug key_averagestabler\ events device_typer[ rematchname
statistics)fnwarmuprepcache start_event end_event_ estimate_msn_warmupn_repeatpisetimesresr filtered_eventss rW fp8_benchr s> D JJ KKJu}}V LE **"""6K
t 4I
1X
JJ **959K 1c&./0H1c#+,-H 8_
BGx Q A5::##$#7 KQ?DXO!!!!5IO NN + + 0 0
x A KKM N ! ! # &&7 8 L ! !
+.{+F G+F41Q^^A
+F G
" **U
"CIIlIIann $$-EQS$TU
#!!Z__4
HH 0%**=
#
O OOQ Q Q
II(#. JO RO 87
H
*
sP 'P 'P!;A8P>3P&;AP> P8 2P>"Q) Q
Q&
P5 0P>>
Q
Fc 4 SSK Jn U" [ 5 " XX#5 $ )Nr ) may_distort_benchmarking_result)$torch._inductor.runtime.benchmarkingr _do_bench_using_profiling)r r r is_vetted_benchmarkingr s rW do_bench_using_profilingr 0 s " U *+D E
C r{ c
( U( d
SSK Jn U" 5 [ 5 nUR 5 n[ U5 nU " 5 UR 5 [ R " [ S5 [ R US9nUR SS9n UR SS9n
U R 5 [ S5 H n UR 5 U " 5 M U
R 5 UR 5 U R U
5 S- n [ S[ X - 5 5 n
[ S[ X,- 5 5 n[ U
5 H
n U " 5 M UR 5 [ R R! [# [ R R$ U5 /S 9 n[ U5 H n UR 5 U " 5 M UR 5 S
S
S
5 [& R) S 5 [& R) WR+ 5 R- S S
S95 [/ UR1 5 Vs/ s H7 nUR2 [# [4 U5 :X d M# UR6 S:w d M5 UPM9 sn5 n[9 U5 U- S:w a [; SU[9 U5 U5 e[9 U5 U- n[/ [= U5 VVs/ s H u nnUU- S:w d M UPM snn5 nUR? 5 UR+ 5 n[& R) S5 [& R) UR- S
S95 [A S U 5 5 S- U- n[& R) SU5 U$ ! , ( d f GN= fs snf s snnf )r r )may_ban_benchmarkingr r Tr r r7 r Nr r r r z Context SynczWFailed to divide all profiling events into #repeat groups. #%s events: %d, #repeats: %szprofiling time breakdown)r c 3 8 # U H oR v M g 7fr r r s rW r ,_do_bench_using_profiling.. s
A=%%%=r r r )!r r rX upperrY r rP r r r r r r r r r r rO r r r r r r\ r r r[ r rR RuntimeError enumerate _build_treesum)r r r r r r device_type_upperdevice_interfacer r r r r r r r r r num_event_per_groupr
actual_eventsr s rW r r H s "M .K#))+/ <D " KKJuyy ME #((t(. s >odaAEos )rR r~ expandr r )seq1seq2s rW sympy_dotr. s6 t9D !! ! <<>c$o>> ??r{ c b U Vs0 s H n[ U5 U_M snR 5 $ s snf r )r values)r$ rT s rW uniquer1 s+ !bBqE1Hb ! ( ( ** !s ,c
[ U [ R 5 ( d [ U[ R 5 ( a4 [ [ R " U 5 [ R " U5 5 $ [ U [
5 ( a [ U[
5 ( d$ U S[
U 5 SU S[
U5 35 e[ X5 $ )Nz: , )r} r~ Exprr_ sympifyr r runtime_ceildiv)numberdenoms rW rj rj s &%**%%E5::)F)Fu}}V,emmE.BCC fc " "z%'='= ("T&\N"UG2d5k]; = 6 ))r{ c t U c g[ U 5 R S5 S n0 SS_SS_SS _S
S _S S
_SS_SS _SS_SS_SS_SS_SS_SS_SS _S S _S S _S!S"_SS#S$S%S&.EnUR [ UR 5 5 Vs0 s H o3U_M sn5 [ U [ 5 ( a U $ S'X! 3$ s snf )(Nz*i8rp r booli1
float8e4nvfp8e4nvfloat8e5fp8e5 float8e4b15fp8e4b15
float8e4b15x4
fp8e4b15x4
float8_e4m3fn float8_e5m2float8_e8m0fnuu8float4_e2m1fn_x2r fp16bfloat16bf16float32fp32float64fp64int8i8int16i16int32i32int64i64u16u32u64)uint8uint16uint32uint64*)r splitupdatelistr0 r} )key dtype_strtysr s rW _type_ofrf sW
{Cs#B'I i G z
w $ D 6 F 6 6 ! " # $ % & ' ( / C4 JJd3::< 01 01 012 S#&&3@a/?,@@ 2s B5c Z U Vs/ s H n[ R " U5 PM sn$ s snf )z
Gets the shape and stride of a tensor. For non-symbolic tensors, this is
trivial. But for symbolic tensors, we need to map from SymIntNode into
sympy.Expr.
)r~ r5 lstr s rW convert_shape_to_inductorrj s% '* *cEMM! c ** *s (c p [ U [ R 5 ( a U R R $ U $ )z
Convert SymInt to sympy.Expr, leave int as is.
Unlike sympy.sympify() which converts int to sympy.Integer,
this function preserves int as int and only converts SymInt to Expr.
)r} rP r2 nodeexprvals rW convert_symint_to_exprrp s( #u||$$xx}} Jr{ c SSK Jn [ U [ 5 ( a U $ [ U [ R
5 ( a [ U 5 $ UR R R R U SS9$ )zD
Like convert_shape_to_symint, but operates on a single expression.
r7 VN)hint)
virtualizedrs r} r r~ r graphsizevars shape_envcreate_symintnode)r rs s rW convert_to_symintrz sk a
!U]]++
F !!++==ad=K
r{ c D U Vs/ s H n[ U5 PM sn$ s snf )zn
Takes a list of shapes from Inductor and converts them into symints (or just
ints if all shapes are static).
)rz rh s rW convert_shape_to_symintr| . s" +. .#Q a # .. .s c N [ S U R R 5 5 $ )z%
Does this op overload have aliasing
c 3 < # U H oR S Lv M g 7fr )
alias_infor r( s rW r is_view..< s F1EA||4'1Es )any_schema argumentsops rW is_viewr 8 s F1E1EF FFr{ c gNFr )r s rW r A s r{ c ^ U R S:w a g[ U R [ R R
5 ( d U R [ R L d g[ [ R R
U R 5 nU[ R L d [ U5 ( a [ U4S jU R 5 5 $ [ R R UR ; =( d T" U5 $ )z
Do all uses of this op have torch.Tag.pointwise or return True for optional `is_pointwise_fn`
Uses in views ops will follow the views uses
call_functionFc 3 < ># U H n[ UT5 v M g 7fr )is_pointwise_use)r uis_pointwise_fns rW r #is_pointwise_use..R s KA#A77s )r r} targetrP _ops
OpOverloadr getitemr r r usersTag pointwisetags)user r s ` rW r r ? s vv 3::uzz 4 455xGWGW9W
%**''
4F
!!!WV__KKKK 99 &++ - H 1HHr{ list[Any]c ^^ [ R R 5 m/ mSUU4S jjnTR " U /[ [ R
X1U45 Q76 n[
U R R 5 S:X a3 [ U R R S R 5 S:X a U4nTR U5 [ R R 0 T5 nUT4$ )Nc ` > TR U 5 TR S[ T5 35 $ )Narg)append placeholderrR )r g
graph_argss rW add_tensor_arg)gen_gm_and_inputs..add_tensor_arg] s, # }}s3z?"3 455r{ r7 r Tensor)r torch.Tensorr r5 )
rP fxGraphr r% r rR r returnsr r outputr4 )r r kwargsr rl gmr r s @@rW gen_gm_and_inputsr W s A%'J6 6
?? u||^F^L D FNN " "#q( &&q).. /8 ;wHHTN b! $B
z>r{ c t U S:X a g [ U 5 nUR 5 ( a UR 5 g g Nr )rY rQ r r s rW r r o s7
/7$$&&$$& 'r{ c [ U5 [ R " S5 [ R " 5 n[ U5 H nU " U6 n[ U5 M [ R " 5 nWc eXt-
$ )Ni9 )r rP manual_seedtime perf_counterr )modelexample_inputsr r t0r resultt1s rW timedr w sk d
B
5\'F
B
7Nr{ c [ R " [ U5 Vs/ s H n[ XX%5 PM sn5 n[ R " U5 U- n[ X- S 5 UR
5 $ s snf )Nz.6f)rP r r r medianprintr ) r r r repeatbaseliner r timingstooks rW print_performancer se ll>CFmLmue 4mLG << 5 (D T_S !# 99; Ms A3c F ^ [ X5 " 5 m[ XU4S j5 g)zKReplace obj.method() with a new method that returns a precomputed constant.c > T $ r r )r s rW r #precompute_method.. s r{ N)rO setattr)objmethodr s @rW precompute_methodr s
S
!
#F C (r{ c , U H n[ X5 M g)zFReplace methods with new methods that returns a precomputed constants.N)r )r methodsr s rW precompute_methodsr s #& r{ c 8 [ X: 5 [ X: 5 -
$ r )r r( r) s rW cmpr s qu:AE
""r{ c [ U [ 5 ( a U /U- $ [ U 5 S:X a [ U 5 " U S /5 U- $ U $ )Nr7 r )r} r rR r )rT sizes rW pad_listliker sD !SsTz
1v{Aw!v%% Hr{ c @ [ U 5 S:X a / $ SS jn[ XS9$ )Nr c [ U [ 5 ( a U $ SSKJn [ X5 ( d eU R 5 $ )Nr7 )rF )r} r schedulerrF get_name)elemrF s rW sort_func tuple_sorted..sort_func s4 dC K0$2222}} r{ rc )r rl r r )rR sorted)rT r s rW tuple_sortedr s$
1v{ ! ##r{ PRV) covariantc 2 \ rS rSr\SS j5 rSS jrSrg) CachedMethodi c g r r )r s rW clear_cacheCachedMethod.clear_cache s ),r{ c g r r selfr r s rW __call__CachedMethod.__call__ s r{ r N)r r r None)r P.argsr P.kwargsr r )r r r r staticmethodr r r r r{ rW r r s , ,Dr{ r c ^ U R nSU S3mSU 0n[ SU ST ST S3R 5 U5 [ R " U 5 " X! S3 5 nS
U4S jjnXCl U$ ) N___cacher z def zC_cache_on_self(self):
try:
return self.zy
except AttributeError:
pass
rv = fn(self)
object.__setattr__(self, "z%", rv)
return rv
_cache_on_selfc B > [ U T5 ( a
[ U T5 g g r r delattrr rc s rW r "cache_on_self..clear_cache s 4 D# r{ r r r r r execlstripr wrapsr )r r ctxwrapperr rc s @rW
cache_on_selfr s
;;D tfF
C *C
F E "' (+e ,
FH oob!#n&=">?G & Nr{ c [ U 5 $ )zU
Variant of cache_on_self for properties. The only difference is the type signature.
)r )r s rW cache_property_on_selfr s r{ c ^ SU 4S jjnU$ )Nc >^ ST SU R S3mSU 0n[ ST ST ST S3R 5 U5 [ R " U 5 " US 5 nS U4S
jjnX2l U$ ) Nr r r r z def inner(self: Any, *args: P.args, **kwargs: P.kwargs) -> RV:
args_kwargs = (args, tuple(sorted(kwargs.items())))
if not hasattr(self, "z2"):
object.__setattr__(self, "z%", {})
cache = self.z
try:
return cache[args_kwargs]
except KeyError:
pass
rv = fn(self, *args, **kwargs)
cache[args_kwargs] = rv
return rv
innerc B > [ U T5 ( a
[ U T5 g g r r r s rW r .wrapper..clear_cache s tS!!c" "r{ r r )r r r r rc
class_names @rW r 'cache_on_self_and_args..wrapper s :,a }F3 Rj ' (+e ,//2e 4 !U #
$ )
, #CL1 # ( r{ )r FN_TYPE[P, RV]r r r )r r s ` rW cache_on_self_and_argsr s
$ $ $L Nr{ c
SSK Jn [ U [ 5 ( ay [ R
" [ R U Vs/ s H? n[ US5 ( d M UR ( d M) UR R PMA sn[ 5 5 $ [ XR 5 ( a U R $ [ 5 $ s snf )Nr7 irrl )
r r} rb r r r or_r rl originsr# r? )
node_scheduler rl s rW aggregate_originsr $ s -&& LL *
*D4( "-1YY " !! )
L
M?? 3 3 $$$|
s C
C
+C
c ` [ U 5 nUS:X ag S nU Vs/ s HA nUR S:X d M SUR ; d M' UR S c M9 U" U5 PMC nn[ [ U5 5 nGOUS:X a / nU H nUR S:X d M S nSnSUR ; a UR S S nO$SUR ; a UR S S nS nU( d Mi [ US
[ 5 ( a UR US
U- 5 M UR US
R U- 5 M [ [ U5 5 nO:US :X a. U Vs/ s H oDR S:X d M UR PM" nnO[ eS R S
/U- 5 $ s snf s snf )N
original_atenc . U R S nSn[ U[ R R 5 ( a UR
R nU$ [ U[ R R 5 ( a [ UR 5 5 nU$ )Nr r )
r r} rP r r _overloadpacketr HigherOrderOperatorr r )originr rc s rW get_origin_meta_str2get_fused_kernel_name..get_origin_meta_strA su "KK 8MC-)>)>??#33<< J M5::+I+IJJ-,,./Jr{ r rP r source_fn_stackr fwd_source_fn_stackbackwardr7
inductor_noder fused) r r r r r# r} r r r r NotImplementedErrorjoin)r descriptive_names all_originsr r sources source_fnsuffixs rW get_fused_kernel_namer% : s $M2KO+ &
%yyO+
( 6;;.
( O,
( '%
G,- g %!FyyO+ $ 3 & ,= >r BI*fkk9 & ,A B2 FI 'F ilC00NN9Q<9NN9Q<#8#86#AB "" G,- o -&1
&1FYY/5QKFKKk
"! 88WI' ((G
<
s" F&F&F& F&!F+8F+c ^ ^ ^! [ U 5 nU Vs/ s H o3R S:X d M UPM nn[ R " [ 5 n[ R " [ 5 nSm U( a [ S U 5 5 n[
U5 S:X ac US R m [ T S5 ( d0 [ T R 5 VV s0 s H u pX_M n
nn U
T l UR U 4S jS9 U GHo n S U R ; a U R S b U R S n Sn
[ U [ R R" 5 ( a [% U R& 5 n
OB[ U [ R R( 5 ( a [% U R+ 5 5 n
U
( a Xm R- U R* 5 S
U R ; a< U R S
S R* n
X] R- U R* 5 GM& U R R/ S 5 S :X d GMH X[R* R- U R* 5 GMr T b S
OSnUR0 SU SSR3 UR5 5 5 SSR3 UR5 5 5 S3nUR0 S3/n[7 UR9 5 5 HA u nnUR- UR0 SU SSR3 [7 U5 5 35 MC T Gb SSK J m UR- UR0 S35 [ 5 n/ n[ U T R> 5 ( Gd SSK J!n S)U 4S jjnS*S jm!S+U!4S jjnU GH n [ U S 5 ( a
U RD c M$ [ U RD S 5 ( a U RD RF b U RD RF H nUR* U; a M URI UR* 5 UR RK UR* 5 nUc MZ U" UUR* 5 u nnUR- UR0 S U S U" U5 S!U S35 M [ U RD S"5 ( d GM+ U RD RL c GME U RD RL HW n UR RK U R* 5 nUc M- U" UU R* 5 u n n UR- S#U - 5 MY GM U H0 n UR- UR0 SU RO S$S%9 35 M2 UR- UR0 S&S'R3 U5 35 US(R3 U5 4$ s snf s sn nf ),a
Retrieves metadata information for a kernel.
Args:
node_schedule (Union[Sequence[BaseSchedulerNode], ExternKernel]):
Either a sequence of BaseSchedulerNode objects or an ExternKernel instance.
wrapper (PythonWrapperCodegen):
An instance of PythonWrapperCodegen, used to define the code comment format.
Returns:
tuple[str, str]:
A tuple containing two strings:
- The first string represents the kernel's metadata.
- The second string represent the kernel's detailed metadata.
r Nc 3 8 # U H oR v M g 7fr )rv )r ns rW r &get_kernel_metadata.. s "CNq77Nr r7 r )_inductor_kernel_metadata_node_to_idx_mapc " > TR U $ r )r* )r( single_graphs rW r %get_kernel_metadata.. s l T TUV Wr{ r r from_nodepartitioner_tag is_backwardzTopologically SortedUnsorted z Source Nodes: [r3 z], Original ATen: []z" Source node to ATen node mapping:z z => r
z Graph fragment:rr c R > [ U TR 5 ( aF [ U R TR 5 ( a! U R R R nO U R nUc UnO UR
n U R
5 nX44$ ! [ a S n X44$ f = fr )r} TensorBoxdata
StorageBox origin_noder
get_layoutr )bufferrw_namer8 r layoutr s rW get_buffer_info,get_kernel_metadata..get_buffer_info s fbll33
KK 9 9 #)++"2"2">">K"("4"4K &"D&++D" # . . 0F |# + " !F|#"s B B&%B&c j SSR U Vs/ s H n[ U5 PM sn5 S3$ s snf )N[r3 r3 )r r )shaperT s rW stringify_shape,get_kernel_metadata..stringify_shape s1 499e%a@@% U c gT" U R 5 nT" U R 5 nU R nS[ U R U U U S3$ )Nr ")r strider r" r )r< shape_annotationstride_annotationdevice_annotationrB s rW stringfy_layout,get_kernel_metadata..stringfy_layout sl > &5fkk&B%C '6v}}'E&F!'-}}o! FLL123C2D()*;))>??-778M5::+I+IJJ-,,./"'..tyy9 $)) #))K(+00C & &tyy 1
YY]],
-
> 99 % , ,TYY 7 *6)A%zH ??
1XJ&6tyy ATATAV7W6X Y 99%7%<%<%>?@ C
$OO,,NOP &~';';'= >
u s=/diiu
6N5O P
!? GOO#44D!EF%/\ "
- 99 &
$J
$UX
$-
$(
A
# q-00AMM4I 1=='22q}}7J7J7V ]] 0 0 66Y. $!
aff-!"!7!7!?!> $- $)8)H Q"))# *;<
2- #< #D $ $??#3t'7'7PT'7'U&VW
#
GOO#4Jsxx
?S>T!UV TYY 01 11I X #Ws WWWc [ U 5 n [ U 5 nU ( ak U R 5 nUR HB nU( a U" U5 ( a M XB; d M UR U5 U R U5 MD U ( a Mk U$ )zJReturns the set of nodes whose values depend on those within initial_queue)rb r# rS r r[ r )
initial_queue skip_filter
dominated_setrl users rW dominated_nodesrt sx
'M }-M
"JJD{4 0 0( !!$' $$T* - r{ c Z ^^ SSK Jm SUU 4S jjm [ U5 u p#U Vs/ s H nT " U5 ( d M UR PM nn[ U 5 u pcU Vs/ s H nT " U5 ( d M UR PM nn[ [
R " / UQUQ76 5 $ s snf s snf )Nr7 r
c l > [ U TR 5 ( a T" U R 5 $ [ U TR 5 ( a T" U R 5 $ [ U TR 5 =( a= [ U TR
TR TR TR 45 ( + $ r ) r} r5 r6 r7 r@ ComputedBuffer InputsKernel InputBufferTemplateBuffer)r( r is_unrealized_nodes rW r{ *gather_origins..is_unrealized_node$ s a & &%aff- - a ' '%aff- -!RYY'
!!!!
1
-
r{ )r( r@ r r: )r r r$ r r# itertoolschain)
r r kwargs_flattenr ro kwargs_origins args_flatten args_originsr r{ s
@@rW gather_originsr s
" %V,N-;W^c?QRU?Vkckk^NW"4(OL+7S.is_neg_leadC s: tUYY ' VC Na,? VDIIaLTVDV
r{ c v > [ U [ R 5 ( a [ U R 5 S:X aT T" U R S 5 ( a: T" U R S 5 ST" U R S R S 5 3$ SR [
TU R 5 5 $ T" U 5 $ )Nr r7 r z - z + )r} r~ r rR r r r )rm r
sympy_str_muls rW
sympy_str_add sympy_str..sympy_str_addH s dEII & & 499~"{499Q<'@'@' ! 56c- RS HYHYZ[H\:]9^__ zz#mTYY"?@@ & &r{ c > [ U [ R 5 ( aJ T" U 5 ( a ST" U R S 5 3$ SR [ TU R 5 5 $ T" U 5 $ )N-r7 z * )r} r~ r r r r )rm r sympy_str_atoms rW r sympy_str..sympy_str_mulS sa dEII & &4 >$))A,7899 zz#ndii"@AA!$' 'r{ c > [ U [ R 5 ( a U R $ [ U [ R [ R
45 ( a ST" U 5 S3$ [ U [ [ [ [ 45 ( aC U R R SSR [ [ U R 5 5 S3$ [! U 5 $ )N()r3 )r} r~ Symbolr r r rc r` ra rb funcr r r sympy_strr r )rm r s rW r !sympy_str..sympy_str_atom^ s dELL ) )99
uyy%)) 4
5
5}T*+1- -
(H M
N
Nii(()499SDII5N+O*PPQR Rt9 r{ )rm r r r: rm r r r r )rm r r r r s @@@@rW r r < s.
' ' ( ( r{ c SSK Jn [ R ( a9 [ UR
SS 5 =n( a UR S:w a [ U 5 $ [ R " 5 $ )Nr7 rr current_node
index_expr)
ru rs ri compute_all_boundsrO interpreterr rf rg unknown)r rs fx_nodes rW get_bounds_index_exprr k sN !!
~tD
DW
D NNl *5!!""$$r{ c U S S:H $ )Nr rj r )prefixs rW prefix_is_reductionr y s !9 r{ c D U [ R :w d e[ XSSS9$ )1
Used to generate an integer-nonnegative symbol.
Tinteger nonnegative)re SIZErd )r rb s rW sympy_index_symbol_with_prefixr } s' TYY vDd CCr{ c b U =( d [ R =( a [ R $ r )ri debug_index_assertsassert_indirect_indexing)checks rW generate_assertr s /V// TV5T5TTr{ c D U S S:w d e[ R " U SSS9$ )r r r Tr )r~ r r s rW sympy_index_symbolr s) 7c>> <<d ==r{ c SS jn[ R " U 5 R UR 5 VVs0 s H
u p4X2" X45 _M snn5 $ s snnf )z
When the passed replacement symbol v is a string, it is converted to a symbol with name v that
have the same replaced expression integer and nonnegative properties.
c [ U [ R 5 ( d e[ U[ 5 ( a* [ R " UU R
U R S9$ U$ )Nr )r} r~ r4 r r r is_nonnegative)replaced replacements rW to_symbol sympy_subs..to_symbol sV (EJJ//// k3 ' '<< ++ $ 3 3
r{ )r r r zUnion[sympy.Expr, str]r sympy.Symbol)r~ r5 xreplacerZ )rm replacementsr kr s rW
sympy_subsr sh +A == ' '(4(:(:(<=(<IaO (<= =s A
c [ U [ R 5 =( d- [ U [ R 5 =( a U R $ r )r} rP r2 r _has_symbolic_sizes_strides)r( s rW is_symbolicr s3 a & 1ell#E(E(Er{ c & [ S U 5 5 $ )Nc 3 8 # U H n[ U5 v M g 7fr )r r s rW r "any_is_symbolic.. s ,t!{1~~tr r )r s rW any_is_symbolicr s ,t, ,,r{ )z,aten._fused_moving_avg_obs_fq_helper.defaultz7aten._fused_moving_avg_obs_fq_helper_functional.defaultz fbgemm.dense_to_jagged.defaultz%fbgemm.jagged_to_padded_dense.defaultrun_and_save_rng_staterun_with_rng_statezaten._local_scalar_densezaten._assert_scalarc SSK Jn U R R HH n[ U5 ( a Us $ UR
R
S5 =nc M7 U" U5 ( d MF Us $ g )Nr )r) ro )%torch.fx.experimental.symbolic_shapesr) rv rU is_cudagraph_unsafe_fx_noder rW )r r) rl ro s rW %get_first_incompatible_cudagraph_noder sW L &t , ,K99==' 'C 49Ns9S9SK r{ c [ [ [ U R R 5 5 5 nUR
S:X d eU$ )z$Get the output node from an FX graphr )nextiterreversedrv rU r )r last_nodes rW output_noder s6 T(288>>234I <<8 ## # r{ c U R R SS9n[ S U 5 5 n[ U 5 R S n[ U[ 5 ( a UOU4n[ S U 5 5 nX%- $ )Nr r c 3 # U HX n[ UR R S 5 [ R 5 ( d M= UR S R
v MZ g7fro N)r} r rW rP r r )r rl s rW r "get_all_devices.. sC 9%D diimmE*ELL 9 % %s [ H n U R 5 M g)z
Clear all registered caches.
N)r r r s rW clear_cachesr ) s " "r{ c # ` # [ R R U 5 n U[ R U ' Sv Uc! [ R R U S5 gU[ R U ' g! Uc! [ R R U S5 f U[ R U ' f = f7f)a Thread-safe env var set/restore using atomic C-level lookups.
We avoid mock.patch.dict(os.environ, ...) because it internally calls
os.environ.copy(), which iterates all env var keys then fetches values in
separate steps. That approach is not atomic and can race with background threads
(e.g. Triton async compilation) modifying the environment, causing KeyError,
so we use os.environ.get() for individual keys which is an atomic C-level lookup.
N)osenvironrW rS )rc r olds rW _set_envr 1 sy
**..
C"
3
; JJNN3 % !BJJsO ; JJNN3 % !BJJsOs B.A2 8B.29B++B.c # ` ^# [ 5 SSKJn U" [ R " US95 m [ ST5 [ R ST5 U" [ R R TS5 5 n[ SU5 Sv [ U [ 5 ( a [ U 5 S:X d S 5 e[ R R U5 ( a{ [ R " U5 nU R! U Vs0 s HH nS
U; d M U[ R R# [ R R XF5 5 _MJ sn5 SSS5 SSS5 U( a^ [% 5 ( a- [& R( R+ 5 ( a
[- 5 [. R0 " T[% 5 U4S jS 9 [ 5 gs snf ! , ( d f N= f! , ( d f N= f! [2 a [ R5 S
T5 e f = f! [ 5 f = f7f)z
Contextmanager that provides a clean tmp cachedir for pt2 caches.
Optionally, pass a dict as 'cache_entries' to get a list of filenames and sizes
generated with this cache instance.
r )normalize_path_separator)dirTORCHINDUCTOR_CACHE_DIRzUsing inductor cache dir %stritonTRITON_CACHE_DIRNz!expected empty cache_entries dictz.lockc . > [ R STUS9$ )Nz*Failed to remove temporary cache dir at %s)exc_info)r warning)r pathr inductor_cache_dirs rW r fresh_cache..t s S[[@& % 6A 6r{ )
ignore_errorsonerrorz(on error, temporary cache dir kept at %s)r torch._inductor.cpp_builderr
tempfilemkdtempr r r r r r r} dictrR existslistdirra getsize
is_windowsrP rJ rQ r shutilrmtree Exceptionr )
cache_entriesr deleter
triton_cache_dirfilesfr s @rW fresh_cacher' F s ND1(2B2Bs2KL'
/1C
D II35G H 7 /: ,.>? mT22 }-2W4WW2ww~~&677 "
+; <%,, */ ).A#*!#3 !V277??277<<@P3T#U U).
@ E$ || 6 6 8 8&( MM" )l
5 @? E
DD >@RS
ss +H. G: A G)A9G=
G AG GG)A-G: H.G
G& "G))
G73G: :"H H H++H.)reversec U R n[ [ U 5 5 n[ [ X2SS95 nU( d [ [ U5 5 $ U$ )NTrc r( ) __getitem__r rR rb r r )seqr( gettera_rsort_idxs rW argsortr0 sC
__F
C/C F3D9:H HX&'' Or{ c F ^ SU 4S jjn[ U5 VVs/ s H> u pEU[ U[ R 5 ( a UR R
OU4PM@ nnn[
U[ R " U5 US9nU VVs/ s H u pGUPM nnnU$ s snnf s snnf )Nc ~ > U u p#Uu pESU4S jjnU" X5: 5 ( a gU" X5: 5 ( a gX$: a gX$: a gg)Nc R > [ U [ 5 ( a U $ TR U SS9$ )NT)size_oblivious)r} r:
evaluate_expr)rm rx s rW evaluate*argsort_sym..cmp..evaluate s+ $%% **4*E Er{ r r7 r )rm z%Union[bool, torch.SymInt, sympy.Expr]r r: r )r( r) a_idxa_valb_idxb_valr6 rx s rW r argsort_sym..cmp sN F
EM " " EM " "
= =r{ r* )r( tuple[int, sympy.Expr]r) r= r r ) r r} rP r2 rl rm r r
cmp_to_key) rx r, r( r rb r exprsr r s ` rW argsort_symr@ s 4 n
$FC
Z5<< 8 8affkka@$
5i 2 23 7 IE %
&fccF
& M
's ABB c r U [ R :X a g[ R " SU S9R 5 $ )Nrv r r )rP r^ r element_sizerB s rW get_dtype_sizerD s-
;;r ' 4 4 66r{ c \ rS rSr% S\S' Srg) LineContexti r contextr Nr r r r r r r r{ rW rF rF s
Lr{ rF c * \ rS rSr% S\S' S\S' Srg)ValueWithLineMapi r r z list[tuple[int, LineContext]]line_mapr NrH r r{ rW rJ rJ s
J++r{ rJ c \ rS rSrSrSSS jjr\R SS j5 rSS jr S S jr
S S jr S S jr S S jr
S S
jrS S jrS S jr S S
jrS!S"S jjrS!S#S jjrS!S#S jjr S$ S%S jjrS&S jrS S jrS'S jrS(S jrSrg))IndentedBufferi c / U l Xl g r )_lines_indent)r initial_indents rW __init__IndentedBuffer.__init__ s GI % r{ c # \ # U R n Xl S v X l g ! X l f = f7fr )tabwidth)r rV prevs rW set_tabwidthIndentedBuffer.set_tabwidth s% }} ! $M MDMs
,
! ,),c [ 5 nSn/ nU R H n[ U[ 5 ( a
U" 5 nUc M$ O5[ U[ 5 ( a UR X$R 45 MX Un[ U[ 5 ( d eUR U5 UR S5 USUR S5 - -
nM [ UR 5 U5 $ )Nr7 rR ) r
rP r} DeferredLineBaserF r rG r writecountrJ getvalue)r bufr linemaplilines rW getvaluewithlinemap"IndentedBuffer.getvaluewithlinemap s j
13++B" .//t< B ,,::/ dC(( (( IIdO IIdO
TZZ%% %A 88r{ c 6 U R 5 R $ r )rc r r s rW r^ IndentedBuffer.getvalue s '')///r{ c [ 5 nU R H n[ U[ 5 ( a
U" 5 nUc M$ O[ U[ 5 ( a M< Un[ U[
5 ( d eUR
S5 ( a UR US S 5 M UR U5 UR S5 M UR 5 $ )N\r rR ) r
rP r} r[ rF r endswithr\ r^ )r r_ ra rb s rW getrawvalueIndentedBuffer.getrawvalue s j++B" .//t< B ,, dC(( ((}}T"" $s)$ $ $ ||~ r{ c 8 U R R 5 g r )rP clearrf s rW rn IndentedBuffer.clear s r{ c , [ U R 5 $ r )r: rP rf s rW __bool__IndentedBuffer.__bool__ s DKK r{ c : SU R U R - - $ )Nr2 )rQ rV rf s rW r IndentedBuffer.prefix s dllT]]233r{ c & U R S5 g )NrR writelinerf s rW newlineIndentedBuffer.newline s t r{ c [ U[ 5 ( a U R R U5 g [ U[ 5 ( a9 U R R UR U R
5 5 5 g UR 5 ( a. U R R U R
5 U 35 g U R R S5 g Nr )r} rF rP r r[ with_prefixr stripr rb s rW rw IndentedBuffer.writeline s dK ( ( KK t $
.
/
/ KK t / /
> ?
ZZ\\ KK $++- 7 8 KK r "r{ c 8 U H nU R U5 M g r rv )r linesrb s rW
writelinesIndentedBuffer.writelines" s D NN4 r{ c L ^ ^ [ R SUU 4S jj5 nU" 5 $ )Nc 3 ># T=R T -
sl S v T=R T - sl g ! T=R T - sl f = f7fr rQ )offsetr s rW r "IndentedBuffer.indent..ctx) s8 LLF "L
' & & s A4 AA Ar Iterator[None])
contextlibcontextmanager)r r r s `` rW indentIndentedBuffer.indent( s$ " " '
# ' u r{ c . U =R U-
sl g r r r r s rW do_indentIndentedBuffer.do_indent3 r{ c . U =R U- sl g r r r s rW do_unindentIndentedBuffer.do_unindent6 r r{ c [ U[ 5 ( a [ S5 nUR HR n[ U[ 5 ( a M U( d M# [ U[
U5 [
UR 5 5 -
5 nMT [ R " U5 ( a SnUR HV n[ U[ 5 ( a U R R U5 M5 [ R X[ U5 S 5 MX g [ R " U5 nU( a UR 5 nU( d g UR 5 nUR! S5 H nU R U5 M g )Ninfr rR )r} rM floatrP rF minrR r mathisinfr rw r textwrapdedentrstripr` )r
other_coder} r rb r s rW spliceIndentedBuffer.splice9 s j. 1 15\F"))!$ 44 TS 5G)G HF * zz&!!")) dK00KK&&t,",,TF
3FG * " 4J ' . . 0
#**,J %%d+q! ,r{ c [ U R S9nU R Vs/ s H
o1" U5 PM snUl U$ s snf N)rR )rM rQ rP )r r r rb s rW r IndentedBuffer.mapS s8 DLL9-1[[9[Td4j[9
:s =c @ [ U 5 SU R 5 S3$ )Nr r )r r^ rf s rW __repr__IndentedBuffer.__repr__X s t* Qt}} /q11r{ c U R UR :X d e[ U R S9nUR U R 5 UR UR 5 U$ r )rQ rM r rP )r otherr s rW __add__IndentedBuffer.__add__[ sK ||u}},,, DLL9 t{{# u||$
r{ c XR ; $ r )rP )r new_lines rW containsIndentedBuffer.containsc s ;;&&r{ )rQ rP rV Nr )rR r r r )rV r r r )r rJ r r r r r r: )rb z)Union[LineContext, DeferredLineBase, str]r r )r z3Sequence[Union[LineContext, DeferredLineBase, str]]r r r )r r r 'contextlib.AbstractContextManager[None])r r r r F)r zUnion[IndentedBuffer, str]r} r: r r )r zCallable[[Any], Any]r rM )r r r rM )r z)Union[DeferredLineBase, LineContext, str]r r: )r r r r rV rS r r rX rc r^ rk rn rq r rx rw r r r r r r r r r r r r{ rW rM rM s H& ! !9(0 ( !4 #!H!
! EJ"4"=A"
"4
2'r{ rM c 6 ^ \ rS rSrSU 4S jjrSS jrSrU =r$ )FakeIndentedBufferig c " > [ TU ] 5 g r )superrS )r __class__s rW rS FakeIndentedBuffer.__init__h s
r{ c V US:X a [ R X5 $ [ SU S35 e)Nr zTried to call self.z on FakeIndentedBuffer. This bufferis currently used on TritonTemplateKernel to prevent actualwrites to the body without explicitly specifying the body with`TritonTemplateKernel.set_subgraph_body(name)`)object__getattribute__r )r r s rW r #FakeIndentedBuffer.__getattribute__k s9 ; **46 6!$ (=
=
r{ r r )r r r r )r r r r rS r r
__classcell__r s @rW r r g s
r{ r c # # [ R [ R p S v Xs[ l [ l g ! Xs[ l [ l f = f7fr )r stdoutstderr)initial_stdoutinitial_stderrs rW restore_stdout_stderrr v s9 %(ZZN@
!/
CJ
CJs A> AAAc h \ rS rSrSrS
S jrSS jrSS jrSS jrSS jr SS jr
SS jr SS
jr S r
g )r[ i z.A line that can be 'unwritten' at a later timec > UR 5 ( d SnXl g r{ )r} rb r~ s rW rS DeferredLineBase.__init__ s zz||D r{ c [ e)zJReturns either self.line or None to indicate the line has been 'unwritten'r rf s rW r DeferredLineBase.__call__ !!r{ c [ e)z3Returns a new deferred line with the same conditionr r~ s rW _new_lineDeferredLineBase._new_line r r{ c @ U R U U R 35 $ r r rb )r r s rW r| DeferredLineBase.with_prefix s ~~ 455r{ c T U R U R R 5 5 $ r )r rb r rf s rW r DeferredLineBase.lstrip s ~~dii . . 011r{ c > U R U R U 5 $ r r )r r s rW r+ DeferredLineBase.__getitem__ s ~~dii .//r{ c , [ U R 5 $ r )r: rb rf s rW rq DeferredLineBase.__bool__ s DII r{ c , [ U R 5 $ r )rR rb rf s rW __len__DeferredLineBase.__len__ s 499~ r{ )rb N)rb r )r zUnion[str, None])rb r r r )r r r r )r r )r zUnion[int, slice]r r r r r )r r r r r rS r r r| r r+ rq r r r r{ rW r[ r[ s- 8
""620 r{ r[ c D ^ \ rS rSrSrSU 4S jjrSS jrS S jrSrU =r $ )
DelayReplaceLinei z6At end of codegen call `line.replace(key, value_fn())`c < > [ TU ] U5 Xl X l g r )r rS rc value_fn)r rc r rb r s rW rS DelayReplaceLine.__init__ s
r{ c j U R R U R U R 5 5 $ r )rb replacerc r rf s rW r DelayReplaceLine.__call__ s# yy 4==?;;r{ c D [ U R U R U5 $ r )r rc r r~ s rW r DelayReplaceLine._new_line s $-->>r{ )rc r )rc r r zCallable[[], str]rb r r )rb r r r )
r r r r r rS r r r r r s @rW r r s @!
<? ?r{ r c [ U [ R 5 ( a U nO [ R " [ 5 U 5 n[ R
" U5 n[ R R ( aF UR c eUR S: d UR S:X a [ R S5 ggUR S:X a SOSnUR nXC: a [ R S X4S
.S 9 gg) N
z6GPU arch does not support max_autotune_gemm mode usageFTrJ rr D z,Not enough SMs to use max_autotune_gemm mode)min_sms avail_sms)extra)
r} rP r rX r createversionrq majorr r r multi_processor_count)index_or_devicer propr r s rW
is_big_gpur s /5<<00 lno> " "6 *D
}}zz%%% ::>TZZ2 - KKP QKK5(bbG**I : %>
r{ c [ R R 5 ( a( [ R R 5 R $ [ R
R S5 R $ )NrH )rP rJ rQ get_device_propertiesgpu_subslice_countrH r r r{ rW get_max_num_smsr sI yy yy..0CCC :: + +F 3 I IIr{ c [ R R 5 ( d g[ R R [ R R 5 5 n U R
S:H $ )zEReturns true if the device is a NVIDIA B200, otherwise returns false.Fr )rP rH rQ r r r )device_propertiess rW
using_b200r sM :: " " $ $
889R9R9TU " "b ((r{ c [ R R 5 ( a
[ 5 $ [ R R 5 n [ 5 U b U -
$ S-
$ )zFHandle experimental carveout if set otherwise return hardware SM countr )rP rJ rQ r r
_get_sm_carveout_experimental)carveouts rW get_num_smsr sM
yy xx557H H,@ HHa HHr{ c SSK JnJn Uc
[ 5 nUR S5 nX - [
- nU" UUUUR " 5 S9$ )zKBuilds and returns a WorkspaceArg for the device side TMA workspace buffer.r7 )r8 WorkspaceZeroModeF)r] zero_moder
outer_name)codegen.commonr8 r r from_boolTMA_DESCRIPTOR_SIZE unique_name)num_tma_descriptorsr num_programsr8 r r r s rW get_tma_workspace_argr
sU @"} !++E2I -0C CD ++- r{ c U R U; a! [ R SU R U5 [ U R R
5 =( a+ U R U; =( a [
U R 5 $ )NzDNot using template since dtype %s is not in allowed layout dtypes %s)r r r is_gpur r r )r< allowed_layout_dtypess rW _use_template_for_gpur sf ||00 R LL !
v}}!!" & LL 1 1 & v}} %r{ c U R 5 [ R R 5 R S5 Vs/ s H oR 5 PM sn; $ s snf NrQ )r ri max_autotune_gemm_backendsr` r} backendrT s rW _use_autotune_backendr P ==?!<<BBDJJ3O Oa O A c U R 5 [ R R 5 R S5 Vs/ s H oR 5 PM sn; $ s snf r )r ri max_autotune_conv_backendsr` r} r s rW _use_conv_autotune_backendr r r ) enable_int32
enable_float8check_max_autotunec SSK JnJn [ R [ R
[ R /nU( a> [ R [ R
[ R [ R /nU( a/ UR [ R [ R /5 [ U R R 5 =( a [ X5 =( d/ U R R S:H =( a U R U; =( ak [ R" =( d [ R$ =( d U( + =( a/ [' S5 =( a U" U R UR( 5 $ )Nr7 )BackendFeaturehas_backend_featurer TRITON)r r! r" rP r rJ rL rT extendrD rE r r r r r ri max_autotunemax_autotune_gemmr TRITON_TEMPLATES)r< r r r r! r"
layout_dtypess rW use_triton_templater) s D]]ENNEMMBM
u{{S
e 1 153D3D EF v}}))* A)&@
O
""e+M
0M
P
VF$<$<
VDV@V
P
"( +
P
~/N/N O r{
output_layout
add_guardsc ^^^^^^ SSK Jn SSKJm S U4S jjmS UU4S jjnS
UUU 4S jjm SUUU4S jjm SU4S jjm U" 5 =( a$ [ U4S
jU 5 5 =( a U" U 5 $ )u.
Return True iff *all* supplied tensors satisfy the CUDA TMA constraints
that Triton relies on today.
* https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html
A tensor is accepted when:
* 1 ≤ rank ≤ 5 (cuTensorMapEncodeTiled)
* dtype in _TMA_SUPPORTED_DTYPES (CUtensorMapDataType enum)
* Base pointer 16-byte aligned
* Exactly one contiguous ("inner") dim with stride 1
* All "outer" dims have 16-byte aligned strides
* Inner dim size × itemsize is a multiple of 16
* For 1-byte dtypes (e.g. FP8), inner dim ≥ 32
r )has_triton_tma_devicer7 rr c X > TR R R U [ 5 $ r )rv rw statically_known_multiple_of
TMA_ALIGNMENT)
expr_bytesrs s rW _aligned can_use_tma.._alignedG s ww < U c gU R nU R nU R nT" U R 5 ( d gT" XU5 $ )NTF)r rF r r )r< sizesstridesr r3 _is_tma_compatibles rW _is_tma_compatible_layout.can_use_tma.._is_tma_compatible_layoutJ sG > --
&&!%%88r{ c > U R 5 nU R 5 nU R 5 nU R 5 TR R
; a gU R
5 =nb UR S:X a T" XU5 $ T" XU5 $ )NFrJ )get_size
get_stride get_dtyper rv unaligned_buffers
get_devicer )r r6 r7 r m_devicers r8 _is_tma_compatible_xpus rW _is_tma_compatible_matrix.can_use_tma.._is_tma_compatible_matrixW sw
,,.
::<17744 4 &H 3
8N)%%@ @!%%88r{ c > [ U 5 nUR nUS: d US: a gU[ ; a gT( aK T
R R R U 5 nT
R R R U5 nOjU Vs/ s H( nT
R R R
U5 PM* nnU Vs/ s H( nT
R R R
U5 PM* nn[ U5 V Vs/ s H4 u pT
R R R US5 ( d M2 U PM6 n
n n[ U
5 S:w a gU
S n [ U5 H u pX:X a M T" X- 5 ( a M g X[ n T" X- 5 ( d gUS:X a, T
R R R U S5 ( d ggs snf s snf s snn f )Nr7 r Fr T)
rR itemsizert rv rw
guard_int_seq
symbolic_hintr statically_known_equalsstatically_known_geq)r6 r7 r rankrG sizes_i strides_ir str r inner_idx inner_dimrs r3 r, s rW r8 'can_use_tma.._is_tma_compatiblee s
5z>> !8tax - - gg&&44U;G((66w?IBGH%Qqww''55a8%GHFMNg))77;gIN
#9-
-ww 77A>
-
u:?!H y)EA~BM** * & ,-- q=!1!1!F!FyRT!U!U; IN
s /G US nTR R R U5 nTR R R US5 ( d gSnU HT nTR R R U5 nTR R R Xu5 ( d MT g g)Nr r7 Fl T)rv rw rI rJ statically_known_gt) r6 r7 r last_stridelast_stride_hint
MAX_UINT32r size_hintrs s rW rB +can_use_tma.._is_tma_compatible_xpu s bk 77++99+Fww 778H!LL
D((66t# U H
nT" U5 v M g 7fr r )r r rC s rW r can_use_tma.. s ?h)!,,h )r2 Union[int, sympy.Expr]r r: )r< Optional[Layout]r r: )r r@ r r: )r6 Sequence[sympy.Expr]r7 zSequence[_IntLike]r torch.dtyper r: )torch.utils._tritonr. ru rs r )
r+ r, matricesr. r9 rs r3 r8 rC rB s
` @@@@@rW can_use_tmarc 2 s " : X 9 9 9 9 /#/#/ /
/ /b##
* 5 ?h? ? 5 %m 4r{ )r, c [ R R ( a U OS n[ S U 5 5 =( a, [ X#US.6=( a [ R R
$ )Nc 3 Z # U H! n[ UR 5 5 S :H v M# g7f)r N)rR r< )r r s rW r *use_triton_tma_template.. s 5HqC
"Hs )+r* )ri r
enable_template_tma_storer rc enable_persistent_tma_matmul)r+ r, rb r< s rW use_triton_tma_templateri sK %mm E E]4F 5H 55 7 J O 7 MM 6 6r{ c f [ X US.6( d gSSKJn SSKJn U" 5 =( a U" 5 $ )Nr* Fr )%has_triton_tensor_descriptor_host_tmar7 is_datacenter_blackwell_arch)ri ra rk codegen.cuda.cuda_envrm )r+ r, rb rk rm s rW !use_triton_blackwell_tma_templatero s5 # : IC 1 2 U7S7UUr{ c X; =( a X; $ r r )scale_option_ascale_option_b
scaling_typess rW use_triton_scaling_templatert s
* N~/NNr{ )maxsizec f [ R R S5 SL$ ! [ a gf = f)zCheck if CuTeDSL is importable; cache the result for reuse.
Call ensure_cute_available.cache_clear() after installing CuTeDSL
in the same interpreter to retry the import.
cutlassNF importlibutil find_specr r r{ rW ensure_cute_availabler| s3 ~~'' 2$>> #
00c f [ R R S5 SL$ ! [ a gf = f)zCheck if NVIDIA Universal GEMM (cutlass_api) is importable; cache the result for reuse.
Call ensure_nv_universal_gemm_available.cache_clear() after installing cutlass_api
in the same interpreter to retry the import.
cutlass_apiNFrx r r{ rW "ensure_nv_universal_gemm_availabler s3 ~~''
6dBB r} c f [ R R S5 SL$ ! [ a gf = f)a3 Check if nvMatmulHeuristics is importable; cache the result for reuse.
nvMatmulHeuristics provides performance model-based kernel selection
for NVIDIA GEMM operations.
Call ensure_nvmatmul_heuristics_available.cache_clear() after installing
nvMatmulHeuristics in the same interpreter to retry the import.
nvMatmulHeuristicsNFrx r r{ rW $ensure_nvmatmul_heuristics_availabler s4 ~~''(<=TII r} c [ 5 ( d g[ S5 ( d gSSKJn [ UR
R 5 ( d gU" 5 ( d g[ R /n [ X)5 ( d g[ R ( d [ R ( d g[ XUS9( d g[ S X4 5 5 ( a gU( a U( a gUc gUc Ub gg)a
Returns True if we can use the blackwell kernel for grouped mm.
Required conditions:
1. CuTeDSL backend is enabled
2. CuTeDSL is available
3. We are on a blackwell arch
4. The dtype is bf16
5. Max autotune or max autotune gemm is enabled
6. A, B, and the output are 16B aligned
7. We are not using dynamic shapes
8. A is 2d
9. B is 3d
10. Offsets are provided
11. Bias and Scale are not provided
FCUTEDSLr7 rl )r+ c 3 8 # U H n[ U5 v M g 7fr )
is_dynamicr rT s rW r 3use_blackwell_cutedsl_grouped_mm..0 s
1.Q:a==.r T)r| r rn rm r r r rP rJ r ri r% r& rc r )
mat_amat_br< a_is_2db_is_2doffsbias scale_resultrm r( s
rW use_blackwell_cutedsl_grouped_mmr s 2 ! " " + +C &--$$ % % ' ) )^^$M 7 7 6#;#; u6 :
15.
111 g | <3 r{ c r SSK Jn UR R R X- U- SS9nUS:: d U[
R R : a gSSKJ n [ R R ( a g[ R [ R [ R /n[! X5 =( a9 [
R" =( d [
R$ =( a [' S5 nU( a; U" 5 ( d/ [( R+ S [
R R, 5 gU$ )
Nr7 rr r fallbackr F)try_import_cutlassCUTLASSzFailed to import CUTLASS lib. Please check whether _inductor.config.cutlass.cutlass_dir %s is set correctly. Skipping CUTLASS backend for now.)ru rs rv rw optimization_hintri rw cutlass_backend_min_gemm_sizecodegen.cutlass.utilsr rP r rq r rJ rT r r% r& r r r cutlass_dir) r< r r( r rs gemm_sizer r( r s rW use_cutlass_templater ? s 221519r2JIA~V^^%Q%QQ9
}} ]]ENNEKK@M f4 -
# U H
nT" U5 v M g 7fr r )r dimr s rW r 1use_nv_universal_gemm_template.. s
C]c $S ) )]r\ c 3 p ># U H+ oR 5 TR R ; v M- g 7fr )r rv r? )r trs s rW r r s&
O>N::<17744 4>Ns 36T)r r r| r r ru rs aot_compilationr r rP r rq ri r% r& r r ) r< r r( r r r r r
dims_to_checktensors_to_checkrs r s @@rW use_nv_universal_gemm_templater ` s < P " " - / / * *
}}V#u}}'8'8 6#;#;
1IM} Q
C]
CCC ~ %
O>N
OOO r{ c [ R R R 5 nUS:X a gU R 5 UR S5 Vs/ s H o"R 5 PM sn; $ s snf )z8Check if CUTLASS should be used for the given operation.ALLTrQ )ri rw cutlass_enabled_opsr r` r} )op_name enabled_opsrT s rW _use_cutlass_for_opr sY ..44::
3;
AI ! # $ $Q '
RZ1_ % %a (
! !! $ " ,, < 5FF#8>IK ''=s ( F5F<Fc T [ R R U 5 R $ r )rP rH r gcnArchNamer s rW _rocm_native_device_arch_namer s :: + +F 3 ? ??r{ c SS K n SSKJnJn SSKJn [ R R U R 5 nXAX#4$ ! [ a SS jnSS jn " S S5 nS n N&f = f) Nr )gen_ops_librarygen_ops_preselected)CKGemmOperationc / $ r r r r{ rW r *try_import_ck_lib..gen_ops_library/ Ir{ c / $ r r r r{ rW r .try_import_ck_lib..gen_ops_preselected2 r r{ c \ rS rSrSrg)*try_import_ck_lib..CKGemmOperationi5 r N)r r r r r r r{ rW r r 5 s r{ r )r r ) ck4inductor(ck4inductor.universal_gemm.gen_instancesr r ck4inductor.universal_gemm.opr r r dirname__file__r )r r r r package_dirnames rW try_import_ck_libr sh
''//+*>*>? -@ QQ s ;A A$#A$c J [ R ( d [ R ( d g[ R R
( d gU R R S:w a g[ U R 5 n[ R R Vs0 s H o"R S5 S U_M sn=( d UR S5 S U0nUR 5 [ R R - Vs/ s H nX2 PM nnU( d gU R [ R [ R [ R" 4; a g[% 5 u n nU( d [& R) S5 gU[ R l gs snf s snf )NFrH :r z,Please pip install Composable Kernel packageT)ri r% r& rP r rq r r r r archr` rY ck_supported_archr r rJ rL r r r ck_dir)r< native_archr requested_archsrequested_supported_archsck_package_dirnamer s rW use_ck_templater < s@ 6#;#; ==
}}V# 0
>K39;;3C3CD3Cawws|A)3CD # q!;IO
!%%'&++*G*GG!GA G ! %
||EMM5>>5==II"3"5 1a BC+FKK + E!s FF c SSK Jn [ S5 =( a> [ U 5 =( a, UR R
R
X- U- SS9S: $ )Nr7 rr CKr r r ru rs r r rv rw r r< r r( r rs s rW use_ck_gemm_templater a sP d# K F # K
GG . .quqy2 . F Jr{ c SSK Jn [ S5 =( a> [ U 5 =( a, UR R
R
X- U- SS9S: $ )Nr7 rr CKTILEr r r r r s rW use_ck_tile_gemm_templater k sP h' K F # K
GG . .quqy2 . F Jr{ c < [ S5 =( a [ U 5 $ )Nr )r r r< s rW use_ck_conv_templater u s %d + G0GGr{ c [ R =( d [ R =( a U R R S:H $ r )ri r% r& r r r s rW _use_template_for_cpur y s2 7v 7 7 &
--
%&r{ c SSK Jn [ UR U5 ( d eUR R nUR R
n[
U 5 =( al UR 5 [ R :H =( aD [ U5 S:H =( a/ [ U5 S:H =( a US US :H =( a US S:H n[ XUSS9=( a# UR R 5 =( d U$ )Nr7 )rA r F)require_constant_mat2)
r rA r} r< r rF r r> rP rL rR use_cpp_gemm_template
is_contiguous)r< mat1mat2rA mat1_size mat1_stride mat1_each_batch_is_contiguouss rW use_cpp_bmm_templater s dkk6 * ** *
I++$$K f% " NN
- "
^q
"
" " ^y|
+ "
^q
" !t5 Q !!#D'Dr{ c
SSK Jn SSKJn SSKJn SSKJn
[ U 5 ( a [ S5 ( d g[ R R ( d gUR 5 [ R [ R 4; n [ R" [ R$ [ R& [ R [ R /n U
" UUU ( a U R( OS UUS9u ppp[+ X45 ( a g[- X'R. 5 ( a UR1 5 nU " UR 5 5 u nnU" S U
UUUR 5 UR 5 U[3 5 U( + US
9
nS S jnU R( U ; =( aT US L=( aI U" U5 =( a: [- X'R4 5 =( a UR7 5 =( d U( + $ )
Nr7 r
)create_micro_gemm)*get_gemm_template_output_and_compute_dtype)mm_argsCPPF) out_dtypemat2_transposed use_4x2_dim
micro_gemm) input_dtype input2_dtype output_dtype num_threadsuse_ref q_group_sizec N U R 5 U R 5 S S:H $ )Nr r7 )
freeze_layoutr= rT s rW is_last_dim_stride12use_cpp_gemm_template..is_last_dim_stride1 s" ||~b!Q&&r{ )rT r@ r r: ) r r codegen.cpp_micro_gemmr
codegen.cpp_utilsr kernel.mm_commonr r r ri cppweight_prepackr> rP r[ rP rL rJ halfr has_free_symbolsr} BaseView unwrap_viewparallel_num_threadsr7 is_module_buffer)r< r r r r is_woq_int4r r r
r r int8_gemmr( r r( r r r r r s rW r r s 9M) ( (0Ee0L0L :: $ $ U[[%**$==I]]ENNEJJ UZZXM") "+&,,' # A!T $ $$ !@ AQROL!" NN$^^%!(* ! J'
% C d " C % C
t]] + C
"
"
$
A,A(A r{ c ~ [ R =( d [ R ( + =( d [ S5 $ )NATEN)ri r% r& r r r{ rW use_aten_gemm_kernelsr, s- 7v 7 7 ' v &'r{ c b \ rS rSr% \R
" S5 rS\S' S
S jrS
S jr S S jr
Sr g ) DebugDirManageri r r prev_debug_namec @ [ [ R 5 U l g r )r r. counterr rf s rW rS DebugDirManager.__init__ s ../r{ c [ R R R U l U R SU R
3U l U R [ R R l g )N_tmp_)rP _dynamori debug_dir_rootr/ r new_namerf s rW __enter__DebugDirManager.__enter__ sM $}} 3 3 B B //0dggY?
.2mm
+r{ c [ R " U R 5 U R [ R
R l g r )r r r7 r/ rP r5 ri r6 )r r s rW __exit__DebugDirManager.__exit__ s*
dmm$.2.B.B
+r{ )r r7 r/ Nr )r r r r ) r r r r r} r] r1 r rS r8 r; r r r{ rW r. r. s& ooa G0<
Cr{ r. c ^ SSK Jn [ 5 mSU4S jjn[ R R USU5 [ R R 5 U " U0 UD6nS S S 5 W[ T5 4$ ! , ( d f N= f)Nr7 r< c ( > TR U 5 g r )r[ code source_codess rW save_output_code*run_and_get_code..save_output_code s r{ rB r@ r r r )
rv r= r# r patchr rP r5 resetrb )r r r r= rB r rA s @rW run_and_get_coderG so
%$.LL
=*<>N O
T$V$
P 4 % %%
P Os 'A77
Bc UR SS5 n[ U /UQ70 UD6u pE/ nU HU nUR [ R " SU[ R
5 5 U( d MA U Vs/ s H oSS PM nnMW XF4$ s snf )N remove_quoteFz '''.*?'''r )rS rG r$ r findallDOTALL) r r r rI r rA kernelsr@ r s rW run_and_get_kernelsrN
s ::ne4L+B@@@FG rzz,bii@A <29:'a|'G:G ? ;s -Bc * ^ SU 4S jjn[ U5 $ )Nc R > T" 5 n U R 5 R 5 U $ r )r r )r r s rW run_with_backward1run_fw_bw_and_get_code..run_with_backward
s!
r{ )r r )rG )r rQ s ` rW run_fw_bw_and_get_coderS
s
- ..r{ c t ^^ SSK Jn / mSU4S jjmS U4S jjn[ R R USU5 [ R R UST5 [
R R 5 U " U0 UD6nSSS5 SSS5 T$ ! , ( d f N= f! , ( d f T$ = f)
zLGet the inductor-generated code, but skip any actual compilation or running.r7 r< c ( > TR U 5 g r r r? s rW rB "get_code..save_output_code
s D!r{ c > " S S5 nU R ( a U R 5 OU R 5 u p#T" UR 5 U( a T" UR 5 U" 5 $ )Nc , \ rS rSrSrSS jrSS jrSrg) @get_code..patched_compile_to_module..DummyModulei"
z4This is empty to replace the generated triton modulec g r r rf s rW rS Iget_code..patched_compile_to_module..DummyModule.__init__%
s r{ c g r r r s rW callEget_code..patched_compile_to_module..DummyModule.call(
s r{ r Nr r r r r r r )r r r r r rS r^ r r r{ rW DummyModulerZ "
s F
r{ ra )r codegen_with_cpp_wrappercodegenr )r ra wrapper_code kernel_coderB s rW patched_compile_to_module+get_code..patched_compile_to_module!
s[ 04/?/?D ) ) +T\\^ " ++, [ . . /} r{ compile_to_modulerB NrD )r r= r r )rv r= r rE r rP r5 rF )r r r r= rf r rB rA s @@rW get_coderi
s $ L" ,
.0I
-);=MN
O
ON
s# "B( 'BB(
B% !B((
B7c [ U /UQ70 UD6nS[ U5 s=:: a S:: d O S[ U5 35 eUS $ Nr7 r z%expected one or two code outputs got r )ri rR )r r r rA s rW get_triton_coderl C
sQ B000L L! &Q &
/L0A/BC & ?r{ c [ U /UQ70 UD6u p4S[ U5 s=:: a S:: d O S[ U5 35 eUS $ rk )rG rR )r r r r rA s rW run_and_get_triton_codern M
sU 'r;D;F;OA L! &Q &
/L0A/BC & ?r{ c ^^^ SSK Jm SSKJn UR m/ mSUUU4S jjn[
R R USU5 U " U0 UD6nS S S 5 UT4$ ! , ( d f WT4$ = f)Nr r< rD c h > T" U 0 UD6 U S n[ UT5 ( d eTR U5 g )Nr )r} r )r r rv r= graph_lowerings real_inits rW fake_init-run_and_get_graph_lowering..fake_initb
s: 4"6"Q% //// u%r{ rS r` )torch._inductor.graphr= torch._inductor.output_coderE rS r rE r ) r r r rE rs r r= rq rr s @@@rW run_and_get_graph_loweringrw Y
sv 4; ((IO& &
?J BT$V$
C ? ""
C B ? ""s A
A/c # # SSK Jn UR U n [ R " X5 UR U ' Sv X2R U ' g! X2R U ' f = f7f)zs
Override the lowering of aten_op with override_fn.
The first argument of override_fn is the original lowering fn.
r )loweringN)torch._inductorry loweringsr partial)aten_op override_fnry orig_fns rW override_loweringr n
sY ) )G.&/&7&7 &M7#
&-7#g7#s A"'A A"A A"c ^ ^^ SSK Jn UR mSUUU 4S jjn[ R R
R
USU5 $ )zf
Add hook functions to be called at the beginning and end of Scheduler.__init__.
Used for unit tests.
r ) Schedulerc F > T" X5 T" X5 nT( a T" X5 U$ r r )r rU outr post_fnpre_fns rW r (add_scheduler_init_hook..wrapper
s% y i' I %
r{ rS )r r rU r r r )torch._inductor.schedulerr rS unittestr rE r )r r r r r s `` @rW add_scheduler_init_hookr
s> 4 G == % %iW EEr{ c [ R ( a [ R U 5 g[ R U 5 g)z
Warnings that will be actionable for PyTorch developers, but not
end users. Allows us to easily disable them in stable releases but
keep them on for nightly builds.
N)ri developer_warningsr r info)msgs rW developer_warningr
s$ C
r{ c [ R R S5 n U S- [ [ R 5 : aV [ [ R U S- 5 S: a3 [ R U S- S S:w a [ R U S- $ [ R H) nUR S5 ( d M U[ S5 S s $ g! [ a NJf = f)a
An experimental API used only when config.benchmark_kernel is true.
The benchmark name is only available at codegen time. So we can not
directly call it in benchmark_all_kernels which is run after codegen.
The function assumes the argument after --only is the benchmark name.
It works for torchbench.py/hugginface.py/timm_models.py. But for ad-hoc
scripts, this function may return None.
There are 2 flavors of --only argument we need handle:
1. --only model_name
2. --only=model_name
z--onlyr7 r r z--only=N)r argvr rR
ValueErrorr )rb r s rW get_benchmark_namer
s
hhnnX& !Gc#((m #CHHS1W%&*q!!$+88C!G$ $ xx >>) $ $s9~'( (
s BC
C"!C"c & [ S U 5 5 $ )Nc 3 * # U H oS :H v M g7fr7 Nr r s rW r is_ones..
%u!Avu r rZ s rW is_onesr
%u% %%r{ c & [ S U 5 5 $ )Nc 3 * # U H oS :H v M g7f)r Nr r s rW r is_zeros..
r r r r s rW is_zerosr
r r{ c & [ S U 5 5 $ )Nc 3 # U HI n[ U[ R 5 ( d M$ UR [ R " S 5 :H v MK g7f)r N)r} rP r r )r r s rW r is_cpu_device..
s9 D dELL ) + u||E**s
#A*Ar )inputss rW
is_cpu_devicer
s r{ c [ U [ R 5 ( d S5 eU R ( a [ R
$ [ R $ )Nz8only support sympy.Expr as input to get_sympy_Expr_dtype)r} r~ r4 r rP rV rN rn s rW get_sympy_Expr_dtyper
s@ c5:: & & B & ~~{{}} r{ c / # U ( a. [ R R " U0 UD6 nUv S S S 5 g S v g ! , ( d f g = f7fr )rP r r )should_profiler r r s rW
maybe_profiler
s;
^^
#
#T
4V
4G 5
4 5
4s (A=A
A Ac p [ R R n U S: a [ R " 5 n U $ Nr7 )ri r threadsrP get_num_threads)r s rW r& r&
s+ jj G{'') Nr{ c SSK Jn U " 5 nUR S[ R R
( a S5 $ S5 $ )Nr7 )get_backend_options
num_stagesr r )runtime.triton_helpersr rW rP r rq )r optionss rW get_backend_num_stagesr
s2 ;!#G ;;|%--*;*;Q CC CCr{ c N [ U [ R R R R
S:H S9nUb U$ SSKJnJn [ R R 5 =( a! [ R R 5 S: nU [ R [ R [ R 4; d e[ R " U5 R R# S5 ( a SSKJn U" 5 nU [ R [ R 4; a U( a U" X5 $ [ R R R R
S:X a U" [ R U5 $ U" [ R U5 $ U [ R [ R 4; a U( a U" U 5 $ [ R R R R
S:X a U" [ R 5 $ U" [ R 5 $ )z
We don't want to throw errors in this function. First check to see if the device is in device_info.py,
then fall back to the inaccurate triton estimation.
tf32)is_tf32r )get_max_simd_tflopsget_max_tensorcore_tflops)rv r
clock_rate)max_clock_rate)r rP backendsrH matmulfp32_precisiontriton.testingr r rQ get_device_capabilityr rJ rL inspect signature
parametersrW torch._utils_internalr )r ds_topsr r SM80OrLaterr sm_clocks rW get_device_tflopsr
s
u~~**11@@FJG M**))+
0P0P0R W 1K
U]]ENNEMMB BB B,-88<<\JJ8!# U]]ENN3 3 ,U= = >> % % 4 4 >,U]]HE E&u}}h? ? U]]ENN3 3 ,U3 3 >> % % 4 4 >,U]]; ;&u}}5 5r{ c SSK Jn U " 5 $ )Nr
get_dram_gbps)r r r s rW get_gpu_dram_gbpsr s
, ?r{ c x SSK Jn U R R R S5 R SS5 $ )Nr r max_shared_mem)triton.runtimer r r r rW r s rW get_gpu_shared_memoryr & s. % == 4 4Q 7 ; ;.N s) Jcz!UZZ(<_ r# maybe_get_sizemaybe_get_strider@ TypeErrorr )r r r s rW r r Q s
bmmR[[":K:KRYYW
0 0 2 8b99=M""$*> > > Ayy)) =d1gY GH H r{ c \ rS rSrSrSrSrg) Placeholderie KERNEL_NAMEDESCRIPTIVE_NAMEr N)r r r r r r r r r{ rW r r e s K *r{ r c v SSK Jn [ R " SSS9 n[ R
" 5 n[ R
" 5 n[
U[ U5 S9R " U6 [ SUR 3US9 [ UR US9 [ R " 5 n[ X5 U " UR 5 S S S 5 [ R " 5 U-
n U" UR 5 UR R 5 UR 5 [ S UR 3US9 [ UR US9 UR! 5 UR! 5 :H n
[" R% S
UUR& U
U 5 S S S 5 g ! , ( d f N= f! , ( d f g = f) Nr7 )stable_topological_sortrl zutf-8)modeencoding)r fake_modezBefore:
)filezAfter:
zZ%s, save before/after graph to %s, graph before/after are the same = %s, time elapsed = %s)pattern_matcherr r NamedTemporaryFileior
r^ rZ propagater rv r nowr] lint recompiler^ r r r ) r r inpr r r& before_ioafter_io
start_time time_elapsedr s rW pass_execution_and_saver o sE 9 $ $
KKM ;;=R#3C#89CCSI
"(($1-
bhhY' \\^
#B
, N - ||~
2 )
#!,
bhhX& H$5$5$7 7 h
FF
+
-
,
s% BF*2FC F*
F' #F**
F8c SSK Jn [ XR 5 =( a [ U R UR
5 $ )z:
Check if input buffer is a multi-outputs template buffer
r7 r
)r r r} CppTemplateBufferr< MultiOutputLayout input_bufr s rW is_multi_outputs_templater s7 i!5!5 6 :"..< r{ c SSK Jn [ XR 5 =( a7 [ U R
5 S:H =( a [
U R
S 5 $ )zD
Check if input buffer is a output of multi-outputs template buffer
r7 r
r )r r r} MultiOutputrR r r r s rW #is_output_of_multi_outputs_templater sJ 9nn- ; !Q & ; %i&6&6q&9 :r{ c U c gSSK Jn [ XR 5 =( a: [ XR 5 ( + =( a US L =( d U R
UL =( Gd_ [
U 5 UR L =( Ga@ [ [ R R S5 =( a; U R
[ R R R R :H =( d [ [ R R S5 =( a; U R
[ R R R R :H =( df [ [ R R S5 =( a; U R
[ R R R R :H $ )NFr7 r
all_to_all_singleall_gather_into_tensorreduce_scatter_tensor)r r r} _CollectiveKernel _WaitKernel op_overloadr FallbackKernelr rP r torchrecr defaultr r rl r r s rW
is_collectiver sM | 4--. 34 0 0 3
4Z
14++r1
T
b'''
**,?@ U$$ (:(:(L(L(T(TT
**,DE E$$99%%<<DDE
**,CD Y$$ (:(:(P(P(X(XX/r{ c < SSK Jn [ U 5 UR L $ Nr7 r
)r r r r )rl r s rW is_waitr s : ''r{ c SSK Jn [ X5 ( a [ S U R 5 5 $ [ U R 5 =( a US L =( d U" U 5 $ )Nr GroupedSchedulerNodec 3 8 # U H n[ U5 v M g 7fr )contains_collectiver s rW r &contains_collective.. s @. s :\=##\r )r r r} r r" r rl )r# r s rW r' r' s4 >%..:U\\:::uzz""r{ c SSK Jn [ U[ R R
5 ( a U/n[ XR 5 =( a U R U; $ r )r r r} rP r r r r r s rW is_fallback_opr* sF "ejj++,,T d - - . I43C3Cr3IIr{ c @ X!U R R 5 $ r ) defining_opr )buf_name name_to_bufname_to_fused_nodes rW buf_name_to_fused_snoder0 s! ( 3 ? ? H H J KKr{ c gr r r# s rW r r ur{ c U" U 5 ( a g UR U 5 U R H- n[ UR X#5 nXa; a M [ UUUUUS9 M/ g )N criteria_cb)r[ unmet_dependenciesr0 r find_recursive_deps_of_node)r# collected_node_setr. r/ r6 depdefining_op_for_deps rW r8 r8 sf 55!'' 5 HHk
4 # #
(r{ c gr r r2 s rW r r r3 r{ c
U" U 5 ( a g UR U 5 U R 5 H nUR H nUR c eUR R 5 S:X a M2 UR R 5 U; a MR X6R R 5 nXq; a Mu [ UUUUUS9 M M g )NOUTPUTr5 )r[ get_outputsr rl r find_recursive_users_of_node)r# r9 r. r/ r6 ors user_ops rW r@ r@ s 55!
GGD99( ((yy!!#x/yy!!#+==(););)=>G, ("" '
!r{ c j [ R R R ( a SOSnX-
U-
$ )zaComputes the number of inputs to the aot fw graph which have fixed addresses (params and buffers)r r )rP
_functorchri functionalize_rng_ops)dynamo_gm_num_inputsaot_fw_gm_num_inputsnum_rng_seed_offset_inputss rW num_fw_fixed_argumentsrI 4 s3
$
$
:
: 69S SSr{ c SS jnSn/ nU R R H8 nUR S:X d M U" U5 ( a UR U5 US-
nM: U[ [ [
U5 5 5 :X d e[
U5 $ )z6
Infers which inputs are static for a backwards graph
c SU R ; =( a; SU R ; =( a% SU R ; =( a SU R ; $ )Ntangentsbwd_seedbwd_base_offset
bwd_rng_stater r s rW is_saved_tensor'count_tangents..is_saved_tensorD sH aff $
.!&&(
.!/
. qvv-
r{ r r r7 )rT r5 r r: )rv rU r r rb r rR )fx_grP arg_countstatic_arg_idxsr( s rW count_tangentsrU ? s
IO
ZZ
44= q!! &&y1 NI d5_)=#> ? ?? ? r{ c > \ rS rSr% S\S' SS jr\S S j5 rSrg)
BoxedBooliX r: r c U R $ r )r rf s rW rq BoxedBool.__bool__\ s zzr{ c @ [ U [ 5 ( a SU l U $ gr )r} rW r r s rW disableBoxedBool.disable_ s c9 % % CIJr{ r Nr )r r r zUnion[BoxedBool, bool]) r r r r r rq r r[ r r r{ rW rW rW X s K r{ rW c # ^ ^# SSK Jn UR m S SU U4S jjjn[ R R USU5 S v S S S 5 g ! , ( d f g = f7f)Nr7 r9 c : > TR U5 T" XX#XE5 $ r rV )r kernel_namere re gpucpp_definition kernel_listorig_define_kernels rW
define_kernel.collect_defined_kernels..define_kernelm s' ;'! {c
r{ rd )NTN)r r: r_ r re r re
Optional[str]r` r: ra rf r r )codegen.wrapperr: rd r rE r )rb r: rd rc s ` @rW collect_defined_kernelsrh g s 5-;; #'(,
"
&
/ - P
Q P Ps AA2A! A2!
A/+A2c U S- $ )N __original__r r s rW get_cloned_parameter_buffer_namerk ~ s
. r{ c U [ ; $ r )rN r s rW r r s
Y r{ c : [ R R SL$ )z,Check if we're running on ROCm/HIP platform.N)rP r rq r r{ rW is_rocmrn s == D ((r{ c 0 U S:g =( a [ U 5 $ )NrI )r r s rW device_need_guardrp s U? -vf~-r{ c h U [ R :X aD [ R R 5 ( a! [ R R 5 S: $ U [ R :X a$ [ R
R 5 ( a gU [ R [ R 4; $ )N)r r T)rP rJ rH rQ r rJ rV r: rB s rW ,needs_fallback_due_to_atomic_add_limitationsrr sq 5::#:#:#<#<zz//1F:: %.. UYY%;%;%=%=ejj111r{ c U R [ R R R [ R R R
4; a Uc gU R [ R R R :X a SOSnUS U4; =( Gd& U=( a [ U5 =( a [ U5 =( d U R [ R R R :H =( ap US:H =( ad U=( a[ US:H =( aO [ R R =( a. [ R R =( d
[ 5 S:g =( dJ X:H =( a# U[ R [ R 4; =( d [ R " 5 $ )NFr[ r r r7 )overloadpacketrP r atenscatter_reduce_scatter_reducescatter_r rr ri r fallback_scatter_reduce_sumdynamic_threadsr& r: rV $are_deterministic_algorithms_enabled)r r
self_dtype src_dtypesrc_device_type
src_is_tensor reduce_tys rW use_scatter_fallbackr s] "" IINN * *EIINN,I,I J K " ++uyy~~/F/FFE
tY // 8 8
H'
H&,,#? ??RVWW r{ c [ U R R 5 ( d g[ R =( d [ U 5 $ r )r r r ri assume_aligned_inputsr )
example_inputs rW should_assume_input_alignedr s5 -&&++ , , ' ' K+<]+KKr{ c X [ R R R 5 n U ( d [ R
" 5 $ U R ( a U R R ( d [ R
" 5 $ U R R nUR 5 $ r ) rP _guardsTracingContexttry_getr nullcontextr rx suppress_guards)tracing_contextrx s rW #maybe_get_suppress_shape_guards_ctxr sv
mm22:: [ U TT5 u pT" U 5 n[ U5 ( a [ R " X5 U$ r )copy_misaligned_inputsrR rP _foreach_copy_)
new_inputs old_tensors new_tensorsr inputs_to_checkr mutated_input_idxss rW r )align_inputs_from_check_idxs..runE
sD #9 );$
J { :
r{ )r list[InputType]r r )rR )r r r r s ``` rW align_inputs_from_check_idxsr =
s(
?q Jr{ c X SU R 5 ; a SnO;[ S [ U R 5 U R 5 5 5 5 S- n[ R
" X4S5 R
5 n[ R
" X R 5 U R 5 5 $ )Nr c 3 6 # U H u pUS -
U- v M g7fr r )r rA rF s rW r )clone_preserve_strides..[
s T:S f$:Ss r7 r )r r r rF rP
as_stridedclone)rT needed_sizer: s rW clone_preserve_stridesr U
s AFFH}
T#affh
:ST TWX X
a
6
<
<
>F FFFHahhj 99r{ c T / n/ nUSLnU H nX n[ U[ R 5 ( d S[ U5 35 eUR 5 [
- ( d MW [
U5 X' U( d Mm Xb; d Mt UR U5 UR X 5 M X44$ )z
Clones misaligned tensors which we inferred were aligned. Returns a tuple of [old_tensors], [new_tensors] for every
cloned tensor which is in `return_pair_idxs`.
Nz Expected tensors only, but got: )r} rP r r data_ptr ALIGNMENTr r )r check_inputs_idxsreturn_pair_idxsr r ret_pair_definedr _inps rW r r a
s ')K&(K (t3
}$ --
.tDzl ;
- ==?Y & & 24 8JM A$9""4("":=1 ##r{ c / nU HV nX n[ U[ R 5 ( d M( UR 5 [ - S:X d ME UR U5 MX [
U5 [
U5 :w a U$ U$ )zO
We require all inputs to be aligned, so introduce a copy for any
that aren't.
r )r} rP r r r r rR )r static_input_idxsaligned_static_input_idxsrb r s rW remove_unaligned_input_idxsr
sp !# eU\\ * *0@90LQR/R % , ,S 1 ! $%->)??(( r{ c SSK Jn [ R " [ R 5 R
nUR R R nUR R R R n[ R ( a& UR R R X5 gUR R R X:* 5 ( a gUR ( a. UR R R U S: 5 ( a gU" U 5 =( a U" U 5 U:* $ )Nr7 rr Tg@x DF)ru rs rP iinforT r rv rw rX rx has_hintri assume_32bit_indexing check_leqr r )r rs int_maxrX r s rW expr_fits_within_32bitr
s kk%++&**G **Iww ))22H
## ""1. ww--al;;
77 1 1!d( ; ; A; 29Q<722r{ c 6 ^^^ [ R R R 5 nUb UR b [ UR 5 S:X d e[
U 5 mUR c eUR H nUc UR R S 5 M# Sm[ R R R 5 =n( a UR mSUU4S jjmUR R [ U4S jU 5 5 5 M g g g )Nr Fc r > Tc [ U 5 $ T( a TR U 5 $ TR U 5 $ r )r deserialize_symexprevaluate_symexpr)r fakify_first_callrx s rW map_expr4set_tracing_context_output_strides..map_expr
s7 ( "1v
( ( < # U H
nT" U5 v M g 7fr r )r r r s rW r 5set_tracing_context_output_strides..
s 5u!(1++ur\ )r r r z,Union[float, int, SymInt, SymFloat, SymBool])
rP r r r output_stridesrR r r r r )r compiled_graphrG r? r r r rx s @@@rW "set_tracing_context_output_stridesr
s mm**224Gw 5 5 A7))*a///).9 ,,888#22E}&&--d3$)! --66>>@@3@(+(=(=%9 9 &&--5u55 3 Br{ c 4 [ R b [ R $ [ R " 5 ( d g[ R R 5 ( a g SSKJn U [ R R S5 : $ ! [ a gf = f)NFr REMOTE_CACHE_VERSIONz.pytorch/remote_cache:fx_graph_memcache_version)
ri fx_graph_remote_cache is_fbcoderP _utils_internalis_fb_unit_test torch._inductor.fb.remote_cacher ModuleNotFoundErrorjustknobs_getval_intr s rW should_use_remote_fx_graph_cacher
s
##/+++ ,,..H 5#8#8#M#M8$ s "B
BBc 2 [ R " SSU 5 $ )Nz
[^a-zA-Z0-9_]r )r subr s rW normalize_namer
s
66"C ..r{ ztl.int1z
tl.float8e4nvz tl.float8e5z
tl.float8e4b8ztl.float8e5b16ztl.uint8)ztl.boolztl.float8_e4m3fnztl.float8_e5m2ztl.float8_e4m3fnuzztl.float8_e5m2fnuzztl.float8_e8m0fnuztl.float4_e2m1fn_x2z^.*[.]c j [ R S[ U 5 5 n[ R X5 $ )z"Convert torch.dtype to triton typetl.)_triton_type_rer r _triton_type_mappingrW )r triton_type_names rW triton_typer
s+ &**5#e*= # #$4 GGr{ c [ R X 5 nUR SS5 n[ [ U5 n[ U[ R 5 ( d eU$ )Nr r )_torch_triton_mappingrW r rO rP r} r )r
adjusted_type type_namer s rW triton_type_to_torchr sM )--e;M %%eR0Iy)I i - -- - r{ c U R ( + =( a U R 5 UR 5 :H =( a U R 5 UR 5 :H =( a U R UR :H =( a U R UR :H =( ae U R 5 R
5 UR 5 R
5 :H =( a! U R 5 UR 5 :H $ r ) is_mkldnnr rF r r untyped_storager r r6 r s rW is_same_tensorr s NN < IIK5::< ' < KKMU\\^ + <
JJ%++ % <
KK5<< ' <
" + + -1F1F1H1Q1Q1S S <
!U%9%9%; ;r{ c U R =( a U R 5 UR 5 :H =( a U R UR :H =( as U R UR :H =( aS [ R
R R U 5 [ R
R R U5 :H $ r )r r r r rP r mkldnnr r s rW is_same_mkldnn_tensorr s P IIK5::< ' P JJ%++ % P
KK5<< ' P
II % %d +uyy/?/?/H/H/O O r{ c g)N)r isnan logical_not logical_andsignbitand_leltgegteqner
xorr r r{ rW boolean_opsr " s r{ c * \ rS rSr% S\S' S\S' Srg) OpDtypeRulei6 r3 type_promotion_kindOptional[torch.dtype]override_return_dtyper NrH r r{ rW r r 6 s 8800r{ r zdict[str, OpDtypeRule]op_dtype_propagation_rulesc ( [ X5 [ U ' g r )r r )r r r s rW #register_op_dtype_propagation_rulesr ? s
(3( t$r{ zOrderedSet[str]op_requires_libdevice_fp64c . [ R U 5 g r )r r[ r s rW #register_op_requires_libdevice_fp64r L s ""4(r{ c SSK Jn U ( d$ UR R 5 R n U S:X a [
R $ U S:X a gU S:X a [
R $ [
R $ )Nr rr r rI rJ ) r rs rv get_current_device_or_throwr ri cpu_backend xpu_backend cuda_backend)r rs s rW get_current_backendr P s_ - gg99;@@ e!!! !!!"""r{ c U [ R [ R 4; a= [ R R
( a [
5 S:X a [ R $ U $ )z"Maybe upcast [b]float16 to float32r
)rP r rJ ri r
codegen_upcast_to_fp32r rL rB s rW upcast_compute_typer" _ s@ %-- 00 MM 0 0 !X -}} Lr{ KeyTypeValTypec v \ rS rSrSrSS jrSS jrSS jrSS jrSSS jjr SS jr
SS
jr SS jr SS jr
S
rg)
ScopedDictin z
A dictionary-like object that allows for scoped updates. It maintains
an original dictionary and a set of new items that can override
the original items within the scope. The original dictionary is
unmodified.
c Xl 0 U l g r
original_dict new_items)r r) s rW rS ScopedDict.__init__v s *13r{ c \ XR ; a U R U $ U R U $ r r* r) r s rW r+ ScopedDict.__getitem__z s, .. >>#& &!!#&&r{ c X R U' g r )r* )r rc r s rW __setitem__ScopedDict.__setitem__ s #sr{ c H XR ; =( d XR ; $ r r- r s rW __contains__ScopedDict.__contains__ s nn$A/A/A(AAr{ Nc t XR ; a U R U $ U R R X5 $ r )r* r) rW )r rc r s rW rW ScopedDict.get s2 .. >>#& &!!%%c33r{ c [ U R 5 nU R H nX R ; d M US-
nM U$ r )rR r) r* )r r( r s rW r ScopedDict.__len__ s< "" # A***Q r{ c # # U R S h vN U R H nXR ; d M Uv M g N-7fr r( )r r s rW __iter__ScopedDict.__iter__ s8 %%%% A*** &s AA A
Ac R [ U R =( d U R 5 $ r )r: r) r* rf s rW rq ScopedDict.__bool__ s D&&8$..99r{ c [ er r r s rW __delitem__ScopedDict.__delitem__ s !!r{ r- )r) Mapping[KeyType, ValType])rc r# r r$ )rc r# r r$ r r )rc r r r: r )rc r# r Optional[ValType]r rB r )r zIterator[KeyType]r )rc r# r r )r r r r r rS r+ r0 r3 rW r r: rq r? r r r{ rW r& r& n s5 4'
$B4
:"r{ r& )frozen_defaultc . ^ SU4S jjnU c U$ U" U 5 $ )Nc 0 > [ R " U STS9$ )NT)kw_onlyr ) dataclasses dataclass)r r s rW wrapir_dataclass..wrap s $$S$vFFr{ )r rl r rl r )r r rI s ` rW ir_dataclassrK s G { 9r{ c [ R R R 5 n U b' U R ( a U R R
$ g r )rP r r r fw_metadatabw_donated_idxs)r s rW get_donated_idxsrO s= mm22::>+R 0D
5') ww||L)UCH8<2::??H#34(D Kr{ c N SSK Jn [ X5 =( a U R S L$ )Nr ) SubgraphPythonWrapperCodegen) torch._inductor.codegen.wrapperr r} partition_signatures)r r s rW #is_codegen_graph_partition_subgraphr Q s' L 79 5 ( ( 4r{ c [ R R R R =( d [
R S L=( a$ [ R R R $ r )rP r ri r
cudagraphs&_unstable_customized_partition_wrapperr graph_partitionr r{ rW is_using_cudagraph_partitionr Z sN
%%00 F 1 9 9 E 1 //
0
01r{ c SSK Jn UR R R U S5 ( a; UR R R U S5 ( a [ R $ [ R $ )Nr7 rr l i ) ru rs rv rw statically_known_ltrK rP rT rV )r rs s rW dtype_from_sizer a sX ww++ e
''
/
/h
?
?{{{{r{ )r rJ c n U S:X a( [ R R R 5 $ SU ; a gg)z3
Returns True if the device supports MKL-DNN BF16.
r rJ TF)rP r r _is_mkldnn_bf16_supportedr s rW is_mkldnn_bf16_supportedr o 3 eyy 99;; + r{ c n U S:X a( [ R R R 5 $ SU ; a gg)z3
Returns True if the device supports MKL-DNN FP16.
r rJ TF)rP r r _is_mkldnn_fp16_supportedr s rW is_mkldnn_fp16_supportedr { r r{ c
x U Vs/ s H n[ [ U5 5 PM nnU HS n[ U5 [ U5 :X d e[ U5 H' u pR[ X5 [ [ U5 5 5 X5' M) MU / nUR SR S [
X5 5 5 5 [ U5 [ U5 S- - [ U5 S-
- nUR SU- 5 U H3 nUR SR S [
XC5 5 5 5 M5 SR U5 $ s snf )N|c 3 4 # U H u pS X S 3v M g7fr2 Nr )r hrl s rW r tabulate_2d.. s H3G41AaWA,3G r r7 r c 3 4 # U H u pS X S 3v M g7fr r )r r rl s rW r r s H7Gtq!Cl7Gr rR )rR r r r r r r r )elementsheadersr widthsrowr r total_widths rW tabulate_2dr s #*
+7ac#a&k7F
+3x3w<''' cNDAFIs3q6{3FI #
E LLH3w3GHHIf+Vq 1S[1_EK LL{"#
SXX Hs37G HHI 99U ,s D7c # # [ U R 5 5 [ UR 5 5 - nU H6 nU R U5 nUR U5 nUUb UOUUb UOU4v M8 g7f)a
Zip two dictionaries together, replacing missing keys with default values.
Args:
dict1 (dict): The first dictionary.
dict2 (dict): The second dictionary.
d1_default (Any): the default value for the first dictionary
d2_default (Any): the default value for the second dictionary
Yields:
tuple: A tuple containing the key, the value from dict1 (or d1_default if missing),
and the value from dict2 (or d2_default if missing).
N)r# rY rW )dict1dict2
d1_default
d2_defaultall_keysrc value1value2s rW zip_dictsr sp ( %**,'*UZZ\*BBH 3 3
(Fj (Fj
s A1A3c v SS jn SS jnU R S[ R R 5 nU R 5 n U( aq U" U SS5 U" U SS5 U" U S[
R R ( + 5 U" U SS 5 U" U S
[ R R 5 U" U S S 5 U R S
[ R R 5 nU R S[ R R 5 nUS:X a U( a [ S5 eU $ )a
Ensures the configuration is internally consistent for standalone AOTInductor.
If `aot_inductor_mode.compile_standalone` is set to True in the provided
`config_patches` (or falls back to the global config), this function ensures
that the following configs are also enabled:
- `aot_inductor.package_cpp_only`
Args:
config_patches (dict[str, Any]): A dictionary of user-provided config
overrides for AOTInductor compilation.
Returns:
dict[str, Any]: The possibly-updated `config_patches` dictionary.
c U R U[ [ U5 5 nUc X U' g U( d X2:w a [ SU SU S35 eg g )NzInvalid config: =z3 when aot_inductor_mode.compile_standalone is True.)rW rO ri r config_patches config_name config_valuer s rW patch_config2maybe_aoti_standalone_config..patch_config sY ""; 0LM =*6; '50 ";-q >qr
1r{ c U R U[ [ U5 5 nX2:w a [ R SUU5 X U' g )NzDOverriding: %s=%s when aot_inductor_mode.compile_standalone is True.)rW rO ri r r r s rW force_patch_config8maybe_aoti_standalone_config..force_patch_config sB ""; 0LM KKV
'3{#r{ z$aot_inductor_mode.compile_standalonez aot_inductor.package_cpp_onlyTz aot_inductor.embed_kernel_binaryz#aot_inductor.emit_multi_arch_kernelz+aot_inductor.model_name_for_generated_files
aoti_modelzaot_inductor.link_libtorchz aot_inductor.dynamic_linkageFz"aot_inductor.cross_target_platformz$aot_inductor.package_constants_in_sowindowszconfig.aot_inductor.package_constants_in_so is not supported for windows cross-compilation. Please use config.aot_inductor.package_constants_on_disk_format = binary_blob.)r dict[str, Any]r r r r r r )rW ri aot_inductor_modecompile_standalonecopyrP r rq test_configs use_libtorch aot_inductorcross_target_platformpackage_constants_in_sor )r r r r r r s rW maybe_aoti_standalone_configr sk " & 58 HK
3&
358
3HK
3
3 (++. 33
$((*N^%DdK^%GN Au}}GXGXCX
I<
( , ,
>+I5Q * . .,11
- 0 0.33
).E
]
r{ c [ R R ( a) [ R R S:X a [ S5 e[ R R ( a0 [ R R
S:X a [ S5 eSnSnX!4$ [ R R S:X a SnSnX!4$ U S:: a gSn[ R " 5 ( + nX!4$ ) z
Decide whether we should mmap weights, and whether to store the weights with .so.
If force_mmap_weights or package_constants_on_disk_format == "binary_blob" configs are set, respect the config.
Returns tuple (use_external_weights, use_mmap_weights).
binary_blobzconfig.aot_inductor.package_constants_on_disk_format = binary_blob and config.aot_inductor.force_mmap_weights cannot both be True.r zKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi 5w)FF)ri r force_mmap_weights package_constants_on_disk_formatr r r ) consts_sizeuse_mmap_weightsuse_external_weightss rW determine_aoti_mmap_flagsr s .. @ @M Q
J
-- 4 4 A ]
$ #55
;;}L # #55m# !++-- 11r{ c SSK Jn U R R nUc g[ U[
5 ( d [
S5 eUS:X a g[ R " SU5 ( d [
S5 eg)zD
Validates if a model name is suitable for use in code generation.
r rh Tz4Invalid AOTI model name: Model name must be a stringr z^[a-zA-Z_][a-zA-Z0-9_]*$zVInvalid AOTI model name: Model name can only contain letters, numbers, and underscores) rz ri r model_name_for_generated_filesr} r r r r )ri
model_names rW is_valid_aoti_model_namer 8 sn
'$$CCJ j# & &OPPR 88/ < < d
r{ c < U( a [ U 5 $ [ U 5 $ r )r) r( )rT
unbacked_onlys rW get_free_symbolsr S s $Q''A r{ c * 0 [ R ES[ R R S[ R R [
R 5 5 0En [ R " 5 ( a [ R " S5 U S' U $ )z9
Get a base environment for running Python subprocesses.
PYTHONPATHTORCH_CUSTOM_PYTHONPATHr6
PYTHONHOME) r r rW r| r r r ri r sysconfigget_path)envs rW python_subprocess_envr Z so
** bjjnn %rzzsxx'@
C % . .v 6L Jr{ c . \ rS rSr% SrS\S' S\S' Srg)CUDAGraphWrapperMetadataiu z
Metadata for Customized CUDAGraphWrapper.
Currently assumes there is 1 dynamo graph and will extend to
multiple graphs in the future.
r num_partitionspartition_indexr Nr r r{ rW r r u s r{ r .c $ \ rS rSr% SrS\S' Srg)CUDAGraphWrapperi Nz Optional[CUDAGraphWrapperType]r r )r r r r r r r r r{ rW r r s
.2G
+2r{ r c U [ l g r )r r )r s rW !set_customized_partition_wrappersr s 5<*2r{ c H ^ U R R nU R R / UQU R R QU R R 5 nU R R n[
R " X45 u p4SS jnU Vs/ s H: nU" U5 ( a( [ R R R USS9OUPM< nnSS jmSU4S jjnU Vs/ s H
og" U5 PM nn[
R " X45 u pX4$ s snf s snf ) Nc [ U [ R R R 5 =( a3 [ U [ R R R
5 ( + $ r )r} rP r r r@ GeneratorStater s rW
_is_tensor_ir(snode_args_kwargs.._is_tensor_ir sH !U__ / / 6 67
u!!00A
=
r{ F) guard_shapec , [ R " XUS9$ )Nr )rP r )r r r s rW _tensor"snode_args_kwargs.._tensor s {{4V< [ U [ R 5 ( d U $ T" U R 5 U R U R
5 nU$ r )r} rP r r r r )r r r s rW to_real_tensor)snode_args_kwargs..to_real_tensor s: !U\\**Haffh2
r{ r )r r )r r r r ) rl r fill_non_provided_args
constant_argsr pytreer$ rP r r ir_node_to_tensortree_unflatten) r# r r flat_argsflat_args_pytree_specr r( r r s @rW snode_args_kwargsr s :: D :: , ,*$*))*
D ZZ
F'-':':D>'J$I
A ,,QE,B
= -66Iq "II6((JLD <% 7s
AD,D c SSK Jn U R nUR R ( a( UR UR R S- 5 nUR S5 $ )Nr7 rr r )primals_r
fwd_rng_staterO rL )ru rs r rv removeprefixr )r: rs dep_names rW is_nonfreeable_buffersr sN xxH ww||(( );< I r{ c x [ X S3- 5 nUR 5 sSSS5 $ ! , ( d f g= f)z,Load a template file and return its content.z .py.jinjaN)openread)r template_dirr& s rW
load_templater s+
lvY // 0Avvx
1 0 0s +
9c U R n[ U[ R R [ R R
45 ( d S[
U5 35 e[ R ( d g[ [ R R R R [ R R R R /5 nX; a g[ [ R R R /5 n[ U[ R R
5 ( a X; $ [# U 5 ( + $ )zLDecide whether fallback for a node. This is only used in inductor lite mode.z6Expected OpOverload or HigherOrderOperator, but found F)r r} rP r r r r ri fallback_by_defaultr# r ru _assert_scalarr lift_fresh_copy higher_order triton_kernel_wrapper_functionalr! )rl r "skip_fallback_due_to_dynamic_shape
fallback_hopss rW should_fallback_by_defaultr
s
[[F &&
(F(FG O ?V ~NO % % *4 IINN ) ) 1 1 IINN * * 2 2
*& 3 @ @AM &%**8899&&&t, ,,r{ ) z-torch.ops._c10d_functional.all_reduce.defaultz.torch.ops._c10d_functional.all_reduce_.defaultz9torch.ops._c10d_functional.all_gather_into_tensor.defaultz8torch.ops._c10d_functional.reduce_scatter_tensor.defaultz4torch.ops._c10d_functional.all_to_all_single.defaultz6torch.ops._c10d_functional_autograd.all_reduce.defaultzBtorch.ops._c10d_functional_autograd.all_gather_into_tensor.defaultzAtorch.ops._c10d_functional_autograd.reduce_scatter_tensor.defaultz=torch.ops._c10d_functional_autograd.all_to_all_single.defaultc U [ ; $ )z0Check if an operation is a collective operation.)COLLECTIVE_OPS)r s rW is_collective_opr
s
n $$r{ c p [ R " 5 ( a SSKJn U $ / $ ! [ a / s $ f = f)Nr tlx_only_cuda_options)ri r )torch._inductor.fb.tlx_templates.registryr r r s rW r r s<
W( ( I s & 55c X- S-
U- U- $ )z(Round x up to the nearest multiple of y.r7 r )rT ys rW _round_upr " s
UQY1 !!r{ c SSK JnJn U" US5 ( a UR UR 4$ [ U5 S: Ga2 U" US U S 5 ( a U" US S5 ( d% U" US S5 ( a, U" US U S 5 ( a UR UR 4$ U" US U S 5 ( a U" US [ U S S5 5 ( d2 U" US U S 5 ( a6 U" US [ U S S5 5 ( a UR UR 4$ U" US [ U S S5 5 ( a6 U" US [ U S S5 5 ( a UR UR 4$ U[ R :X a SOSnU[ R :X a U[ R :X a [ U S S5 [ [ XS - S5 S5 - n [ U S S5 [ [ XS - S5 S5 - n
U" X)5 ( d
U" X*5 ( a UR UR 4$ U[ R :X Ga
[ R" R$ ( d [ U S S5 [ [ XS - S5 S5 - n [ U S S5 [ [ XS - S5 S5 - n
U" X)5 ( d
U" X*5 ( a UR& UR 4$ g [ U S S5 U- U S - n [ XS - S5 U S - n
U" X)5 ( d
U" X*5 ( a UR& UR 4$ g )
z2
Core implementation for scale/swizzle inference.
r )r6 SwizzleTyper7 r rs rr rN rF NN)torch.nn.functionalr6 r
TensorWise
NO_SWIZZLErR RowWiserj BlockWise1x128BlockWise128x128rP rH rD r
BlockWise1x16SWIZZLE_32_4_4rF r rq
BlockWise1x32) mat_size
scale_size scale_numel mat_dtype scale_dtypeeq_fnr6 r K_multiplierexpected_numel_aexpected_numel_bs rW _infer_scale_swizzle_implr* ' sg =
[! %%{'='=== :! *Q-! - -%
1
q2I2I *Q- # #jmXa[(I(I && (>(>> >
*Q-! - -jmWXa[#%>?? *Q-! - -jmWXa[#%>?? --{/E/EE E A S 9 : :u qM78A;4@
@
//1G1GG G "U%;%;;1L E***{e>Q>Q/Q$Xa[#6 LA;. 3Q:
%Xa[#6 LA;. 3Q:
/ /53W3W ,,k.H.HH H e***}} (!c :Y {2B7> )!c :Y {2B7> [33u8 8 #00+2L2LLL8 'x{B 7, FRS T &|qk'A2 FRS T [33u8 8 #00+2H2HHH r{ c [ U R S U R S 4[ UR 5 UR 5 U R UR S S9$ )aR
Infer the scaling type and swizzle mode from matrix and scale tensor shapes/dtypes.
This function determines how scale factors are laid out relative to the matrix:
- TensorWise: Single scale for entire tensor
- RowWise: One scale per row
- BlockWise1x128/128x128: Block-scaled with float32 scales
- BlockWise1x32: MXFP8 with float8_e8m0fnu scales (swizzled on NVIDIA)
- BlockWise1x16: NVFP4 with float8_e4m3fn scales (swizzled)
Args:
mat: The matrix tensor (FP8 or FP4)
scale: The scale factor tensor
Returns:
Tuple of (ScalingType, SwizzleType) or (None, None) if unrecognized
r r7 c
X:H $ r r r s rW r %infer_scale_swizzle.. s 16r{ r! r" r# r$ r% r& )r* rA r numelr )matscales rW infer_scale_swizzler2 w sO ( %))A, ! -%KKM))KK!
r{ c r ^ SSK Jm U R 5 nUR 5 nU( a
US US 4nU( a& [ R " [
R US5 OSnSU4S jjn[ [ U5 S: a
US US 4OUS S4[ U5 UU R UR US9$ )z
Infer the scaling type and swizzle mode for IR nodes (used during graph lowering).
This is the IR-compatible version of infer_scale_swizzle, using symbolic
size comparisons via V.graph.sizevars.statically_known_equals.
r rr r7 c N > TR R R X5 $ )z5Compare values using symbolic equality when possible.)rv rw rJ )r( r) rs s rW symbolic_eq+infer_scale_swizzle_ir..symbolic_eq s ww 77==r{ r r. )r( r r) r r r: ) r rs r< r r r r! r* rR r r )r0 r1 transposer! r" r# r5 rs s @rW infer_scale_swizzle_irr8 s .||~H !J QK!- DN)""8<<Q?STK> %/28}/A(1+x{+QR UVGW$ ))KK
r{ r )ry r r r )r r r r: ) d )r Callable[[], Any]r r r r r r )r9 r: F)
r r; r r r r r r: r r r )r z"Union[Optional[torch.device], str]r torch.device)r$ zIterable[sympy.Expr]r r )r, r_ r- r_ r r )r$ z Iterable[_T]r zValuesView[_T])r7 r] r8 r] r r] )rc r r r )ri z"Iterable[Union[int, torch.SymInt]]r zlist[sympy.Expr])ro Union[int, torch.SymInt]r r] )r r] r r= )ri z Iterable[Union[int, sympy.Expr]]r z list[Union[int, torch.SymInt]])r torch._ops.OpOverloadr r: )r r5 r z'Callable[[torch._ops.OpOverload], bool]r r: )r r r r r r r z&tuple[GraphModule, list[torch.Tensor]])rH )r r r r )r7 rH )
r Callable[..., Any]r
Sequence[Any]r r r r r r )r r r g ?rH )r r? r r@ r r r r r r r r r r )r r r r r r )r r r r r r )r( r r) r r r )rT zUnion[int, Sequence[int]]r r r
Sequence[int])rT ztuple[_T, ...]r zlist[_T])r z!Callable[Concatenate[Any, P], RV]r zCachedMethod[P, RV])r r r z*Callable[[FN_TYPE[P, RV]], FN_TYPE[P, RV]])r 0Union[Sequence[BaseSchedulerNode], ExternKernel]r zOrderedSet[Node])r Sequence[BaseSchedulerNode]r z8Literal[True, 'torch', 'original_aten', 'inductor_node']r r )r rB r r: r tuple[str, str]r )rp zIterable[torch.fx.Node]rq z Optional[Callable[[Any], bool]]r OrderedSet[torch.fx.Node])r zSequence[IRNode]r zdict[str, IRNode]r rE r )r r r zValueRanges[Any])r r r r: )r re rb r r r )r r: r r: )r r r r )rm r r zdict[sympy.Expr, Any]r r )r( r r z,TypeGuard[Union[torch.SymInt, torch.Tensor]])r r r r: )r torch.fx.GraphModuler zOptional[torch.fx.Node])r rF r r5 )r rF r zOrderedSet[torch.device]r )r r r r )rc r r r r r )NNT)r" zOptional[dict[str, Any]]r rf r# r: r r )r, r@ r( r: r list[int])rx r+ r, z.Sequence[Union[int, torch.SymInt, sympy.Expr]]r( r: r rG )r r` r r r r )r zUnion[int, torch.device]r r: r )r r r r< r
Optional[int]r r8 )r< rA r zlist[torch.dtype]r r: )r r r r: )
r< rA r r: r r: r r: r r: )rb r@ r+ r^ r, r: r r: )rb r@ r+ rA r, r: r r: )rq r6 rr r6 rs zlist[ScalingType]r r: )r r r r r< rA r r: r r: r
Optional[Any]r rI r rI r r: )
r< rA r r r( r r r r r: r )r< rA r r r( r r r r r@ r r@ r zOptional[IRNode]r zOptional[_IntLike]r r: )r r r r: r )
r r r( r r r r r r r: )r r r( r r r r r: )r r r( r r r r rG )r r r r )r zQtuple[Optional[str], Callable[[], list[Any]], Callable[[], list[Any]], type[Any]])r< rA r r: )r< rA r z Union[ReinterpretView, Buffer]r r@ r r: )FTFN)r< rA r r@ r r@ r r: r r: r( r: r rH r r: )r Callable[P, _T]r r r r r ztuple[_T, list[str]])r r? r ztuple[Any, list[str]])r rJ r r r r r r )r rJ r r r r r r )r rJ r r r r r z tuple[Any, list[GraphLowering]])r} r? r~ r? r r )r r? r z Optional[Callable[..., Any]]r r )r r r r )r rf )rZ r@ r r: )r zSequence[torch.Tensor]r r: )ro r r r` )r r: r r r r r z
Iterator[Any])r r` r r )r r r r: )r r r r )r z
Iterable[Any]r r: )
r r? r r4 r r@ r r r r )r z"Optional[Union[Buffer, Operation]]r r: )rl z Optional[Union[Node, Operation]]r z!Optional[torch._ops.OperatorBase]r r: )rl z"Optional[Union[IRNode, Operation]]r r: )r# rF r$ z-Optional[Callable[[BaseSchedulerNode], bool]]r r: )r# rF r r: )rl zOptional[Operation]r z?Union[torch._ops.OpOverload, Collection[torch._ops.OpOverload]]r r: )r- r r. r r/ r r r ) r# rF r9 z MutableSet[BaseSchedulerNode]r. zdict[str, SchedulerBuffer]r/ z dict[str, BaseSchedulerNode]r6 zCallable[[Any], bool]r r )rF r rG r r r )rR rF r r )rb r r r )r r r r )r rf r r: )r r r r: )r r` r r: )r r> r rf r| r` r} r` r~ r r r: r r: )r rC r r )r r r r: )r r r r: )r r )r rJ r r r r r ztuple[_T, str])r Sequence[InputType]r zOptional[ShapeEnv])r Callable[[list[InputType]], _T]r rA r zOrderedSet[int]r rL )rT r r r )r r r rA r zOptional[OrderedSet[int]]r z-tuple[list[torch.Tensor], list[torch.Tensor]])r rK r rA r rA )r r r r: )r r@ r rE r r )r r` r r )r r r r` )r6 r r r r r: )r ztuple[str, ...])r r r r3 r r r r )r r r r )r rf r r )r r` r r` )r zOptional[type[Any]]r r: r r )r zOptional[list[int]])r rQ )r ztorch._ops.OperatorBaser rD )r
torch.fx.Noder r: )rl rB r r: )r r: r r: )r r r r` )r r r r: )r zSequence[Sequence[T]]r z Sequence[T]r r )
r rA r rA r ValType | Noner rN r zEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None])r r r r )r r r ztuple[bool, bool])rT r* r r: r zOrderedSet[sympy.Symbol])r zdict[str, str])r CUDAGraphWrapperTyper r )r# rF r z tuple[list[Any], dict[str, Any]])r: r; r r: )r r r r/ r r )rl rM r r: )r r )rT r r r r r )r! ztuple[Any, Any]r" ztuple[Any, ...]r# r r$ r` r% r` r& zCallable[[Any, Any], bool]r #tuple[Optional[Any], Optional[Any]])r0 r r1 r r rP r )r0 r> r1 r> r7 r: r rP (
__future__r rS r rG enumr ry r r r} r r r r r r r r r r r r r r collections.abcr r r r r r r
r r r
typingr r r r r r r r r r r r r typing_extensionsr r r r r~ rP torch.utils._pytreer _pytreer $torch._inductor.analysis.device_infor torch._inductor.runtime.hintsr !torch.fx.passes.regional_inductorr! torch.utils._dtype_abbrsr" torch.utils._ordered_setr# r$ r% OPTIMUS_EXCLUDE_POST_GRADr r( r) r* r+ r, r- r. pathlibr/ r0 r1 r2 torch._prims_commonr3 torch.fxr4
torch.fx.noder5 r r6 r r8 rg r: dependenciesr; rv r= r r> r? r@ rA rB rC output_coderE r rF rG rN rL r rX torch._dynamo.device_interfacerY torch._dynamo.utilsrZ torch.autogradr[ torch.autograd.profiler_utilr\ (torch.fx.passes.graph_transform_observerr] torch.fx.passes.shape_propr^ torch.utils._sympy.functionsr_ r` ra rb rc torch.utils._sympy.symbolrd re torch.utils._sympy.value_rangesrf rg r ri runtime.runtime_utilsrj r6 _IS_WINDOWS getLoggerr r rl r r4 VarRangesr r InputTypegetenvXPU_KERNEL_FORMATGPU_KERNEL_BIN_EXTSr r r1 r r[ rP r\ rR r] rT rV r rJ rL rN rD rE float8_e4m3fnuzfloat8_e5m2fnuzrt r rx rz r Functionr rH r r r r r r r% r. r1 rf rj rp rz r| r r r r r r r r r r r r r FN_TYPEr r r r r r% rn rt r r r r r r r r r r frozensetro r r r r r r r r r r r' clear_on_fresh_inductor_cacheclear_inductor_cachesfresh_inductor_cacher0 r@ rD rF rJ rM r r r[ r r r r r r
r r r r) rc ri ro rt r| r r r r r r r r r r r r r r r r r r r r, r. rG rN rS ri rl rn rw r r r r r r r r r r& r r r r r r r r r r# r Enumr r r r r r r r' r* r0 r8 r@ rI rU rW rh rk r rn rp rr r r r r r r r r r r r r r r r r rZ r compiler r r r r r r r r r r r r" r# r$ r& rK rO rQ r\ r^ rc rm r rv r r r r SUPPORTED_MKLDNN_DEVICESr r r r r r r r r r PartitionFnTyperO r r r r r r r
r r
r r r* r2 r8 )r r s 00rW r sv "
C B $ $ ? : E 0 / ; ($
>> //C$"/,5!$TT,=
+ CL
D 0 % 2 K 0 8 D = llg%
!
T] UZZ' ( U5<<ell:; < Ebii(I7 S
! "
2<
2 . ( {Q'A-+2B XDX XB5
L ENN
d# $ "G GX #(
!
4 #( __ _
_ ! _
_D ;@
+ * " *+A * * #AL+ +++ "/ )/#/G @OI
I<I
I0 *8+0' ' !
( %'
!
) ' #
$ cN TT"
;sAv&*
+E8WQU^ E: ) ++/+\ C ,4).4)O4) 4)nW2CW2
!W2 W2x 48*0 (G
G$5G G:, ^ % DU >2 - $
$ '&
< !# I " " "( .2 7+7 7
7 7 7v !5 $ " 49
( $$ 7$ $
$N Q7 7*
, , ,
S' S'l
@ @ @ ?' ? 8 J J ) )I #' (
+<
#
< :>RW||&6|KO| |@ BG&,:> BG
V
V&,
V:>
V
V O O O %O
O Q Q Q > > >
> >
>
> > >
>B P " E EE E E
E
E E E
EPJ CO,) , EF !)?B ( . 5( 5(p @ @ R R:"JH& 8@F : ""& "&= =
= = =
=
= =
=@' C C"&& & & &$ & 2: /(V & 2: ## :#$#* .
..@.. .$ IMF F)EFF* B&& D D '6 '6T Q .0(#K(*$)) *
"-
4A
HK
F 1 " -1#
*#)#
#L( @D Q Q< Q
Q#J
JGJ
JLL .LDRLL *=
5
,
5
'
< *= 5 , 5
'
:T 2
,! )
.2$&$!$ $ $
$
$
$N HB L
' &2:2: *" ( % 0 : 37$ $$$ 0$ 3 $< $ $$3N!3B :&/ '#)* $%
+?*D*D*FG*F$!*FG **Y'H &
1 1 1
68 2 7
8 1
/9l O :) # # )
)
-" 01 -"` D)t *499 4 42Y5!H"J4
1 *
& "&!%
$
$
K
FRj&2R6 6
d# $ 38$ ./ @ 3 3 *:); &= F
"-L
%
"
M M M M M
M &
M )M` *(B " " " " ) " "{ Hs $z