L i9ddlZddlZddlmZejdefdZejddeeefdeeeffdZ ejdefdZ ejdefdZ ejdefd Z ejdefd Z ejdefd Zejddefd Zejdefd ZejdefdZejdefdZy)N)Anyreturnc, ddl}y#t$rYywxYw)NrTF)triton ImportError)rs Y/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/torch/utils/_triton.pyhas_triton_packager s s  fallbackc ddl}td|jjdddD\}}||fS#t$r|cYSwxYw)Nrc32K|]}t|ywN)int).0vs r z%get_triton_version..sOSVOs.)rtuple __version__splitr)r rmajorminors rget_triton_versionrsTOV-?-?-E-Ec-J2A-NOO uu~ s8; A A cddl}|jjxr6|jjdk\xr|jj S)Nr r)torchcuda is_availableget_device_capabilityversionhip)rs r_device_supports_tmar#sK  ! " JJ , , .& 8 " !! !c^trtr ddlm}m}yy#t $rYywxYw)Nrcreate_1d_tma_descriptorcreate_2d_tma_descriptorTF)r r#$triton.tools.experimental_descriptorr'r(rr&s r has_triton_experimental_host_tmar*&s9  !      s  ,,cZtrtr ddlm}yy#t$rYywxYw)NrTensorDescriptorTF)r r#triton.tools.tensor_descriptorr-rr,s r%has_triton_tensor_descriptor_host_tmar/7s9  !     s  **c.txs tSr )r/r*r$rhas_triton_tmar2Gs 0 2 X6V6XXr$cNtruddl}|jjr3|jj dk\r|j j r|jjr ddlm }m }yy#t$rYnwxYw ddl m }y#t$rYywxYw)Nrr)&experimental_device_tensormap_create1d&experimental_device_tensormap_create2dTmake_tensor_descriptorF)r rrrr r!r"xputriton.language.extra.cudar4r5rtriton.languager7)rr4r5r7s rhas_triton_tma_devicer;Ls JJ # # % 002f<MM%% YY # # %      B   s$7B B  B B B$#B$ctrsddl}|jjr3|jj dk\r|j j r|jjr ddlm }yy#t$rYywxYw)Nrrr6TF) r rrrr r!r"r8r:r7r)rr7s rhas_triton_stable_tma_apir=lsv JJ # # % 002f<MM%% YY # # % B    s7A?? B  B ctsyddlmdtdtfd}dtdtfd}dtdtfd}||||d dtffd }|S) NFr)get_interface_for_devicedevice_interfacercP|jjjdk\S)N)Workerget_device_propertiesrr@s rcuda_extra_checkz$has_triton..cuda_extra_checks"&&<<>DDIIr$c:ddl}d|jjvS)Nrcpu)triton.backendsbackends)r@rs rcpu_extra_checkz#has_triton..cpu_extra_checks0000r$cy)NTr1rEs r _return_truez has_triton.._return_truesr$)rr8rHmtiac|jD](\}}|}|js||s(yy)NTF)itemsr)device extra_checkr@r?triton_supported_devicess r is_device_compatible_with_tritonz4has_triton..is_device_compatible_with_tritonsG#;#A#A#C  FK7? ,,.;?O3P r$)r torch._dynamo.device_interfacer?rbool)rFrKrMrTr?rSs @@r has_tritonrWsx  GJ3J4J1#1$1 st!  d , --r$c^ddlm}ddlm}|jj }||S)Nr) make_backend)driver)triton.compiler.compilerrYtriton.runtime.driverrZactiveget_current_target)rYrZtargets rtriton_backendr`s%5, ]] - - /F  r$cddlm}t}|d|j}t j |j djjS)Nr) triton_key-zutf-8) %torch._inductor.runtime.triton_compatrbr`hashhashlibsha256encode hexdigestupper)rbbackendkeys rtriton_hash_with_backendrmsR@G \N!GLLN+ ,C >>#**W- . 8 8 : @ @ BBr$))rr) functoolsrftypingrcacherVr rrrr#r*r/r2r; lru_cacher=rWr`strrmr1r$rrss D sCxeCHo d  $     t   YYY t>T4$ .D..D     C#CCr$