)`i DdZddlZddlZddlZddlmZddlZddlm cm cm Z ddl Z ddlmZddlmZmZmZmZdefdZdedejfd Zd Zejd e jdefd ZGd deZej dfdeedeeejfdedefdZ!dS)a3 Copyright (c) 2025 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)Union)ir) AddressSpaceNumericPointerTypereturnctjdduo tjdduS)Ncutlassz cutlass.cute) importlibutil find_specm/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/cute_dsl/utils.pyis_cute_dsl_availablers=  ++47 A N $ $^ 4 4D @rdtypectjtjtjtjtjtjtjd}||S)N)float16bfloat16float32 float8_e5m2 float8_e4m3fnfloat8_e8m0fnu float4_e2m1fn)r Float16BFloat16Float32 Float8E5M2 Float8E4M3FN Float8E8M0FNU Float4E2M1FN)r dtype_maps rget_cutlass_dtyper$$sC?$?) -!/ -I U rctt|jd}tjtjtjtjtjtj tj tj tj tj tjtjtjtji}|||}|t'|d|S)zE Return the corresponding torch.dtype per the given DSL type Nz is not supported by torch)getattrtorch__name__lowerr TFloat32rrrrrrrrr rFloat8E4M3B11FNUZfloat8_e4m3fnuzget TypeError) cutlass_dtype torch_dtypetorch_type_maps rcutlass_to_torch_dtyper21s%!7!=!=!?!?FFK %-%.E-e1!5#8N$((77 =DDDEEE rdevicecJtj|jSN)r'cudaget_device_propertiesmulti_processor_count)r3s r get_num_smr9Hs : + +F 3 3 IIrceZdZdZejjdfdejfdZdefdZ dZ dZ d Z e dejfd Ze deefd Ze d Zddd dedefdZdZdefdZdZdS)_Pointera_Runtime representation of a pointer that can inter-operate with various data structures, including numpy arrays and device memory. :param pointer: The pointer to the data :type pointer: int or pointer-like object :param dtype: Data type of the elements pointed to :type dtype: Type :param mem_space: Memory space where the pointer resides, defaults generic :type mem_space: _cute_ir.AddressSpace, optional :param assumed_align: Alignment of input pointer in bytes, defaults None :type assumed_align: int, optional :ivar _pointer: The underlying pointer :ivar _dtype: Data type of the elements :ivar _addr_space: Memory space of the pointer :ivar _assumed_align: Alignment of the pointer in bytes :ivar _desc: C-type descriptor for the pointer :ivar _c_pointer: C-compatible pointer representation N mem_spacec||_||_||_||jdz|_n||_d|_d|_t|j|jzdksJd|jddS)Nrzpointer must be z bytes aligned)_pointer_dtype _addr_spacewidth_assumed_align_desc _c_pointerint)selfpointerrr< assumed_aligns r__init__z_Pointer.__init__ds   $  "'+"2D  "/D  4=!!D$771<<< Bt2 B B B=<<<)rFr?rArOs r__str__z_Pointer.__str__s+EDM**EEE$2BEEEErc*|Sr5)rrrOs r__repr__z_Pointer.__repr__s||~~r)r( __module__ __qualname____doc__r]rgenericrJrFrPrSrVr[propertyrrrRrrrarrgrmstrrrrtrrrr;r;Os0,4+@+H   (     ,BsBBBB   !!!   27   X tG}X  X ,0DGGGsGWGGGG   FFFFFrr;valuer<c,t|tr|}njt|tjr1tj|tjj}| Jdntdt|t||||S)aCreate a pointer from a memory address :param dtype: Data type of the pointer elements :type dtype: Type[Numeric] :param value: Memory address as integer or ctypes pointer :type value: Union[int, ctypes._Pointer] :param mem_space: Memory address space, defaults to AddressSpace.generic :type mem_space: AddressSpace, optional :param assumed_align: Alignment in bytes, defaults to None :type assumed_align: int, optional :return: A pointer object :rtype: Pointer .. code-block:: python import numpy as np import ctypes from cutlass import Float32 from cutlass.cute.runtime import make_ptr # Create a numpy array a = np.random.randn(16, 32).astype(np.float32) # Get pointer address as integer ptr_address = a.ctypes.data_as(ctypes.POINTER(ctypes.c_float)) # Create pointer from address y = make_ptr(cutlass.Float32, ptr_address) NzPointer address is Nonez;Expect int or ctypes.POINTER for value but got type(value)=)rI) rirFrLr;castrNr{r.type)rr{r<rI address_values rmake_ptrrsJ%   E6? + +  E6?;;A ((*C(((( Ld5kk L L    M5)= Q Q QQr)"rwrL functoolsimportlib.utilr typingrr cutlass._mlir.dialects.cute_mlirdialectscuter]r' cutlass._mlirrcutlass.cute.typingrrrrboolrrzrr$r2cacher3rFr9r;rxrrrrrs  ............ DDDDDDDDDDDDt S W]    . Ju|JJJJJ ]]]]]w]]]F+2 0R0R =0R fo% &0R0R  0R0R0R0R0R0Rr