o j9:j @sUddlmZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl m Z m Z mZmZmZddlmZmZddlmZddlmZddlmZmZmZmZmZmZmZmZmZm Z m!Z!ddl"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(dd l)m*Z*ddl+Z+dd l+m,Z,m-Z-m.Z.ddl/m0m1m2Z3ddl4m5m6Z7ddl8Z9ddl:Z9ddl;m6mm?Z?dd l@mAZAdd lBmCZCddlDmEZEddlFmGZGddlHmIZIddlJmKZKmLZLddlMmNZNmOZOmPZPmQZQmRZRmSZSddlTmUZUmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_ddl`maZaddlbmcZcddldmeZeddlfmgZgmhZhmiZimjZjddlkmlZlddlmmnZnmoZoddlpmqZqmrZrmsZsmtZtmuZuddlomvZvmwZwmxZxmyZymzZzm{Z{ddl|m}Z}ddl~mZmZmZmZddlmZdd lmZmZdd!l6mZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZdd"lmZmZmZerdd#lTmZdd$l`mZdd%lmZdd&lmZdd'lmZdd(lmZdd)l6mZneZd*ed+<z ddlZejZd,ZWneydZd-ZYnwe&d.Ze d/Ze d0Ze d1Zee,BZd*ed2<eeBe,BZd*ed3<e9jje9jjBZd*ed4<eeZeje jd5d6Ze9jjZ e!ed7eed7fd8d9ee!eeed7fd7d8d9fdBfZd*ed:<dd?d@Zejd,dAGdBdCdCZʐddGdHZːddLdMZ̐ddQdRZ͐ddSdTZΐddZd[Zgd\Zgd]Z dddbdcZҐddedfZ dddgdhZedddkdlZe -dddndlZ -dddqdlZՐddudvZ֐ddydzZאdd|d}Zؐdd~dZِdddZڐdddZېdddZܐdddZݐdddZddZGdd9d9Zed-dAGdddZeGdddeZdddZeGdddeZeGdddeZedededededededdZded< ,dːdddZeGdddeZde-dfddd„Ze ee,ee,gefZd*ed<GddńdeZGddDŽdeZGddɄdeZeGdd˄deZeGdd̈́deZeGddτdeZdddфZdddӄZ , -  - dϐddd܄ZdddބZdddZeGdddeZeGdddeZeGdddeZeGdddeZeGdddeZeGdddeZeGdddeZeGdddeZGdddeZeGdddeZeGdddeZeGdddeZdddZdddZGdddZeGdddeZGdddeZGdddeZGdddeZGdd d eZ Gd d d eZ eGd d d eZ GdddeZ ed-dAGdddeerZ ed-dAGddde eZGddde ZGdddeZGdddeZeGdddeZeGdddeZed-dAGdddeZejd,dAGd d!d!ZGd"d#d#eZGd$d%d%eZeeBeBeBeeeBeBeBBZGd&d'd'ZGd(d)d)eZGd*d+d+eZGd,d-d-eZGd.d/d/eZGd0d1d1eZ Gd2d3d3eZ!dՐd7d8Z"ed-dAGd9d:d:eZ#Gd;d<d<e#Z$Gd=d>d>e$Z%ed-dAGd?d@d@e#Z&ed-dAGdAdBdBe&Z'GdCdDdDe'Z(GdEdFdFe&Z)GdGdHdHe Z*GdIdJdJe&Z+GdKdLdLe+Z,GdMdNdNe+Z-GdOdPdPe&Z.GdQdRdRe&Z/GdSdTdTe&Z0GdUdVdVe&Z1GdWdXdXe&Z2GdYdZdZe2Z3Gd[d\d\e)Z4Gd]d^d^e&Z5Gd_d`d`e&Z6Gdadbdbe'Z7Gdcdddde&Z8Gdedfdfe&Z9Gdgdhdhe&Z:Gdidjdje&Z;ed-dAGdkdldlZGdqdrdre=Z?eGdsdtdteZ@Gdudvdve&ZAGdwdxdxeAZBGdydzdzeAZCd֐d}d~ZDGddde=ZEejGdddeZFGdd7d7eFZGGdddeFZHed-dAGdddeZIdאddZJed-dAGddde&ZKed-dAGddde&ZLdؐddZMed-dAGddde&ZNGddde=ZOGdddeZPeGdddePZQeGdddePZReGdddePZSeGdddePZTGddde=ZUGdddeUZVGdddeUZWGdddeUZXdِddZYdِddZZdڐddZ[dS() annotationsN)Callable GeneratorIterableIteratorSequence)AbstractContextManager nullcontext)Enum)partial) AnycastClassVarLiteraloverload SupportsFloat SupportsInt TYPE_CHECKING TypeAliasTypeVarUnion) assert_neverNeveroverride ParamSpecSelfTypeIs)patch)ExprIntegerSymbol)identity)GraphModuleSerializer)can_auto_functionalize)metricsget_free_symbols)FakeScriptObject)get_opaque_obj_repris_opaque_value)compute_required_storage_lengthis_boolean_dtype(is_contiguous_for_memory_format_or_falseis_float_dtypemake_channels_last_strides_for StrideType) &_remove_effect_token_unbacked_bindingscompute_unbacked_bindings free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNodehas_free_unbacked_symbols IterateExprsrebind_unbackedresolve_unbacked_bindingsShapeEnvSymTypes)Node OrderedSet)_disable_current_modes)CleanDivFloorDivModModularIndexing)SymT)config dependencies)BackendFeature CodegenSymbolget_scheduling_for_deviceindex_prevent_reorderingKernel)Depextract_free_symbols#extract_input_node_reduction_rangesextract_read_writesSymbolUsageCollectorOpsHandler var_builder)LoopBody) OpCounterCSE OpCountResult ReductionType StoreMode) benchmarker)DeviceProperties ReductionHint)argsort argsort_sym cache_on_selfcache_on_self_and_argsceildivconvert_shape_to_inductorconvert_shape_to_symintdeveloper_warningdo_bench_using_profilingdtype_from_sizeget_dtype_sizeget_kernel_metadataGPU_ALIGN_BYTES ir_dataclass is_dynamicis_gpu sympy_dotsympy_index_symbolsympy_index_symbol_with_prefix sympy_product sympy_substensor_is_aligned)opsOpsValueV) SympyBoolean)Argument) IntLikeType)CUTLASSTemplate)PythonWrapperCodegen) GraphLowering)IndentedBufferrrvTF_P_T_U_V_IntLike_NumLike _OpOverloadsz prefix TensorBoxr IRNode _NodeOrNodesxobjectreturnTypeIs[int | Integer]cCst|ttfSN) isinstanceintrrrY/home/nk/hobo-godmode/plappi-mvp/.venv/lib/python3.10/site-packages/torch/_inductor/ir.py _is_staticr)frozenc@s>eZdZUded<ded<ded<ded<d ed <d ed <d S)GraphPartitionSignatureOrderedSet[sympy.Symbol] symbol_inputsz0dict[str, IRNode | sympy.Expr | TorchBindObject] input_nodes list[IRNode] output_nodeszdict[str, bool]input_deallocationboolskip_cudagraph list[str]constant_namesN__name__ __module__ __qualname____annotations__rrrrrs  r node_or_nodes_NodeOrNodes | NoneNonecsdfdd |dS)Nnodesrrrc s|durdSt|ttfr|D]}|qdSt|tr*|D]}|q!dSt|ttttt j j j t ttttf sFJdt|ddS)NzFound zE, which is not a supported top level IR node. See [Note: Inductor IR])rlisttupledictvalues ExpandView DynamicScalar AssertScalarrsympylogicboolalgBooleanrrEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputtype)rnode_check_tensorboxrrrs4    z%validate_ir.._check_tensorbox)rrrrr)rrrr validate_irs rnamestrCallable[..., OpsValue]cs(tts Jtdfdd }|S) Nargsrkwargsrrqcstt|i|Sr)getattrrprrrrrfnzops_wrapper..fn)rrrrrrq)rrr)rrrrr ops_wrappersrorder Sequence[int]&Callable[[Sequence[_T]], Sequence[_T]]cs(tt|tt|dfdd }|S)Nindex Sequence[_T]rc0ttks JfddttDS)Ncg|]}|qSrr.0i)r inv_orderrr z4inverse_reorder..reindex..lenrangerrrrreindexz inverse_reorder..reindexrrrr)rziprrrrrrrinverse_reordersrcsdfdd }|S)Nrrrcr)Ncrrrr)rrrrr%rz1same_reorder..reindex..rrrrrr#rzsame_reorder..reindexrrrrrr same_reorder"srreindex1&Callable[[Sequence[_U]], Sequence[_V]]reindex2&Callable[[Sequence[_T]], Sequence[_U]]&Callable[[Sequence[_T]], Sequence[_V]]csdfdd }|S)Nrrr Sequence[_V]c |Srrrrrrrr. z fuse_reindexing..reindex)rrrrr)rrrrrrfuse_reindexing*sr)rrD)rrrrDseq#Sequence[int | torch.SymInt | Expr] shape_envShapeEnv | NonecCs4|dus tdd|Drt|}|St||}|S)z1 Convert strides to fill order (argsort) Ncs |] }t|ttjfVqdSrrrrrrsrrr >z!get_fill_order..)allrZr[)rr sorted_idxrrrget_fill_order8s  rSequence[int | Integer]cs0ddt|Dfddtt|D}|S)z Convert stride order to fill order For channel last format, stride order = [3, 0, 2, 1] and fill order = [1, 3, 2, 0] cSi|]\}}||qSrrridxposrrr Mz+stride_order2fill_order..cg|]}|qSrrrlookuprrrNz+stride_order2fill_order..) enumeraterr)r fill_orderrrrstride_order2fill_orderFsrcCs>t||}ddtt|D}t|D]\}}|||<q|S)z) Convert strides to stride order cSg|]}dqSrrr_rrrrYz$get_stride_order..)rrrr)rrroutrelemrrrget_stride_orderRs  rreplace_symbols_with_hintsrcCdSrrrrrrrir_node_to_tensor_r torch.TensorcCrrrrrrrrcs IRNode | Nonetorch.Tensor | Nonecs|durdS|rtjjjntfdd|D}t|r,fdd|jD}nt |}| }| }t |}t |}tjjjtj||||d}Wd|S1s_wY|S)Ncg|]}|qSrrrshape_fnrrryrz%ir_node_to_tensor..crrrrrrrr|r)sizestridedtypedevice)rrgraphsizevarsoptimization_hintr!get_sizeis_storage_and_layout get_layoutrFlexibleLayoutcontiguous_strides get_dtype get_devicer`rsuppress_guardstorch empty_stridedzero_)rrrrrrtrrrris.   valueSequence[_T] | NoneSequence[_T | None] | NonecCst|tr |s dgS|Sr)rrr-rrrmay_convert_to_optionalsr1/IRNode | OutputSpec | torch.device | None | str str | NonecCsbt|ts |dur |St|tjr|jSt|ttfr!t|St d|dt|j ddS)Nzget_device_type(: )) rrr)rrr OutputSpecget_device_typer'rrrrrrr7s   r7"IRNode | torch.device | None | strcCstt|}|dvrtt|ddkrdSdS|dus"t|}dur$dSddlm}t|ts5Jt|t||S)N)cpucudaxpu_backendtritonTFrD)TritonScheduling) r7rrErIcodegen.tritonr>rr issubclass)rrdevice_schedulingr>rrr is_tritons  rBcCs t|dkS)Nr9)r7rrrris_cpurrCBuffer | TensorBox alignmentrcst|tr|dust|st|rdStjfdd|ddD}tt |ddt |dd}t||}t j j |S)NFc3s"|] }tt|dVqdSrN)rEqrArrErrrs z-is_aligned_realized_tensor..rD)rrmaybe_get_strider3 get_strider!rAndOrrGLerrrrguard_or_false)rrEaligned_stridesaligned_last_dim is_alignedrrHris_aligned_realized_tensors    $ rSstrides1Sequence[_IntLike]strides2shapecCsnt|t|krt|t|ksJt|||D]\}}}tjj|dr&qtjjt||s4dSqdS)zP Returns true if the strides are equal, ignoring dimensions of size 1 . rDFT) rrrrrrstatically_known_leqrOrrG)rTrVrWdims1s2rrrsignificant_strides_equals$r\tensorstridesSequence[int | torch.SymInt]cCst|s|Stddt||Dr|St|||s"|St|\}}g|j}t|D]\}}t j j |drE||||<q3t |j|j|j||j|j}tt||dS)a Tries to match the strides of the tensor to those in the meta_strides. Strides of insignificant dimensions - size 0 or 1 - will be updated. If there are real stride differences (NHWC vs NCHW), or the tensor is not realized, then the input will be returned cs$|] \}}tjj||VqdSrrrrrstatically_known_equalsrrZr[rrrr  z2try_match_insignificant_strides..rDdatalayout)r"rrrKr\r!as_storage_and_layoutrrrrrrrX FixedLayoutrrroffset is_pinnedrReinterpretView)r]r^storage old_layout new_striderr new_layoutrrrtry_match_insignificant_stridess.     rqgmtorch.fx.GraphModulecCsD|jjddd}ddt|jD|jd<ddlm}||dS)NoutputoprcSsg|]\}}|qSrr)rrr rrrr sz.gm_original_output_strides..user_visible_output_idxs)record_original_output_strides)r find_nodesrrmetatorch._inductor.compile_fxrx)rr output_noderxrrrgm_original_output_stridess   r}inputsSequence[IRNode] list[Expr]cCs@t}|D]}|t|ddO}|t|ddO}qt|S)NF unbacked_only)r=r&r!rKr)r~sym_varsinprrrget_symbolic_inputss rcCsHt|tr|j}t|tr|}t|tr|j}t|tr"|SdSr)rrrfBaseView unwrap_view StorageBoxBufferget_namerrrr try_get_names   rc@s2eZdZUdZeZded<dZded<ej ddZ d ed <ej ddZ d ed <ej ddZ d ed<ej ddZ ded<ej ddZded<eejdddZeejdddZedddZdddZdd%d&Zdd'd(Zdd*d+Zdd,d-Zdd.d/Zdd1d2Zdd4d5Zdd6d7Zddd;d<Z 8ddd@dAZddCdDZ ddFdGZ!ddIdJZ"ddLdMZ#ddOdPZ$ddRdSZ%ddTdUZ&ddWdXZ'ddZd[Z(e)dd]d^Z*dd`daZ+ddbdcZ,ddedfZ-dddidjZ.ddldmZ/ddodpZ0ddqdrZ1ddtduZ2ddwdxZ3ddzd{Z4dd|d}Z5dd~dZ6dddZ7dddZ8ddddZ9dddZ:dddZ;dddZdddZ?dddZ@ ddddZAdddZBdddZCdddZDdddZE ddddZFdddZGdddZHdddZIdddZJdddZKdddZLdddZMdddZNeOre)dddZPdSdS)rzBase class for all intermediate representation (IR) nodes in TorchInductor. Note: This is an abstract base class. Most methods raise NotImplementedError and must be overridden by concrete subclasses. zClassVar[OrderedSet[Any]]_current_originsNzClassVar[int | None]_current_stream_idxF)initOrderedSet[Any]originslist[str] | None tracebacktorch.fx.Node | None origin_nodedict[str, Any]r int | None stream_idxOrderedSet[Node]rGenerator[None, None, None]ccs.tj}||Bt_z dVW|t_dS|t_wr)rr)roldrrrcurrent_origins:s  zIRNode.current_originsccs*tj}|t_z dVW|t_dS|t_wr)rr)rrrrrcurrent_stream_idxDs zIRNode.current_stream_idxrrcCst|tttttfSr)rComputedBuffer InputsKernel InputBufferrlTemplateBufferrrrris_realized_nodePszIRNode.is_realized_nodecCs t|Srrcreateselfrrrwrap_for_lowering] zIRNode.wrap_for_loweringattrrr-r rcCst|||dSr)r __setattr__)rrr-rrr_post_init_setattr`szIRNode._post_init_setattrcCsZt|j}|d||dtjrtnd|dd|di|d|jdS)Nrrrrr)r=rrrEdebug_ir_tracebackr format_stackr)rrrrr __post_init__fs    zIRNode.__post_init__OrderedSet[str]cCtdd|DS)Ncs|]}|jVqdSrrrdeprrrrrz(IRNode.get_read_names..r= get_readsrrrrget_read_namesqzIRNode.get_read_namescC|jSr)rrrrr get_tracebacktzIRNode.get_tracebackcCrrrrrrrget_origin_nodewrzIRNode.get_origin_nodeOperation | NonecCrrrrrrrget_defining_opzrzIRNode.get_defining_oplist[Subgraph]cCgS)z'Return subgraphs contained in this noderrrrr get_subgraphs}zIRNode.get_subgraphscCst}|j}t|tr|}|jrt|g}|D]:}t|dr*|jr*||jqt j j j di |jg}t|ts>q|D]}t j j j |d}|rR||q@q|S)N stack_trace postToPre)r=rr ExternKernelrrhasattrraddr) _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrrrpre_grad_nodes node_namerrrrget_stack_tracess6       zIRNode.get_stack_tracesTshorten Sequence[str]cCsdt|dd}|rt|dkr|ddd}|s!|gSg}|D]}|d||d7}|d q'|g|S) Nzorigins=r@=z...zstack_traces = { })rrrappendsplit)rrrstack_trace_strrrrr common_reprs    zIRNode.common_reprlinesSequence[object] multilinecCsbt|t||}ttt|}|r&td|}t|jd|dSt|jd|dS)Nz, z( z )(r5)rrmaprindentjoinrr)rrrr new_linesrrr str_helpers zIRNode.str_helper torch.dtypecCrrrrrrrr&rzIRNode.get_dtypetorch.dtype | NonecC z|WStyYdSwr)r&NotImplementedErrorrrrrmaybe_get_dtype   zIRNode.maybe_get_dtypeLayoutcCtdt|d)Nz#get_layout() is not implemented by !rrrrrrr#rzIRNode.get_layout Layout | NonecCrr)r#rrrrrmaybe_get_layoutrzIRNode.maybe_get_layoutr6cC|Sr)r#rrrrget_output_speczIRNode.get_output_specOutputSpec | NonecCrr)rrrrrrmaybe_get_output_specrzIRNode.maybe_get_output_speccCst|tS)z4True for single tensor output (excludes MultiOutput))rrrrrrrhas_tensor_outputzIRNode.has_tensor_outputSequence[Expr]cCr)Nz!get_size() is not implemented by rrrrrrr!rzIRNode.get_sizeSequence[_IntLike] | NonecCrr)r!rrrrrmaybe_get_sizerzIRNode.maybe_get_size)_IntLike | sympy.Rel | Sequence[_IntLike]cCrrr!rrrrrWz IRNode.shapercC t|Sr)rmr!rrrr get_numelrzIRNode.get_numelcCtjjt|dSNrrrrrstatically_known_truerrGrrrrris_zero_elementszIRNode.is_zero_elementsr3cCtdt|)a) If the IRNode refers to data which has not been materialized (e.g., it is a Pointwise/Reduction that could potentially have more compute fused into it), realize the IRNode into physical memory, ending the possibility of fusing into it, but allowing, e.g., multiple users to access the data without having to recompute. Check StorageBox.realize for a particularly notable implementation. TODO(ezyang): I think, in principle, every IRNode should have an implementation of this, and most of the time no-op is OK, but you really do have to audit each IRNode for this, so for now, raise an error if it's not implemented. Note that some code in graph.py will catch this thrown error and suppress it with a warning. zrealize NYI on rrrrrrealizeszIRNode.realizewriterIndentedBuffer | NonecCr)Nzcodegen_reference NYI on rrr rrrcodegen_referencezIRNode.codegen_referencetorch.device | NonecCrrrrrrrr'rzIRNode.get_device torch.devicecC|}|dus J|Srr'rrrrrget_device_or_error zIRNode.get_device_or_errorcCdSNFrrrrrhas_exceeded_max_reads rzIRNode.has_exceeded_max_reads$Callable[[Sequence[Expr]], OpsValue]cCtt|jrrrrrrrr make_loaderrzIRNode.make_loader Callable[[Sequence[Expr]], Expr]cCrrrrrrr make_indexerrzIRNode.make_indexerrUcCrrrrrrrrKrzIRNode.get_stridecCrr)rKrrrrrrJrzIRNode.maybe_get_stridecCrrrrrrrrrzIRNode.get_namecCrr)rrrrrrmaybe_get_name!rzIRNode.maybe_get_namecCs(z |tjjvWStyYdSwr)rrrr graph_inputsrrrrris_input_buffer's  zIRNode.is_input_buffer thresholdcCrrrrr"rrrhas_large_inner_fn-rzIRNode.has_large_inner_fnusersrcCrrrrr%rrr mark_reuse0rzIRNode.mark_reusecCrrrrrrr realize_hint3rzIRNode.realize_hintcCrrrrrrrr6rzIRNode.unwrap_viewcCrrrrrrr freeze_layout9rzIRNode.freeze_layoutrr allow_paddingcCrrrrrr*rrrfreeze_layout_with_stride_order<z&IRNode.freeze_layout_with_stride_ordercCrrrrrrrrfreeze_layout_with_fill_orderArz$IRNode.freeze_layout_with_fill_orderrcCrrrrrrrrfreeze_layout_with_same_orderDrz$IRNode.freeze_layout_with_same_order exact_stridescCrrrrr2r*rrr freeze_layout_with_exact_stridesGr-z'IRNode.freeze_layout_with_exact_stridesdependencies.ReadWritescCrrrrrrrget_read_writesLrzIRNode.get_read_writesOrderedSet[Dep]cC |jSrr6readsrrrrrOrzIRNode.get_readscCrr)rrrrrr num_readsRrzIRNode.num_readsr~cCrrrrrrrget_storage_numelUrzIRNode.get_storage_numelrrcCrrrrrrrrget_free_symbol_usesXr-zIRNode.get_free_symbol_usescCrrrrrrrget_reduction_type]rzIRNode.get_reduction_typecCrrrrrrrget_reduction_size`rzIRNode.get_reduction_sizecCrrrrrrr is_externcrzIRNode.is_externcCrrrrrrris_no_opfrzIRNode.is_no_oprcCrrrrrrrconstant_to_deviceirzIRNode.constant_to_devicecCrrrrrrrget_mutation_nameslrzIRNode.get_mutation_namescCrrrrrrrget_operation_nameorzIRNode.get_operation_namecCrrrrrrrget_inputs_that_alias_outputrrz#IRNode.get_inputs_that_alias_outputcCrrrrrrrrwrz IRNode.dtype)rrrr)rrrrrrrrrr)rrr-r rrrrrr)rrrrrrrrT)rrrr)TT)rrrrrrrrrr)rrrr)rrrr6)rrrrrr)rr)rrrrrr3rr r rrrrrrrrrrrrUrrr"rrrr%rrrFrrr*rrrrrrrrrUrrr2rUr*rrrrr5rr7rrrr~rrrrrrrrrr)Qrrr__doc__r=rrr dataclassesfieldrrrrr staticmethod contextlibcontextmanagerrrrrrrrrrrrrrrr&rr#rrrrr!rpropertyrWrrr r r'rrrrrKrJrrr!r$r'r(rr)r,r/r1r4r6rr;r<r>r?r@rArBrCrDrErFrrrrrrr%s           !                                          c@seZdZd9ddZd:ddZd;d d ZdddZd?ddZ d@ddZ dAddZ dAddZ dBd d!Z dCd#d$ZdDd&d'ZdEd)d*ZdFd,d-ZdGd/d0Z 1dHdId3d4ZdJd6d7Zd8S)K OperationrrcCsd|_i|_dSr)operation_name_config_patchesrrrrr}s zOperation.__post_init__rcCtrrrrrrr'rzOperation.get_devicercCt|dsJ|jSNr)rrrrrrrzOperation.get_origin_nodercCrw)Nr)rrrrrr get_originsryzOperation.get_originsrcCrw)Nr)rrrrrrget_stream_idxryzOperation.get_stream_idxrcCs|jdusJ|jSr)rsrrrrrEryzOperation.get_operation_namercCr)zHGet config patches for this operation (e.g., coordinate_descent_tuning).rtrrrrget_config_patcheszOperation.get_config_patchespatchescCs ||_dS)z&Set config patches for this operation.Nr|)rrrrrset_config_patches zOperation.set_config_patchesrcCrrrrrrrrArzOperation.is_externcCrrrrrrrrBrzOperation.is_no_opr5cCrurrvrrrrr6rzOperation.get_read_writesrcCs ||vSr)rrrrrr is_user_ofrzOperation.is_user_ofrcCr)Ncsrrrrrrrrrz+Operation.get_read_names..rrrrrrrzOperation.get_read_namesr7cCr8rr9rrrrrrzOperation.get_reads list[Buffer]cCrurrvrrrr get_outputsrzOperation.get_outputsrcCtSrr<rrrrget_unbacked_symbol_defsrz"Operation.get_unbacked_symbol_defsFrcCr)a When unbacked_only=True: Returns the unbacked symbols which are required to be in scope in order to successfully perform codegen for this buffer. For example, a buffer that corresponds to an extern kernel call that takes i0 as an argument would return {i0} here. This is used to generate necessary dependencies that ensure we actually bind i0 in codegen before you try to use it. Note that this is NOT transitive; in particular, if this buffer takes in as input another buffer with dynamic shape (e.g., (i0,)), we will not report it here, because you will already have a dependency on that buffer, which will eventually have a dependency on i0 if necessary. When unbacked_only=False: Similar to `unbacked_only=True` but including all free symbols instead of only free unbacked symbols. r<r=rrrr>szOperation.get_free_symbol_usesrcCr)z Gets extra global memory size needed by this buffer. Some algorithms (e.g. group gemm) may require extra global memory in the generated code. rrrrrrget_workspace_sizeszOperation.get_workspace_sizeNrIrWrK)rr)rrr\)rr)rrrrrRrd)rrrrrJrerrrrr_rhrf)rrrrr'rrzr{rEr}rrArBr6rrrrrr>rrrrrrr{s(                 rrc@s*eZdZUded<ded<ded<ded<ed dQdRddZdSddZdTddZeZdUddZ dVddZ dWddZ dWd d!Z e dXd&d'ZeejfdYd*d+ZedZd-d.Zd[d0d1ZedTd2d3Zd\d]d7d8ZdQd^d:d;Zd_d?d@Zd`dBdCZdadDdEZdbdGdHZdWdIdJZdcdLdMZdddOdPZd4S)eLoopsrrrrCallable[..., Any]inner_fnrUrangesFrrrrcs,tjgfdd|jD|RS)Nc3|]}t|VqdSrr%rerrrrz-Loops.get_free_symbol_uses..)r=unionrinner_fn_free_symbolsr=rrrr>s zLoops.get_free_symbol_usesnamesrrcsFdjjdtjgfdd|DdjgS)N'cs g|] }|dt|qS=)rrrrrrrs z!Loops._to_str.. origin_node=)rrrrr inner_fn_strr)rrrrr_to_strs z Loops._to_strcC |dSNrrrrrr__str__rz Loops.__str__rcCrrrrrrrr'rzLoops.get_devicercCrrrrrrrrrzLoops.get_origin_nodercCrrrrrrrr!rzLoops.get_sizecCrrrrrrrget_pointwise_sizerzLoops.get_pointwise_sizerr rrcOsN|dd}|dd}||i|}|d||d|p |jt|S)Nrr)poprrrr)clsrrrtbrrrrrs    z Loops.createrrCcsfddt|DS)Ncs*g|]\}}|dkrtjjnt|qSrD)rSZerorl)rnrrrrrz Loops._index..)r)rrrrr_index s z Loops._indexrTc Cstt}t|2ttdd|j|| WdWdS1s0wYWddS1s@wYdSNallow_indexingT) rSrr MockHandlerset_ops_handlerrrr$r inner_fn_argsgetvalue)r opcounterrrrinner_fn_opcounts  RzLoops.inner_fn_opcountSequence[Sequence[_IntLike]]cCs||jfSr)rrrrrrrrzLoops.inner_fn_argscCstjj|jg|RSr)rrKernelFormatterHandler ir_to_stringrrrrrrr s zLoops.inner_fn_strNr"rcCs&|durd}t|tj}|j|kSr)maxrErealize_opcount_thresholdrnum_opsr#rrrr$&s zLoops.has_large_inner_fnOrderedSet[Symbol]cCs||j}t|j||dSNr)rrrMr)rrrrrrr,s zLoops.inner_fn_free_symbolssymbolr rcCsR||j}t|}t|||Wd|jS1s!wY|jSr)rrrPrrrrusages)rrrhandlerrrrcollect_inner_fn_symbol_usage1s   z#Loops.collect_inner_fn_symbol_usager7cCsvttdd*|r t|||jWdSt||jWdS1s4wYdSr) rrr$r?rOrr!r@r:rrrrr:s$zLoops.get_readscCt|jSr)r=r read_buffersrrrrrHrzLoops.get_read_namesrcCrr)rrrrrrrr;KrzLoops.num_readscCr)Nz+get_reduction_size() is not implemented by rrrrrrr@NzLoops.get_reduction_sizer3cCr)Nz+get_reduction_type() is not implemented by rrrrrrr?SrzLoops.get_reduction_typercCr)Nz+constant_to_device() is not implemented by rrrrrrrCXrzLoops.constant_to_devicer_rh)rrrrr\rWrKrS)rr rr rr)rrUrrCrr)rrTrrrr]rrrr)rr rrrerJrfrUri) rrrrr]r>rr__repr__r'rr!r classmethodrrnrCINDEXrr\rrrr$rrrrr;r@r?rCrrrrrsB                   rrExpr | Sequence[Expr]rrrqcCs"|jr ttd|Std|S)Nnanr)is_floating_pointrpconstantfloat)rrrrr nop_loader_fn^s rc@sLeZdZdddZdddZeZdd d Zdd d Zd ddZd!ddZ dS)" PointwiserrcCs|r tt|jdS|jSNr)rr rrrrrrrrgszPointwise.make_loaderrcCrrrrrrrrnrzPointwise.__str__Sequence[sympy.Expr]cCrrrrrrrr@srzPointwise.get_reduction_sizer3cCrrrrrrrr?vrzPointwise.get_reduction_type output_nameindexer!Callable[[Sequence[Expr]], Never]varsrrcCs"|}t|p d||||SNunnamed)rrpstorerrrrloaderrrr store_outputyszPointwise.store_outputrrrcCs.|}ttd||}t||j||jdSFMove this to a given device. Requires that all reads are to constants.override_devicerrrr)rrrConstantBufferrrrrrrrrrrCszPointwise.constant_to_deviceNrYr\)rrrU)rr3rrrrrrri) rrrrrrr@r?rrCrrrrres      rc@s6eZdZUded<dZded<dd d ZdddZdS)Scatterroutput_indexerNrV scatter_moderrrrcCs6|}ttd||}t||j||j|j|jdS)rr)rrrrrr) rrrrrrrrrrrrrrCszScatter.constant_to_devicerr3rrrrr cCs6|}|dur d}tj|||||||jdS)Nr)mode)rrprrrrrrrrs zScatter.store_outputri)rr3rrrrrr )rrrrrrCrrrrrrs    r logical_ormaximumminimummulr bitwise_xor)anyrminprodsumdotxor_sumz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNreduction_typearg_break_ties_leftCallable[..., object]csRtvrtSdvrdfdd }|Sd kr"dd d }|Std )Nargmaxargminatuple[object, object]brtuple[OpsValue, OpsValue]c s|\}}|\}}dkrt||}nt||}t||}trCt||}t||} t|t|| }t|t|| }rKt||nt||} t|t|| }t|||t|||fS)Nr) rpltgteqr-ner logical_andwhere) rra_valuea_indexb_valueb_indexmaskequala_isnanb_isnantierrrrrargmax_combine_fns&       z3get_reduction_combine_fn..argmax_combine_fnwelford_combine#tuple[OpsValue, OpsValue, OpsValue]c SsR|\}}}|\}}}||}||} || } ||| |||||| | fSrr) rra_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_wrrrwelford_combine_fns   z4get_reduction_combine_fn..welford_combine_fnzunknown reduction_type=)rrrrrr)rrrrrr)rr)rrrr rrr rget_reduction_combine_fns rcs\eZdZUded<ded<ded<ded<dkd d ZeZeddldmfdd ZdnddZdoddZ dpd d!Z dqd#d$Z drd&d'Z dldmd(d)Z dsd-d.Ze /dtdud:d;Zedvd>d?Zeejd/fdwdBdCZedxdFdGZedxdHdIZedydMdNZedzdPdQZe /dtd{dWdXZed|d`daZed}dedfZe /dtd~dgdhZeddidjZZS) ReductionrUreduction_rangesrUrr src_dtyperYreduction_hintrrcCr)N)rrrrrrrrrrzReduction.__str__Frrrcs(ttjfdd|jDBS)Nc3rrr%rrrrr rz1Reduction.get_free_symbol_uses..)superr>r=rrr= __class__rrr>szReduction.get_free_symbol_usesrcCrr)rrrrrr@rzReduction.get_reduction_sizer3cCrr)rrrrrr?rzReduction.get_reduction_typerrrrreduction_varsSequence[Symbol]rc Cs8t|j|j|j|||}t|pd|||dSr)rp reductionrrrrstore_reduction)rrrrr"r-rrrr%s zReduction.store_reductionrcCt|jt|jSr)rrrrrrr index_length#rzReduction.index_lengthSequence[Sequence[Expr]]cCs$||j}||jtj}||fSr)rrrrCR0_INDEX)rrrindexrrrr&s zReduction.inner_fn_argscCs.||j}||jtj}t|j|||dSr)rrrrCr)rMr)rrrr*rrrr+s  zReduction.inner_fn_free_symbolsrrrc Cs>|}ttd||}t||j||j|j|j|j t j dS)rrrrrrrrrr) rrrrrrrrrrrYDEFAULTrrrrrC2szReduction.constant_to_deviceN dst_dtyperCallable[_P, OpsValue]rReductionType | Literal['scan']reduction_numelr input_nodertuple[ReductionHint, _IntLike]c # Cs|t|g} tjj| stjdfStjj|} tjjt|} |dkp5tj|t j  o5|dvo5t j } |dkr?tjdfSt |} | j}d}| r`tjtjj|dd}tjtjj|dd}ndd d}|}| dkr|| | }|dkrytj|fS|durt|trttddt|\}}Wdn1swY|dur|durtjjt||}| |krtd|||||tjdfStj|fS| |ks| |ddkrtjdfSt||||||dkr|nd|tjd}ddd}||\}}|r||\}}t|dkrtjdfSt !|"|#\\}}}d}d}|D].}tjj$||} tjj%| |t&|'}!t(dd|!D}"|"rK|d7}q"|d7}q"||kr^tj|| | fStj)|| | fS)NrDscanrr T)inner_reductionFreduction_numel_hintr numel_hintrcSrNrDr)r6r7rrrinner_reduction_splitsoz4Reduction.num_splits..inner_reduction_splitsrzUse previous IRNode's range and reduction_ranges instead of split. current ranges: %s, current reduction ranges: %s, current split: %d, new ranges: %s, new reduction ranges: %srIrrr+rrtuple[Sequence[Expr], bool]c s|}|dus Jtdt|||d|d}|}|jdus%Jdd|jD}g}d}t|jdddD]7t fd d |Drq| j j t jjvrqt jjj }t|jd d}|t|jd d|krqd }q:||fS) NrrrrrgrfcSs&g|]}t|trt|tjs|qSr)rrrNumberrrrrrrs zBReduction.num_splits..get_read_indices..FcSrrrrrrrz@Reduction.num_splits..get_read_indices..keyc3s|] }|jjvVqdSr)rr2r?mdrrrzAReduction.num_splits..get_read_indices..rT)r'rr$r&r!r6 range_varssortedr:rrrrrrrname_to_bufferrrg decide_layout) rrcb read_writesrGindiceschangedbuforiginal_striderrDrget_read_indicess:   z.Reduction.num_splits..get_read_indicesrcss |] }|dkp |dkVqdS)rrDNrrrrrrrz'Reduction.num_splits..)r6rr7rrr)rrrr;)*rmrrrrall_unbacked_explicitly_hintedrYr,r  has_featurerGREDUCE_TO_SINGLE_ELEMENTrEsplit_reductionsrXrmulti_processor_count functoolsr choicesreduction_split_factorINNERrrrrr$rN!replace_backed_symbols_with_hintslogrrrrFindex_vars_squeezer!r@simplify_with_ranges stride_hintsrkeysrOUTER)#rr-rrrrrr0r1exprsr6r7 should_splitpropsnum_smmin_elements_per_threadr9outer_reduction_splitsr new_rangesnew_reduction_rangesextracted_numel_hintrrQrMrNr r"ranges1 num_outer num_innerrjr^outerrrr num_splitsAs             #       zReduction.num_splits.fn..cSsg|]}t|qSr)rrrrrrrrz>Reduction._unroll_reduction_fn..fn..)rWreduce itertoolsproductr) combine_fnrrsrrrs  z*Reduction._unroll_reduction_fn..fnrrr*rcs*dd|D}||t|tjfS)NcSg|]}t|qSrrexpandrrrrrrzDReduction._unroll_reduction_fn..value_fn..)rp index_exprr)int64rr*) flatten_indexrrrrs sz0Reduction._unroll_reduction_fn..value_fncs |dSr8rr)rrrr@s z0Reduction._unroll_reduction_fn..N)rrUrr )rrUr*rUrr)rrrr guard_int_seqr_fixed_indexerr$r%)rrrrr)rxrrrrrsr_unroll_reduction_fns   zReduction._unroll_reduction_fnrrc  stjjtdkrBd#fdd } | d| d| d| dd vs.Jd d$fdd } tj||| t|dSdkrcdvrRd$fdd } nd$fdd } tj|| |dStt rt t j krt|dkszt |jrdkrtj||||dS||||| \} }d%fdd }||}|tjkr| }|dkr| dusJttddt| \}}Wdn1swY|dusJ|dusJ||||||| S|dkr<||||||| }d}t jjrt|trd&d d!}||}|r:t|jts*Jt|j|jjd|_|_ ||_!|_"|Stt||||d"}|S)'zy Create a reduction node. May split the reduction to multiple layers to expose more parallelism. rvalrrbool | float | intcsTtjkr t|Sjrt|tsJt|t|St|ts&Jt|t|Sr) r)rrrrrrrrrr-rrpy_cnst2s z!Reduction.create..py_cnstrD)rrrrz* not supported for zero-dimension tensors!rrrqcstSrrprr)r-rrtypes_to_initsrrconst_fnHz"Reduction.create..const_fnrrycs tdSrrrrrrrVrzReduction.create..fncddD}||S)NcSg|]}tjjqSrrrrr rrrr\rz0Reduction.create..fn..rrreduction_index)rrrrr[ rrcs$tr|S|dkrt|tjS|Sr8)rrrE min_num_split)r)r0rr_maybe_increase_splits  z/Reduction.create.._maybe_increase_splitrINrTcur_noderComputedBuffer | NonecSsd|}t|dkr dStt|}|tjjvrdStjj|}t|ts'dS|j dus0J|Sr8) rrnextiterrrrrIrrrfr?)r read_namesbufnamerOrrr_find_split_reductions     z/Reduction.create.._find_split_reductionr+)rrrr)rrrrq)rrrr)rrrr)#rrrrsimplifyrmrrrrrrrEunroll_reductions_thresholdrirrrprYr,rrr$rN!create_multilayer_existing_rangescreate_multilayerr=mix_order_reductionrrfrr _split_size_original_inner_fn_original_ranges_original_reduction_ranges)rrr-rrrrrrr1rrrhintrrrhrir split_reductionrr)r-rr0rrrrrs             zReduction.creater_NumLike | Sequence[_NumLike]c Cs|dvrt|r tdSt|rdSt|jS|dvr0t|r$tdSt|r*dSt|jSt|r6dnd}t|r>dnd}||||||||f|||ftd|fd |S) N)rrz-infF)rrinfTrrD)rrrrrwelford_reduceronline_softmax_reduce)r-rr+r)iinforr)rrzeroonerrrdefault_accumulators2    zReduction.default_accumulatorcCs|dkrdSt||S)Nrr)rrrrrrr default_value s zReduction.default_valuerr~r7cCsP|dkr|S|dkr|dkr|tjkrtjS|dkr&|dkr&|tjkr&tjS|S)NrIi)rYra OUTER_TINY)rr7rrrr_multilayer_second_step_hints z&Reduction._multilayer_second_step_hintrcCs|durdStjj||sdS|zt|Wn ty&YdSw|}t |ddD]\}}tjj|drC|Sq3dS)z If we are reducing over the full tensor, and it is non-dense in the last dimension, reindex so we reduce over the dense dimension. initially just handle complete reduction case NrIrD) rrrrrbrr rhrrKr)rr0r1r^rrrrr$check_for_split_dense_dim_reindexing%s$   z.Reduction.check_for_split_dense_dim_reindexingrr block_sizedefaultrc sV||}t|g|tjjtt |d d fdd } | S) Nrrr#rrrqcsl|\}|^}||dfdd }r3t}tt|t|}t||S|S)NrrqcsgSrrr)rMr new_indexrrrbodyarzCReduction._multilayer_wrap_loader..wrapper_fn..body)rrq)rcrprr}masked)rrreduction_blockr index_dtyperrrr need_maskr0r)rMrr wrapper_fnZs    z5Reduction._multilayer_wrap_loader..wrapper_fn)rr#rr#rrq) rViewdynamic_reshape_indexerrrrrrrrGrA) rrrr0rrrr1 dense_indexrrrr_multilayer_wrap_loaderEs z!Reduction._multilayer_wrap_loader4Callable[[Sequence[Expr], Sequence[Expr]], OpsValue]original_rangesoriginal_reduction_rangesrhSequence[Integer]ri@Callable[[Sequence[sympy.Expr], Sequence[sympy.Expr]], OpsValue]csNtddDsJdt|t|t|d fd d }|S) Ncs|]}|dkVqdS)rDNrr?rrrryzDReduction._multilayer_wrap_loader_existing_ranges..z8Only enabled for numel_hint == 1, found original_ranges= merged_indexrnew_reduction_indexrrqcs:|dt}|td}|t|t|Sr)rr)rr original_idxrrrrrrrs zEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fn)rrrrrrq)rrrr)rrrrrhrirrrr'_multilayer_wrap_loader_existing_rangesps  z1Reduction._multilayer_wrap_loader_existing_rangesrr list[Integer]c  s|tjtjfvr |ntj} t|| ||||| | } | | d fdd }tj j t |}| | || } ||dt|ksDJtt|||||t|d| || d S) a Break a large reduction up into multiple smaller reductions recursively rrUrrrqcsg||Srrrintermediate_loaderrrintermediate_fnsz;Reduction.create_multilayer_helper..intermediate_fnNr+)rrUrrUrrq)r)float16bfloat16rrrr rrrrrr rmrrr)rrr-rrrrrhrirrrintermediate_dtype intermediaterr7rrrcreate_multilayer_helpersD z"Reduction.create_multilayer_helperc  Csdt|} t| |d|} |||} |||| || | | }|||||||g||| g||| S)rrD)rmr@rrr)rrr-rrrrrrrr1r0rrrrrrrs2   zReduction.create_multilayerc Cs8||||||} ||||| ||g|||| d| S)rrI)rr) rrr-rrrrrhrirrrrrrrs( z+Reduction.create_multilayer_existing_rangesr\r_rrSrU) rr3rrrrr"r#rrrfrr(rir)rrr-rrrrr.rrUrrUrr/r0rr1rrr2) rrqrrUrrrrrrr)rrr-rrrrrrrrrrrUrrYr1rrrrrrrrr)rr~r7rrrYrrY)r0r~r1rrr)rrrrUr0r~rr~rr~rrr1rrr) rrrrrrrhrrirrr)rrr-rrrrrrrrrrhrrirrrUrr~rrYrr)rrr-rrrrrrrrrrrUrr~rrYr1rrr)rrr-rrrrrrrrrrhrrirrrUrrYrr) rrrrrrr]r>r@r?r%r'rrrCrnrprrrYr,rrrrrrrrrr __classcell__rrr rrs^          - + Q     *  ? -rrrSequence[int] | Nonerjrrcsdfdd }|S) 1A closure containing math to read a given elementrrrrcsddur t|tksJt|tksJ}t|D]\}}}|dkr/|||}q |Sr8)rr)rresultrstszrjrrrrr*s z_fixed_indexer..indexerN)rrrrr)rrrjrrrrr#s r INNER_FN_TYcs2eZdZUded<dfdd Zd ddZZS)!MultiOutputReductionr output_indexrrr-r inner_fns#INNER_FN_TY | Sequence[INNER_FN_TY]rrrrrUrrrYc sXtrftdkrd} nd fdd } tj||| |||||d | |_dS) NrDrrr reduction_idxrtuple[OpsValue, ...]cstfddDS)Nc3s|]}|VqdSrrrrrrrrrSrz@MultiOutputReduction.__init__..loader..)rrrrrrPsz-MultiOutputReduction.__init__..loaderr+)rrrrrr)callablerr__init__r) rrr-rrrrrrrrr rrr<s     zMultiOutputReduction.__init__rr3rrrrr"r#rr c CsXt|j|j|j|||}t|ttfsJt |||j }t |p&d|||Sr) rpr$rrrrrrrrrr%)rrrrr"rr-rrrr%as  z$MultiOutputReduction.store_reduction)rrr-rrrrrrrrrUrrrrYrr) rr3rrrrr"r#rr )rrrrrr%rrrr rr9s %rc@s"eZdZeejdfdddZdS)OnlineSoftmaxReductionNrrr-rrrrrrr num_outputrrrYr1rrSequence[TensorBox]c s<tfddt|D} | D]} | q| S)z> Create the reduction disregarding splitting. c 3s.|]}ttd| VqdS)rN)rrrr output_idxrr-rrrrrrrrs" z0OnlineSoftmaxReduction.create..)rrr ) rrr-rrrrrrr1resultsr,rrrrts  zOnlineSoftmaxReduction.create)rrr-rrrrrrrrrrrrrYr1rrr)rrrrrYr,rrrrrrss  rc@s<eZdZeejfdddZedddZedddZ dS)WelfordReductionrrrrrSequence[Callable[..., Any]]rrrrrUrrYrrc sFdvsJtjjrJdtjjt}dfdd } |d kr5| d } | d } | d } | | | fS|d kr_dfd d dkrTd | d | d fStfddDSt j d |d\} }t j krw| |d kr| |SfddtdD}|D]}|q|S)N)rrzGwelford reduction usage is explicitly disabled, please check you configrrrrcs&dfdd }tj|tdS) Nrrrrqcs tSrrr)rrrrrsz8WelfordReduction.create..const..inner_fnrrrrrqrrr)rr)rrrrrconstsz&WelfordReduction.create..constrrDrrcs&dfdd }tj|tdS) Nrrrrqcr)NcSrrrr rrrrrzKWelfordReduction.create..copy..inner_fn..r)rr)rrrrrrz7WelfordReduction.create..copy..inner_fnrrr)rr)rrrrrrcopysz%WelfordReduction.create..copyrc3|]}|VqdSrrr)rrrrrz*WelfordReduction.create..)rr0cs*g|]}tt| qSr)rrrr)rrrrrrrrrr s z+WelfordReduction.create..r)rrrr)rrrr)rEmtiadisable_welford_reductionrrrrrrmrrrprYr,rrr )rrrrrrrrr0rmeanm2weightrrrr,r)rrrrrrrrrrsZ     zWelfordReduction.createrrcCr)N)rrrrrrrrr r:zWelfordReduction.default_valuerr~c  s&ttjjttd } | r:|dkr:dfd d } j||dt | dd t | d d f|d|dSt d t |t fdd|Dg|g||} | D]} | qbdddtjjt|} | |}t |t fdd| D|gd|S)rrrrrrr-rrrqcs t|Srr)rrr-rrrr5 s z4WelfordReduction.create_multilayer..constantr0rD)rrrrrrrrc 3s&|]}j|ddVqdS)rrN)r)rr)rrr0rrrrrM s  z5WelfordReduction.create_multilayer..rrrrcSs|g||Srr)rrrrrrintermediate_loader_fn` szBWelfordReduction.create_multilayer..intermediate_loader_fnc3s |] }t|dVqdS)rN)r rr)rrrrp  N)rrrrr-rrrq)rrrrrrrrq)rmrrrrrrrGrArr r@rrrr r r)rrrrrrrrrrr intermediatesrr7r)rrrrr0rrrr sb         z"WelfordReduction.create_multilayerN)rrrrrrrrrrrrUrrYrrr)rrrrrrrrrrrrUrr~rrYrr) rrrrrYr,rrnrrrrrrrs { rcseZdZUded<ded<ded<ded<ded <d ed <d ed <ded<eddIdJfdd ZdKfdd ZdLd#d$ZdMd%d&ZdNd'd(Z dNd)d*Z dNd+d,Z dOd-d.Z dPd0d1Z dIdJd2d3Zeejfd4d5dQd=d>ZedRdGdHZZS)SScanr scan_rangesr=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rxzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rrYrrrtuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]rFrrrrcDttjfdd|jDBtjfdd|jDBS)Nc3rrr%rrrrr rz,Scan.get_free_symbol_uses..c3rrr%rrrrr r)rr>r=rr rr=r rrr> s zScan.get_free_symbol_usesrc0t|jt|jt|jksJtdSr)rrr rrrrr rrr "zScan.__post_init__rr3r%Callable[[Sequence[_IntLike]], Never]rr scan_varsr#r csR|||tfdd|jD}t|j|j|}t|p d|||jS)Nc3|]}|VqdSrrrrrrrr rz'Scan.store_reduction..r) rrrrpr3r rxrr)rrrrrrrrrrr% s zScan.store_reductioncCr)Ncustomrrrrrr? rzScan.get_reduction_typecCrr)r rrrrr@ rzScan.get_reduction_sizecCrrrrrrrr! rz Scan.get_sizecCrrrrrrrr rzScan.get_pointwise_sizecCr&r)rrr rrrrr' rzScan.index_lengthrcC.||j}||jtj}|||}|fSr)rrr rCr)rrrr*rrrrr   zScan.inner_fn_argscC8||j}||jtj}|||}t|j||dSr)rrr rCr)rrMrrrrr*rrrrr   zScan.inner_fn_free_symbolsT)can_fallback_to_atenrr+tuple[Callable[[Sequence[Expr]], Any], ...]axisrrSequence[TensorBox | None]c sg d dd g tjtjs$dgtStdkr9tjtjs9dgtStjj} | t } ttksNJ| t | drg fddt tDS|jdd | d\} t | dkrtjjduptotdkotdk} | s|rdgtSd} nt d fd d  f ddt tD}|D]}|q|S)NrDc&g|]}tj||dqSrrrrrrr rrrrr zScan.create..r)rrrrpointwise_rangesr rx scan_numelz3.3.0rr scan_indexrrcHt|tks Jt|tksJg|d||dSrr)rr))rr'r rrr  zScan.create..reindexcsBg|]}t d|| |d qS)) rrr rrrrr rxrrrrrr$) rxrr rrr'rrr  scan_typerrrr s* )rrr)rrr)rrrrSrGSCANrTUPLE_REDUCTIONrrrmrrrNrrprr)versionhip has_tritontriton_version SplitScanr )rrr rrrrxrrrrr(rpsupports_splitrrr) rrxrr rrr'rrr r-rrr sV         z Scan.createrrrrr'r(rr2c s*d fdd } tj|||| ||d|d S) Nrrrrrqcs$g|d||dSrrrrrrrr7 s$z#Scan.num_splits..wrapper_fnr3)rr-rrrrrr0)rrrrrrq)rrp) rrrrrr'r rxr(rrr6rrp* s zScan.num_splitsr_rrI) rr3rrrrrr#rr rUrSrfr)rrr r rrrrrrrxr rrYrrrr rr )rrrrrrrrr'rr rrxr r(rrr2)rrrrr]r>rr%r?r@r!rr'rrrrYr,rrprrrr rr| s6           brc@s eZdZdS)r4N)rrrrrrrr4G sr4cseZdZUded<ded<ded<ded<ded <d ed <d ed <ded<ded<edd;d<fdd Zd=fdd Zd>d"d#Zd?d$d%Zd@d&d'Z d@d(d)Z d@d*d+Z dAd,d-Z dBd/d0Z d;d.c3rrr%rrrrre r)rr>r=rr8rr=r rrr>] s zSort.get_free_symbol_usesrcrr)rrr8rrrrr rrri rzSort.__post_init__rr3rrrrr"r csV|||tfdd|jD}t|j||j|j}t|p"d|||j S)Nc3rrrrrrrru rz'Sort.store_reduction..r) rrrrpsortr r9r:rr)rrrrr"rrrrrr%m s zSort.store_reductioncCr)Nr;rrrrrr?{ rzSort.get_reduction_typecCrr)r8rrrrr@~ rzSort.get_reduction_sizecCrrrrrrrr! rz Sort.get_sizecCrrrrrrrr rzSort.get_pointwise_sizecCr&r)rrr8rrrrr' rzSort.index_lengthr(cCrr)rrr8rCr)rrrrrr rzSort.inner_fn_argscCrr)rrr8rCr)rrMrrrrrr rzSort.inner_fn_free_symbolsrr'tuple[Callable[[list[Expr]], Any], ...]rrr c  s<g d dd g tjtjs$dgtStjj} | t } t j j r8t j j } nd} t j j oF| t| | } | sPdgtSttksZJ| t| drs fddttDSd  fd d  f d dttD}|D]}|q|S)NrDrcr!r"r#r$r%rrr r&zSort.create..rr sort_indexrrcr*rr+)rr=)rr'r8rrr r,zSort.create..reindexcsDg|]}ttd|| | d qS)) rrr rrrrr8rrrr9r:r)rrr7r$) r:rr rrr'rrrr8r9rrr s,)rrr=rrr)rrrrSrGSORTrrrrmrEr=decompose_sort_opspersistent_reductionsrrrNrr )rrr rrrr9r:rrr sort_numelis_persistent_kernel max_rblockrrr) rr:rr rrr'rrrr8r9rr s4      z Sort.creater_rrI) rr3rrrrr"rrr rUrSrfr)rrr r rr<rrrrr9rr:rrrYrr rr )rrrrr]r>rr%r?r@r!rr'rrrrYr,rrrrr rr7L s0          r7cCs(z t|ddWdStyYdSw)NFfreezeT)rhrrrrrr" s   r"cCs@zt|dd\}}|r||WStyYdSwNFrD)rhshould_pad_strides pad_strides is_contiguousr)r_bufferrgrrr is_contiguous_storage_and_layout s  rKrEwant_contiguous stride_orderSequence[int | Integer] | Noner*r2tuple[StorageBox, Layout]c Cst|trt|j|||||dSt|tr)t|j|||||d\}}||jfSt|trc|r[|r?||s>Jn|durK|j ||dn|durW|j ||dn| t||fSt|t rvt|j|d\}}||j fSt)z Try to simplify x into a StorageBox and a Layout. allow_padding only affect how we apply stride_order. When allow_padding is True, we have the freedom to add padding when applying the stride_order. rErLrMr*r2Nr*rD)rrrhrfrr#rr)rIr,r4rJrlrgr) rrErLrMr*r2r rgbufferrrrrh sR      rhcCs2zt|dd\}}||WStyYdSwrF)rhis_stride_orderedr)rrMrJrgrrr"is_stride_order_storage_and_layout; s   rTrcCsrt|ttfr t|jSt|tr*|j}tjj |j t |j t }t|jp)|St|tr7|tjjvSdSr)rrr is_unalignedrfrlrgrrrrstatically_known_multiple_ofrjrdrrfrrunaligned_buffers)rrghas_unaligned_layoutrrrrUE s   rUc@seZdZUded<eddEdFdd ZdGd d ZdHddZdIddZe dJddZ dKddZ dLddZ dMddZ dNd d!ZdOd#d$ZdPd(d)ZdQd*d+ZdRd-d.ZdSd/d0ZdTd2d3ZdQd4d5ZdQd6d7ZdUd9d:ZdVdd?ZdXdBdCZdDS)YrrrfFrrrrcC |j|Srrfr>r=rrrr>[  zBaseView.get_free_symbol_uses*Callable[[Sequence[Expr]], Sequence[Expr]]cCstd|)Nzmake_reindexer NYI on rvrrrrmake_reindexer_ rzBaseView.make_reindexerrc&|j|dfdd }|S)Nrrrrcrrrrinnerrrrrf rz&BaseView.make_indexer..indexer)rrrr)rfrr]rrrr_rrb  zBaseView.make_indexerrcr^)Nrrrrqcrrrrr_rrro rz$BaseView.make_loader..loaderr)rfrr]rrrr_rrk rbzBaseView.make_loaderrcC |jSr)rfr&rrrrrt rzBaseView.dtypercCrdrrfr#rrrrr#x rzBaseView.get_layoutrcCrdrrfr'rrrrr'{ rzBaseView.get_devicercCrrrrrrrr~ rzBaseView.get_origin_nodercCrdrrfrrrrrr rzBaseView.get_namercCrrrrrrrr rzBaseView.get_pointwise_sizer%rrcCrYrrfr'r&rrrr' rzBaseView.mark_reusecCrdrrfrrrrrr rzBaseView.has_exceeded_max_readsr3cCrdrrfr rrrrr  rzBaseView.realizecCs|jdSrrfr(rrrrr( rzBaseView.realize_hintr~cCrdrrfr<rrrrr< rzBaseView.get_storage_numelcCrdrrfrArrrrrA rzBaseView.is_externcCs$t|jts Jt|j|jSr)rrfrris_module_bufferrrrrrn s zBaseView.is_module_bufferrcCrdrrfrrrrrr rzBaseView.get_read_namesr7cCsFttddt||jWdS1swYdSr)rrr$rOrr!r:rrrrr s$zBaseView.get_readscCs"|}t|tr|j}t|ts|Sr)rrrf)rrrrrr s   zBaseView.unwrap_viewrrcCs2|}ttd||}t||||dSr)rrrrrr&r!rrrrrC szBaseView.constant_to_deviceNr_rrr\rZrYrOrPrWrKr\rSr^rRrUrIrgrJrerHri)rrrrr]r>r]rrrqrr#r'rrrr'rr r(r<rArnrrrrCrrrrrW s2                    rc@sFeZdZUded<eddd Zedd d Zdd dZdddZ dS)rrrrrnew_sizerUrcCs tjj}dd|D}|}dgt|t|t|}t|t|ks(Jtt|D]_}||dkrE||dus>J||||<q.||dusTtjj||rUq.t|st|s||}||}|duskJ|dusqJ||}|j |dddksJd|d|d |q.|S) zReplace `-1` with correct sizescSrzrr{rrrrr rz.ExpandView._normalize_size..NrIrfallbackzBroadcast failed in ExpandView(, z) on dimension ) rrrrr!rrris_size_one_or_falser5r )rrqrold_sizerv1v2diffrrr_normalize_size sB   zExpandView._normalize_sizerc Cs|||}t|rWt|\}}t|t|j}|dksJtjjg|}t|j |jD]\}}| t j j |s<|ntjjq-t|j|jt|||j|j} t|| dSt||dS)Nrre)rfr)rzr"rhrrrrrrrrrrrrrurirrrrjrkrlr) rrrqrmrnskiprorrrprrrr s,      zExpandView.createcCrrrrrrrr! rzExpandView.get_sizer\cs6|}|jt|tdfdd }|S)NrrrcsRt|d}t|tksJttD]}|dkr&tjj||<q|Sr8)rrrrrr)rractualr{rrr s  z*ExpandView.make_reindexer..reindexrrrr)r!rfr)rtargetrrr|rr] s   zExpandView.make_reindexerN)rrrqrUrrU)rrrqrUrrrSrp) rrrrrnrzrrr!r]rrrrr s  '  rc@sFeZdZUded<eddd Zedd d ZdddZdddZdS) PermuteViewrdimsrrrrrcs||}t|ttt|ksJt|r>t|\}tjjfdd|Dfdd|Dj j }t ||dSt ||dS)Ncg|]}j|qSrrrrnrrr& rz&PermuteView.create..crrrrrrrr' rre)rfr) _map_neg_dimsr=rrr"rhrirrrjrkrlr)rrrrmrprrrr s    zPermuteView.create list[int]csfddDS)Ncs$g|]}|dkr |nt|qSrr+)rrYrrrr1 s$z-PermuteView._map_neg_dims..r)rrrrrr/ zPermuteView._map_neg_dimsrcsDt||jttt|jksJ|jfdd|jDS)Ncrrrrrrrr8 rz(PermuteView.get_size..)r=rrrrrfr!rrrrr!3 s   zPermuteView.get_sizer\cs`ddt|jDfddtt|jDtttt|jks'Jd fdd }|S) NcSrrr)rrrnrrrr= rz.PermuteView.make_reindexer..crrrrinvrrr> rz.PermuteView.make_reindexer..rrrcsfddDS)NcrrrrrrrrD rz?PermuteView.make_reindexer..reindex..rrrrrrA z+PermuteView.make_reindexer..reindexr~)rrrrr=)rrrrrr]: s zPermuteView.make_reindexerN)rrrrrr)rrrrrSrp) rrrrrrrr!r]rrrrr s    rc@s8eZdZeddddd Zedd dZdddZdS) SqueezeViewNrYrrrYrrc sLt|rzt|\}}g}g}dur*ttsJtdkr(t|jks*Jtt|j|j D]3\}\}} durOt j j |sN|||| q3|kr^|||| q3|dksfJdq3t|j|j|||j|j} t|| dSdurt|dd|DS|dksJt|fddt|DS)NrrDzexpected squeezed size to be 1recSsg|] }tjj|s|qSr)rrrrrurrrrro s  z&SqueezeView.create..csg|] \}}|kr|qSrrrrrrrrrw )r"rhrrrrrrrrrrrrrurrirrrjrkrlrrr!) rrrYrmrnrqrorrrrprrrrK sF       "zSqueezeView.createrr>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]]cs@dd|D}ddt|Dt|d fdd }||fS) NcSsg|]}|dkr|qSrrrrrrr} rz(SqueezeView.squeezer..cSsg|] \}}|dkr|qSrrrrrrr~ rrrrtuple[Expr, ...]csTt|tksJ|dtjjg}t|D]\}}|||<qt|S)N )rrrrrr)rrrrlengthnot_onerrr s " z%SqueezeView.squeezer..reindex)rrrr)rr)rrqrrrrsqueezery s zSqueezeView.squeezerrfr rcCtd)Nzuse SqueezeView.create())AssertionError)rrfrrrr rzSqueezeView.__init__)rrrYrrr)rrrr)rfr rr)rrrrrrnrrrrrrrI s - rc@sXeZdZUded<ded<dddZdd d Zdd d ZeZedddZ dddZ dS) GenericViewrrr\rrcCrr)rrrrrr] szGenericView.make_reindexerrcCsBddtt|jD}t||}ddtt|d|S)NcSg|]}ttj|qSr)rlrCrrrrrrr  z+GenericView.reindex_str..zlambda rtr4)rrrrrrrr)r index_old index_newrrr reindex_str s  zGenericView.reindex_strcCs$||jd|jd|gS)Nsize=zreindex=)rrfrrrrrrr szGenericView.__str__rrrqrcCs||t||dS)Nrfrr)r)rrrqrrrrr szGenericView.createcCrrrrrrrr! rzGenericView.get_sizeNrpr\)rrrqrrr\rrrS) rrrrr]rrrrrr!rrrrr s     rc@sbeZdZdZedddZeedd d ZedddZ e dd ddZ e dd!ddZ dS)"rz This class handles tensor reshaping by computing appropriate index transformations to map the new shape back to the original storage layout. rrrrcCs<t|}t|}tjjjj}|t|dr||}|Sr)rr|rrrrr evaluate_exprLt)rrrrrrhandle_negative_index s   zView.handle_negative_indexrrrqrcsrtts Jt|\tjjr|St t dkp.t t dk}t |}dd d dfd d }dvrVdfdd }|t |dS|ra|t St|si||St|dd\}}|j} tjj} tjj| } tjj} ddlm} | | | | |d}|durdd|D}t|j|j||j|j}t||dS||S)NrrrrqrrorrlcSs6t|dd\}}t|j|j|||j|j}t||dS)NT)rLre)rhrirrrjrkrl)rrqrormrnrprrrcreate_reinterpret_view s z,View.create..create_reinterpret_viewrc sPz}|t|dWSty't|}|tYSw)a Handle the case where view is not possible with current strides. Try dynamic_reshape_indexer first; if it fails with unbacked symbols (guard_or_false can't resolve comparisons), fall back to making the tensor contiguous. r)rrr4rrequire_contiguousr$r%)rrrrrqrvrr"handle_unbacked_or_dynamic_reshape s    z7View.create..handle_unbacked_or_dynamic_reshaperr tuple[int, ...]cstdgtSr)rrr)rvrr fake_reindex rz!View.create..fake_reindexrFrD)_compute_stridesize_obliviouscSs(g|]}t|dr |jjnt|qSr)rrexprrrrrrrr szView.create..re)rrrqrrorrrlrrrr)rr rr)rrrresolve_negative_sizer!rrrrstatically_known_list_equalsrr3rKrr$r%r"rhrto_symints_or_intstorch._subclasses.fake_implsrrirrrjrkrl)rrrqunbacked_symbols_in_sizesrIrrrmrn old_strideold_size_symintold_stride_symintnew_size_symintrnew_stride_symintrorprrrr sZ    z View.creatervtuple[list[Expr], list[Expr]]cCsdd|D}dd|D}t|}tt|D]}||dkr3tjj||<tt|t|||<nqtj j t|t|||fS)NcSg|] }tjj|qSrrrrrrrtrrrr6 z.View.resolve_negative_size..cSrrrrtrrrr7 rrI) rrrrrOner?rmrrrr check_equals)rvrqrrrrr2 s  zView.resolve_negative_sizeNrU dense_dimrrc Cs\z ||||}W|Stttfy-t|g}|||}|||}t||}Y|Swr)_dynamic_reshape_indexerrr4 IndexErrorrmr)rrvrqrrflatrrrrrrC s    zView.dynamic_reshape_indexerr\cstjjjdfdd }dd tt|Dtt|}t|}|d uo3|t|d ko3t|d k}|rF|d us b. For unbacked symbols, guard_or_false returns False, so we fall back to divisibility checks. rrIrD) rrGrGtrrrrrVr4)rr)rOrr compare_sizes_ sz4View._dynamic_reshape_indexer..compare_sizescSrr)rlrCVIEWrrrrr rz1View._dynamic_reshape_indexer..NrDrrrcsHt|tksJt|tftt|tfddDS)Nc3rr)rnrt replacementsrrr rzAView._dynamic_reshape_indexer..reindex..)rrrrr)r view_exprrrr s$z.View._dynamic_reshape_indexer..reindex)rrrrrrr~)rrrrrOrrrrrrrrrrrrBrreverseinsert)rvrqrr stack_new stack_oldreordering_dense_dimold_dimsize_oldvarsize_newvar2 size_new2divisormodulus dense_exprrr)rOrrrrT s $            zView._dynamic_reshape_indexer)rrrrrr)rrrqrrr)rvrrqrrrr)rvrUrqrUrrrr)rvrrqrrrrr\) rrrrkrnrrrrrrrrrrrr s m rcseZdZUdZded<d3fdd Zd4d d ZeZd4d d Zd5ddZ d6ddZ e d7ddZ d8ddZ d8ddZd9ddZd:dd Zd;d!d"Zd3d#d$Zed %dd?d.d/Zd@d1d2ZZS)Arlz*Pretend our storage has a different layoutrrgrrcs2tt|jtrt|d|jdSdS)Nrf)rrrrfrrrrrr rrr s  zReinterpretView.__post_init__rcC||j|jgSr)rrfrgrrrrr s zReinterpretView.__str__cCrdrrgrrrrr rzReinterpretView.get_namercC|jjSr)rgrrrrrr' rzReinterpretView.get_devicercCrrrrrrrr rzReinterpretView.get_origin_nodercCrr)rgrrrrrr rzReinterpretView.dtypercC t|jjSr)rrgrrrrrr! rzReinterpretView.get_sizecCrr)rrgrrrrrrK rzReinterpretView.get_stridercdfdd }|S)NrrrrqcsFj}t||}jjjjkr!t|jjjS|Sr)rgrrploadrrrfto_dtype_bitcast)rr tmp_loaderrrrr s z+ReinterpretView.make_loader..loaderrrrrqrrcrrrr szReinterpretView.make_loaderrcCrdr)rgrrrrrrrzReinterpretView.make_indexercCrrrgrrrrr#rzReinterpretView.get_layoutcCrrrrrrrr) rzReinterpretView.freeze_layoutFrrrcCs*t|jj|t|jj|Bt|jj|BSr)r&rgrrrjr=rrrr>s   z$ReinterpretView.get_free_symbol_usesNr r cCs@tjjj|j|jj|jj|jj|dur|j ntjjj |jj dSr) rrr wrapper_codecodegen_reinterpret_viewrfrgrrrj writelinerr rrrr sz!ReinterpretView.codegen_referencercCrr8rrrrrr;%rzReinterpretView.num_readsrIr\rWrKrOrSrYrZrPr_rhrrVrf)rrrrkrrrrrr'rrqrr!rKrrr#r)r]r>r r;rrrr rrl s,             rlc@sXeZdZUdZded<edd d Zdd d ZeZe dddZ dddZ dddZ dS) DtypeViewz(Pretend our storage has a different typer target_dtyperr new_dtyperrcCsHt|rt|\}}t|j||j|j|j|j}t||dSt ||dS)Nre)rfr) r"rhrirrrrjrkrlr)rrrrmrnrprrrr/s   zDtypeView.creatercCrr)rrfrrrrrr>rzDtypeView.__str__cCrr)rrrrrrCr~zDtypeView.dtypercCrdrrfr!rrrrr!GrzDtypeView.get_sizercsjdfdd }|S)Nrrrrqcst|jjjSr)rprrrfrrr`rrrrMsz%DtypeView.make_loader..loaderrrfrrcrrrrJs zDtypeView.make_loaderN)rrrrrrr\rOrSrY) rrrrkrrrrrrqrr!rrrrrr)s     rc@s2eZdZdZedd d Ze ddddZdS) SliceViewzqView that represents a slice along a single dimension. Corresponds to tensor[..., start:end:step, ...]. rrrYrstartendrtuple[int, int]cstjj||tdd||fDrtjtjntdd||fDr0tjtjnjj dfd d dfdd }||dd}|||}||fS)zz Normalize start and end such that both are in the range [0, x.get_size()[dim]] and start <= end. cs|]}t|VqdSr)r3rtrrrrdrz0SliceView.normalize_start_end..css*|]}t|tr|tjtjVqdSr)rrhasrMinMaxrtrrrrgs rrlowerrupperrcs<||r|n||}||r|}|S||}|Sr)statically_known_geqrX)rrr clamped_lower clamped_full)max_funcmin_funcrrrclampss z,SliceView.normalize_start_end..clamprrr Expr | intcs$|dur|S|}|||Sr)r)rrrr)rrdim_sizerr clamp_wrap~s  z1SliceView.normalize_start_end..clamp_wraprN)rrrrrrrr) rrrrrrrrrr) rrrrr!rrrr evaluate_min evaluate_max)rrrYrrrr)rrrrrrrnormalize_start_endYs"   zSliceView.normalize_start_endrDTsteprrc stttsdksJzdkr"|dkr"dkr"|WSWn ty,Ynwt||r?|||\}t|d<t |r}t |\}}t|j } | | <t |j |j| |j|j |j} t|| dSd fdd } t|| d S) NrlrDrerrrcsDt|tksJd|dt|}||<|S)Nz wrong ndim r)rrrrYrqrrrrrs$z!SliceView.create..reindexrr~)rr|rr TypeErrorrr!rr@r"rhrrirrrjrkrlr) rrrYrrrrrmrnrorprrrrrs8       zSliceView.createN) rrrYrrrrrrr)rDT)rrrYrrrrrrrrrrr)rrrrkrrrrrrrrSs 1rc@sFeZdZUded<ded<dddZdd d Zdd dZdddZdS) BaseConstantrrrrrrcCrNrrrrrrr!rzBaseConstant.get_sizercCrrrrrrrr'rzBaseConstant.get_devicercCrrrrrrrrrzBaseConstant.get_origin_noder7cCrrr<rrrrrrzBaseConstant.get_readsNrSrWrKre)rrrrr!r'rrrrrrrs    rc@sDeZdZUded<ded<ded<dd d Zdd d ZdddZdS)Constantr r-rrrrrrcr)NrrrrqctjjSr)rprr-rrrrrrrz$Constant.make_loader..loaderrrrcrrrrzConstant.make_loaderr3cCrrrrrrrr rzConstant.realizercCt|j|j|dS)N)r-rr)rr-rrrrrrCrzConstant.constant_to_deviceNrYrUri)rrrrrr rCrrrrrs   rc@s:eZdZUded<ded<ded<dd d Zdd d ZdS)IndexingConstantr rrrrrrrcr)Nrrrrqcrr)rpr}rrrrrrrrz,IndexingConstant.make_loader..loaderrrrcrrrrrzIndexingConstant.make_loaderrcCr)N)rrr)r rrrrrrrCrz#IndexingConstant.constant_to_deviceNrYri)rrrrrrCrrrrr s  r cCspd}d}ttt||D](\}}|dkrq tjj||s)tjj||s)dS|td|9}||9}q dSNrDFT) reversedrrrrrrrbrr)rrWexpected_strideexpected_stride_maxryrrris_contiguous_strides_for_shapes rcCs tj|jSr)rEpadding_alignment_bytesitemsizerrrrget_align_for_dtype rrc@s2eZdZdZdddZdddZ ddd dZdS)r6zxAbstract base for Layout, MultiOutputLayout, NoneLayout. Represents the memory layout of the output of an Operation.rrcCrrrrrrrr'rzOutputSpec.get_devicercCrrrrrrr storage_sizerzOutputSpec.storage_sizeFrrrcCrrrr=rrrr>r-zOutputSpec.get_free_symbol_usesNrWrfr_rh)rrrrkr'rr>rrrrr6s   r6c@s8eZdZdZdeddfdMddZedNddZejdOddZedNddZ e jdOddZ edPddZ e jdQddZ dRd d!Z e Z dSd"d#Z dTd%d&ZdUd'd(ZedVd,d-ZdUd.d/ZdWd2d3ZdUd4d5ZedXd7d8ZdYd9d:ZdUd;d<ZdZd>d?Zd[dAdBZd\dEdFZdPdGdHZed d]d^dKdLZdS)_rzo Layout base class Carries tensor meta-information including offset and whether it is pinned. NrFrrrrrrrSequence[Expr] | NonerjrrkrrrcCs|dur t|}||_||_t|t|ks!Jd|d|tdd|Ds,J||_||_||_||_ |j rE|jj dksGJddSdS)Nr , stride=css|] }t|ttfVqdSr)rrrrrrrr7z"Layout.__init__..r9zOnly CPU tensors can be pinned) r$r%rrrr_size_stride_offsetrkr)rrrrrrjrkrrrr(s $ zLayout.__init__cCrrrrrrrrAr~z Layout.sizer-cC ||_dSrrrr-rrrrErcCrrrrrrrrIr~z Layout.stridecCrrrrrrrrMrcCrrrrrrrrjQr~z Layout.offsetcCrrrrrrrrjUrrc Csd}|jdkr d|j}|jjdurdnd|jj}d}|jr'd|j}t|jd|jj|d|jd|jd |j||d S) Nrrz , offset=:z , is_pinned=z('z', z, size=rr5) rjrrrkrrrrr)rrjdevice_index_str is_pinned_strrrrrYs"   "zLayout.__str__cCrrrrrrrr'irzLayout.get_devicercCsPtjtjt|jt|j|j|j|j dWdS1s!wYdS)N)rr pin_memory) rr fake_moder)r*r`rrrrrkrrrr get_examplels$zLayout.get_examplecCst|j|jSr)rrrrrrrrIvrzLayout.is_contiguousrWrUr^cCsVt|}|dvs|ddkrdSt|t||D]\}}}|dkr(||kr(dSqdS)N)rrDFT)rrr.)rWr^ndimleftrightrrrris_channels_last_contiguousysz"Layout.is_channels_last_contiguouscCsJt|jtttt|j|jD]\}}}|dkr"||kr"dSqdSr )rrr r$r%rr)rr'r(rrrr is_transposedszLayout.is_transposedrrcstjtks JddtjD}fdd|D}fdd|Dddd }|d gt}ttD] }||||<q=ttd D]'}||||d k}t|tsrtjj j ||||d kd d }|rwdSqPd S)NcSs*g|]\}}tjjj|dddkr|qS)rrrrD)rrrrr )rrrYrrrrs z,Layout.is_stride_ordered..crrrrrrrrrcrrrrrrrrrarrrrcst|fdd|DS)Ncg|]}|qSrr)relement sorted_arrrrrrzDLayout.is_stride_ordered..sorted_indices..)rH)r+rr.rsorted_indicessz0Layout.is_stride_ordered..sorted_indicesrIrDTrF)r+rrr) rrrrrrrrrr _shape_envr)rr non_1_indicesrr0stride_orderedrrr)rrrrSs*  zLayout.is_stride_orderedcCs:dgtttdt|jd}t|g|}||SNrrD)rr rrrrSr.rrris_channels_last_stride_ordereds" z&Layout.is_channels_last_stride_ordered in_stridescst|}t|dkr |Stjst||r|St}t|dr)|j ddr)|St ddt ||D }tjs>|r>|SttjdrHtjjnddfd d r`tfdd|Dr`|St|}t|}ddtt|D}d||d<d} t|ddddD]?\} } || d} || || } t| ttjfo| tjko| |dkpt| tjotj}| || <|rt| |||| <d} q| s|Stjd7_|S)z The padding does not change stride order but makes sure all strides larger than the threshold are multiple of align. rrzdislike_paddingFcsrrrrrrrrrz&Layout._pad_strides..r1Nrsympy.Expr | intrrcs4durdSt|tjsdStfdd|jDS)NFc3s|]}|VqdSr)is_unbacked_symintrrrrrrzILayout._pad_strides..contains_unbacked_symints..)rrrrr2rr:rrcontains_unbacked_symintss  z6Layout._pad_strides..contains_unbacked_symintsc3rrrr)r<rrrrcSrrrr rrrrr z'Layout._pad_strides..rD)rT)rr8rr)rrrEpad_channels_lastrr)rrget_current_noderrzrrrvchainpad_dynamic_shapesrr1rrrrrrrrrpadding_stride_thresholdrr^r$num_comprehensive_padding)r6rraligncurrent_fx_noderhrMr new_stridespaddedrankrprev_idxrrequire_paddingr)r<rr _pad_stridessZ        zLayout._pad_stridescCs>t|ts Jt||jdusJ||j|j|j|_dSr)rr$rrrJrrrrrrrH szLayout.pad_stridescCstjot|tSr)rEcomprehensive_paddingrr$rrrrrGrzLayout.should_pad_stridesricCs<t|tr|S|r|t|j|j|j|j|j|j Sr) rrirGrHrrrrrjrkrrrras_fixeds zLayout.as_fixedrcCs(tjsJdt|jd|S)Nzconvert z to FixedLayout first)r$rrrrLrrrrrr$s zLayout.make_indexerotherrcCsRt|to(|j|jko(|j|jko(|j|jko(|j|jko(|j|jko(|j|jkSr)rrrrrrrjrk)rrMrrr__eq__*s       z Layout.__eq__cCt|j|j|jSr)r*rrrjrrrrr5rzLayout.storage_sizerrcCs$t|j|t|j|Bt|j|BSr)r&rrrjr=rrrr>8s   zLayout.get_free_symbol_uses)rrrrrrrrrjrrkrrrrSr-rrrrTr-rrrr\rX)rrrR)rWrUr^rUrr)rrrr)r6rrrrrrrrIrrirZ)rMrrrr_rh)rrrrkrrrqrsetterrrjrrr'r$rIrnr)r*rSr5rJrHrGrLrrNrr]r>rrrrrsN               #  N     rc@seZdZdZdddZdS)riz A Tensor layout we cannot changerrcCrO)r)rrrrjrrrrrFrzFixedLayout.make_indexerNrZ)rrrrkrrrrrriCsricseZdZdZdZdDddZedEd d ZedFd dZedGddZ edHddZ edIddZ e dJddZ e jdKddZ e dJdd ZejdKd!d Ze dLd#d$ZejdMd%d$Z dNdOd(d)Z dNdPd+d,ZdQd-d.ZdRd/d0ZdSd2d3ZdTd7d8Z 9 dUdVfdBdC ZZS)Wr$z| A Tensor layout that we are allowed to change Assumption: layout change should NOT add or remove free symbols FrricCst|S)z Compute what the strides would be if this layout were frozen, without actually modifying the layout. This is used for speculative stride computation during Triton template code generation. )rdeepcopyrLrrrr!get_fixed_layout_without_freezingTsz0FlexibleLayout.get_fixed_layout_without_freezingsizesrrcCsNt|dkrgStjjg}t|ddD] }|||dqtt|S)NrrDrI)rrrrr rr)rVreversed_stridesrrrrr%^s   z!FlexibleLayout.contiguous_stridesrcCs\ttt|t|ksJ||ftjj}dgt|}|D] }|||<|||}q|S)z Create a stride based on the order the dimensions should be filled in. In this format, channels last would be: [1, 3, 2, 0] N)r=rrrrr)rVr next_strider^rrrr fill_orderedgs$zFlexibleLayout.fill_orderedrcCs0ttt|t|ksJt|}t||S)z Create a stride based on the sorted order of a permuted range. In this format, channels last would be: [3, 0, 2, 1] )r=rrrr$rY)rVrrrrrr3xs zFlexibleLayout.stride_ordered memory_formattorch.memory_formatcCsP|tjkr t|tS|tjkrt|tS|tjkr t|St d|t )aq Create a stride based on a memory format. Memory format is translasted into a stride order, so channels_last is the same as: FlexibleLayout.stride_ordered(sizes, [3, 0, 2, 1]) This interface does not support memory_format `torch.preserve_format` which should be used to deduce a format from another source z>stride_ordered_for_memory_format, unsuppored memory_format: %s) r) channels_lastr$r3NHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatr%r\rr)rVrZrrr stride_ordered_for_memory_formats      z/FlexibleLayout.stride_ordered_for_memory_formatrrUcCsDt|t|ks Jtjj|}ttt||jd}t ||S)z Create a stride that has the same stride order as given stride For example, if given stride is [1000, 1, 100, 10], the fill order should be [1, 3, 2, 0] rB) rrrrrguarding_hints_or_throwrHr __getitem__r$rY)rVrrrrr same_ordereds  zFlexibleLayout.same_orderedcCrrrrrrrrr~zFlexibleLayout.sizer-rcC|d|||_dS)Nr)!assert_free_symbol_uses_unchangedrrrrrr  cCrrrrrrrrr~zFlexibleLayout.stridecCre)Nr)rfrrrrrrrgrcCrrrrrrrrjr~zFlexibleLayout.offsetcCre)Nrj)rfrrrrrrjrgr*rcCsH||j|}|r|r|||j|j}t|j|j|j||j|jSr) r3rrGrJrrirrjrk)rrr*rorrras_stride_orders zFlexibleLayout.as_stride_orderr2cCs>|}|r|r|||j|j}t|j|j|j||j|jSr)rGrJrrrirrjrk)rr2r*rorrras_exact_stridess zFlexibleLayout.as_exact_stridescCD||j|}|r|||j|j}t|j|j|j||j|jSr) rYrrGrJrrirrjrk)rrrorrr as_fill_orderzFlexibleLayout.as_fill_ordercCrjr) rdrrGrJrrirrjrk)rrrorrr as_same_orderrlzFlexibleLayout.as_same_order$dict[tuple[str, bool], sympy.Symbol]cCs<i}dD]}dD]}||f}ttt|||||<qq|S)N)rrrjTF)r=r&r)rinitial_free_symbolsrrrCrrrget_initial_free_symbol_usess z+FlexibleLayout.get_initial_free_symbol_usesrrr6cCsFdD]}|j||f}tt||}||ks Jd|d|qdS)Nroz)Expected free symbols unchanged, but got z vs )rpr=r&)rrr-rold_free_symbolsnew_free_symbolsrrrrfs z0FlexibleLayout.assert_free_symbol_uses_unchangedNrrrrrrMrNrkcs@|r t||}nt|}tj|||||d||_dS)Nrk)r$rYr%rrrqrp)rrrrrMrkr^r rrrs  zFlexibleLayout.__init__rR)rVrrr)rVrrrrr)rVrrrrr)rVrrZr[rr)rVrrrUrrrSrPrTrQr_)rrr*rrri)r2rUr*rrri)rrrri)rrUrri)rrn)rrr-r6rrr) rrrrrrrMrNrkrrr)rrrrkrrUrnr%rYr3rardrqrrSrrjrhrirkrmrqrfrrrrr rr$KsJ                r$csLeZdZdZdfdd Zdd d Zdd d Zed ddddZZ S)NonOwningLayoutz,Is a view into the storage of another tensorviewBaseView | TensorBoxrrcs,|}t|j|j|j|j||_dSr)r#rrrrrrrv)rrvrgr rrr,s zNonOwningLayout.__init__rcC |Sr)rLrrrrrr6rzNonOwningLayout.make_indexerrcCs4|jj}|dkr dSddlm}tjj||S)NrTrD) ALIGNMENT) rvr#rjutilsryrrrrrV)rrjryrrrmaybe_guard_aligned9s  z#NonOwningLayout.maybe_guard_alignedFrrcCsVt|jtsJ|jj}t|tsJt||j}t|ts%Jt||j|Sr) rrvrlrfrrrrgr>)rrbox input_bufferrrrr>As  z$NonOwningLayout.get_free_symbol_uses)rvrwrrrZrRr_rh) rrrrkrrr{r]r>rrrr rru)s ruc@eZdZdZdS)CommBufferTypesymm_memN)rrrSYMM_MEMrrrrrMrcs4eZdZUdZded<ded<d fdd ZZS) CommBufferLayoutax A layout that signifies the buffer is a comm buffer. In terms of striding, the layout is identical to `FixedLayout`. Buffers with this layout do not participate in in-place reuse - it can be neither the source nor the target for in-place reuse. For detailed motivation and usage of this layout, see NOTE [lowering-time collective optimization]. rcomm_buffer_typer group_namergFlexibleLayout | FixedLayoutcsJt|tr |n|}tj|j|j|j|j|j |j d||_ ||_ dS)Nrrrrrjrk) rr$rLrrrrrrrjrkrr)rrgrrfixedr rrr`s zCommBufferLayout.__init__)rgrrrrr)rrrrkrrrrrr rrQs  rc@sdeZdZUded<ejdddZded<ejdddZded <dd d ZdddZ dddZ dS) NoneLayoutrrcCdgSrrrrrrr@~rAzNoneLayout.default_factoryrrcCrrrrrrrr@rArrrcCrrrrrrrrrzNoneLayout.storage_sizer6cC|SrrrrrrrLrzNoneLayout.as_fixedcCrrrrrrrr'rzNoneLayout.get_deviceNrfrQrW) rrrrrlrmrrrrLr'rrrrrss    rcseZdZd#fdd Zed$dd Zejd%d d Zd&ddZd'ddZd(ddZ e d)d*ddZ d+ddZ d,d!d"Z ZS)-MutationLayoutSHOULDREMOVErrrrcs@t|||d||_|}tj |dSr) rrrr&r!r get_bufferrrrrmark_buffer_mutated)rrrr rrrs z#MutationLayoutSHOULDREMOVE.__init__rcCr8r) real_layoutrrrrrrrz!MutationLayoutSHOULDREMOVE.strider-rcCrrrrrrrrrrcCrxr)rrrrrrrrz'MutationLayoutSHOULDREMOVE.storage_sizercs2dfdd |j}t|tsJt||S)Nrr rcsBt|tr |jSt|tr|St|tr|jS|Sr)rrrrr MutableBoxrf)r unwrap_viewsrrrs      z;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_views)rr rr )rrrr)rrrrrrs z%MutationLayoutSHOULDREMOVE.get_bufferrcCs|j}t|ts J|Sr)rrgrr)rrgrrrrs z&MutationLayoutSHOULDREMOVE.real_layoutFsrcdst unsafe_aliasrc Cs|tj|t|tr|j}||sAt j | | | ddt||Dd}t|ttfs>J|j}|t|dsNJ|t|jjts]Jt|jjt||j_|jS)NcS g|] \}}tjj||qSrrrrrcheck_equals_and_simplifyrrrrrrrz;MutationLayoutSHOULDREMOVE.realize_into..rrf)r rrrrrrrrfr(rrr'r&rrr!rrrrgr$rr)rrrrrrrr realize_intos*   z'MutationLayoutSHOULDREMOVE.realize_intorcCrrrrrrrrLrz#MutationLayoutSHOULDREMOVE.as_fixedrcCrdr)rrrrrrrrz'MutationLayoutSHOULDREMOVE.make_indexer)rrrrrS)r-rrrrT)rrrPr_)rrrrrrrr)rrrZ)rrrrrqrrSrrrrrrLrrrrr rrs      'rcsleZdZUded<ded<d]fdd Zd^d d Zd_d dZd`ddZdaddZdbddZ e dcddZ ddddZ dedd Z dfd"d#Zdgd%d&Zdhd'd(Zdid*d+Zdjd-d.Zd]d/d0Z 1dkdld5d6Zdmd7d8Zdnd:d;Z 1dkdod=d>Zdjd?d@ZdpdBdCZdqdrdGdHZd]dIdJZdsdLdMZdsdNdOZdtdQdRZed 1dkdudUdVZ dvdWdXZ!dwdYdZZ"djd[d\Z#Z$S)xrr3rr6rgrrcst|dddSrx)rrrrr rrrs zBuffer.__post_init__rcCrxr)r#rrrrrrrzBuffer.make_indexerrcCs|jsJ||jSrrrrrrrryzBuffer.get_nametorch.Tensor | torch.SymIntcCs&t|jtr |jStt|jjr)rrgrr$rrrrrrrr$s  zBuffer.get_examplercCrxr)rr'rrrrr'rzBuffer.get_devicercCrrrrrrrrrzBuffer.get_defining_oprcCr8r)r#rrrrrrrz Buffer.dtypercCg|jSr)r#rrrrrr!rzBuffer.get_sizercCrr)r#rrrrrrK rzBuffer.get_stridercCr8r)r#rjrrrr get_offsetrzBuffer.get_offsetrcCs"t|jtr |jStt|jjr)rrgrrrrrrrrr#s zBuffer.get_layoutcCrrrrrrrrrzBuffer.get_output_specrcCrr)rrrrrr<rzBuffer.get_storage_numelrcCr8r)r#rkrrrr get_is_pinnedrzBuffer.get_is_pinnedcCs0t|jtrt|jts|j|_dSdSdSr)rrgrrurLrrrrr)s zBuffer.freeze_layoutFrrr*cC0t|jts Jt|j|jj||d|_dSNrQ)rrgr$rrhr+rrrr,%sz&Buffer.freeze_layout_with_stride_ordercC,t|jts Jt|j|j||_dSr)rrgr$rrkr.rrrr/+z$Buffer.freeze_layout_with_fill_orderrcCrr)rrgr$rrmr0rrrr1/rz$Buffer.freeze_layout_with_same_orderr2cCrr)rrgr$rrir3rrrr43s z'Buffer.freeze_layout_with_exact_stridescCrrrrrrrr;rzBuffer.is_zero_elementsrcs*r ttdSdfdd }|S) Nrrrrrqcs}tjp d||Sr)rrprrrrrrrrCsz"Buffer.make_loader..loaderr)rr rr&rcrrrr>szBuffer.make_loaderNr r cCrrrr rrrr IrzBuffer.codegen_referencecCrrrrrrrrJLrzBuffer.decide_layoutrcCt|jtr |jjgSdSr)rrgrurvrrrrrrFO z#Buffer.get_inputs_that_alias_outputcCrr)rrgrrrrrrrrDTrzBuffer.get_mutation_namesrcCst|gSr)r=rrrrrrYrzBuffer.get_read_namesrrcCrrr<r=rrrr>\zBuffer.get_free_symbol_usescCrrr<rrrrrbrzBuffer.get_unbacked_symbol_defscCrrrrrrrr erzBuffer.realizecCrrrrrrrshould_allocatehrzBuffer.should_allocaterIrZr\)rrrWrLrOrS)rrrTrPrQrfrRr_r`ra)rrrr)r2rr*rrrrYrrVrjrJrhrrU)%rrrrrrrr$r'rrqrr!rKrr#rr<rr)r,r/r1r4rrr rJrFrDrr]r>rr rrrrr rrsL                         rc@s0eZdZd ddZd ddZejZdd d Zd S)OperationBufferrrcCs|gSrrrrrrrprzOperationBuffer.get_outputsrrcCrrrrrrrrsrzOperationBuffer.get_defining_oprcCst|t|dSr)rrrrrrrrry zOperationBuffer.__post_init__NrrrrrI)rrrrrrrrErrrrrrms   rc@eZdZdddZdS)rrrcCrr8rrrrrr;rzInputBuffer.num_readsNrf)rrrr;rrrrr~rc@r~) DonatedBufferaY Represents a donated buffer which is a saved tensor that is not alias to any fwd inputs, fwd user outputs, and bwd outputs. We generally cannot inplace reuse the input tensor memory during backward since it might be used in another function. However, donated buffer can be inplace reused during backward to save memory. N)rrrrkrrrrrrrc@s.eZdZUdZded<d ddZdd d ZdS)rNrrrrcr)Nrrrrqcs,}ttjj||Sr) r#rrprrrr constant_namerrrrrrrs z*ConstantBuffer.make_loader..loaderrrrcrrrrszConstantBuffer.make_loaderrrrcCsttj|||jdSNrrg)rrrrrrrgrrrrrCsz!ConstantBuffer.constant_to_devicerYri)rrrrrrrCrrrrrs   rc@sLeZdZdddZed ddd d ZddddZdddZdddZd S)NoneAsConstantBufferrr7cCrrr<rrrrrrzNoneAsConstantBuffer.get_readsFrrrcCrrr<r=rrrr>rz)NoneAsConstantBuffer.get_free_symbol_usesNr r rcCs tjjjSr)rrrrnone_strr rrrr rz&NoneAsConstantBuffer.codegen_referencer6cCs tddSNr)rrrrrrrz$NoneAsConstantBuffer.get_output_speccCrrrrrrrrrz&NoneAsConstantBuffer.has_tensor_outputrer_rhrrVrQrR) rrrrr]r>r rrrrrrrs   rc@sBeZdZUded<ed dddd ZddddZdddZd S)rrrFrrrrcC t|j|Sr)r&rr=rrrr> z*ShapeAsConstantBuffer.get_free_symbol_usesNr r rcCstjj|jSr)rrrrcodegen_sizevarrr rrrr rz'ShapeAsConstantBuffer.codegen_referencecCrrrrrrrrrz'ShapeAsConstantBuffer.has_tensor_outputr_rhrrVrR)rrrrr]r>r rrrrrrs  rcsheZdZUdZded<dZded<dZded <dZd ed <dZd ed <dZ d ed<e j d[ddZ e e j d[ddZd\ddZd]ddZd^ddZd_dd Zd`d"d#Zed dadbd'd(Zdcfd*d+ Zddd,d-Zded/d0Zdfd2d3Zdgd5d6Zedhd8d9Z  didjd>d?Ze dkdldIdJZdmdLdMZdmdNdOZ d\dPdQZ!dddRdSZ"dddTdUZ#dndYdZZ$Z%S)orzb Represents a buffer that is computed during kernel execution rather than being an input. rrfFzClassVar[bool]_force_realizeNrrCallable[..., Any] | NonerrrrrIterator[None]c cs|jdusJ|jdusJ|jdusJ|jdusJt|jts+Jt|j|j}|j}z3t|j |j |j|j|j|j |j |j d}||_t|j |j |j|_|j|dVW||_||_dS||_||_w)Nr+)rrrrrrfrrrgrrrrrriget_default_sizes_body clear_cache)rold_datarnnew_datarrrwith_original_inner_fns>   z%ComputedBuffer.with_original_inner_fnccs*tj}z dt_dVW|t_dS|t_wNT)rr) old_valuerrr force_realizes zComputedBuffer.force_realizer3cCs(|jdur|jSt|jdr|jjSdS)z Returns self.name if it exists, otherwise returns the name of the data node if that exists. If neither exist, returns None. Nr)rrrfrrrrget_computed_buffer_names  z'ComputedBuffer.get_computed_buffer_namercCrdrrfr;rrrrr; rzComputedBuffer.num_readsr7cCrdrrfrrrrrrrzComputedBuffer.get_readsrcCrdrrorrrrrrzComputedBuffer.get_read_namesr5cCst|jttttfstjtttdSt t dd,|j r7t ||j|jWdSt ||jWdS1sKwYdS)Nr:writes index_exprsrT)rrfrrr7rrF ReadWritesr=rrr$r?rOget_store_functionrr@r!rrrrr6s& $zComputedBuffer.get_read_writesrrrcCs6|j||j|B}|r|||O}|Sr)rgr>rfhas_store_functionr6)rrrrrrr>+s z#ComputedBuffer.get_free_symbol_usesrcs<|s|jtjjvr|dkr|js|jSt Sr) r?rrrrmutated_buffersr;rrfrrrr rrrEs   zComputedBuffer.make_loadercCst|jttttfSr)rrfrrr7rrrrrrPrz!ComputedBuffer.has_store_functionCallable[..., None]cCs`|}t|jtttfrt|jj |j |St|jt s'Jt |jt|jj |j |Sr)r#rLrrrfrrr7r r%rrrrrarrrrSs z!ComputedBuffer.get_store_functionlist[int] | Nonecst|jtrYt|j|j\\}}|j }t dd|Ds&Jfdd|D}|rYt|jt t frA|j |n|fdd|D}ddlm}|||SdS) al If our layout is still flexible, try to determine the stride order based on stride orders of reads. TODO(jansel): A better algorithm here would look at downstream consumers of this value and try to do global graph-level layout optimization. This is also something just begging to be autotuned. css"|] }t|tjtjfVqdSr)rrFStarDep MemoryDepr?rrrrj  z0ComputedBuffer.get_fill_order..cs.g|]}t|tjrt|jddDqS)cSsi|] }|dkr|tjjqSrrrvrrrroz..)rrFrrnrr?)r"rrrns  z1ComputedBuffer.get_fill_order..csg|] }tjj|qSrrrrrr_rr)rMrrrysrDpick_loop_orderN)rrgr$rFr]rfrr@r6r:rrr7r schedulerrr!)r index_varsr r:stride_lengthsrr)rMr"rr[s*     zComputedBuffer.get_fill_orderrcCs6t|jtr|}|r||dS|dSdSr)rrgr$rr/r)r.rrrrJs  zComputedBuffer.decide_layoutMtuple[tuple[list[Expr], list[Expr]], LoopBody, tuple[list[Expr], list[Expr]]]c Cstj||dd\}}ttd|t| | r"|n|dd|g|R}Wdn1s8wYg}g}g}g}| D]+\}} ||dvrb|rWJ| || | qI||dvsjJ| || | qI||f|||ffS)NqrrrDr) rFr]rr@rrrr'rRrr?itemsr) rr var_rangesrr reduce_vars index_size reduce_sizerrrrrrs2       z%ComputedBuffer.get_default_sizes_bodyextra_indexing_constraints'tuple[dict[Any, Any], list[Any]] | Nonerecompute_sizes_body_func5tuple[tuple[list[Expr], list[Expr]], LoopBody | None]cs\\}}}\}}|r|||f|||f\\}}}\}}g|j|durut|tr4t|dks6J|\}} t|tsEJt|t| tsPJt| t dd| Ds[J|j } | |kshJ| |ffdd| D} | 7g| t j tjs|dfdd } ||} tt ptj } | || || \}}}| || || \}}}tj||dd\\}}}t|||||g|||}||f|fS)an This is a main place where we do loop transformations in a backend-agnostic way. Here we: 1) Remove any 1 dimensions 2) Fuse contiguous dimensions together 3) Reorder dimensions based on stride orders Optional argument extra_indexing_constraints can be used to append additional indexing expressions to existing ones derived from buffer's body. This can be useful to fuse scheduler nodes with compatible ranges, e.g. (s0*s1*...,) and (s0, s1, s2, ...) on CPU by preventing indexing simplifications and obtaining index/reduce ranges for the scheduler node compatible with other nodes. Optional argument recompute_sizes_body_func can be used to recompute sizes and body on the default body. This can be useful to append additional loop transformations. Nrcs|]}t|tVqdSr)rr)rfrrrrrz6ComputedBuffer.simplify_and_reorder..csg|]}|vr|qSrrr)index_formulasrrrz7ComputedBuffer.simplify_and_reorder..x_varsSequence[sympy.Symbol] support_varsrVrsimplify_loopsrrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]]c s||\}}}dkr8tdkr8ttt}||ddkr8fdd|D}t|}t|}||}|rTtjj ||t ||\}}} t ||} n|} || |fS)NrrrcrrrrrVrrrrzUComputedBuffer.simplify_and_reorder..simplify_and_reorder..) _apply_loop_reorderingr?rrrrrrrrr_simplify_loopsrJr) rrrVrnewsizesreindex0rrr_prunerr memory_addrsrrrsimplify_and_reorders&     zAComputedBuffer.simplify_and_reorder..simplify_and_reorderpr) rrrrrVrrrrr)rindexing_exprsrrrrrrrrrget_write_exprsrrrrSrGPREFER_STORE_LOOP_ORDERextendget_read_exprsrir7rEloop_ordering_after_fusionrFindex_vars_no_squeezerR)rrrrrrrrextra_indexing_rangesextra_indexing_exprexpected_var_rangesrrshould_merge_loops iter_ranges iter_reindexr  reduce_rangesreduce_reindex iter_varsrrrrrsx   6    z#ComputedBuffer.simplify_and_reorderrrrrVrrlist[sympy.Expr] priority_idxrc sddlm}|dur g}z*fdd|D}t|t|kr)t|dtks+Jtt|||}Wn tyVtjrLt dt t |tt t}Ynwfdd|Dt |t|fS) zU Shuffle the order of loops around to hopefully improve performance. rDrNcsg|] }tjj|qSrrr)rrrrrWz9ComputedBuffer._apply_loop_reordering..rz%Did not simplify complex index: %s %scrrrrrrrrgr)rrrrr  ExceptionrErr\warningrrrrr)rrrVrrrr^rr)rrVrrrBs,     z%ComputedBuffer._apply_loop_reorderingrcCrdr)rfrrrrrrjrz!ComputedBuffer.get_pointwise_sizecCrdrrfr@rrrrr@mrz!ComputedBuffer.get_reduction_sizecCrdrrfr?rrrrr?prz!ComputedBuffer.get_reduction_typecCrdr)rfrrrrrrBsrzComputedBuffer.is_no_opcCrrrrrrrrvrzComputedBuffer.should_allocaterrrcCrY)rrfrCrrrrrCyr[z!ComputedBuffer.constant_to_device)rrrUrfrerJrdr_rhrYrR)rr)rrrI)rrNN)rrrrrrr) rrrrrVrrr rrrrrSri)&rrrrkrrrrrrrorprrnrrr;rrr6r]r>rrrrrJr\rrrrr@r?rBrrCrrrr rrsR       "       ' "  '    rc@s2eZdZUdZded<ded<ded<ded<dS) FinalizeCodegenResultzNStructured result from TemplateBuffer._finalize_codegen for external backends.rsourcerimports call_preamble call_argsN)rrrrkrrrrrr~s  rcseZdZdZ   dKdLfdd ZedMddZdNddZdOddZdPdQdd Z dRd"d#Z dSd%d&Z dTd'd(Z  dUdVd-d.Z dTd/d0ZdWd2d3ZdXd7d8ZedYddZdIdJZZS)[rz Base class for template operators that support epilogue and prologue fusion. Subclasses: TritonTemplateBuffer (built-in Triton templates), HelionTemplateBuffer (Helion kernels), etc. Nrgr6r~rmake_kernel_renderrmutated_inputsIterable[IRNode] | Noneallowed_prologue_inpsOrderedSet[str] | None named_inputsdict[str, IRNode] | Nonerrcstjd|dt|_|_tj_ tj i_ i_ i_ |r-t|ni_|_g_|durYjd}t|tsJJt||fdd|D_|p]t_d_d_dS)Nrrcg|] }ttd|qSrMutationOutputrrrOrrrrrrz+TemplateBuffer.__init__..)rrrunwrap_storager~rrrrregister_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsrmutation_outputsrrrr'r=rallow_epilogue_fusionallow_prologue_fusion)rrgr~rrrr  first_inputr r'rrs.     zTemplateBuffer.__init__rcCst|jtr td|jS)Nz1Multi-output templates do not have a single dtype)rrgMultiOutputLayoutrr#rrrrrrs  zTemplateBuffer.dtyper5cCs |jddS)NT normalize)rOrrrrr6rzTemplateBuffer.get_read_writesr4rOrderedSet[dependencies.Dep]cst}|jD]9}t|ttfsJt||tjts$Jtjjd fdd }|t j | d|d j O}q|S) z(Build read dependencies from all inputs.r Sequence[Any]r*rr cs$t|dksJt|Sr)rrprrrr inp_indexerrrdummysz4TemplateBuffer._read_deps_from_inputs..dummyrr3Nrr6r*r6rr ) r=r~rrlrrrgrrrFrOr!r:)rr4r:inp_rawr9rr7r_read_deps_from_inputss   z%TemplateBuffer._read_deps_from_inputsFcst|jtr%ttj|tddddg}tj | ||tdddS|| dfd d }tj ||d|d }|j| |O_|S)aExtract read/write dependencies for this TemplateBuffer. When the layout is MultiOutputLayout (multi-output templates), the buffer itself has no data layout, so we cannot build an indexer. Instead, synthesize a trivial write dep and derive read deps from the named tensor inputs (``_named_inputs``). For single-output templates with a concrete layout, fall through to the standard path. rr) var_namesrN)r:rrrGrrr6r*rr c"t|dksJt|dSNrfakerrprrrrrrr9z1TemplateBuffer.extract_read_writes..dummyr3r:)rrgr2r=rFrrrrrr<r#rrOr!r:)rr4rr9depsrrBrrOs,   z"TemplateBuffer.extract_read_writesrcCstjjSr)rrrrrrrr@rz!TemplateBuffer.get_reduction_sizer3cCrrrrrrrr?rz!TemplateBuffer.get_reduction_typecCrrrrrrrr rzTemplateBuffer.should_allocaterrr9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]cCs|gfdfSrr)rrrrrrr s z#TemplateBuffer.simplify_and_reordercCs t|jtS)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrgr2rrrris_multi_outputs_templater[z(TemplateBuffer.is_multi_outputs_templatercCrr)rrrrrget_allowed_prologue_inpsrz(TemplateBuffer.get_allowed_prologue_inps hook_outputsdict[str, str]FinalizeCodegenResult | NonecCr)aCalled after epilogue/prologue subgraph codegen with rendered hook outputs. ``hook_outputs`` maps placeholder keys (e.g. ````, ````) to the Triton code generated by Inductor for each fused subgraph. Return a ``FinalizeCodegenResult`` to provide custom source code and call metadata, or ``None`` to use the default codegen path. Nr)rrHrrr_finalize_codegen!s z TemplateBuffer._finalize_codegenrrrcCsNt|trt|jtr|jSt|}t|tr|j}t|jtr%| |S)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrrf MultiOutputr realize_inputrrgr$r))rrrrrrrealize_template_input/s   z%TemplateBuffer.realize_template_input)direct_alias_at_leafon_tensor_leafon_non_tensor_leaf template_buf structuredrrOdict[int, IRNode] | NonerPFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | NonerQCallable[[int], None] | Nonetuple[TensorBox, ...]cs4itd fdd t|gS) zLWalk a structured output tree, creating MultiOutput nodes for tensor leaves.rtrrMlist[tuple[type, int]]rlist[TensorBox]c st|ttfr$g}t|D]\}}||g|t||fq |St}t|tjrqr<|vr.walkN)rtrrMrXrrY)rvcountr)rrRrSrOrPrQrrbrbuild_multi_outputs;s z"TemplateBuffer.build_multi_outputs)NNN)rgr6r~rrrrrrrr r!rrrOrd)r4rrr5r_)r4rrr5rSrUrRr)rrrrrrErJ)rHrIrrJ)rrrr) rRrrSrrOrTrPrUrQrVrrW)rrrrkrrqrr6r<rOr@r?rrrFrGrKrrNrgrrrr rrs6 2     &      rcsReZdZ  ddfdd Zed ddfdd Zd ddZd!ddZZS)"TritonTemplateBufferNrgrr~rrCallable[_P, _T] | NonerrrrrrcsBtj|||||d|jdusJ|j|ji|_d|_d|_dS)a NOTE:[TritonTemplates with multiple outputs] We want the ability for TritonTemplates to output multiple tensors. Triton kernels have no notion of outputs and this is done by creating tensors that are then mutated by the kernel. Currently our STORE_OUTPUT codegen doesn't support creating multinode outputs for triton templates. We work around this by creating an extra input buffer during the lowering and we mark them as mutated inputs. )rrN)rrrr+ subgraph_inps subgraph_outs)rrgr~rrrr rrris zTritonTemplateBuffer.__init__Frrrcst|}|jr |jng}|jr|jng}|D]%}t|tjr)|t||qt|t r7|||q|dus=Jq|D]}t|t rP|||q@|dusVJq@|Sr) rr>rkrjrrrupdater&r)rrresrkrjrr r rrr>s    z)TritonTemplateBuffer.get_free_symbol_usesrcC |g|jSrr.rrrrrrz TritonTemplateBuffer.get_outputsrcCsd|jd}|S)NzTritonTemplateBuffer(layout=r5r)rr rrrrszTritonTemplateBuffer.__str__r) rgrr~rrrirrrrrrr_rhrr\) rrrrr]r>rrrrrr rrhhs rhcs|eZdZdZd'fd d Zd(ddZd)ddZd*ddZd)ddZd)ddZ d+ddZ d,d!d"Z d)d#d$Z d-d%d&Z ZS). ChoiceCallera1 Represents a possible choice used in autotune_process.py. During autotuning, self.benchmark() is first called to get benchmark result, and if this choice is selected, self.output_node() is called to get the output_node. Children classes: TritonTemplateCaller, CUTLASSTemplateCaller. rrrrrgr descriptionrrcsPt||_||_||_||_d|_d|_i|_d|_ d|_ i|_ i|_ dSr) rrrrgrrqfailed_benchmark_with_cudagraphsrrr decompositiondecomposition_kwargsconfig_patches)rrrrgrqr rrrs  zChoiceCaller.__init__rr r rrcsP||jrtfddStjrtfddStjd|iddS)NcSrrralgorrrr@z(ChoiceCaller.benchmark..crwrrrrxrrr@rzr r) to_callablersrWbenchmark_gpu_with_cuda_graphrE/profile_bandwidth_with_do_bench_using_profilingrb benchmark)rr rrrxrr~s zChoiceCaller.benchmarkcCrurrvrrrr call_namerzChoiceCaller.call_namercCrurrvrrrrr{rzChoiceCaller.to_callablecCr)z Hash key for the underlying kernel. By default, we assume there are no runtime params, so kernel hash key defaults to choice caller's hash key. )hash_keyrrrrkernel_hash_keyszChoiceCaller.kernel_hash_keycCrurrvrrrrrrzChoiceCaller.hash_keyrcCrurrvrrrrr|rzChoiceCaller.output_node6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]cCsiS)zRInformation returned here is logged to the autotune log file when that is enabled.rrrrr info_dictrzChoiceCaller.info_dictcCr)Nunsupported_choicerrrrrautoheuristic_idrzChoiceCaller.autoheuristic_idcCs d|_dS)z Mark the choice as failed so that it can be removed later. Useful for when we decouple compilation and tuning. TN)rrrrrr mark_faileds zChoiceCaller.mark_failed) rrrrrgrrqrrr)rr r rrrr\)rr)rr)rrrI)rrrrkrr~rr{rrr|rrrrrrr rrps        rpc@r)TritonTemplateCallerBaserr cCrurrvrrrrget_make_kernel_renderrz/TritonTemplateCallerBase.get_make_kernel_renderN)rr )rrrrrrrrrrrcs~eZdZdZd)fdd Zed*ddZed+ddZ d,d-ddZe j d.ddZ d/d d!Z d,d0d#d$Z d1d'd(ZZS)2MultiTemplateBufferaG Represents a Buffer with multiple backing implementation choices. Choices can be TritonTemplates or ExternKernels. During scheduling if there is a potential epilogue we will benchmark each of the choices with the epilogue to determine an implementation. Otherwise, the fastest base choice will be chosen. rgrr~rchoice_timings_fn1Callable[[int | None], dict[ChoiceCaller, float]]unfiltered_choiceslist[ChoiceCaller]rrrrcsJtj||d|d||_i|_||_||_tdd|D|_i|_dS)N)rgr~rrcss0|]}t|tpt|tjjjo|jVqdSr)rrr)rselect_algorithmExternKernelCallerhas_out_variant)rchoicerrrrs  z/MultiTemplateBuffer.__init__..) rr_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders)rrgr~rrrr rrrs zMultiTemplateBuffer.__init__rcCr)z^ Are all possible choices TritonTemplates or Extern Kernels with out variants )rrrrroutput_plannable#sz$MultiTemplateBuffer.output_plannablecCrr)rrrrrrX*r~zMultiTemplateBuffer.choicesN hint_overriderdict[ChoiceCaller, float]cCs$||jvr |||j|<|j|Sr)rr)rrrrrchoice_timings.s  z"MultiTemplateBuffer.choice_timingscallerrrccsZt|tjjjsJt||j|jksJ|j}||_z dVW||_dS||_wr) rr)rrTritonTemplateCallerrrgrr)rrrenderrrrswap_as_triton_caller5s  z)MultiTemplateBuffer.swap_as_triton_callercCsRt|tjjjsJt|||jjksJ| |jj ks"J| |_ dSr) rr)rrrrr!rgrrKrrr)rrrrrfinalize_as_triton_callerCs z-MultiTemplateBuffer.finalize_as_triton_callertuple[ChoiceCaller, float]cCs&|j|d}t||jd}|||fS)N)rrB)rrr)rrtimings min_choicerrrget_min_choiceKs  z"MultiTemplateBuffer.get_min_choicecallers*dict[int | None, TritonTemplateCallerBase]cCs0|D] \}}||j|<q|jd|_dS)z;Finalize with multiple callers for different hint overridesN)rrrr)rrrrrrrfinalize_as_triton_callersRsz.MultiTemplateBuffer.finalize_as_triton_callers) rgrr~rrrrrrrrrrR)rrr)rrrr)rrrr)rrrr)rrrr)rrrr)rrrrkrrqrrXrrorprrrrrrrr rrs      rcs2eZdZdfdd ZdddZdddZZS)CUTLASSTemplateBufferrgrr~rrCallable[_P, _T]workspace_sizertemplatervsupports_epilogue_fusionrrrcs&t|||||_||_||_dSr)rrrrr)rrgr~rrrrr rrr^s  zCUTLASSTemplateBuffer.__init__cCs|jdur|jSdSrrrrrrrmrz(CUTLASSTemplateBuffer.get_workspace_sizecCs$|D] }t|ddqdSr)rrprr)rrtrrremulate_store_fnps z&CUTLASSTemplateBuffer.emulate_store_fn)rgrr~rrrrrrrvrrrrrfrI)rrrrrrrrrr rr]s rcs,eZdZdfd d Zdfdd ZZS)CppTemplateBufferrgrr~rrrrrvrr rrcs&t|||||_||_d|_dSr)rrrroutputs)rrgr~rrrr rrrvs zCppTemplateBuffer.__init__cspt|jtr3t|jtsJt|j|jd}t|ts#Jt||j}t|ts1Jt||St Sr) rrgr2rrrrrrr#)r first_outputrgr rrr#s   zCppTemplateBuffer.get_layout) rgrr~rrrrrvrr rrrP)rrrrr#rrrr rrus rcs0eZdZdZ ddfdd ZdddZZS)CuteDSLTemplateBufferz Buffer for CuteDSL (CUTLASS Python DSL) template kernels. Similar to other template buffers but specialized for CuteDSL operations. Nrgrr~rrrrr rrrrcst||||_|_g_|dur?tjdts'Jtjdjd jfdd|D7_dSdS)Nrcr"r#r$r&r'rrrrz2CuteDSLTemplateBuffer.__init__..) rrrrrrr~rrr')rrgr~rrrr r'rrs"zCuteDSLTemplateBuffer.__init__rcCrrrrrrrrrz!CuteDSLTemplateBuffer.get_outputsr) rgrr~rrrrr rrrrr)rrrrkrrrrrr rrs  rcsPeZdZdZ     d"d#fdd Zd$ddZd%ddZ d&d'd d!ZZS)(NVUniversalGemmBufferz Buffer for NVIDIA Universal GEMM kernels. Unlike CuteDSL templates which use Jinja templates, this generates simpler Python code that directly calls the cutlass_api library. rNrgrr~rkernelr accumulator_typevariantrr scale_type_a Any | None scale_type_bswizzle_type_aswizzle_type_brrc sjtj||dd||_||_|g|_||_||_||_||_| |_ | |_ |j j |j j d|_|j|_dS)N)r) kernel_namemin_cc)rrrrrrrrrrrmetadatarrkernel_metadata_make_kernel_renderr) rrgr~rrrrrrrrr rrrs zNVUniversalGemmBuffer.__init__cCr)z#Return the workspace size in bytes.rrrrrrr~z(NVUniversalGemmBuffer.get_workspace_sizercCrrrrrrrrrz!NVUniversalGemmBuffer.get_outputsout_noderrtuple[Any, Any]c sddlm}ddlm}g}|jD]}t|tr|j}t|tr#|j}| |qt |j }|||||j |j |j|j|j|j|j|jd fdd}|fS)z Create a kernel renderer for code generation. Returns (kernel, render) tuple where: - kernel: NVUniversalGemmKernel object with call_kernel() method - render: function that returns source code string r)NVUniversalGemmKernel) Placeholder) rrr|rrrrrrrrcsSr)rr render_kernelrrrrz9NVUniversalGemmBuffer._make_kernel_render..render)Btorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelrtorch._inductor.utilsrr~rrrfrrr KERNEL_NAMErrrrrrrr) rrrrrrrrrrrrrs2       z)NVUniversalGemmBuffer._make_kernel_render)rNNNN)rgrr~rrr rr rr rrrrrrrrrrrrrfrr)rr rrrr) rrrrkrrrrrrrr rrs ! rr#Sequence[IRNode | Sequence[IRNode]]TypeIs[Sequence[IRNode]]cCstdd|DS)Ncsrrrrrrrrrrz#is_node_sequence..)r)rrrris_node_sequence rrc@szeZdZUded<d!ddZd"d d Zd#d dZed$ddZe d%ddZ d&ddZ d'ddZ e d d(d)ddZd S)*rrr~rrrrcCs |j|}t|ts J|Srr~rrr)rrinputrrr input_names zInputsKernel.input_namer5csttj}tj|jD]#}t|tr |fdd|Dq t|tr&q | | q ttjfdd| D}tj ||tdS)Nc3|] }|VqdSrrrtrrrr rFz/InputsKernel.get_read_writes..c3rrrr&rrrr'  r) r=rFrLrr~rrrlrrrrr)rr:rrrrrr6s    zInputsKernel.get_read_writesr7cCr8rr9rrrrr1rzInputsKernel.get_readsrrcCs~t|tr|j}t|tr|j}t|trt|tst|}t|tr)||St|t r0|St|t tfs=Jt ||Sr) rrrfrrrlrrMunwrap_storage_for_inputTorchBindObjectrrrrrrrr4s      z%InputsKernel.unwrap_storage_for_inputlist[IRNode | Sequence[IRNode]]cCs@g}|D]}t|trdd|D}nt|}||q|S)NcSrzr)rrrrrrrNrz/InputsKernel.unwrap_storage..)rrrrr)r~ inputs_newrrrrr(Gs   zInputsKernel.unwrap_storagercCrrrrrrrrATrzInputsKernel.is_externcCrr8rrrrrr;WrzInputsKernel.num_readsFrrcCsNttj}|jD]}t|tr|||O}q |D] }|||O}qq |Sr)r=rr r~rrr>)rrrr inner_inprrrr>Zs   z!InputsKernel.get_free_symbol_usesN)rrrrrdrer)r~rrrrRrfr_rh)rrrrrr6rrrrnr(rAr;r]r>rrrrrs       rc@s eZdZd ddZd ddZdS) NopKernelrrcCrrrrrrrrBirzNopKernel.is_no_opr7cCrrr<rrrrrlrzNopKernel.get_readsNrRre)rrrrBrrrrrrhs rc@s^eZdZdZeddd Ze ddddZed ddddZed ddZ d!ddZ d S)" ConcatKernelzn There isn't actually a real kernel for concat, we just change the storage for the upstream data. r~rrYrrrc Csz|d}|d}t|d}dg}||g}d|kr)t|ks,JJtdt|D]Z}||} |||t| t|ksLJ|||ksVJ|||ks`Jtt|D]} | |krw|| | | || <qftjj || | | || <qf|||q3t |} t jrt| ||dj} tt|D]!}||} t| r| } t| trt| j| jrt|} nqtdd|D}tjjjd}|durt|trtdd|Drt|} tdd|D}|dusJtdt|||| |d gd }t|}g}t |D]{\}}t|t!t"fs'Jt#||$|t%j&||||||dd }t|t'sEJt#|t|j(tsSJt#|j(|j(|t|j)t!rf|j)*}n|j)}t|tr|+r|}durt,|j#rt-|s||.qt|dkrtj/|t0j1rtj2|tj3||_4|5|j(|_(tj6||S) z6 Create the concat kernel from inputs rrDcsrr)r"rtrrrrrz&ConcatKernel.create..FcssB|]}d|jvo|jdjtjdp|jdjtjdVqdS)rrZN)rzrIr)r\r^rargrrrrs   css"|] }t|o |jVqdSr)r"r#rkrtrrrrs N)rrrrrkrrgr~r)7r'r&rr!rrrrrrrrr$r%rErKrrJrr"r#rrir)rrr.r current_noderrrrrrrrrrrrr~rfrr!rirhrErSrGFOREACHregister_operation_listr)rr(r*)rr~rYrrrq offsets_start offsets_endr input_sizern output_striderrgany_input_is_storage_and_layout fx_node_argsrk concat_kernelrop_namesrr}input_unwrappeddevrrrrvs                 zConcatKernel.createNrrrrrcCst|tr ||j|St|ttfsJt|t|jtrPt|jjt r*|jj s,dS|dur2dSt | t | kr@dSt ddt| | DSt|jdoct|jjtoct|jt S)NFTcsr`rrarcrrrrrdz=ConcatKernel.can_realize_into_without_copy..rg)rrcan_realize_into_without_copyrfrrrrrgrirrrKrrrr$ExternKernelAlloc)rrrrrrrs*      z*ConcatKernel.can_realize_into_without_copyFrrcCs t||Sr)rr>r=rrrr>rz!ConcatKernel.get_free_symbol_usesc Cst|tst|rt|\}}t||d}t|ts Jt|t|tr,||j|St|trL| t |jds=J| ||rLt ||j_ |jStj|||ddt||Dd}|||S)NrergcSrrrrrrrr3rz-ConcatKernel.realize_into..r)rrlr"rhrrrrfrr rrrurgrrr'r&rrr!)rrrrmrgpwrrrrs,        zConcatKernel.realize_intocCrrrrrrrr:rzConcatKernel.should_allocate)r~rrYrrrr)rrrrrrr_rh)rrrrrrrR) rrrrkrrrr]r>rrrrrrrps w# !rcseZdZUdZdZded<ejedZ ded<dZ d ed <dZ d ed <dZ d ed <eje dZded<dZded<dZded<ejedZded<dZded<ejedZded<eje dZded<       ddfd$d% Zdd'd(Zdd*d+Zdd,d-Zdd.d/Z ddd3d4Zdd5d6Zddd7d8Zdd9d:Zdd;d<Zdd>d?Ze ddCdDZ!e"ddJdKZ#e"ddMdNZ$e"ddOdPZ%e"ddQdRZ&e"   SdddZd[Z'e" Sddd]d^Z(e" Sddd`daZ)e"ddbdcZ*e"ddddeZ+e"ddfdgZ,e"ddhdiZ-ddjdkZ.ddldmZ/dddqdrZ0ddsdtZ1ddvdwZ2dddydzZ3dd{d|Z4dd}d~Z5dddZ6dddZ7dddZ8dddZ9e:d SddddZ;dddZS)rz A class that represents Kernels which are not directly lowered to Inductor Loop Level IR, such as custom operators, or aten operators which we fallback to. rr6 constant_argsrrrNReinterpretView | None output_viewr3python_kernel_namecpp_kernel_name Iterable[str]ordered_kwargs_for_cpp_kernel_OpOverloads | None op_overloadzlist[dict[str, Any]] | Nonearg_propertieszdict[str, dict[str, Any]]allarg_propertiesz dict[str, dict[str, Any]] | Nonekwarg_propertiesz"dict[sympy.Symbol, pytree.KeyPath]unbacked_bindingszlist[MutationOutput]r.rrgr6r~rdict[str, Any] | Nonerrc sttj|||d||_|r|ni|_||_| |_||||| |_| i|_ g|_ t j j|_i|_dSNr)rrrrrrset_cpp_kernel_nameset_python_kernel_namercollect_arg_kwarg_propertiesrr.rrrrfx_noder) rrrgr~rrrrrrrr rrrZs"     zExternKernel.__init__rcCrnrrorrrrrzrzExternKernel.get_outputsrcCrrr<rrrrr}rz%ExternKernel.get_unbacked_symbol_defscCst|jtjjrdd|jjjDn ddtt|j D|_ t|jtjjr1dd|jjjDni|_ t|jtjjrW|j sJdd|jjjD|_ dd|jjjD|_ dSg|_ dS)NcSs$g|]}|js|j|j|jdqS))rrr) kwarg_onlyr real_typerrtrrrrsz=ExternKernel.collect_arg_kwarg_properties..cSsg|]}iqSrrrrrrrr cSsi|] }|j|j|jdqS))rr)rrrrtrrrrrz=ExternKernel.collect_arg_kwarg_properties..cSg|]}|jr|jqSrrrrtrrrr cSsg|]}|jr|qSr)rrtrrrrs )rrr)_ops OpOverload_schema argumentsrrr~rrr schema_kwargsrrrrrs*    z)ExternKernel.collect_arg_kwarg_propertiescCs$t|jtr||dSdSr)rrgr$apply_constraintr)rrrrrJs  zExternKernel.decide_layoutwrapperrwrcCsZt||\}}|r|||s|}|r+ddlm}|||dd}|||dSdS)NrD)'set_kernel_post_grad_provenance_tracingT)rA)re make_commenttry_get_kernel_namerrwrite_provenance_debug_handle)rrr origin_str_detailed_origin_strr debug_handlerrrcodegen_comments  zExternKernel.codegen_commentcCrurrvrrrrrcodegenrzExternKernel.codegencCs||_tjjrt|jtjjsdS|j}|jdurh|j dkra|j dkr+|j ddn|j dd}ddlm}|d|j i}tdd |Dd d }|d krX|d |}d |d|_dS|jj|_dSdS)Natenr.rr inductor_fallback_opszaten.css*|]}|drt|ddVqdS)rrDN) startswithrrrrrrs(z3ExternKernel.set_cpp_kernel_name..rDr_vz at::_ops::z::call)rrrr cpp_wrapperrrr)r r  namespace _overloadnamerrreplacetorchgen.aoti.fallback_opsrrrr r)rrropnamer version_infolatest_versionrrrrs.        z ExternKernel.set_cpp_kernel_namecCsd||_|dur dS|j}|durdSt|tjjr"d|j|_dS|jddd|j|_dS)Nztorch.ops.higher_order.._ops..ops.r) rrrr)r HigherOrderOperatorrrr$)rrrrrrrsz#ExternKernel.set_python_kernel_namecCsddlm}|}r|jntjj}tjjr|jStjj r=t tjj |s-Jttjj |j dur4dStjj |j |S|jS)NrD) CppWrapperCpu)codegen.cpp_wrapper_cpur,r'rrrr device_type fx_wrapperrr!rrrget_c_shim_func_name)rr,drrrrrs  z ExternKernel.try_get_kernel_namercCrr)rrrrrget_kernel_namerzExternKernel.get_kernel_namerrrcCs:tj||||||d}||S)N)rrrrrr) rrr'r&rr!rrr )rrrrr copy_input szExternKernel.copy_inputrrrr ftuple[Any, list[Any], list[Any], Callable[[Any, Any], Any], dict[sympy.Symbol, pytree.KeyPath] | None]cs||d}t|\}gg}g}g}|D]} | tdr:tjjjj| dd} d|| || qt drhd|| | j j } | j j dkrZ| dus\J|t jj| qtdrd|| || jqtdrd|| q d|| || qdfd d } fdd|D}|D] } t| rt| ddqg}|D]}} t| ts| tjjvr|tjj| qt| ts| tjjvr|tjj| qt| tr || qt| tr|| jqt| t jj j r>| j j } | j j dkr0| dus2J|t jj| q|t!| q| ||\}}||i|}d}tj"j}rtj#j$%d}t&}tj#j't j(j)j*urz|d}t+tj#}|t,|tj#|Wdn 1swYt-|||}t|t.t/fs|gn|}|D]+}t|t j0r|j1rt2j3sd}tjj#j$%dd}r|d|}|tj_4q|||| |fS)aPartition kernel args into tensor and non-tensor, realize tensor inputs, re-run fake tensor propagation with the realized strides, and return (example_output, tensor_args, non_tensor_args, unflatten_args, unbacked_bindings). unflatten_args(new_tensor_args, new_non_tensor_args) reconstructs the original (args, kwargs) tree from replacement lists. rrN)rFr:Tnew_tensor_argsrnew_non_tensor_argsrtuple[list[_T], dict[str, _T]]csdg}t|}t|}D]}|r|t|q |t|q t|}|dg|difS)Nrr)rrrpytreetree_unflattenr)r5r6r it_tensorsit_non_tensors is_tensorr)args_flat_is_tensor args_specrrunflatten_argsUs z3ExternKernel.process_kernel..unflatten_argscr,rrMrtrrrrcrz/ExternKernel.process_kernel..rDrrDzEsparsity not handled. Please file issue for sparse inference weights.rz Found from : )r5rr6rrr7)5r8 tree_flattenrrrrrrcreate_symintnoderGeneratorStaterrrr)r:default_generators clone_stateOpaqueObjectStater-rr"rhrrr constantstorchbind_constantsr get_valueropaque_example_valuerirrr#rrzrr r_higher_order_opseffects with_effectsr0r7r1rrrZ is_sparserEgraph_partitiondisable_cudagraphs_reason)rrrr binded_args args_flat tensor_argsnon_tensor_argsreal_non_tensor_argsrr device_indexr?r example_argsnew_args new_kwargsexample_outputrr node_meta_valctxexample_out_lir,msgrr)r=r>rrprocess_kernels                     zExternKernel.process_kernelrlc Cspt|ts Jt|t|tr|S|}tj|}|dus$J| }|dur]d|j vr]t|tt t fr]t|j tr]t|j dtjdsSt|j dtjdr]|t|n|tj|dd\}}|d}||}tjj||}tjj||} tjj||} t|| | } || krt d| | |t!t|j"t#|$|%|| | dd d S) z In order to pass this to an extern kernel we need a ReinterpretView not a View. This allows us to avoid some unneeded copies. Nrrrrrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFrre)&rrrrlrrrrrrrrzrrrgr$r,r)r\r^r1r.r!r)rFr]rrr^ stride_vars offset_varrjr\rrrfrirr&) rr x_unwrap_viewrOx_unwrap_view_fx_node index_argsrrGrr^rjexpectedrrrconvert_to_reinterpret_viewsn       z(ExternKernel.convert_to_reinterpret_viewcCs8|durtSt|ttjjjtfrt|dSt|t r@t t j tj|j||dWdS1s;wYt|trG|St|trR||jSt|trct||j|dSt|tr|t|rz||WStyYnwt|tr||St|t tt!fr|S|"|S)Nr;)rrre)#rrrrrrrrrrr>rrradd_tensor_constantr)r]r-r&r'rrrMrfrlr#rr r"rrhrr NonTensorObjrr3rrrrrMs@            zExternKernel.realize_inputcCsDt|rt|dkr|S|D] }|dkr|Sq||Sr4)r"rrKr3)rrrrrrrequire_stride1#s  zExternKernel.require_stride1Frrr2rr*rc s|dus dus JdvrsStrttrS|rGt|o-tj }tdd|r@t t j j jn||dStddd|dStttfr}|re|srr}tjr}dur{tSSt}trt|}trtdt|tr|r||srt|jrSttr|r|sLjrtjrɈSttrtjtrtjtst}rt|drt|jtsz!|j_|r|j ||dWSr|j!|dWSWn t"yYnwd} } durOt j j fd d t#t$D} | D] } t%j&j'(| d d qA|)tdd||d|rkt|siJS| r| durxduszJt%j&j'*| tSS) N)rrDTF)rErLrMr*rPzHthe MutationLayoutSHOULDREMOVE's real layout shouldn't be FlexibleLayoutrfrQcs4g|]}|dr|dr|qS)rr)rbrr!rr2rrrrrsz0ExternKernel.require_strides..rrD)+rr"rr#r$rTr3rrhrrrrrrbrirurSr\r!rqrrrrrrfrrlrrrrhrequire_stride_orderrequire_exact_stridesrrrr)rloweringslice_r3r|) rrrr2r*use_current_stride_ordermutation_layoutrr expanded_dims orig_sizerYrrlrrequire_strides-s                 zExternKernel.require_stridesrUcCs|j|dd|D|dS)NcSs$g|]}t|tjr|jjn|qSr)rr)SymIntrrrrrrrsz6ExternKernel.require_exact_strides..)r2r*ru)rrr2r*rrrrnsz"ExternKernel.require_exact_stridesrcCs|j|||dS)N)rr*rw)rrrr*rrrrmsz!ExternKernel.require_stride_ordercC ||tSr)rmr]rrrrrequire_channels_lastr[z"ExternKernel.require_channels_lastcCrxr)rmr_rrrrrequire_channels_last_3dr[z%ExternKernel.require_channels_last_3dcCs,ddd}||r |S||t|S)Nrrrrc Ss@z|}Wn ttfyYdSw|tjjvotjj|jSr)rAttributeErrorrrrrrH is_mkldnn)rrrrris_mkldnn_tensors  z9ExternKernel.require_contiguous..is_mkldnn_tensorrrrrrnr$r%r!)rrr}rrrrs  zExternKernel.require_contiguouscCs||t|Srrrrrrrequire_contiguous_stridessz'ExternKernel.require_contiguous_stridescCrrrrrrrrrzExternKernel.apply_constraintcCst|ts Jt|t|tst|}|jsJdt|}t|j}||krStd|j||t ||D]}|j|d}| ||vrJ||n|j|dq7|S)Nz/ExternKernel.arg_properties should not be emptyzv%s has %d unprovided positional arguments. Will check if they are in the keyword arguments or will use default values.rr) rrrrrrr\rrrr)rrrn_args n_pos_argsrarg_namerrrfill_non_provided_argss(   z#ExternKernel.fill_non_provided_argsrrrc Cstjjrog}d}|r"|jr"t|jt|ksJddd|jD}t|jD]E\}}|durF|dus5J|||}|rC|dnd}nt|j|}|jr_|t|jkr_|j|dnd}| tjj ||q'|Sdd|jDS)NzDnames passed to codegen_const_args does not match self.constant_argscSsi|]}|d|qSr)rrrrrr8rz3ExternKernel.codegen_const_args..rcSrrrrrrval_to_arg_str)rrrrrrKrz3ExternKernel.codegen_const_args..) rrrr!rrrrrr~rrr) rrrname_to_arg_propertiesrrproptype_rrrrcodegen_const_args,s2  zExternKernel.codegen_const_argscCstjjr|jdur|g|j|j|j}d}n|j}d}g}t|D]4\}}tjjrN|j r6|t |j ks:Jd|j | d}| tjj ||q$| tjj |q$|rb|||S)NFTz-Invalid access to ExternKernel.arg_propertiesr)rrrr!rrr~rrrrrrrrrrr)rr~need_codegen_constant_argsrrrrrrr codegen_argsMs&zExternKernel.codegen_argsrcKsT||vr ||S||jvr|j|S|j|}dur#|dSt|d)zGiven an argument name, queries for values in (in order): 1. any provided kwargs for this function. 2. the class self.kwargs member. 3. any available default arguments in self.allarg_properties.Nrz not in self.allarg_properties)rrrr)rrrrrrrget_kwargs_valuegs    zExternKernel.get_kwargs_valueskip_outcCstjjrO|jdurt|jdkrgSg}|jD]5}|r |dkr q||}t|t r0| |q|j dus7J|j |i d}| tjj ||q|Sdd|jD}|S)Nrr rcSs(g|]\}}|dtjj|qSrr)rkrrrrrsz/ExternKernel.codegen_kwargs..)rrrr!rrrrrrrrrrrrrr)rrrrrrrrrcodegen_kwargsts$     zExternKernel.codegen_kwargscCsT|jdur&|jj}t|dd}|dd}|ddd}|d|}|Sd}|S) Nrunknown_namespacer)r*rrDr unknown_op)rrrr$rsplit)rr op_namespaceop_namerrr get_op_names   zExternKernel.get_op_namec Cstjr=tjjs?t|dkrdStjj|}tjj| }| }| d| d|d|d|d dSdSdS)Nrzassert_size_stride(rtr5) rE size_assertsrrrr!rmr!rcodegen_shape_tuplerKrrr)rrrrrrrrcodegen_size_assertss z!ExternKernel.codegen_size_assertsc Csttjr6tjjs8|}|tjjv}|}|r(|d|dt d|ddS|d|d|ddSdSdS)Nzassert_alignment(rtr5z # buffer z (op: z) is assumed to be not aligned) rEalignment_assertsrrrr!rrWrrrf)rrralignedrrrrcodegen_alignment_assertss z&ExternKernel.codegen_alignment_assertscCs@tjjrtjjr dS||}|d|d|ddS)zc Track outputs of fallback operators if config.test_configs.track_memory_lifecycle Nz track_tensor(z, 'z')) rE test_configstrack_memory_lifecyclerrrr!"write_memory_track_allocation_oncerr)rrrrrrcodegen_memory_trackings z$ExternKernel.codegen_memory_tracking'tuple[list[Sequence[Expr]], list[Expr]]cCs|}|}|gg|fS)zD get output sizes and strides, for template_codegen )r!rK)rrrrrrget_group_strides zExternKernel.get_group_stridetuple[Expr, Sequence[Expr]]c stjj|}|}fdd|D}ddtt|Dttt||jdd}ddt |DfddttD}fd d|D| }|}tjj ||g\}}} t d \} t t|fd d|D} tt|| }|t|fS) zC Manually get canonicalization of the output index cr,r)r rt)rrrrrz-ExternKernel.canonicalize..cSsg|] }td|qS)r1)rkrrrrrrT)rCrcSrrrrrrrrrz-ExternKernel.canonicalize..crrrrrrrrrcrrrr)rrrrrccrrrrt)add_varrrrr)rrrrr!rKrrrHrcrrrrQrrrnrr|r) rrVr^ index_orderrrr new_sizesrrr  replacementr)rrrrr canonicalizes$   zExternKernel.canonicalizercCsP|rtnt}t||}|jD]}|||O}q|jD]}|||O}q|Sr)maybe_free_unbacked_symbolsmaybe_free_symbolsrr>rrr)rrmaybe_get_symbolsrrrrrr>s   z!ExternKernel.get_free_symbol_usescsPtdd}d|g}|fddtD7}|dj|S)Nrzpython_kernel_name=cs$g|]}|jdt|jqSr)rr)rrmrrrrsz(ExternKernel.__str__..r)rrlfieldsrrr)rrrrrrrs   zExternKernel.__str__rNNNNrN)rr3rgr6r~rrr6rrrrrr3rr3rrrrrrrrrIr)rrwrr3rrrrwrrrr3rr)rr3rrrUr\)rrrr)rrrr rr rr4)rrrrlr)NNF) rrrrr2rr*rrrr_)rrr2rUr*rrr)rrrrr*rrr)rr6rrrr6)rrrrrr)rrrr rr )rrrr)rr)rrrh)?rrrrkrrrlrmrrrrrrrrrrrrr.rrrrrJrrrrrr2rnr3rrarhrMrkrurnrmryrzrrrrrrrrrrrrrrr]r>rrrrrr rr>s           %   %    ! G #  $        $ !    !  rcsBeZdZdddZ       dd fdd Zd!ddZZS)"ExternKernelOutrrwrrcC||dSr)generate_extern_kernel_outrrrrr rzExternKernelOut.codegenrNrgrr~rrr6rrrrrr3rrrrc s`||} t| tsJt| td|| ||pid|||| tj||_ tj |dSr) r(rrrrrrrrr)rr*) rrgr~rrrrrrrunwrapped_inputsr rrr s  zExternKernelOut.__init__rcCrrrrrrrr)rzExternKernelOut.should_allocaterr)rgrr~rrr6rrrrrr3rr3rr6rrrrrR)rrrrrrrrrr rrs rceZdZd fdd ZZS) RandomSeedsrfrrrrrcsFttj}tjt|tj|gdg|j|j|ggddtj j ddS)Nr<zaten.randint.low_outzat::_ops::randint_low_out::call)rgr~rrrr) r)rr~rrrirrrrandintlow_out)rrfrlimitsr rrr.s  zRandomSeeds.__init__)rfrrrrrrrrrrrrr rr-rcsJeZdZdddZ      dd fdd Zd!ddZd"ddZZS)#rrrwrrcCrr)generate_extern_kernel_allocrrrrrBrzExternKernelAlloc.codegenrNrgr6r~rrr6rrrr3rrrrc sp||} tdd| DsJtd|ttt| ||pid|||| g|_tj ||_ tj |dS)NcsrrrrrrrrQrz-ExternKernelAlloc.__init__..) r(rrrr rrrrrrr)rr*) rrgr~rrrrrrrr rrrEs"  zExternKernelAlloc.__init__rcCrrrrrrrrerz!ExternKernelAlloc.should_allocatecCrurrvrrrrrhrz"ExternKernelAlloc.apply_constraintr)rNNNrN)rgr6r~rrr6rrrr3rr3rr6rrrrrRrI)rrrrrrrrrrr rrAs  rcsJeZdZdZdfd d Zdd d ZdddZdddZdddZZ S)r%zP An output buffer that represents the mutation of a pre-existing buffer rgr6 mutated_noder mutating_noderrrrcsDtjd|d|}tj||g|_||_tj||_ dSr) rrrrrrrmutation_namesrr)r)rrgrrmutated_node_namer rrrqs  zMutationOutput.__init__cCrr)rrrrrr{rzMutationOutput.get_defining_oprcCrr)rrrrrrD~rz!MutationOutput.get_mutation_namesrcCrrrrrrrrrzMutationOutput.should_allocatercCs |}dddd|DDS)NcSg|]}|dur|qSrrr&rrrrs z7MutationOutput.get_mutation_buffers..css|] }tj|VqdSr)rrrtry_get_bufferrrrrrrFz6MutationOutput.get_mutation_buffers..)rD)rrrrrget_mutation_bufferss z#MutationOutput.get_mutation_buffers)rgr6rrrrrrrrrjrRrr) rrrrkrrrDrrrrrr rr%ls  r%cs`eZdZUdZiZded<edd d Zedd d Zdfdd Z dddZ dddZ Z S) TMADescriptorad An IR node representing a generic host-side TMA descriptor in the Triton API Mostly useful for user-defined Triton kernels relying on host-side TMA; but can, in principle, be used for Inductor's Triton templates, too. See TMADescriptorExperimental and TMADescriptorStable for the two implementations (the old API and the new API) zdict[Any, TMADescriptor]_CACHEr]rtma_metatuple[str, tuple[Any, ...]]rcCsTt|dksJ|ddkrt|g|dRS|ddks Jt|g|dRS)Nrr experimentalrDr9)rTMADescriptorExperimentalTMADescriptorStable)rr]rrrr _create_impls  zTMADescriptor._create_implcCs2t||f}||jvr||||j|<|j|Sr)r[rr)rr]rrCrrrrs   zTMADescriptor.creater~r6rrcsVtdtt||dttt|t|d||_ t j ||_ t j |dS)Nre)rrrurlr#r rrrr]rrrr)rr*)rr]r~rr rrrs zTMADescriptor.__init__rrwcCrr)generate_tma_descriptorrrrrrrzTMADescriptor.codegencCrr)r]rrrr get_tensorrzTMADescriptor.get_tensor)r]rrrrr)r]rr~r6rr6rrrrH) rrrrkrrrrrrrrrrrr rrs     rcs&eZdZdZ ddfd d ZZS)rz the new host-side TMA Descriptor API: (the ones obtained via create_{1d,2d}_tma_descriptor calls). See also TMADescriptorStable for the new API. Nr]rrlist[int | torch.SymInt] block_dims element_sizerrrcst|dvsJt|t|ksJ|dur|j}||_||_||_t|j|_|g}g|j|j|j}tj|||ddS)N)rDrr]r~r) rr&rrrrrGrr)rr]rrrr~rr rrrs*   z"TMADescriptorExperimental.__init__r) r]rrrrrrrrrrrrrkrrrrr rrs rcs"eZdZdZdfdd ZZS) rz the new host-side TMA descriptor API (the ones obtained via TensorDescriptor.from_tensor). See also TMADescriptorExperimental for the old API. r]r block_shapercs||_tj||g|ddS)Nr)rrr)rr]rr rrrs  zTMADescriptorStable.__init__)r]rrrrrrr rrsrcs,eZdZ ddfdd ZdddZZS)SubgraphBufferNrgrrrrrrsexample_inputs list[Any] subgraph_namerrvrc sbtd||||_||_tj||_tj|tj |j|||_ t |j s,Jt |j }|D]}||j j|j<|j j|jq3dd|D|_ddlmm} t|j Kdddd} i| |pii} | | |j j|jWdn1swY|r|j jD]} | |qWddSWddS1swYdS)NcSg|]}|jqSrr)rsym_varrrrr z+SubgraphBuffer.__init__..rFATEN) max_autotunemax_autotune_gemmmax_autotune_gemm_backends)rrrrrrrrr)rr* make_subgraphsubgraphrr~rr graph_input_namesr sym_inputstorch._inductor.configrrEset_graph_handlerrrun operationsrr) rrgrrrrrrvrsym_inpinductor_config base_patchesmerged_patchesrvr rrrs:     "zSubgraphBuffer.__init__rrwrrcCsRGddd}t|jsJdd|jD}|||jg|j||jgdS)Nc@r)z,SubgraphBuffer.codegen..CodegenGraphrrxcSs||_|j|_dSr)rr)rrrrrr8s z5SubgraphBuffer.codegen..CodegenGraph.__init__N)rrx)rrrrrrrr CodegenGraph7rrcSg|]}|qSrr rr,rrrr=rz*SubgraphBuffer.codegen..)rr~'codegen_subgraph_with_flattened_outputsrrr)rrr outer_inputsrrrr6s zSubgraphBuffer.codegenr) rgrrrrrrsrrrrrvrrrrrrrrrrr rrs.rcseZdZdZd.ddZd/ddZed0d d Zd1ddZd2ddZ e d d3d4fdd Z d5ddZ d6fd#d$ Z ed7fd&d' Zd8d)d*Zd9d,d-ZZS):UserDefinedTritonKernelz> A user-defined triton kernel (e.g. via @triton.jit). r(tuple[Kernel, Any, list[str], list[str]]csddlm}ddlm}||jg}g}g}t|retdr0|fddj Dn tds7J|j tdrRj D] }| j j|qEn td sYJ|jj}j |||fS) Nr) Autotuner)kernel_side_table restore_idxc3s|] }jj|VqdSr)r arg_namesrrrrrXrzBUserDefinedTritonKernel.get_kernel_and_metadata.. restore_value reset_idx reset_to_zero)triton.runtime.autotunerr*torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrrrrrrrrrrconfigs)rrrrrestore_value_argsreset_to_zero_argsrrrrget_kernel_and_metadataKs,           z/UserDefinedTritonKernel.get_kernel_and_metadatarcCstjsdS|jjs dSt|jjdkrdSt|jdksJt|jdt s(dSt|jdj t s3dSt|jdj j t s?dSt|jdj j j t sLdStdd|jdj j j jDs^dSdS)aE For kernels like @triton.jit def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offs < n_elements x = tl.load(in_ptr0 + offs, mask=mask) y = tl.load(in_ptr1 + offs, mask=mask) tl.store(out_ptr + offs, x + y, mask=mask) @torch.compile def fn(a, b): out = torch.empty_like(a) grid = (triton.cdiv(a.numel(), 1024),) add_kernel[grid](a, b, out, a.numel(), BLOCK_SIZE=1024) return out.relu() We can potentially fuse the relu epilogue into the add_kernel. We do this by pruning the `out` tensor allocation and directly writing the relu-output. FrDrcsrrFrr?rrrrrz.T)rE*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrfrrrrrrrrrrks$ z)UserDefinedTritonKernel.can_fuse_epiloguerrwrcCs|j|ddS)N)epilogue_fusion_codegenrrrrrrzUserDefinedTritonKernel.codegenr tuple[ComputedBuffer, str]cCs |||S)zQ epilogue_fusion: (fused epilogue node, modified kerel src code) r )rrr rrrcodegen_with_epilogue_fusions z4UserDefinedTritonKernel.codegen_with_epilogue_fusion!tuple[ComputedBuffer, str] | Nonec s$ddlm}\}}}}|||j||j|\}} } } fddjD} |rNtjj j dks6Jt t jj j j } | | vsFJ|\}}|| | <dd|jDdd|jD}tfd d |D}g}g}g}g}t| ttd | D]w\}}||vr|rq||||t|tr||||qt|ttttjfr|||t|q||vr|d |tq|d ur |r|d |tq| | qt!dt|d|"|||j#|||||| | d$j%j d d S)YOverrides the parent member. See https://github.com/pytorch/pytorch/issues/151692r)triton_version_uses_attrs_dictcsi|]}||qSrrrrrrrrrz4UserDefinedTritonKernel._codegen..rDcSrrrrrrrrrrz4UserDefinedTritonKernel._codegen..cSrr) is_constexprnumrrrrrrc3s|]}|VqdSrrr)rrrrrz3UserDefinedTritonKernel._codegen..rrINzUnsupported arg type: r4T) arg_typesraw_argsraw_keys triton_meta inductor_metar=roriginal_fxnode_name)&rrr!define_user_defined_triton_kernelrgridrrrrLrrrrparamsr=rvr?rrrepeatrrrr r&rrrrrrrrrgenerate_kernel_callr'r)rrr rrrrrnew_namerrextra_launch_args named_argsmutable_arg_nameepilogue_computed_bufferr  constexprsconstexpr_namesrrraw_keys_filteredraw_args_filteredrrr)rrrrs                z UserDefinedTritonKernel._codegenFrrcst|t|j|BSr)rr>r&r r=r rrr>s z,UserDefinedTritonKernel.get_free_symbol_usescCrrr<rrrrr rz0UserDefinedTritonKernel.get_unbacked_symbol_defsrrr r tma_descriptor_metadatar kernel_argscsg}i}g}D]0\}} t| tr1t| } ||vr't| ||} || | ||<q || | ||<q t |dksCJ|d _ t|t sUJt |tdtj d|t|||_|_\} } } } t| dszJfdd| jD_ddlm}m}t | dkr| djni}ddl}| j_|j_ |j_!_"|| i||_#fddj#j$j%D_&fddj&D_'t(j)*dS) Nrrrcsg|]}|vr|qSrrrr.rrrLrz4UserDefinedTritonKernel.__init__..)identify_accessed_tensorsidentify_triton_storescs(g|]}t|jtr|jqSr)rrrrrrCr/rrrgs cs g|] }ttjd|qSr#)r%rrr&rrrrm)+rrrrrrMrrrrr'rrrrrrrrr rrrrrr0r1rastr kernel_srcparse kernel_astr r.rrLrr r.rrrr*)rrr r-r.r~rrrrr,rrr r0r1autotuned_kwargsr4r )r.rrr#s`           z UserDefinedTritonKernel.__init__r5cstjstSfddjjjDtjjj}ddt |j Dt j t fddjjjDt fddjjjDt d}|S)Ncs i|] }|jj|jqSr)rr.r)rformal_arg_deprrrr{r3z;UserDefinedTritonKernel.get_read_writes..cSsi|] \}}|j|qSr)rr)rr9 mut_outputrrrrs cg|]}|qSrrenamer) read_renamesrrrz;UserDefinedTritonKernel.get_read_writes..cr;rr<r) write_renamesrrrr?r)rErrr6rrLr:rrrr.rFrr=)rformal_arg_writesrLr )r>rr@rr6ss2    z'UserDefinedTritonKernel.get_read_writesrcC t|jSr)rr.rrrrrrz#UserDefinedTritonKernel.get_outputsrcCrrrrrrrr'rz"UserDefinedTritonKernel.get_device)rrrRr)rrwr rrr)rrwr rrrr_rhr) rrr r r-rr.rrrrdrrW)rrrrkrrrrrrr]r>rrr6rr'rrrr rrFs"  8   g P %rcsJeZdZdZdddZdd d Zdd d ZdddZdfdd ZZ S)InplaceBernoulliFallbackE This needs to be a custom class to handle mutation properly rrwrrc Cstdd|jDs Jdd|jD\}tjjr5||d|ddtt |j d|j dS||d|ddtt |j d|j dS)Ncsrrrrrrrrrz3InplaceBernoulliFallback.codegen..css|] }tt|VqdSr)r rr rrrrrrrrtz, NULL)r5) rr~rrrr!rr2rrreprrending)rrrrrrrs,,z InplaceBernoulliFallback.codegenrcCrrrrrrrrrz(InplaceBernoulliFallback.should_allocatercC |dgSrrrrrrrDrz+InplaceBernoulliFallback.get_mutation_namesrcCrrr<rrrrrrz1InplaceBernoulliFallback.get_unbacked_symbol_defsrrrrrr csVtjdt|d||g||dtj|tj ||_ tj |dS)Nrr) rrrr'r(rrrrrr)rr*)rrrrr rrrs  z!InplaceBernoulliFallback.__init__rrRrjr)rrrrrr rr rrrrkrrrDrrrrrr rrCs    rCcs\eZdZdZd ddZd!d d Zd"d d Zd#ddZd$fdd Ze d%d&ddZ Z S)'InplaceCopyFallbackrDrrwrrcCs |\}}}||||dSr)rcodegen_device_copy)rrrr non_blockingrrrrszInplaceCopyFallback.codegenrcCrrrrrrrrrz#InplaceCopyFallback.should_allocatercCrGrrHrrrrrDrz&InplaceCopyFallback.get_mutation_namesrcCrrr<rrrrrrz,InplaceCopyFallback.get_unbacked_symbol_defsrgr6r~rrr6csJtjd|||dddtj|dtj||_tj|dS)Nz aten.copy_aoti_torch_copy_)rrr) rrrrrrrr)rr*)rrgr~rr rrrszInplaceCopyFallback.__init__FrrrrMcs6fdd||fD}|f}tt|d||}|S)Ncr,rr@rrArrrrz.InplaceCopyFallback.create..r)rKrr')rrrrMr~rrrrArrs zInplaceCopyFallback.createrrRrjr)rgr6r~rrr6rrr_)rrrrrMrrrK) rrrrkrrrDrrrrrrrr rrKs    rKc@sBeZdZdZdddZdd d Zdd d ZdddZdddZdS)MutatingFirstArgExternKernelrDrrwrrcCsXt|jsJgdd|jDtt|j}||dd|d|jdS)Ncss|]}|VqdSrrrrrrrrz7MutatingFirstArgExternKernel.codegen..rrtr5) rr~rrErrr2rrF)rrargrefsrrrrs z$MutatingFirstArgExternKernel.codegenrcCrrrrrrrrrz,MutatingFirstArgExternKernel.should_allocatercCrGrrHrrrrrDrz/MutatingFirstArgExternKernel.get_mutation_namesrcCrrr<rrrrrrz5MutatingFirstArgExternKernel.get_unbacked_symbol_defscCrrrrrrrhas_side_effectsrz-MutatingFirstArgExternKernel.has_side_effectsNrrRrjr) rrrrkrrrDrrQrrrrrOs   rOcr) ResizeStorageBytesvariablerrqrrrcst|ts Jdtjdt|d||g|fdtj | tj ||_ tj |d|_d|_t|tttfsFJt|tjj|j dS)NzTODO: dynamic shapesr)rz"inductor_ops.resize_storage_bytes_z&torch::inductor::resize_storage_bytes_)rrrrrr'r(rrrrrr)rr*rrrrrrnever_reuse_buffersrrf)rrSrqr rrrs   zResizeStorageBytes.__init__)rSrrqrrrrrrr rrRrrRcs(eZdZd fdd Zd d d ZZS) SetSourceTensorKernel self_tensorrstorage_tensorrrcs|tj|||gdtjjjjdt |t t t fs$Jt |tjj|jtjj|tjj||}tt|d||tt|d||g|_dS)Nz!torch.ops.aten.set_.source_Tensor)rrr)r)rrr#r)rprset_ source_TensorrrrrrrrrrTrrfrr'r%rr.)rrVrWrr rrr.s"  zSetSourceTensorKernel.__init__rcCs|d|dgSr4rHrrrrrFBrz2SetSourceTensorKernel.get_inputs_that_alias_output)rVrrWrrrrj)rrrrrFrrrr rrU-srUcsReZdZdZd!ddZd"d d Zd#d d Zd$ddZdddd%fdd ZZ S)&ScatterFallbackz This needs to be a custom class to handle mutation properly. This class handles both aten.scatter_ and aten.scatter_reduce_. It also handle the case `src` being a scalar properly. rrwrrcCrr)generate_scatter_fallbackrrrrrMrzScatterFallback.codegenrcCrrrrrrrrPrzScatterFallback.should_allocatercCs"|jd}t|ts J|gSrr)rrrrrrDSs  z"ScatterFallback.get_mutation_namesrcCrrr<rrrrrXrz(ScatterFallback.get_unbacked_symbol_defsNTru include_selfrrrrrYrrrrur3r]c st|t_jrfdd|||fD}|f} nfdd||fD}||f} tjdt|d|| ||dt|ddg|dt j | t j _t j dS) Ncr,rr@rrrrrjrz,ScatterFallback.__init__..cr,rr@rrrrrmrrr\rur])rrr)rr src_is_tensorrrrr'r(rrrrrrr)rr*) rrrrYrrrur]tensorsrr rrr[s&   zScatterFallback.__init__rrRrr)rrrrrYrrrrrrur3r]rrrrJrrr rrZFs     rZcsJeZdZdZdddZdd d Zdd d Zd ddZd!fdd ZZ S)"IndexPutFallbackzQ This needs to be a custom class to handle mutation and indices properly rrwrrcCrr)generate_index_put_fallbackrrrrrrzIndexPutFallback.codegenrcCrrrrrrrrrz IndexPutFallback.should_allocatercCrGrrHrrrrrDrz#IndexPutFallback.get_mutation_namesrcCrrr<rrrrrrz)IndexPutFallback.get_unbacked_symbol_defsrtorch._ops.OpOverloadrrrMrrr6 accumulater c s|_dd|D}fdd||g|D}d}tjdt|d||fd||dtj dtj _ tj dS) NcSrrrrrrrrrz-IndexPutFallback.__init__..cr,rr@rtrrrrraoti_torch_index_put_outrzaten.index_put_)rrrr) rMrrrr'r(rrrrrr)rr*) rrrrMrrc valid_indicesr_rr rrrs   zIndexPutFallback.__init__rrRrjr) rrbrrrMrrr6rcr rrrJrrr rr`s    r`c@s$eZdZeddd Zdd dZdS) DeviceCopyrrrrrMrrc Cs(|}|dus J|s:t|tjjvr:tdd|Dr:tj j s:tjj r5tj |tj || |Stj |tj |td|f}t|}d}|r\|}t|jog|jdkog|}|jdkost|jos|}|rt|rd|_tt|||||d||g|S)Ncss|] }|tjjvVqdSr)rrrrHr?rrrrrFz$DeviceCopy.create..zDeviceCopy in input programr9Trt)r'rArrrrrrrrE aot_inductoruse_runtime_constant_foldingr!add_device_inforCrarrr!rKrirr"r#rkrfrir&rM) rrrrMx_devicerris_destination_pinnedis_source_pinnedrrrrsL          zDeviceCopy.createrrwrcCs\|}t|dks J|jr||d|j|ddS||d||ddS)NrrrD)rrrrLr )rrrrrrrszDeviceCopy.codegenN)rrrrrMrrrr)rrrrrrrrrrrfs /rfcs`eZdZdZdddZd ddZd!fdd Zd"ddZed d#d$ddZ d%ddZ Z S)&DynamicSelectStorageOffseta The result of computing a dynamic selection index is determined as follows: when the index in the select operation is unbacked, the actual index calculation is ambiguous for negative indices (index + size) versus non-negative indices (just index). To resolve this, we allocate an unbacked SymInt to represent the storage offset and decompose the select operation into a call to as_strided, computing the storage offset at runtime with this node. rr7cCrrr<rrrrrrz$DynamicSelectStorageOffset.get_readsrcCrrrrrrrrrz*DynamicSelectStorageOffset.should_allocateunbacked_offset_symbol sympy.Symbolr base_offsetsympy.Symbol | intbase_dim_striderrrcsDtdttddg||_||_||_||_||_ ||_ dSNr9r) rrrr)rrnrrprrrr)rrnrrprrrrr rrrs  z#DynamicSelectStorageOffset.__init__rcC t|jgSr)r=rnrrrrr rz3DynamicSelectStorageOffset.get_unbacked_symbol_defsFrcCrr)r&rr=rrrr> rz/DynamicSelectStorageOffset.get_free_symbol_usesrrwcCs|j||jddS)Nr)codegen_dynamic_select_indexrrrrrr rz"DynamicSelectStorageOffset.codegenrerR)rnrorrorprqrrrqrrqrrrrrr_rhr rrrrkrrrrr]r>rrrrr rrms   rmcs`eZdZdZdddZdddZd fdd Zd!ddZed d"d#ddZ d$ddZ Z S)%DynamicSliceSizeac Computes the output size of a slice call, handling the correct semantics in codegen. We do this for flexible handling for unbacked indices (to not data-dependent error). Slicing has 4 semantics for indices, i.e. x[start:] could be: 1) start < -x.size(0) -> x[0:] # negative out-of-bounds 2) start in [-x.size(0), 0) -> x[x.size(0) + start:] # negative slicing 3) start in [0, x.size(0)) -> x[start:] # standard slicing 4) start >= x.size(0) -> empty slice # positive out-of-bounds If the appropriate semantics are known beforehand, the output size is computed based on the start & end indices. If not (with unbacked indices), a new unbacked symbol is created to represent the output size, and codegen handles computing the correct case. rr7cCrrr<rrrrr& rzDynamicSliceSize.get_readsrcCrrrrrrrr) rz DynamicSliceSize.should_allocateunbacked_size_symbolrorrqrrrcs>tdttddg||_||_||_||_||_ dSrs) rrrr)rrxrrrr)rrxrrrrr rrr, s  zDynamicSliceSize.__init__rcCrtr)r=rxrrrrr< rz)DynamicSliceSize.get_unbacked_symbol_defsFrcCst|j|t|j|Sr)r&rrrr=rrrr>? s  z%DynamicSliceSize.get_free_symbol_usesrrwrcCrr)codegen_dynamic_slice_sizerrrrrG rzDynamicSliceSize.codegenrerR) rxrorrqrrqrrqrrqrr_rhrrvrrr rrw s   rwcsJeZdZdZdddZdddZdfdd ZdddZdddZZ S)rz; The result of a call to aten._local_scalar_dense. rr7cCrrr<rrrrrP rzDynamicScalar.get_readsrcCrrrrrrrrS rzDynamicScalar.should_allocatesymrokeypathpytree.KeyPathrfrrcs<|tdttdd||g||_||_dSrs) r rrrr)rr(rzr{)rrzr{rfr rrrV s  zDynamicScalar.__init__rcCrtr)r=rzrrrrr` rz&DynamicScalar.get_unbacked_symbol_defsrrwcCrr)codegen_dynamic_scalarrrrrrc rzDynamicScalar.codegenrerR)rzror{r|rfrrrrr) rrrrkrrrrrrrrr rrK s   rcs`eZdZdZdddZdddZdfdd ZdddZed ddddZ d ddZ Z S)!rz5 The result of a call to aten._assert_scalar rr7cCrrr<rrrrrl rzAssertScalar.get_readsrcCrrrrrrrro rzAssertScalar.should_allocatescalarrsr`rrcs,tdttddg||_||_dSrs)rrrr)rr~r`)rr~r`r rrrr s zAssertScalar.__init__cCrrrrrrrrQ} rzAssertScalar.has_side_effectsFrrcCrr)r&r~r=rrrr> rz!AssertScalar.get_free_symbol_usesrrwc CstjsdStt|jdd}tjjrdStjjr;d|d}tjj j |j dd}| d|d|j d|d dStjj j|j dd}| d |d | d t|j d| |d dS)NFrzstd::to_string(r5)rzif (!(z()) { throw std::runtime_error("Expected z but received " + z); }zif not (z):z raise RuntimeError(z = None)rEscalar_assertsrrr>rrrr/r!rcodegen_cpp_sizevarr~rr`codegen_python_sizevarrEr)rrr symbol_strsizevarrrrr s& zAssertScalar.codegenrerR)r~rsr`rrrr_rhr) rrrrkrrrrQr]r>rrrrr rrg s   rc@seZdZUded<ded<dS)ExternKernelNoderrzexport_schema.NoderNrrrrrr s  rcseZdZdZ dfdd Zd?ddZd@ddZdAd!d"Ze dBd&d'Z dCd)d*Z dDd,d-Z dDd.d/Z d0d1Zed?d2d3Ze dEd7d8ZedFd:d;ZZS)Gr\z A class that represents a fallback kernel for handling operators that are not directly support by inductor. It currently supports functional ops, view ops, inplace aten ops, and mutating ops that are auto-functionalizable. Nrrgr6rrrUrnontensor_argsr6r?rrrr)dict[sympy.Symbol, pytree.KeyPath] | Nonerrcsxtj|t|t||dd_|pi_t|tjjtjj fs.Jd|dt |d|_ |_ |dur:in|_ jdusDJtjjg_g_tj tjj r[dSdj vrddSj j}tjjj r|j|ddSdd d } |jrtj s| j stdj  jj\} }dfdd } tjj || |D] \} } | | | qdS)NrIF#Fails to create FallbackKernel for r4 not supported_c10d_functionalrrvrbrrcSs2tj|tjjjpt|dotjjj|jvS)N py_kernels)r)_C%_dispatch_has_kernel_for_dispatch_keyr DispatchKey Functionalizerrrurrrhas_functionalize_impl s  z7FallbackKernel.__init__..has_functionalize_implz'NYI: Can't generate FallbackKernel for infotorch._C.Argumentrr rcstjtjrt|ttfsJt|tjr#t|ttfr#J|dur)dSjdur0dSdfdd }t jrO|durK|D]}||qDdSdStjsWJ||dS)Nr,rrrcsLj|jdusJjjr$jtt|d|dSdSr) alias_namesrr alias_infois_writer.r%rr'r,)rrrr add_alias!szPFallbackKernel.__init__..handle_aliasing_and_mutation..add_alias)r,rrr) rrr)ListTyperr library_utilsis_tensor_like_typeris_tensorlist_like_type)rrroptional_tensor_argr)rrhandle_aliasing_and_mutation!s"     z=FallbackKernel.__init__..handle_aliasing_and_mutation)rvrbrr)rrrr rr)!rrruse_runtime_dispatchrrr)r r r+rrr?rrrrr warn_fallbackrrrr _libraryrzmutates_and_returns_first_argrr is_mutabler#rr~r zip_schema)rrgrrUrr?rrschemarrrrrr rrr sV      zFallbackKernel.__init__r5csHt}|jtjjjur"|jD]}t|t r!| t | }q|Sr)rr6rr)_prims rng_primsgraphsafe_run_with_rng_staterrrD with_readrFrr)rrLrr rrr6*!s    zFallbackKernel.get_read_writesrrwcCs|||jt|ddSNr)(codegen_unbacked_symbol_defs_for_outputsrrrrrrrcodegen_unbacked_symbol_defs6!sz+FallbackKernel.codegen_unbacked_symbol_defsrcC>t|dd}rttjjj|}|dusJt|StSrrr8rrrrrr=r`rrresolvedrrrr;!   z'FallbackKernel.get_unbacked_symbol_defsrcstjGdddt|jsJfdd|jD}|||j\}}tjjrDt |j t j j rD|||}ddt|j jj|D}ndd|D}|j||S)Nc@ eZdZUded<dddZdS) z)FallbackKernel.codegen_args..Shimr refrrcSrr)rrrrrrJ!rz2FallbackKernel.codegen_args..Shim.__repr__Nr\)rrrrrrrrrShimF! rcsg|]}|qSrrrtrrrrN!rz/FallbackKernel.codegen_args..cSs"g|] \}}tjj||jqSr)rrrrrr)rparamrrrrrR!scSrrrrtrrrrW!r)rl dataclassrr~r?rrrrr!rrr)r r rrr r rrl)rrUrrrrrrE!s  zFallbackKernel.codegen_argsSequence[torch.Tensor] | Noner\r cCs|r dd|Dnd}|r|sJdd|D}|dSt|tjr%|jSt|tjjtfs2t|r7tdSt|tt frvt dd|D}dd|D}t |d krX|dS|s\dS|D]}t|tjshJt |j rq|Sq^|dSdS) NcSsg|] }t|ts|qSr)rrrrrrrb!rz.FallbackKernel.find_device..cSsg|] }|r|qSrrrrrrrh!rrr9css|] }td|VqdSr)r\ find_devicertrrrrq!s  z-FallbackKernel.find_device..cSsg|]}|r|qSrr)rrrrrrw!rrD)rr)rZrr ScriptObjectr'r)rrr=rrir)rUr\non_torch_bind_tensor_argsdevices device_setrrrrr]!s@     zFallbackKernel.find_devicercCsddlm}||jS)Nr) is_impure)torch._library.utilsrr)rrrrrrQ!s  zFallbackKernel.has_side_effectsrcCsrt|jtjjtjjfsJd|jdt|jdt|jtjjs6d|jvr6|jjj r6t |jr6gS|j S)Nrr4rr) rrr)r r r+rrr rr#rrrrrrF!s   z+FallbackKernel.get_inputs_that_alias_outputcCst|jdks J|jSr8)rrrrrrrD!sz!FallbackKernel.get_mutation_namesc sftdjttsJtjj \} |}fddj D}j}t j js=g||Stdg}|||}dd d t|tjjjrb||d |dj}n|jj}t|dkrjrrjnj}|d j}||g} n fddt|jD} jdusJttjj || idd} t j!"| g||S)a ProxyExecutor Design Note We export the ExternFallbackNodes (for custom ops) into a serialized file and run it with a host side proxy executor to address the ABI problem This is currently only implemented for fbcode. Eventually, we will also make this work for OSS. Detailed design doc can be found at https://docs.google.com/document/d/1wC4DOZFaYym2t1Esz0X5yxlLI3RDnSiyRbUus3bkJ64/edit?usp=sharing z4Extern kernel node added for node %s with target %s.csg|] }j|fiqSrrr2rrrrr!rz.N return_type1torch.TensorType | torch.ListType | torch.JitTypertIRNode | Sequence[IRNode]rexport_schema.ArgumentcSsnt|tjtjfrE|}t|ttfrt|dksJ|d}t|tjr8t|ts+Jtj j tj | ddS|dus>Jtj j ddSt|tj rjt|tjrjt|ts^Jt|tj j dd|Dd St|tjrt|tjr|durtj j tjj ddd St|tsJtj j tjj tj | ddd St|tjrtj j |d Std t|) NrDrr) as_tensorT)as_nonecSsg|] }tj|dqS)r) export_schemaTensorArgumentr)rr rrrr!szZFallbackKernel.export_extern_kernel_node..handle_single_output..) as_tensors)as_optional_tensor)as_intzUnsupported return type )rr) TensorTypeNoneTyperrrrrrtrrrrgetElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError)rrtr rrrhandle_single_output!sR     zFFallbackKernel.export_extern_kernel_node..handle_single_outputrrDcsg|] \}}|j|qSr)r)r return_schemart)rrrr"s)rr~rr)rr)rrrtrrr)#r\rrrrr\rr?r~rrrrrraot_moder"serialize_inputsr)rM torchbind CallTorchBindrreturnsr rrr.rrrrr;rextern_kernel_nodesr) rrordered_kwargsr serializernamed_argumentsrrroutput_argumentsrr)rrrrexport_extern_kernel_node!sN      5     z(FallbackKernel.export_extern_kernel_nodecsj}|dus J|jdkr5t|tjjsJt|tjj r4ddl m }t ||vr4t d|d_n|jdkrHt|tjjsGJt|n tjj rS|tjjv_tjj rt|tjjrjsdfd d jj\}t|fddjD}tfddt||jjD_|jrƈ}jdusJjdusJ| !jfddj|j"rj"nj#n|$tj%t&r'|(|)|*|dS)rNrrrzG%s is missing a c-shim implementation, using proxy executor as fallbackT _quantizedr, torch.JitTyperrcs$t|tjr |St|tjSr)rr)rr NumberTyper is_numberrrrC"s   z)FallbackKernel.codegen..is_numberc3s"|] }j|fiVqdSrrrrrrrP"rz)FallbackKernel.codegen..c3s(|]\}}t|to|jVqdSr)rcomplexr)rrrrrrrU"s  csgSr)rrrrrrr@c"rz(FallbackKernel.codegen..)r,rrr)+rr"rr)r r rrrrr!r%rrr\rrrErgcustom_ops_to_c_shimsr?r~rrvr?rrrr r rrr,generate_fallback_kernel_with_runtime_lookuprrr.generate_fallback_kernelrgrrrrr)rrrrr args_iter exported_argsr)rrrrr"sl                zFallbackKernel.codegenrtrricCsJd}z|}Wn tyYnwt|j|jt|t||dS)NFrt)rkrrirrr_rr)rtrkrrrr]r"s    zFallbackKernel.tensor_to_layoutrc s\tjf}||vrttdtjj}nt}|j|g|Ri|\}}}} } Wdn1s3wYt |t j j rt |t j rddlm} m} m} m}d}| |jr\||}|durd| |}|durt| |dkrt|j|jg|jg|d}t|t|t||t||dStdd|D||st |t jj j!s|t j"j#j$urt d t |t j j rtjj%s׈rt&j'||||| || d }|dur|S|durt(d |||| || d nsJd t)d |||| || d dfdd |g}t |tt*fr|_+|St |t,r(t*|_+|S|g_+|S)z9Create an instance of FallbackKernel from an _OpOverloadsNr)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variantrDrrrr)rgr~rrrrcsrr)rUrrrrr"rz(FallbackKernel.create..r9rhas_unaligned_inputrrrz"Not sure where to find device infortr rMlist[tuple[Any, int]]rcs tttfrtfddttDSttr,fddDSttj rNt  }t j sDsDtsLtjj|j|SttrUSttjr_jjSttjjtfsltrvttdSdusJdtddS)Nc3s,|]}|t|fgVqdSrrrgenerate_outputrMrtrrr"s  zAFallbackKernel.create..generate_output..cs*i|]\}}||t|fgqSrr)rrCrrrrr"rzBFallbackKernel.create..generate_output..rzFallbackKernel output type z is not supported)rrrrrrrrr)rZrLr]rE assume_unaligned_fallback_outputrorrrrWrrrrvrrrrr'r)rr)rtrMrOrrrrpacked)rMrtrr"sR       z.FallbackKernel.create..generate_output)rtr rMrrr )-r*_fused_moving_avg_obs_fq_helper_functionalr rrrrr#r rarr)r r rZtorch._library._out_variantrrrrr rrirrrWrrr_make_out_variant_kernel_namerrrMrrrp higher_orderprintr!ExternKernelMultiOut try_createrr2rrr)rrrrfake_incorrect_kernelscontextr\rUrVr?rrrrrout_oprg out_resultrrrrr"s         + zFallbackKernel.createrrgr6rrrUrrr6r?rrrrrrrrdrrr)rUrr\r6rr rRrj)rtrrri)rrrr rr rr\)rrrrkrr6rrrrnrrQrFrDrrrr]rrrrrr rr\ s,  w    %  y U r\cs>eZdZdZdddZdddZd d d dfddZZS) ComplexViewz9View a complex number as two dtyped numbers or vice versarrcCrrrrrrrr'#rzComplexView.should_allocatercCrGrrHrrrrrF*#r[z(ComplexView.get_inputs_that_alias_outputNrrgr6rrrUrrr6r?rrrrrrc stj|||||||ddS)Nr)rrrrgrrUrr?rrr rrr.#s  zComplexView.__init__rRrjr )rrrrkrrFrrrrr rr ##s   r c@seZdZdZd ddZdS) MemoryCheckKernelz Custom kernel for memory checking that generates direct function calls TODO - the custom op was erroring with str inputs. should be able to custom op directly. rrwrrcCsl||j\}}}t|}t|}|r&|dd|d|d|d}n d|d|d}||dS)z.Override codegen to write direct function callzV# note: dont currently distinguish between buffers returned and dealloc'd in last stepzcheck_memory_step(allocated=z, freed=z, is_final_step=r5N)rrrEr)rr alive_list dead_list is_final_step alive_repr dead_reprcallrrrrK#s zMemoryCheckKernel.codegenNr)rrrrkrrrrrr D#sr c@r) r2rrrrcCrrrrrrrr'a#rzMultiOutputLayout.get_deviceNrW)rrrrr'rrrrr2]#rr2cs`eZdZdddZ dd fdd Zed dd!ddZd"ddZd#ddZd$ddZ Z S)%rLrrwrrcCs,|||js||||dSdSr)codegen_multi_output!skip_size_stride_alignment_checksrrrrrrrf#s  zMultiOutput.codegenFrgr6rrrMlist[tuple[Any, ...]]r rcs>td||gdtj||_tj|||_||_dSr) rrrrrr)rr*rMr )rrgrrMr r rrrl#s   zMultiOutput.__init__rrcCs&|jd}t|tsJ|||Sr)r~rrr>)rrr1rrrr>y#s  z MultiOutput.get_free_symbol_usescCst|jdkot|jdtS)NrDr)rr~rrrrrrr#szMultiOutput.should_allocatercCsdd|jDS)NcSs.g|]}t|trt|dkr|qSr)rr\rrFrrrrrrr#sz.)r~rrrrrF#sz(MultiOutput.get_inputs_that_alias_outputr5cst}|jD]}t|tr|t|q|| d fdd }| }t j p;|dup;t |j }tj||d|d }tj||jtd S) Nrr6r*rr cr>r?rArrBrrr9#rCz*MultiOutput.get_read_writes..dummyrr3rr:)r=r~rrrrFrrr#rr'rErrirrOr!rr)rr:rr9rshould_normalizewrite_rwrrBrr6#s,     zMultiOutput.get_read_writesrr_) rgr6rrrMr r rrrrhrRrjrd) rrrrrr]r>rrFr6rrrr rrLe#s     rLcsDeZdZdZdfd d ZedddZdddZdddZZ S)rzMultiOutput for opaque objects.rgr6rrrMr opaque_valuer rrcstj|||dd||_dS)NTr )rrrK)rrgrrMr r rrr#s zOpaqueMultiOutput.__init__rcCr)NzOpaqueMultiOutput has no dtype)r{rrrrr#rzOpaqueMultiOutput.dtypecCrrrrrrrr#rz#OpaqueMultiOutput.wrap_for_loweringr5cCsVt}|jD]}t|tr|t|qtt|g}tj||tdS)Nr) r=r~rrrrFrrr)rr:rrrrrr6#s  z!OpaqueMultiOutput.get_read_writes) rgr6rrrMr r r rr)rr)rrrd) rrrrkrrqrrr6rrrr rr#s  rc@s$eZdZdZd ddZd d d Zd S)AllocatingMultiOutputzMultiOutput with Inductor-controlled allocation for .out() variant ops. Overrides should_allocate()=True so Inductor allocates the output buffer, and skips tuple-indexing codegen since .out() writes directly into these buffers. rrcCrrrrrrrr#rz%AllocatingMultiOutput.should_allocaterrwrcCs"|js||||dSdSr)r rrrrrrr#s zAllocatingMultiOutput.codegenNrRr)rrrrkrrrrrrr #s r r rbcCs4|j}|jjdd}|j}d|d|d|S)z8Build fully-qualified kernel name for an out-variant op.z::rDz torch.ops.r)r"r rrr#)r nsrrrrrr#srcsTeZdZUdZded<ded<d(fd d Zd)ddZedddd*d&d'ZZ S)+rzMulti-output .out() variant lowering. Subclass of FallbackKernel that emits .out() calls with pre-allocated output buffers. Uses AllocatingMultiOutput child nodes for each output. r out_arg_nameszlist[AllocatingMultiOutput]out_variant_output_nodesrr r rbrrrcs2tj|i|||_g|_t||_||_dSr)rrr r rrr)rr r rrr rrr#s   zExternKernelMultiOut.__init__rrwcCs||||dSr)r generate_extern_kernel_multi_outrrrrr$rzExternKernelMultiOut.codegenNFrrr\rrrUrrVr6r?rrrrrr&Sequence[AllocatingMultiOutput] | Nonec Cs<ddlm} m} m} | |jsdSt|ttfsdS| |} | dur$dS| | }tdd|Ds3dSt |t |kr=dS|t |d||||||| |d }g}t |D]7\}}t |j |jg|jg|d}t||t||fgd }tjs}| s}t|stjj|j||qS||_||_t|trt|St|S) zGCreate an ExternKernelMultiOut if the op has a matching .out() variant.r)rrrNcss|] }t|tjVqdSr)rr)rZrrrrr'$rFz2ExternKernelMultiOut.try_create..r)rrr r r)rgrrM)rrrrr rrrrrr2rrirrrWrr rrErrorrrrWrrrr r)rrr\rrUrVr?rrrrrrr r rrr tensor_outrg multi_outrrrr $sd      zExternKernelMultiOut.try_create) rr r rbr rrr rrr)rrbr\r rrrUrrVr6r?rrrrrrrrr ) rrrrkrrrrr rrrr rr#s   rc@seZdZUdZded<dvddZdwd d Zdxd d ZdyddZdzddZ d{ddZ d|d}ddZ d~dd Z dd!d"Z dd#d$Zdvd%d&Zdd'd(Z )ddd-d.Zdd/d0Zdd2d3Z )ddd5d6Zdd8d9Zdd;d<Zdd=d>Zdd@dAZddCdDZddFdGZdvdHdIZdvdJdKZddNdOZddQdRZd{dSdTZddUdVZ ddWdXZ!e"d )ddd[d\Z#dd^d_Z$ddadbZ%d|ddedfZ&e'ddhdiZ(ddkdlZ)ddmdnZ*ddodpZ+e'ddrdsZ,d{dtduZ-e-Z.dS)rzC TensorBox / StorageBox allow in-place mutation of Tensors rrfrrcCrdrrirrrrr_$rz!MutableBox.has_exceeded_max_readsrcCrdrrfrrrrr'b$rzMutableBox.get_devicercCrdrrrrrrre$rzMutableBox.make_loaderrcCrdr)rfrrrrrrh$rzMutableBox.make_indexerrUcCrdr)rfrKrrrrrKk$rzMutableBox.get_stridercCrdrrgrrrrrn$rzMutableBox.get_nameNr"rcCrYr)rfr$r#rrrr$q$rzMutableBox.has_large_inner_fnr%rrcCrYrrhr&rrrr't$rzMutableBox.mark_reusecCrdrrkrrrrr(w$rzMutableBox.realize_hintcCrdr)rfrrrrrrz$rzMutableBox.unwrap_viewcCrdr)rfr!rrrrr!}$rzMutableBox.is_input_buffercCrdr)rfr)rrrrr)$rzMutableBox.freeze_layoutFrrr*cC|j||Sr)rfr,r+rrrr,$r-z*MutableBox.freeze_layout_with_stride_ordercCrYr)rfr/r.rrrr/$rz(MutableBox.freeze_layout_with_fill_orderrcCrYr)rfr1r0rrrr1$rz(MutableBox.freeze_layout_with_same_orderr2cCr r)rfr4r3rrrr4$r-z+MutableBox.freeze_layout_with_exact_stridesr5cCrdr)rfr6rrrrr6$rzMutableBox.get_read_writesr7cCrdrrrrrrr$rzMutableBox.get_readscCrdrrrrrrr;$rzMutableBox.num_readsr~cCrdrrlrrrrr<$rzMutableBox.get_storage_numelr3cCrdrrrrrrr?$rzMutableBox.get_reduction_typercCrdrrrrrrr@$rzMutableBox.get_reduction_sizecCrdrrmrrrrrA$rzMutableBox.is_externcCrdr)rfrBrrrrrB$rzMutableBox.is_no_oprrcCrYrrrrrrrC$rzMutableBox.constant_to_devicercCrdr)rfrDrrrrrD$rzMutableBox.get_mutation_namescCrdr)rfrErrrrrE$rzMutableBox.get_operation_namecCrdr)rfrFrrrrrF$rz'MutableBox.get_inputs_that_alias_outputcCrdrrjrrrrr $rzMutableBox.realizerrcCrYrrZr=rrrr>$rzMutableBox.get_free_symbol_usesrcCrdrrorrrrr$rzMutableBox.get_read_namesrcCrdr)rfrrrrrr$rzMutableBox.get_defining_opr r cCrYr)rfr r rrrr $rzMutableBox.codegen_referencer6cCrdrrfrrrrrrg$s zMutableBox.layoutrcCrdrrerrrrr#$rzMutableBox.get_layoutcCrdrr rrrrr$rzMutableBox.get_output_speccCrdrrrrrrr!$rzMutableBox.get_sizercCrr)rfrrrrrr$rzMutableBox.dtypecCsnt|jtrt|jdt|jjd}d}|jj}n t|jd}|j}d}|tt||g}d|S)Nrz))r5r)rrfrrrrrr)rline0endlr`rrrrr$s    zMutableBox.__str__rRrWrYrZr[r\rr]r^rIrHr_r`rarbrcrdrerfrgrUrSrirjrhrJrLrVrQrPrO)/rrrrkrrr'rrrKrr$r'r(rr!r)r,r/r1r4r6rr;r<r?r@rArBrCrDrErFr r]r>rrr rqrgr#rr!rrrrrrrrW$sb                                     rc@s>eZdZeed ddZeed ddZed ddZd S) rrfrrcCrrrrfrrrr$rzTensorBox.creatercCrrrr# rrrr$rcCst|tr|Stt|Sr)rrrrr# rrrr$s  N)rfrrr)rfrrr)rfr)rrrrrnrrrrrr$sc@sjeZdZdZdddZdddZdd d Zdd d ZdddZdddZ d ddZ d!ddZ d"ddZ dS)#rz7 StorageBox allow in-place mutation of Tensors rrcCs&t|jttfr|jtjjvSdSr)rrfrrlrrrrr rrrrr!%szStorageBox.is_input_buffercCst|jto|jtjjvSr)rrfrrrrrrHrrrrrn%s zStorageBox.is_module_bufferr3cCst|jr |jSt|jttttfsJt |j|j }|j }|j }|dus1Jt dt||j|jdd|jd|_tj|j|j_tj|j|j|j_||j_||j_|jjj|j_|jjS)NF)rrrrkr=)rrrfrrrrrr7rrrr'rr$r&r!rrrr)rr*rrrr)rrrrrrrr %s4        zStorageBox.realizercCs4t|jttfr|jjdkr|dSdSdS)zL Called on buffers we expect to be forced to realize later. rDN)rrfrrrnontrivial_read_countr rrrrr(,%s  zStorageBox.realize_hintr"rcs^ddlmfdd|D}|sdSt|}t|}t|}||ko.||dko.||kS)Nris_nonfreeable_bufferscs g|] }|stj|qSr)rrrget_dep_size_hintrr% rrr9%s  zCStorageBox.has_accumulated_enough_reads_by_size..Fr)rr& rrrr)rr" size_of_reads total_sizemax_sizemin_sizerr% r$has_accumulated_enough_reads_by_size6%s   z/StorageBox.has_accumulated_enough_reads_by_sizecCs8t|jto|tjkp|ptjduo|tjSr) rrfrr;rErealize_acc_reads_thresholdr$ realize_acc_reads_size_thresholdr, rrrrrI%s  z!StorageBox.has_exceeded_max_readsr%csh|dkr2t|jttfr2t|jr'|jddg}tfdd|Dr'dS|tj kp1| SdS)zj A heuristic to decide if we should realize a tensor that is used multiple times. rDexpsigmoidc3s|]}|jvVqdSr)used_opsrtopcountrrr_%rz5StorageBox.should_realize_on_reuse..TF) rrfrrrCrrr;rErealize_reads_thresholdr$)rr% heavy_opsrr2 rshould_realize_on_reuseU%s  z"StorageBox.should_realize_on_reusecCs||r |dSdSr)r6 r r&rrrr'g%s  zStorageBox.mark_reusecCrdrrrrrrr;k%rzStorageBox.num_readsNrRrUrI)r"rrr)r%rrrr^rf) rrrrkr!rnr r(r, rr6 r'r;rrrrr$s      rc@s*eZdZUded<ded<dZded<dS)Subgraphrrrs graph_moduleNzGraphLowering | Noner)rrrrrrrrrr7 o%s r7 bufferscCs,dd|D}ttdd|Dt|kS)NcS"g|] }t|tr |n|qSrrrlrrrRrrrrw%z(_has_aliased_buffers..csrr)r[r< rrrr|%rz'_has_aliased_buffers..)rr=)r9 rrr_has_aliased_buffersv%sr> csjeZdZUdZdZded<dZded<dZded<dfdd ZdddZ e dddZ dddZ Z S)InvokeSubgraphz. Ir node for the invoke_subgraph HOP. NSubgraph | NonerSequence[IRNode] | Noneoperandsrr7 rrgr2rrcs6tjd||d||_tj||_tj|dSr)rrrrrrr)rr*)rrrB rgr rrr%szInvokeSubgraph.__init__rcCs|jr|jgSgSr)rrrrrr%rzInvokeSubgraph.get_subgraphsr@list[ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput]csddlm}tjj}d}|jd}r5d}|jtj j j ur,|j dtj j j us*Jd}|d|d}n&d}|jtj j j urM|j dtj j j usKJd}|j |d}d d |D}fd d |D}g} t|D]\} } t| tttfr|| | qj| || || qj| }|jdurtjj|j||jd |_t|j|jj|Wdn1swY|jj} d} |D] } t| ts| } nq| dusJt||t| d ddfdd fdd t| D}|_|S)zFor each operand, get a realized input, force it to have the same strides as the subgraph inputs, then use an InvokeSubgraphrD)constrain_to_fake_tensorNeager_input_valsrrrrcSg|]}|jdqSrrzrtrrrr%rz)InvokeSubgraph.create..cr,rr@rtrArrr%rrrrrr)rrB rgrtrindrr:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutputc sdt|ttfr |S|}|dusJtt||||| j | j dt |fgddS)NrTr ) rrrr'rLrir&r!rKr#rjrkr)rtrI r)invoke_subgraphrr create_output%s" z,InvokeSubgraph.create..create_outputcsg|] \}}||qSrr)rrrt)rL rrr%r)rtrrI rrrJ )rorD rrrrrzrrr)rprrOrrK rrrrDrGrrr8 rrr graph_outputsr'r? r2r)rrrB rD r fake_operandsrE rj fx_operands new_operandsroperandrroutsr)rrL rK rr%sh       zInvokeSubgraph.createrrwcCrr)codegen_invoke_subgraphrrrrr&rzInvokeSubgraph.codegen)rr7 rB rrgr2rrrM)rr7 rB rrrC r)rrrrkrrrB rrrrrrrrrr rr? %s      hr? cseZdZUdZdZded<dZded<dZded<dZded <dZ d ed <d-fdd Z d.ddZ e d/ddZ ed0d$d%Zd1d(d)Zd2d+d,ZZS)3 Conditionala IR node representing torch.cond Attributes: predicate: A boolean scalar tensor determining which branch to execute. operands: Input tensors passed to both true and false subgraphs. true_subgraph: Subgraph executed when predicate is True. false_subgraph: Subgraph executed when predicate is False. outputs: MultiOutput nodes representing the conditional's outputs. Nr predicaterA rB r@ true_subgraphfalse_subgraphSequence[MultiOutput] | Nonerrrr7 rgr2rrrrc sj||_||_||_||_t|g|\}}tjd|||d|dur&||_tj ||_ tj |dSN)rrgr~r) rU rB rV rW _split_by_sym_typerrrrrrr)rr*) rrU rB rV rW rgrsym_argsrUr rrr&s zConditional.__init__rcC,g}|jr ||j|jr||j|Sr)rV rrW r subgraphsrrrr4&   zConditional.get_subgraphsrint | torch.SymIntint | sympy.ExprcCst|tr|S|jjSr)rrrr)rrrr _maybe_expr<&s zConditional._maybe_exprrtrue_fnfalse_fnrYlist[MultiOutput]c s|}fddDtjjjd}t|ts Jt|g}|D]}t|tr4| |j dq$| |q$tjjj d}dd d } ||fD]8} | jd urtjj | j || j d| _t| j| jj|| | jj|| j_Wd n1s}wYqJ|jd usJ|jd usJ|jj} |jj} d| fd| ffD]\} }t| rtd| d|qt| t| ksJ| | ftt| | D]5\}\}}||ksJ|||f||ksJ|||f|j|jksJ|||fqtdd|gDttjjjtjjj dd }d us#Jdt|||t d|dfddtt| tjjj dD}|_!ddl"m#}||j \}}}}}||j \}}}}}t$|t$|B}fddt%|D_&|S)zNCreate a Sequence of IRNodes from a conditional statement (see .lowering.cond)cr,rr@rtrArrrN&rz&Conditional.create..rIrrM r fake_tensorsSequence[torch.Tensor]rrcSsPg}t||D]\}}t|tr||q|tjt||ddq|SNFrQ)rrrrrrnrr)rM rf retrtr@rrr_require_exact_strides^&s  z2Conditional.create.._require_exact_stridesNrH rc rd zVOutput aliasing is currently not supported in compiled torch.cond. The outputs of the z% subgraph of torch.cond are aliased: css"|] }t|ts|VqdSr)rrr')rorrrr&s z%Conditional.create..rzcannot determine devicer)rU rB rV rW rgrc svg|]7\}\}}tt|dur|n|dd|Ddd|D|j|jdt |fgqS)NcSrzrrT rb rrrrrr&rz1Conditional.create...cSrzrrl rm rrrr&s r) rLrir'r&rrr#rjrkr)rrrt merged_output) conditionalrrrr&s&   r)-check_input_alias_and_mutation_return_outputscs"g|] }t|j|qSr)r%rgrr)ro rB rrr&r= )rM rrf rg rr)'rMrrrrrrrrr;rrzrr8 rrrrM r> rrrrr'r&r#rjrr8rrrrT r2rtorch._higher_order_ops.utilsrp r=rHr.)rrU rc rd rB rO rN fx_op fake_outputsrj r true_outputs false_outputsrrrt_of_orrp r true_mutated_inputsfalse_mutated_inputsmutated_operand_indicesr)rro rrB rrB&s         $     zConditional.createrrwcCs*|||||jt|didSr)codegen_conditionalrrrrrrrrr&s zConditional.codegenrcCrrrrrrrr&rz$Conditional.get_unbacked_symbol_defs)rU rrB rrV r7 rW r7 rgr2rrrrrM)rr` rra ) rU rrc r7 rd r7 rB rYrre rr)rrrrkrU rrB rV rW rrrrnrb rrrrrrrr rrT &s         rT rr-tuple[list[ShapeAsConstantBuffer], list[Any]]cCs<g}g}|D]}t|tr||jq||q||fSr)rrrr)r non_sym_argsr[ rrrrrZ &s  rZ cseZdZUdZdZded<dZded<dZded<dZded<dZ d ed <d,fdd Z d-ddZ e d.ddZ e d/ddZed0d#d$Zd1d'd(Zd2d*d+ZZS)3 WhileLoopzSThe IR node for while_loop and while_loop_stack_output. It supports input mutation.NrA carried_inputsadditional_inputsr@ cond_subgraph body_subgraphrX rrr7 rgr2rr stack_outputrrrc sr||_||_||_||_tg||\}} tjd|| |d|dur'||_||_t j ||_ t j |dSrY )r r r r rZ rrrr rrrr)rr*) rr r r r rgrr r[ rUr rrr's$  zWhileLoop.__init__rcCr\ r)r rr r] rrrr%'r_ zWhileLoop.get_subgraphscCsrt|s|Sdd|D}t}g}t||D]\}}t||vr*|t|q|t|||q|S)NcSr: rr; r< rrrr8'r= z3WhileLoop._clone_aliased_inputs..)r> r=rr[rrr3r)r unwrapped_buffers seen_buffersroriginal_inputunwrapped_bufferrrr_clone_aliased_inputs0's  zWhileLoop._clone_aliased_inputsr rcCsJt|tr|St|ttfrt|St|trt|Stdt|)NzNYI unsupported output type: )rrrrlrLrrr)r rrr_maybe_wrap_as_tensor_boxJ's   z#WhileLoop._maybe_wrap_as_tensor_boxcond_fnbody_fnrc& sddlm}d*d d }tjjjd }tjjjd } || } d d| D} dd|D} dd| D} fdd|D}t|}||| }fdd|D}||| }||||fD]S}|jdurt| t smJt | tjj |j | |j d|_t|j'|jj| ||urt|jjt| ksJ||jj| |j_Wdn1swYq[|jr|jsJ|jj}|jj}t|rtd|t|dksJ||d}t|ts|tjksJ|t|dksJ|tdksJdd}|dus Jt|t|ksJ||ftt||D]?\}\}}d+dd}||||||||ksOJ||||f||ks_J|||fq!|dushJttjjjtjjj !dd}t||||t"|d||d }|jdurt|jj#tj$j%sJ||jj#| d!}t&|}fd"d|D}t'|}g}g|_(g|_)|rt|dksJd#ttjjj d$D]/\} }!t*t+|!j,|!j-d%d|!.Dd&d|!/Dd'|t0| fg}"|j(1|"|1|"qnVt|D]Q\} }!| |vr/| t|ksJd(t2|}#|j)1t3|#j4|#||1|#qt*t+|!|!|!|!|!5j6d)|t0| fg}"|j(1|"|1|"qt||D]\}$}%|$7tjj8vrrtjj9:|%7q\|S),zcreate the while_loop IR node. stack_output controls whether it stack each iterations' output, which is necessary for training. r)check_input_alias_and_mutation tensor_boxesrrf 'list[int | torch.SymInt | torch.Tensor]rrcSslt|t|ks Jg}t||D]"\}}t|tjr.t|}|tj || ddq||q|Srh ) rrrr)rZr r rrrnr)r rf ri rfknew_tbrrrrj c's    z0WhileLoop.create.._require_exact_stridesrIcSrF rrG rtrrrr'rz$WhileLoop.create..cSrF rrG rtrrrr'rcSrF rrG rtrrrr'rcr,rr@rtrArrr'rcr,rr@rtrArrr'rNrH zOutput aliasing is currently not supported in compiled torch.while_loop. The outputs of the body_fn subgraph of torch.while_loop are aliased: rDz9torch.while_loop is assumed to have at least one operand. lhs_exprsSequence[int | sympy.Expr] rhs_exprsrcSs<t|t|ks Jt||D] \}}tjj||qdSr)rrrrrrr)r r lhsrhsrrr_guard_list_equals'sz,WhileLoop.create.._guard_list_equalsrr)r r r r rgrr rcrrrrq ) all_inputsrrr'rz-NYI: while_loop_stack_output input mutations.rcSrzrrl rm rrrr(rcSrzrrl )rrrrrr(rrzonly carries can be mutated.)rrrrrj)r rrf r rr)r r r r rr);rr r rrrrrr r rrrrr8 rrrrrM r> rrr&r)rr!r'rrrKr8rrrzrr2modulefx GraphModuler=rrr.rLrirrrrrrrr%rgr#rjrr rTr)&rr r r r r r rj fx_carried_inputsfx_additional_inputs fx_all_inputsfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_r cond_outputs body_outputsrrrrvbor r while_loop mutated_idxsmutated_idx_setrmutated_inputs_iter all_outputsrrtr mutated_inputrr r)r rrrU's           "$      zWhileLoop.createrrwcCs.|||j|||jt|didSr)codegen_while_loopr rrrrrrrrr/(szWhileLoop.codegenrcCrrrrrrrr5(rz"WhileLoop.get_unbacked_symbol_defs)r rr rr r7 r r7 rgr2rrr rrrrM)r rrr)r rrr) r r7 r r7 r rr rr rrrrr)rrrrkr rr r r rrrrnr r rrrrrrrr rr &s$           Zr cs@eZdZ ddddfddZdfdd ZdddZZS)rNrrgr6rrrUrrr6r?rrrrrrrc sbtj|||||d|dddlm}||} | dusJ| |_tjj| d|_ |tjj| <dS)Nrr) _get_effect) rrtorch._higher_order_ops.effectsr effect_typerrr effectful_opsrprev_effect_buffer) rrgrrUrr?rrr r r rrrA(s   zEffectfulKernel.__init__r5cs0t}|jdur|jt|j|Sr)rr6r r:rrFrr)rrLr rrr6^(s  zEffectfulKernel.get_read_writesrcCrrrrrrrrQh(rz EffectfulKernel.has_side_effectsrr rdrR)rrrrr6rQrrrr rr@(s  rc@s"eZdZed d d ddZdS) rjFrrrrcCrrr<r=rrrr>m(rz!NonTensorObj.get_free_symbol_usesNr_rh)rrrr]r>rrrrrjl(srjc@sReZdZUded<ded<dddZddd d Zdd dZdddZdddZdS)rrr%FakeScriptObject | torch.ScriptObjectr-rcCrrrrrrrry(rzTorchBindObject.get_nameNr r cCrrrr rrrr |(rz!TorchBindObject.codegen_referencecCrrr0rrrrrJ(rzTorchBindObject.get_valuetorch.ScriptObjectcCst|jtjr |jS|jjSr)rr-r)rreal_objrrrr get_real_obj(szTorchBindObject.get_real_objrcCsZ|}t|r dSt|dsJt|}t|d}dd|D}tt j |dS)Nr__obj_flatten__cSs(g|]}t|tjr||qSr)rr)rZrnumelrtrrrr(s  z1TorchBindObject.get_buf_bytes..) r r)rrr r8rBrWruoperatorr)rreal_script_obj flat_dict flat_elems flat_sizesrrr get_buf_bytes(s zTorchBindObject.get_buf_bytesr\rrV)rr )rr rf) rrrrrr rJr r rrrrrt(s     rc@s0eZdZUdZded<d ddZddd d ZdS)OpaqueValueTypeConstantaIR node for opaque value type constants that appear directly in graph outputs. Unlike TorchBindObject (which references named constants loaded at runtime), this inlines the value's repr into the generated code since value types are reconstructed from their repr. r r-rrcCrBr)rEr-rrrrr(rz OpaqueValueTypeConstant.get_nameNr r cCs0t|j\}}|D] \}}|tjj|<q |Sr)r(r-rrrropaque_value_type_classes)rr obj_repr opaque_typesrr,rrrr (sz)OpaqueValueTypeConstant.codegen_referencer\rrVrrrrkrrr rrrrr (s  r c@s4eZdZUded<ded<d ddZddd d ZdS)rDrrrrrcCrrrrrrrr(rzGeneratorState.get_nameNr r cCrrrr rrrr (rz GeneratorState.codegen_referencer\rrV)rrrrrr rrrrrD(s  rDc@s8eZdZUdZded<ded<dddZddd d Zd S)rGz Represents an opaque object (e.g., ProcessGroup) that is passed through as a graph input. Similar to GeneratorState, this wraps the object with its placeholder name so codegen can reference it properly. rrr r-rcCrrrrrrrr(rzOpaqueObjectState.get_nameNr r cCrrrr rrrr (rz#OpaqueObjectState.codegen_referencer\rrVr rrrrrG(s  rGc@sHeZdZdddZdddZddd d ZedddZedddZdS)_CollectiveKernelrrcCrrrrrrrr(rz!_CollectiveKernel.should_allocatecCrrrrrrrrQ(rz"_CollectiveKernel.has_side_effectsNrr3rcCsRt|jtjjus Jd|j}|dur||_n|jj|_dd|jjD|_ dS)Nz,Setting cpp kernel needs a valid op_overloadcSrrrrtrrrr(r z9_CollectiveKernel.set_cpp_kernel_name..) rrr)r r rr rr r)rrrrrrr(s  z%_CollectiveKernel.set_cpp_kernel_namerrr~IRNode | list[IRNode]rr rc s$tjj|j||g|Ri|\}}}}} Wdn1s!wY| r1J|d| |D]} | tj| q3|d|td||||t |} j fdd| Dj dd| Dd|vrj ttd|dj |ddSdS)Nrrrcr"r#r$r&rrrrr )rz4_CollectiveKernel.create_inplace..cSrrrr rrrr)rr )rrrr#rar rrr'rr8 tree_leavesr.rrrr%) rrr~rr_example_outputrUrVr?r tensor_arginpsrr rcreate_inplace(s@   z _CollectiveKernel.create_inplaceTensorBox | list[TensorBox]%list[MultiOutput] | _CollectiveKernelcsJtjjj||g|Ri|\}}}}} Wdn1s!wY| r1J|d| |D] } t| ts>| q3t|tr||} | dusPJt | d||||fddt |D_ t j |D]\} } t jsyt| stjj| jqnj S|||||t jst|stjjjg_ S)Nrtrcs(g|]\}}t|t|fgqSr)rLr]r)rrr]rrrrrK)sz9_CollectiveKernel.create_out_of_place..)rrrr#rarrr rrr2rrrrErrorWrrr])rrr~rrr\rUrVr?rr rrOr]rr rcreate_out_of_place,)s^      z%_CollectiveKernel.create_out_of_placerRrr) rrr~r rr rr rr) rrr~r rr rr rr ) rrrrrQrrr r rrrrr (s    Ar c2eZdZ ddddfddZdddZZS)_AllReduce_KernelNrrgr6rrrUrrr6r?rrrrrrrc (tj|||||d|d|ddS)Nr+aoti_torch_cpu__c10d_functional_all_reduce_rrrr r rrrj)  z_AllReduce_Kernel.__init__rrwcC2|d||t|jtr||dSdSNz+torch/csrc/inductor/aoti_torch/c/shim_cpu.hinclude_extra_headerrrrgrrrrrrr)   z_AllReduce_Kernel.codegenrr rrrrr rr i)  r cr )_AllReduceKernelNrrgr6rrrUrrr6r?rrrrrrrc r )Nr*aoti_torch_cpu__c10d_functional_all_reducer r r rrr)r z_AllReduceKernel.__init__rrwcCr r r rrrrr)r z_AllReduceKernel.codegenrr rrrrr rr )r r csXeZdZ d"ddd#fddZd$ddZd%ddZed&ddZd'fd d! ZZ S)( _WaitKernelNrrgr6rrrUrrr6r?rrrrrrrc r )Nr+aoti_torch_cpu__c10d_functional_wait_tensorr r r rrr)r z_WaitKernel.__init__rrwcCr r r rrrrr)r z_WaitKernel.codegencCs|jd}t|ts Jt|tr$|jd}t|ts!Jt||gSt|trB|jd}t|tr@|jd\}}|j|gSgSgSr)r~rrr rrLrM)rrrcollr rrrrget_volatile_reads)s       z_WaitKernel.get_volatile_readsrrc Cstjj|||\}}}}}Wdn1swY|r*J|d||t|d||||}|jtt|d||dS)Nrr) rrrr#rarr'r.rr%) rrrr rUrVr?rrrrr create_wait)s(   z_WaitKernel.create_waitr5cs6t}|}|D] }|jt|q |Sr)rr6r r:rrFrr)rrLvolatile_readsvrr rrr6)s z_WaitKernel.get_read_writesrr rr)rrrrrrrd) rrrrrr rr r6rrrr rr )s    r rrcCdt|ttfr t|St|ttfr%ttj}|D]}|t |O}q|St|t j r/t|StSr) rr:rr3rrr=rr rr)rZrrr,rrrr*  rcCr r) rr:rr2rrr=rr rr)rZr rrrr*r rrr r torch.fx.NodecCst|trht|jtrjt|jjtr|jjd|dSt|jjtrl|jjd|t|jjtrEt|jjjtrE|jjjd|dSt|jjtrn|jjj spt|jjj dtrr|jjj dd|dSdSdSdSdSdSdS)Nrr) rrrfrrrrrrLrMr~)rrrrrassign_origin_node *s&   r )rrrr)rrrr)rrrr)rrrr)rrrrrrr)rrrrrr)rrrrr_)rrrrrr)rrrrrr)rrrrrr)r-r.rr/)rr2rr3)rr8rr)rrDrErrr)rTrUrVrUrWrUrr)r]rr^r_rr)rrrsrr)r~rrr)rrrrrrqrN)rrrrrrrr)rrrrrjrrrr~)TFNFN)rrrErrLrrMrNr*rr2rNrrO)rrrMrrrrG)rrUrWrUrr)rrrr)rrrr)r rbrr)r9 rrr)rrrr} )rrrr)rr rr rr(\ __future__rrorrlrWrvloggingr textwraprcollections.abcrrrrrrr enumr r typingr r rrrrrrrrrtyping_extensionsrrrrrr unittest.mockrrrrr torch._export.serde.schema_exportserderrrrrzrtorch._loggingr)torch.fxtorch.utils._pytree_pytreer8torch._dynamo.utilsr!torch._export.serde.serializer"*torch._higher_order_ops.auto_functionalizer#torch._inductorr$rr&"torch._library.fake_class_registryr'torch._library.opaque_objectr(r)torch._prims_commonr*r+r,r-r.r/%torch.fx.experimental.symbolic_shapesr0r1r2r3r4r5r6r7r8r9r: torch.fx.noder;torch.utils._ordered_setr=torch.utils._python_dispatchr>torch.utils._sympy.functionsr?r@rArBtorch.utils._sympy.symbolrCrrErFcodegen.commonrGrHrIrJrKrLrMrNrOrPrQ loop_bodyrR ops_handlerrSrTrUrVruntime.benchmarkingrW runtime.hintsrXrYrZr[r\r]r^r_r`rarbrcrdrerfrgrhrirjrkrlrmrnro virtualizedrprqrrrsrt torch.typesrucodegen.cutlass.templatervcodegen.wrapperrwrrxryrrr= __version__r3r2 ImportErrorrzr{r|r}rr~rrr r r+r getLoggerrr\rrrrrrrrrrrrrr]r_rrrrr1r7rBrCrSr\rqr}rrrrrrrrrrrrrrrrrrr4r7r"rKrhrTrUrrrrrrrlrrrrr rrr6rrir$rurrrrrrrrrrrrrrrhrrPrimitiveInfoTyperprrrrrrrrrrrrrrr%rrrrrrCrKrOrRrUrZr`rfrmrwrrrr\r r r2rLrr rrrrrr7 r> r? rT rZ r rrjrr rDrGr r r r rrr rrrrs  4        4      `        *    #      "     '   XX (# ?/ :&dK   < a^0D$ V)o%_$"Y   : a,?O_]UOO%+!@(?\-29,</5?{ L#kq iE,% Y