)`iNr UdZddlZddlZddlmZddlmZmZmZm Z m Z m Z m Z ddl Z ddlZ ddlZddlmZddlmZddlZddlmZGd d eZGd d eZGd deZdZGddeZGddeZGddeZde jde de jfdZ!de jde de jfdZ"de#de#fdZ$ dde#de#de#d e#de#f d!Z%d"e ddfd#Z&de ddfd$Z'de jde(fd%Z)de jde jfd&Z*d'e e je e je jffde de e je jffd(Z+d)e#de jfd*Z,iZ-ee e e j.fe jfe/d+< dd-e d.e#d/e j.d0e(de jf d1Z0de#de#fd2Z1d3e#d/e j.de jfd4Z2d5e#d/e j.de jfd6Z3d7e e j4e fde j4fd8Z5ej6d/e j.de e#e#ffd9Z7ej6d/e j.de8fd:Z9d;e jde j4ddf d?Z:eeed@krW ddddAd-e dBe edCe e ee fdDe e e e e fdEe e def dFZ; dd-e dBe edefdGZe j4de(f dMZ>d"e#dKe(dLe(d=e j4d>e j4de(f dNZ?d/e j.d"e#dKe(dLe(d=e j4d>e j4de fdOZ@dPe dQe de(fdRZAde(fdSZBddTlCmDZDmEZEde fdUZFd/e j.de(fdVZGd/e j.de(fdWZHd/e j.de(fdXZId/e j.de(fdYZJd/e j.de(fdZZKd/e j.de(fd[ZLd/e j.de fd\ZMde jd]e e e#d^e e j4d_e e j.d-e ddf d`ZNej6daZOGdbdceZPePjQePjRePjSePjTePjUePjVddZWdee ddfdfZXej6d/e j.de(fdgZYde#dhe#de#fdiZZde#dhe#de#fdjZ[ej6d/e j.de#fdkZ\GdldmZ]gdnZ^gdoZ_dpe#de#fdqZ`dre jdpe#de jfdsZa ddre jdpe#due#de jfdvZbdwZcdxee#defdyZd ddzee efd{e ed|e edefd}Zeej6d/e j.fd~ZfdS)a3 Copyright (c) 2023 by FlashInfer team. Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at http://www.apache.org/licenses/LICENSE-2.0 Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. N)Enum)CallableDictIterableOptionalSequenceTupleUnion) TorchVersion) __version__)gen_spdlog_moduleceZdZdZdZdZdS)PosEncodingModerr N)__name__ __module__ __qualname__NONE ROPE_LLAMAALIBId/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/utils.pyrr s DJ EEErrceZdZdZdZdZdZdS)MaskModerr rN)rrr NON_CAUSALCAUSALCUSTOMMULTIITEMSCORINGrrrrr&s%J F FrrceZdZdZdZdS) TensorLayoutrr N)rrrNHDHNDrrrr#r#-s C CCCrr#g+eG?ceZdZdZdS)GPUArchitectureErrorz5Custom exception for GPU architecture-related errors.Nrrr__doc__rrrr'r'5s??Drr'ceZdZdZdS) LibraryErrorz,Custom exception for library-related errors.Nr(rrrr+r+;66Drr+ceZdZdZdS)BackendSupportedErrorz,Custom exception for backend-related errors.Nr(rrrr.r.Ar,rr.x kv_layoutreturnc|jdvrtd|jdkrX|dkr|dS|dkr|dStd||S) N)zx must be 4D or 5Dr3r$r%Invalid kv_layout {}ndim ValueError unsqueezeKeyErrorformatr/r0s r _expand_5dr?GvV-...v{{   ;;r?? " %  ;;r?? "188CCDD D Hrc|jdvrtd|jdkrX|dkr|dS|dkr|dStd||S) N)rr3zx must be 3D or 4Drr$r5r%r6r7r8r>s r _expand_4drBYr@rc|dkrdS|dz }||dz z}||dz z}||dz z}||dz z}||dz z}||dz z}|dzS)Nr rr3 r)r/ns rnext_positive_power_of_2rHksr1uuq AAaKAaKAaKAaKAbLAbLA q5Lr num_tokens num_expertstop_kmax_tile_tokens_dimcd}||z|z}t||z}t|}tt|d|}|S)Ng?rD)intrHminmax)rJrKrLrMimbalance_factornum_tokens_per_experttile_tokens_dims rcalculate_tile_tokens_dimrU|sa(%/K? 58H HII./DEEO#oq113FGGO rpos_encoding_modecttt|s"td|dS)NzInvalid pos_encoding_mode {})hasattrrr<r=)rVs r_check_pos_encoding_moderYs@ ?$5 6 6Q5<<=NOOPPPQQrcttt|s"td|dS)Nr7)rXr#r<r=)r0s r_check_kv_layoutr[s> < + +A-44Y??@@@AArc@|jtjtjfvSN)dtypetorch float8_e4m3fn float8_e5m2r/s r is_float8rcs 7u*E,=> >>rc|tj}tj|jddz|j|j}|d|dd<|S)Nrr r^device)tor_int64zerosshaper^rfcumsum)r/rets r get_indptrrmsU U[A +agaj1nAGAH E E EChhqkkCG Jrpaged_kv_cachecVt|tr%|\}}t||t||fStj|r-t ||}|d\}}||fStdt|)Nr )dimzQUnrecognized paged_kv_cache type {}, expect a single tensor or a tuple of tensor.) isinstancetuplerBr_ is_tensorr?unbindr<r=type)rnr0 paged_k_cache paged_v_caches r_unpack_paged_kv_cacherxs.%(( '5$ } }i 0 0 }i 0 0    ( (  #NI>>'5'<'<'<'C'C$ }m++ _ f f^$$     rn_headsc dtjtj|z}dd|z z}tj|tjdd|z}||krPdd|z z}tj|tjddd||z zzd}tj||g}|S)Nrg@g r g)mathfloorlog2r_powarangecatfloat)ryrGm_0mm_hat_0m_hats rget_alibi_slopesrs TZ '** + ++A $( C #u|Aq1u--..A7{{$(# '5<1qGaK7H3H!#L#LMM Iq%j ! ! 7799r _cache_bufFnamebytesrf zero_initc||f}t|}||d|krO|r"tj|tj|}n!tj|tj|}|t|<|S)Nrre)rgetsizer_riuint8empty)rrrfrkeybufs r_get_cache_bufrs .C ..  C {chhqkkE))  G+e5;vFFFCC+e5;vFFFC 3 Jrc6d|dz zS)Nr ) bit_lengthrbs r _ceil_pow2rs Q""$$ $$rseq_lenct|}d||f}t|}|+tj||tj}|t|<|d|S)Nrange_)rfr^)rrrr_rint32)rrf seq_len_pow2rrs r_get_range_bufrsfg&&L "L " "F +C ..  C {l<ekJJJ 3 xx=r num_qo_headscd||f}t|}|,t||}|t|<|S)N alibi_slopes_)rrrrg)rrfrrs r_get_cache_alibi_slopes_bufrsU *< ) )6 2C ..  C {|,,//77 3 Jrr^ct|trtt|St|tjr|St dt|)Nz-dtype must be a string or torch.dtype, got {})rqstrgetattrr_r^ TypeErrorr=rur^s rcanonicalize_torch_dtypersf% ue$$$ E5; ' '   ; B B4;; O O   rc~|jdkrtdtj|jS)Ncudazdevice must be a cuda device)rur:r_rget_device_capabilityindexrfs rget_compute_capabilityrs6 {f7888 : + +FL 9 99rct|trtj|}|jdkrt d||j|jnd}tj tj |}tj |}tj |tj }||zdzdz dz }|tj S#tj wxYw)a Get GPU memory bandwidth in GB/s for the specified CUDA device. Args: device: torch.device object, e.g., torch.device('cuda:0') Returns: float: GPU memory bandwidth (GB/s) Raises: ValueError: If device is not a CUDA device rz"Device must be a CUDA device, got NrrrDi)rqrr_rfrur:rpynvmlnvmlInitnvmlDeviceGetHandleByIndexnvmlDeviceGetMemoryBusWidthnvmlDeviceGetClockInfoNVML_CLOCK_MEM nvmlShutdown)rf device_indexhandle bus_width mem_clock bandwidths rget_gpu_memory_bandwidthrs&#&f%%{fFfFFGGG$*<#;6<<L O 2<@@6v>> 1&&:OPP *Q.!3d: s +ACC*qkdtype_qdtype_kvc|j|krtd|jd|d|j|krtd|jd|ddS)NzThe dtype of q z does not match the q_data_type z specified in plan function.zThe dtype of k z! does not match the kv_data_type )r^r:)rrrrs r_check_cached_qkv_data_typer0s| w' lag l lw l l l    w( nag n n n n n   rz2.4) device_typesschemafn mutates_argsrrcdS)Nc|Sr]rrbs rz$register_custom_op..Hrrrrrrrs rregister_custom_opr?s {rcdS)Nc|Sr]rrbs rrz"register_fake_op..Nrrrrrs rregister_fake_oprJs {rcdS)Nc|Sr]rrbs rrz$register_custom_op..errrrs rrrRs &{rcdS)Nc|Sr]rrbs rrz"register_fake_op..lrrrrs rrrgs {rcdt|\}}|dkrtjjdkrdSdS)N 12.3sm90sm80)rr_versionrrfmajor_s rdetermine_gemm_backendros5%f--HE1 zzem(F22vvruse_fp16_qk_reductionsuse_custom_maskcD|rdS|tjjkrdS|rdSdS)a Check if the FA3 backend is supported based on the given parameters. NOTE(Zihao): this function is a workaround for the lack of support for certain features in our FA3 backend, and will be removed once the backend is fully supported. Parameters ---------- pos_encoding_mode : int The positional encoding mode. use_fp16_qk_reductions : bool Whether FP16 QK reductions are allowed. use_custom_mask : bool Whether a custom mask is used. dtype_q : torch.dtype The data type of the query tensor. dtype_kv : torch.dtype The data type of the key-value tensor. Returns ------- bool True if the FA3 backend is supported, False otherwise. FT)rrvaluerVrrrrs ris_fa3_backend_supportedrws9<uO0666uu 4rc|rdS|tjjkrdS|rdS|tjtjfvrdS|tjtjfvrdSdS)a; Check if the cutlass backend is supported based on the given parameters. Parameters ---------- pos_encoding_mode : int The positional encoding mode. use_fp16_qk_reductions : bool Whether FP16 QK reductions are allowed. use_custom_mask : bool Whether a custom mask is used. dtype_q : torch.dtype The data type of the query tensor. dtype_kv : torch.dtype The data type of the key-value tensor. Returns ------- bool True if the cutlass backend is supported, False otherwise. FT)rrrr_r`rars ris_cutlass_backend_supportedrso8uO0666uu5&(9:::uE'):;;;u 4rcNt|rt|||||rdSdS)a Determine the appropriate attention backend based on the device and parameters. Parameters ---------- device : torch.device The device to be used. mask_mode : int The mask mode. pos_encoding_mode : int The positional encoding mode. use_fp16_qk_reductions : bool Whether FP16 QK reductions are allowed. use_custom_mask : bool Whether a custom mask is used. dtype_q : torch.dtype The data type of the query tensor. dtype_kv : torch.dtype The data type of the key-value tensor. Returns ------- str The name of the attention backend to be used. fa3fa2)is_sm90a_supportedr)rfrVrrrrs rdetermine_attention_backendrsBB&!! &> '' uurr base_versioncfddlm}||||kS)Nr)r) packagingrparse)rr pkg_versions rversion_at_leastrs;000000   W % %):):<)H)H HHrcBddl}|jdduS)z Check if cuda.cudart module is available (cuda-python <= 12.9). Returns: True if cuda.cudart exists, False otherwise rNz cuda.cudart)importlib.utilutil find_spec) importlibs rhas_cuda_cudartrs+ > # #M 2 2$ >>r)has_flashinfer_jit_cachehas_flashinfer_cubincddl}|jS)Nr)rr )rs rget_cuda_python_versionr sKKK  rcpt|\}}|dkottjjdS)Nrrrrr_rrrs rrrs2%f--HE1 A: F*5=+=vFFFrcpt|\}}|dkottjjdS)N 12.8rrs ris_sm100a_supportedr2%f--HE1 B; G+EM,>GGGrcpt|\}}|dkottjjdS)Nr12.9rrs ris_sm100f_supportedrrrcpt|\}}|dkottjjdS)N z13.0rrs ris_sm110a_supportedr"rrc|t|\}}|dko$|dkottjjdS)N rrrrfrminors ris_sm120a_supportedr ';)&11LE5 B; V5A: V*:5=;Mv*V*VVrc|t|\}}|dko$|dkottjjdS)Nrr rrrs ris_sm121a_supportedr ,r rc(t|rdndS)Nrr)rrs rdetermine_mla_backendr1s&v.. 955E9rexpected_shapeexpected_dtypeexpected_devicec*|r:|jtj|krtd|d|d|j|r(|j|krtd|d|d|j|r(|j|krtd|d|d|jdSdS)NzInvalid shape of z : expected z, got zInvalid dtype of zInvalid device of )rjr_Sizer:r^rf)r/rrrrs rcheck_shape_dtype_devicer5s !'UZ%?%??? P P P P Pqw P P    !'^33 P P P P Pqw P P    1866 S S S/ S S S S     66rcBtSr])rbuild_and_loadrrrget_logging_modulerJs    - - / //rc&eZdZdZdZdZdZdZdZdS)LogLevelrr rrr3r4N) rrrTRACEDEBUGINFOWARNERRORCRITICALrrrrrOs, E E D D EHHHrr)tracedebuginfowarnerrorcriticallvl_strchtt|jdSr])r set_log_level log_level_mapr)r's rr)r)bs*&&}W'='CDDDDDrcL|jdkrdSt|\}}|dkS)NrFr)rurrs rdevice_support_pdlr,fs/ {fu%f--HE1 A:ryc||zdz |zS)z Perform ceiling division of two integers. Args: x: the dividend. y: the divisor. Returns: The result of the ceiling division. r rr/r-s rceil_divr0ns EAI! rc(t|||zS)z'Round up x to the nearest multiple of y)r0r/s rround_upr2|s Aq>>A rcJtj|jSr])r_rget_device_propertiesmulti_processor_countrs rget_device_sm_countr6s : + +F 3 3 IIrc `eZdZdZ d dejdejdedeeedffd Z dS) FP4Tensora@Wrapper class for FP4 tensors. Since PyTorch doesn't natively support FP4, this wrapper contains: - data: uint8 tensor storing the compressed FP4 data, the size of innermost dimension is ceil(original_dim / 2) since each uint8 stores 2 FP4 values - scale: float8_e4m3fn tensor storing the scale factors rNdatascalescale_start_indexoriginal_shape.c |jtjkrtd|j|jtjkrtd|j|jddzdkrtd|jd|dks||jdkr td|d|jd||jdz|jdkr.td|d |jdd|jd||jd d |d d krtd |jd |t j|d dz }|jd |kr)td|jd d|d d|||_||_ ||_ ||_ d|_d S)a3Initialize FP4Tensor. Parameters ---------- data : torch.Tensor uint8 tensor storing the compressed FP4 data scale : torch.Tensor float8_e4m3fn tensor storing the scale factors scale_start_index : int The start token index of the scale factors. This is needed when two kernels (like prefill and decode kernels) are reusing the same scale factor tensor with different offsets. original_shape : Optional[Tuple[int, ...]] The original shape before compression. zdata must be uint8 tensor, got z(scale must be float8_e4m3fn tensor, got rrIz.scale.shape[0] must be a multiple of 128, got zNscale start index must be in the range [0, scale.shape[0]). scale_start_index=z, scale.shape[0]=zTscale start index + data.shape[0] must not exceed scale.shape[0]. scale_start_index=z, data.shape[0]=NzVdata and original_shape must have the same dimensions except the last one. data.shape=z, original_shape=rzIdata last dimension must be ceil(original_shape[-1] / 2). data.shape[-1]=z, original_shape[-1]=z , expected=nvfp4) r^r_rr:r`rjr{ceilr9r:r;r<)selfr9r:r;r<expected_data_dims r__init__zFP4Tensor.__init__s_( : $ $KtzKKLL L ;%- - -U UUVV V ;q>C 1 $ $QQQQ  q $5Q$G$GZ%6ZZINUVZZ  tz!} ,u{1~ = =y%6yyHL ST yyhmhstuhvyy   %z#2#.""555 P"&*PP?MPP !% .*>l>|> > !sCx1 >>>>>>rr8)rrDr rrrrrr3rr4 ) rrDrEr rrrrrr3rr4rErFrGrHrIepilogue_tile_mc d}|dzdkrd}|S)NrErIrrFr)rYshuffle_block_sizes rget_shuffle_block_sizer\s$!! r input_tensorc|dksJd||j\}}t|}|dkrtnt}||zdks Jd|t j|t j}t|D]!}||z}||z} || } ||z| z} ||| <"|S)z Higher-level PyTorch approach to reorder the rows in blocks of size 16 or 32. - We do NOT try to handle custom e2m1 memory usage (i.e. no 'K/2' bytes). - Instead, we purely reorder rows in a standard PyTorch shape [M, K]. r(input_tensor should be a 2D tensor, not rErz+input_tensor.shape[0] must be multiples of r) rprjr\srcToDstBlk16RowMapsrcToDstBlk32RowMapr_rlongrange) r]rYMKr[row_map row_indicesold_row block_idx row_in_blockmapped_row_in_blocknew_rows r get_shuffle_matrix_a_row_indicesrms      " " "G<3C3C3E3EGG # " "  DAq0@@%72%=%=!!CVG ! !Q & & &J6HJJ ' & &+auz222K88''11 !33 %l3003FF& G rrEnum_elts_per_sfcH|jtjks|jtjksJ|dks|dksJ|dksJd||j\}}|dzdksJ|dzdksJt ||}|S)NrErFrr_rIrr3)r^r_rbfloat16rprjrm)r]rYrnrdrergs r#get_shuffle_matrix_sf_a_row_indicesrqs   , , 0Ben0T0T0T T b Or$9$9$9 9      " " "G<3C3C3E3EGG # " "  DAq s7a<<<< q5A::::2<QQK rc\ttdr tjStjS)zFget native fp4 datatype if supported in Torch, otherwise return uint8.float4_e2m1fn_x2)rXr_rsrrrrget_native_fp4_dtypert+s&u())%%{r supported_ccsc  t|}n3#t$r&tdt|jdwxYwgt |D]\}}t |t rtd|d|t |ts*td|dt|jd||fd}|S)a Decorator to mark functions with their supported CUDA compute capabilities. This decorator annotates a function with metadata about which CUDA compute capabilities (CC) it supports. It adds a `_supported_ccs` attribute containing the set of supported compute capabilities and an `is_compute_capability_supported` method to check if a specific compute capability is supported. Parameters ---------- supported_ccs : list or iterable of int A list of supported CUDA compute capability versions as integers (e.g., [75, 80, 86, 89, 90, 100, 103, 110, 120]). These are computed as major * 10 + minor (e.g., SM 8.0 = 80, SM 9.0 = 90). Returns ------- decorator : callable A decorator function that adds compute capability metadata to the decorated function. Attributes Added to Decorated Function --------------------------------------- _supported_ccs : set of int A set of integers representing the supported compute capabilities. is_compute_capability_supported : callable A method that takes a compute capability (int) and returns True if it's supported, False otherwise. Examples -------- >>> @supported_compute_capability([80, 86, 89, 90]) ... def my_kernel_function(): ... pass ... >>> my_kernel_function._supported_ccs {80, 86, 89, 90} >>> my_kernel_function.is_compute_capability_supported(80) True >>> my_kernel_function.is_compute_capability_supported(75) False Notes ----- This decorator is useful in conjunction with the backend_requirement decorator to mark functions with their supported CUDA compute capabilities. Raises ------ TypeError If supported_ccs is not iterable or contains non-integer values. z'supported_ccs must be an iterable, got Nzsupported_ccs[z ] must be an integer, got bool: z] must be an integer, got z: cJt_fd}|_S)Nc|jvSr])_supported_ccs)ccfuncs ris_cc_supportedzHsupported_compute_capability..decorator..is_cc_supported}s,, ,r)setryis_compute_capability_supported)r{r| validated_ccss` r decoratorz/supported_compute_capability..decoratorzs<!-00 - - - - -0?, r) listrrur enumeraterqboolrOappend)ruccs_listirzrrs @rsupported_compute_capabilityr3s6h &&  Td=6I6I6R T T    M8$$!!2 b$   VTQTTPRTTUU U"c"" WWWd2hh>OWWSUWW  R     s 0Abackend_checks common_checkheuristic_funccfd}|S)a Decorator to enforce backend and problem size requirements for kernel functions. This decorator validates that a function is called with a supported backend and compute capability, and optionally validates problem size constraints. It performs runtime checks before executing the function and raises appropriate errors if requirements are not met. If checking overheads are a concern, you can pass a `skip_check` keyword argument to the function to bypass the validation. Parameters ---------- backend_checks : dict A dictionary mapping backend names (str) to requirement checker functions. Each checker function should accept the same arguments as the decorated function and return True if the problem size is supported, False otherwise. Checkers can be decorated with @supported_compute_capability to specify which compute capabilities they support. common_check : callable, optional An optional function that performs additional validation checks common to all backends. Should accept the same arguments as the decorated function and return True if requirements are met, False otherwise. In the case where the kernel function does not have any specific backends, this can be decorated with @supported_compute_capability to specify the function's supported compute capabilities. heuristic_func : callable, optional A function that performs heuristic backend selection when backend is "auto". Must be provided if backend is "auto". Does not do anything if backend is not "auto". Should accept the same arguments as the decorated function. Should return an ordered list of runnable backends with the most preferred backend first. When decorated function is not autotuned, the first backend in the heuristic list will be run. When decorated function is autotuned, the backends in the heuristic list will be autotuned over to find the best backend. Returns ------- decorator : callable A decorator function that wraps the target function with validation logic, and inserts the "skip_check" keyword argument to the function. Attributes Added to Decorated Function --------------------------------------- is_backend_supported : callable Method with signature `is_backend_supported(backend, cc=None)` that returns True if the specified backend is supported, optionally for a specific compute capability (cc). is_compute_capability_supported : callable Method with signature `is_compute_capability_supported(cc)` that returns True if any backend supports the given compute capability. Keyword Arguments Added to Decorated Function --------------------------------------------- skip_check : bool (Defaults to False) If True, the function will not be validated. This is useful for performance-critical code paths. Raises ------ BackendSupportedError If the function is called with an unsupported backend or compute capability. ValueError If the problem size is not supported for the given backend. Examples -------- >>> @supported_compute_capability([80, 86, 89, 90]) ... def _cutlass_check(q, k, v, backend): ... # Validate problem size constraints for CUTLASS backend ... return q.shape[-1] <= 256 ... >>> @supported_compute_capability([75, 80, 86, 89, 90]) ... def _cudnn_check(q, k, v, backend): ... # Validate problem size constraints for cuDNN backend ... return True ... >>> @backend_requirement({ ... "cutlass": _cutlass_check, ... "cudnn": _cudnn_check ... }) ... def my_attention_kernel(q, k, v, backend="cutlass"): ... # Backend invocation ... pass ... >>> # Example with kernel function with no backend requirements >>> @supported_compute_capability([80, 86, 89, 90]) ... def _common_size_check(q, k, v): ... return True ... >>> @backend_requirement( ... backend_checks={}, # Empty backend_checks ... common_check=_common_size_check ... ) ... def backend_agnostic_kernel(q, k, v): ... pass Notes ----- - The decorator automatically extracts compute capability from tensor arguments by finding the first torch.Tensor in args or kwargs. - A `skip_check=True` keyword argument can be passed to bypass validation for performance-critical code paths. - All validation is performed before the wrapped function executes. - Works in conjunction with the @supported_compute_capability decorator to provide fine-grained control over backend and architecture support. c P tjd fd  fd fddtf fd dtdtf fd } fdd t j  f d   _ _| _ _ S) Ncstdj|vrdS|}|dSt|dr||SdS)Nz:Invalid is_backend_supported call: no backend choices for FTr~)r:rrXr~)backendrz req_checkerrr{has_backend_choicess ris_backend_supportedzDbackend_requirement..decorator..is_backend_supporteds&&((  `QUQ^`` .00 5,W5 :4[*KLLK&FFrJJJurcs=tdstdjdSt fdDS)Nr~z.Invalid is_compute_capability_supported call: z8 does not have is_compute_capability_supported decoratorc3bK|])}t|do|V*dS)r~N)rXr~).0checkerrzs r zbbackend_requirement..decorator..is_compute_capability_supported..s[ G%FGGD??CCr)rXr:rr~anyvalues)rzrrrs`rr~zObackend_requirement..decorator..is_compute_capability_supporteds&&(( |-NOO$YI^YYY$CCBGGG$2#8#8#:#:rc|d}s|i|S|vrtd|dj|}|i|o||i|S||i|S)Nrz Backend 'z' is not supported for )rr.r)argskwargsrrrrr{rs r_is_problem_size_supportedzJbackend_requirement..decorator.._is_problem_size_supportedsjj++G'&(( 5#|T4V444n,,+OOO OO)1K'#|T4V44Ud9Uf9U9UU"{D3F333rr1c"tSr])r)rsrrzCbackend_requirement..decorator..has_backend_choices,s'' 'rrc |vSr]r)rrs r has_backendz;backend_requirement..decorator..has_backend0sn, ,rc  |i|sdSg}D]M}|} ||i|r*||r||>#t$rYJwxYw Jd|g|Ri|}|sdS| _dS)NFz#Heuristic function must be providedT)r~rr:suitable_auto_backends) rzrrsuitable_backendsrrrrrwrappers rrzFbackend_requirement..decorator..suitable_auto_backends4s' d0Mf0M0M'u " )  ,W5 "{!':%EEbII:*00999!H"--/T--- ./@ R4 R R R6 R R $ u->G *4s2A AAcd}d}|t|z}|D] }t|tjr|}n!|t |j\}}|dz|z}|S)Nr)rrrrqr_rDrrf)rr capability tensor_argall_argsrrr s r_get_capabilityz?backend_requirement..decorator.._get_capabilityJsJJeFMMOO444H!  eU\22!&JE% 6j6GHH u"RZ%/  rc |dd}|s;j|i|}|t|j}|d}|i|} s t d j rx|dkr!|fi|std jn||s%|rd|nd}t jd |d | d i|st d jnu|st jd | d i|st d jn1|r/ -|ddkr|i|}|g|Ri| |i|S)N skip_checkFrzYInvalid @backend_requirement decorator usage: no backend choices and no common_check for autoz$No suitable auto backends found for z with capability z does not support backend ''z"Problem size is not supported for z% does not support compute capability r) popbindapply_defaultsdict argumentsrr:rr.)rrr bound_argskwargs_with_defaultsrrextrarrrr{rrrr~sigrs rrz7backend_requirement..decorator..wrapper_s L%88J1 H&SXt6v66 ))+++'+J,@'A'A$.229==,_d=f== **,,1E$DtxuBDD'&((&((55&  *>  #8 Vt} V V## 43GZHHDN V @J @ @ @TV"#8#'= ^ ^W ^ ^W\ ^ ^## :9QQ.decorators%%        &       $ 4 4 4 4 4 4 4 4$ (T ( ( ( ( ( ( - - - - - - - -        ,   *   8 )8 )8 )8 )8 )8 )8 )8 )8 )8 )8 )8 )8 )  8 )t(<$2Q/)&9#rr)rrrrs``` rbackend_requirementrs9Vmmmmmmm^ rcvtjtjj|jSr])r_rinitdefault_generatorsrrs rget_default_generatorsrs& JOO : ( 66r)rI)Fr])rE)NN)gr)rr{enumrtypingrrrrrr r r_ torch.versionrtorch.torch_versionr r torch_versionr jit.spdlogrrrr#log2e Exceptionr'r+r.rDrr?rBrOrHrUrYr[rrcrmrxrrrf__annotations__rrrrr^rcacherrrrrrrrrrrrjit.envrrrrrrrr r rrrrrrrrrr r*r)r,r0r2r6r8r`rar\rmrqrtrrrrrrrs   MMMMMMMMMMMMMMMMMM  ,,,,,,<<<<<<))))))d t     4        9        9        I    %, 3 5<    $ %, 3 5<    $$OR"%.1HK0QQQQQQ AAAAAA ??$????%,5< %,elEL.H(IIJ   5< %&    .cel<> DsEL()5<7 8===DI     #(< <@  \    %#%#%%%%C%,$| \ E%+s*:$;       :5<:E#s(O::::  $U\$e$$$$N   |     /4{  FKk        < e!4!444"&  =A $     X   C#./  uS(3-%789        "& X  "& =A $ X  C#./  uS(3-%789   ."& X  5<C$$ $$[ $ k $  $$$$N&& &&[ & k &  &&&&R* L**!* * [ * k * ****ZIcIIIIII ? ? ? ? ?  Gu|GGGGG H HHHHH H HHHHH H HHHHH W WWWWW W WWWWW :%,:3::::  | Xc]+ U[) el+         * 000t^ ^ M M ^!  E3E4EEEE u|         J JJJJJFFFFFFFFT      CC%,%14% \%%%%RNP,14GJ \(P P(PPPPj(,)-ZZh'Z8$ZX&Z ZZZZz 75<777777r