§
    8˜ºh5c  ã                   óº  — d dl Z d dlZd dlZd dlZd dlZd dlZd dlmZ d dlmZ d dl	m
Z
 d dlmZ d dlmZ d dlmZ ej                             ej                             e¦  «        ¦  «        Zej                             ed¦  «        gZej                             ed	¦  «        Zd
gZ e j        ¦   «         d„ ¦   «         Z e j        ¦   «         d„ ¦   «         Z G d„ de¦  «        Zd„ ZddddddœZddddddœZ dZ!d„ Z" G d„ d¦  «        Z# e$d„  e%d¦  «        D ¦   «         ¦  «        Z&de&d<   de&d <   d e&d<   d!„ Z'd"„ Z( G d#„ d$e¦  «        Z) G d%„ d&e¦  «        Z*dS )'é    N)ÚPath)Úknobs)Úcompile_module_from_src)Ú_allocation)Ú	GPUTarget)Ú	GPUDriverÚincludeÚlibÚcudac                  óÜ  — t           j        j        x} r| gS t          j        ddg¦  «                             ¦   «         }d„ |                     ¦   «         D ¦   «         }d„ |D ¦   «         }t          j        d¦  «        }|r!|sd„ | 	                    d¦  «        D ¦   «         }d}|r|d	t          |¦  «        z  z  }|d
z  }n
|dz  }|dz  }t          d„ |D ¦   «         ¦  «        s
J |¦   «         ‚|S )Nz/sbin/ldconfigz-pc                 óJ   — g | ] }d |v ¯|                      ¦   «         d         ‘Œ!S )úlibcuda.so.1éÿÿÿÿ)Úsplit)Ú.0Úlines     úe/var/www/html/services/rag-service/venv/lib/python3.11/site-packages/triton/backends/nvidia/driver.pyú
<listcomp>z libcuda_dirs.<locals>.<listcomp>   s0   € ÐUÐUÐU ¸nÐPTÐ>TÐ>TˆD�JŠJ‰LŒL˜ÔÐ>TÐ>TÐ>Tó    c                 óL   — g | ]!}t           j                             |¦  «        ‘Œ"S © )ÚosÚpathÚdirname)r   Úlocs     r   r   z libcuda_dirs.<locals>.<listcomp>   s&   € Ð1Ð1Ð1 S�BŒG�OŠO˜CÑ Ô Ð1Ð1Ð1r   ÚLD_LIBRARY_PATHc                 óŒ   — g | ]A}t           j                             t           j                             |d ¦  «        ¦  «        ¯?|‘ŒBS )r   ©r   r   ÚexistsÚjoin)r   Údirs     r   r   z libcuda_dirs.<locals>.<listcomp>    sA   € ÐsÐsÐs˜ÅÄÇÂÕPRÔPW×P\ÒP\Ð]`ÐbpÑPqÔPqÑArÔArÐs�ÐsÐsÐsr   ú:zlibcuda.so cannot found!
z!Possible files are located at %s.z:Please create a symlink of libcuda.so to any of the files.z<Please make sure GPU is set up and then run "/sbin/ldconfig"z- (requires sudo) to refresh the linker cache.c              3   ó�   K  — | ]A}t           j                             t           j                             |d ¦  «        ¦  «        V — ŒBdS )r   Nr   )r   r   s     r   ú	<genexpr>zlibcuda_dirs.<locals>.<genexpr>(   s@   è è € ÐSÐSÀd�rŒw�~Š~�bœgŸlšl¨4°Ñ@Ô@ÑAÔAÐSÐSÐSÐSÐSÐSr   )r   ÚnvidiaÚlibcuda_pathÚ
subprocessÚcheck_outputÚdecodeÚ
splitlinesr   Úgetenvr   ÚstrÚany)Úenv_libcuda_pathÚlibsÚlocsÚdirsÚenv_ld_library_pathÚmsgs         r   Úlibcuda_dirsr4      s.  € å œ<Ô4Ð4Ðð "Ø Ð!Ð!åÔ"Ð$4°dÐ#;Ñ<Ô<×CÒCÑEÔE€Dð VÐU¨¯ªÑ):Ô):ÐUÑUÔU€DØ1Ð1¨DÐ1Ñ1Ô1€DÝœ)Ð$5Ñ6Ô6ÐØð t 4ð tØsÐsÐ2×8Ò8¸Ñ=Ô=ÐsÑsÔsˆØ
&€CØð ?ØÐ2µS¸±Y´YÑ>Ñ>ˆØÐKÑKˆˆàÐMÑMˆØÐ>Ñ>ˆÝÐSÐSÈdÐSÑSÔSÑSÔSÐXÐXÐUXÑXÔXÐXØ€Kr   c                  ó.   — t           gt          ¦   «         ¢S ©N)Úlibdevice_dirr4   r   r   r   Úlibrary_dirsr8   ,   s   € åÐ+�L™NœNÐ+Ð+r   c                   ó$   ‡ — e Zd Zˆ fd„Zd„ Zˆ xZS )Ú	CudaUtilsc                 óŒ   •— t          | d¦  «        s-t          t          | ¦  «                             | ¦  «        | _        | j        S )NÚinstance)ÚhasattrÚsuperr:   Ú__new__r<   )ÚclsÚ	__class__s    €r   r?   zCudaUtils.__new__8   s<   ø€ Ý�s˜JÑ'Ô'ð 	>Ý ¥¨CÑ0Ô0×8Ò8¸Ñ=Ô=ˆCŒLØŒ|Ðr   c                 óV  — t          t          t          j                             t
          d¦  «        ¦  «                             ¦   «         dt          ¦   «         t          t          ¬¦  «        }|j
        | _
        |j        | _        |j        | _        |j        | _        |j        | _        d S )Nzdriver.cÚ
cuda_utils©ÚsrcÚnamer8   Úinclude_dirsÚ	libraries)r   r   r   r   r    r   Ú	read_textr8   rG   rH   Úload_binaryÚget_device_propertiesÚcuOccupancyMaxActiveClustersÚset_printf_fifo_sizeÚfill_tma_descriptor)ÚselfÚmods     r   Ú__init__zCudaUtils.__init__=   s‹   € Ý%Ý•R”W—\’\¥'¨:Ñ6Ô6Ñ7Ô7×AÒAÑCÔCØÝ%™œÝ%Ýð
ñ 
ô 
ˆð œ?ˆÔØ%(Ô%>ˆÔ"Ø,/Ô,LˆÔ)Ø$'Ô$<ˆÔ!Ø#&Ô#:ˆÔ Ð Ð r   )Ú__name__Ú
__module__Ú__qualname__r?   rQ   Ú__classcell__©rA   s   @r   r:   r:   6   sG   ø€ € € € € ðð ð ð ð ð
;ð ;ð ;ð ;ð ;ð ;ð ;r   r:   c                 ó¼   — | d         dk    rdS |                       d¦  «        rdS i dd“dd	“d
d“dd“dd“dd“dd“dd“dd“dd“dd“dd“dd“dd“dd“dd“|          S )Nr   Ú*ÚCUdeviceptrÚ
tensordescÚCUtensorMapÚi1Úint32_tÚi8Úint8_tÚi16Úint16_tÚi32Úi64Úint64_tÚu1Úuint32_tÚu8Úuint8_tÚu16Úuint16_tÚu32Úu64Úuint64_tÚfp16ÚdoubleÚbf16Úfp32Úf32Úfp64Ú	nvTmaDesc)Ú
startswith)Útys    r   Ú	ty_to_cpprw   Q   sý   € Ø	ˆ!„u�‚|€|Øˆ}Ø	‡}‚}�\Ñ"Ô"ð Øˆ}ðØˆiðàˆhðð 	ˆyðð 	ˆyð	ð
 	ˆyðð 	ˆjðð 	ˆiðð 	ˆzðð 	ˆzðð 	ˆzðð 	�ðð 	�ðð 	�ðð 	ˆxðð 	�ðð  	�]ð!ð" 	ô#
ð 
r   rj   rf   rm   )rn   rp   rq   rr   rs   Ú	pack_fp16Ú	pack_bf16Ú	pack_fp32Ú	pack_fp64ÚiiiKKppOOOOOc                 óœ  ‡‡‡‡— ˆfd„}ˆfd„Šˆfd„Šˆfd„Š ||                      ¦   «         ¦  «        }d„ t          |¦  «        D ¦   «         }d                     ˆfd„|                      ¦   «         D ¦   «         ¦  «        }t          |z   }g }|                      ¦   «         D ]} ‰||¦  «         Œd„ t          |¦  «        D ¦   «         }t	          |¦  «        d	k    r4d
d
                     d„ |                     ¦   «         D ¦   «         ¦  «        z   nd}	g }
|                     ¦   «         D ]b\  }}|dk    rŒ|t          v r&|
                     t          |         › d|› �¦  «         Œ;|
                     t          |¦  «        › d|› �¦  «         Œcd
                     |
¦  «        }g }|                     ¦   «         D ]‹\  }}|d	         dk    r|                     d|› d�¦  «         Œ+|t          v r|                     d|› d�¦  «         ŒN|dk    r|                     d|› �¦  «         Œm|dk    r|                     d|› �¦  «         ŒŒt          t	          |¦  «        ¦  «        }d}d„ |                     ¦   «         D ¦   «         }d„ |                     ¦   «         D ¦   «         }d„ |                     ¦   «         D ¦   «         }d„ |                     ¦   «         D ¦   «         }|                     d¦  «         dt	          |¦  «        d	k    rd
|z   nd› dd
                     |¦  «        › d|                     ˆfd„|                     ¦   «         D ¦   «         ¦  «        › d|› d |	› d!|                     |¦  «        › d|                     |¦  «        › d|                     |¦  «        › d"t	          |¦  «        d	k    rd
d
                     |¦  «        z   nd› d#�}|S )$Nc                 óò  •— g }d}| D �]V}t          |t          ¦  «        �r(|                     d¦  «        �r‰
r‰
|         nd }|dz  }t          j        d|¦  «        }|                     d¦  «        }|                     d¦  «        }|                     d¦  «        dz   }|€C|                     d|z   ¦  «         t          d|z  ¦  «        D ]}	|                     d¦  «         Œn|                     d	¦  «         t          |¦  «        D ]}	|                     d
¦  «         Œt          |¦  «        D ]}	|                     d¦  «         Œ�ŒA|                     |¦  «         �ŒX‰
r|t          ‰
¦  «        k    sJ ‚|S )Nr   rZ   é   ztensordesc<([^[>]*)\[([^]]*)\]é   ú,rX   rc   rt   rb   )
Ú
isinstancer,   ru   ÚreÚmatchÚgroupÚcountÚappendÚrangeÚlen)Ú	signatureÚoutputÚtensordesc_idxÚsigÚmetar„   ÚdtypeÚshapeÚndimÚ_Útensordesc_metas             €r   Ú_expand_signaturez(make_launcher.<locals>._expand_signature~   sž  ø€ ØˆØˆð ð 	#ñ 	#ˆCÝ˜#�sÑ#Ô#ñ #¨¯ª°|Ñ(DÔ(Dñ #Ø:IÐS� ~Ô6Ð6Èt�Ø !Ñ#�åœÐ!CÀSÑIÔI�ØŸš A™œ�ØŸš A™œ�Ø—{’{ 3Ñ'Ô'¨!Ñ+�à�<Ø—M’M #¨¡+Ñ.Ô.Ð.õ # 1 t¡8™_œ_ð -ð -˜ØŸš eÑ,Ô,Ð,Ð,ð-ð —M’M +Ñ.Ô.Ð.å˜t™œð )ð )�AØ—M’M %Ñ(Ô(Ð(Ð(Ý˜t™œð )ð )�AØ—M’M %Ñ(Ô(Ð(Ð(ñ)ð —’˜cÑ"Ô"Ð"Ñ"à"ÐL n½¸OÑ8LÔ8LÒ&LÐ&LÐ&LÐ&LØˆr   c                 ó‚   •— t          | t          ¦  «        r| D ]} ‰||¦  «         Œd S |                     | ¦  «         d S r6   )r‚   Útupler‡   )r�   r‹   ÚxÚ_flatten_signatures      €r   r˜   z)make_launcher.<locals>._flatten_signature¤   s^   ø€ å�c�5Ñ!Ô!ð 	Øð .ð .�Ø"Ð" 1 fÑ-Ô-Ð-Ð-ð.ð .ð �MŠM˜#ÑÔÐÐÐr   c                 óÆ   •— t          | t          ¦  «        r)d                     t          ‰| ¦  «        ¦  «        }d|› d�S | d         dk    rdS | dv rdS t	          | ¦  «        S )Nr�   ú[ú]r   rX   z	PyObject*©Ú	constexprrt   )r‚   r–   r    Úmaprw   )rv   ÚvalÚ_extracted_types     €r   r    z&make_launcher.<locals>._extracted_type¬   sm   ø€ Ý�b�%Ñ Ô ð 	Ø—(’(�3˜°Ñ3Ô3Ñ4Ô4ˆCØ�s�:�:�:ÐØˆaŒ5�CŠ<ˆ<Ø�;ØÐ+Ð+Ð+Ø�;Ý˜‰}Œ}Ðr   c                 ó  •— t          | t          ¦  «        r)d                     t          ‰| ¦  «        ¦  «        }d|› d�S | d         dk    rdS | dv rdS |                      d¦  «        rdS d	d
dddddddddœ
t          | ¦  «                 S )NÚ ú(ú)r   rX   ÚOrœ   rZ   ÚdÚlÚbÚhÚiÚLÚBÚHÚIÚK)
ro   Úlongr_   ra   r]   rd   rh   rj   rf   rm   )r‚   r–   r    rž   ru   rw   )rv   rŸ   Ú	format_ofs     €r   r±   z make_launcher.<locals>.format_of¶   s·   ø€ Ý�b�%Ñ Ô ð 	Ø—'’'�#˜i¨Ñ,Ô,Ñ-Ô-ˆCØ�s�:�:�:ÐØˆaŒ5�CŠ<ˆ<Ø�3ØÐ+Ð+Ð+Ø�3Ø�=Š=˜Ñ&Ô&ð 	Ø�3àØØØØØØØØØð
ð 
õ �B‰-Œ-ôð 	r   c                 ó   — i | ]\  }}||“Œ	S r   r   ©r   rª   Úss      r   ú
<dictcomp>z!make_launcher.<locals>.<dictcomp>Î   s   € Ð>Ð>Ð>™$˜!˜Q��AÐ>Ð>Ð>r   r¢   c                 ó&   •— g | ]} ‰|¦  «        ‘ŒS r   r   )r   rv   r±   s     €r   r   z!make_launcher.<locals>.<listcomp>Ð   s!   ø€ ÐFÐFÐF¨R˜9˜9 R™=œ=ÐFÐFÐFr   c                 ó   — i | ]\  }}||“Œ	S r   r   r³   s      r   rµ   z!make_launcher.<locals>.<dictcomp>Ö   s   € Ð<Ð<Ð<™$˜!˜Q��AÐ<Ð<Ð<r   r   ú, c              3   ó&   K  — | ]\  }}d |› �V — ŒdS )z&_argNr   ©r   rª   rv   s      r   r$   z make_launcher.<locals>.<genexpr>×   s,   è è € Ð LÐ L±°°B ¨  Ð LÐ LÐ LÐ LÐ LÐ Lr   r�   z argrX   Úptr_infoz.dev_ptrÚ_argÚ_storagert   z*tma_ptrz
  c                 óJ   — g | ] \  }}|d          dk    ¯d|› d|› d|› d|› d�	‘Œ!S )r   rX   zDevicePtrInfo ptr_infoz = getPointer(_argr¸   z); if (!ptr_infoz.valid) return NULL;r   rº   s      r   r   z!make_launcher.<locals>.<listcomp>ò   sU   € ð ð ð áˆAˆrØˆaŒ5�CŠ<ˆ<ð 	f ÐeÐe°aÐeÐe¸1ÐeÐeÈaÐeÐeÐeàˆ<ˆ<r   c           	      ó8   — g | ]\  }}|d k    ¯d|› d|› d|› d�‘ŒS )rt   zCUtensorMap* tma_ptrz = getTmaDesc(_argz); if (!tma_ptrz) return NULL;r   rº   s      r   r   z!make_launcher.<locals>.<listcomp>÷   sO   € ð ð ð Ù\aÐ\]Ð_aØ�ÒÐð 	X˜qÐWÐW°AÐWÐWÀaÐWÐWÐWØÐÐr   c           
      ón   — g | ]2\  }}|t           v ¯t           |         › d |› dt          |         › d|› d�‘Œ3S )ú _argz_storage = z(_argz);)ÚFLOAT_STORAGE_TYPEÚFLOAT_PACK_FUNCTIONrº   s      r   r   z!make_launcher.<locals>.<listcomp>û   s_   € ð ð ð áˆAˆrØÕ#Ð#Ð#õ ˜bÔ!ÐYÐY¨ÐYÐYÕ6IÈ"Ô6MÐYÐYÐTUÐYÐYÐYà#Ð#Ð#r   c                 ó*   — g | ]\  }}|d k    ¯d|› �‘ŒS )r�   z&argr   rº   s      r   r   z!make_launcher.<locals>.<listcomp>   s,   € ÐMÐMÐM™U˜Q ¸2ÀÒ;LÐ;Lˆj�QˆjˆjÐ;LÐ;LÐ;Lr   z&global_scratchaR  
#include "cuda.h"
#include <stdbool.h>
#include <Python.h>
#include <dlfcn.h>

static inline void gpuAssert(CUresult code, const char *file, int line)
{
   if (code != CUDA_SUCCESS)
   {
      const char* prefix = "Triton Error [CUDA]: ";
      const char* str;
      cuGetErrorString(code, &str);
      char err[1024] = {0};
      strcat(err, prefix);
      strcat(err, str);
      PyGILState_STATE gil_state;
      gil_state = PyGILState_Ensure();
      PyErr_SetString(PyExc_RuntimeError, err);
      PyGILState_Release(gil_state);
   }
}

#define CUDA_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); }

typedef CUresult (*cuLaunchKernelEx_t)(const CUlaunchConfig* config, CUfunction f, void** kernelParams, void** extra);

static cuLaunchKernelEx_t getLaunchKernelExHandle() {
  // Open the shared library
  void* handle = dlopen("libcuda.so.1", RTLD_LAZY);
  if (!handle) {
    PyErr_SetString(PyExc_RuntimeError, "Failed to open libcuda.so.1");
    return NULL;
  }
  // Clear any existing error
  dlerror();
  cuLaunchKernelEx_t cuLaunchKernelExHandle = (cuLaunchKernelEx_t)dlsym(handle, "cuLaunchKernelEx");
  // Check for errors
  const char *dlsym_error = dlerror();
  if (dlsym_error) {
    PyErr_SetString(PyExc_RuntimeError, "Failed to retrieve cuLaunchKernelEx from libcuda.so.1");
    return NULL;
  }
  return cuLaunchKernelExHandle;
}

static void _launch(int gridX, int gridY, int gridZ, int num_warps, int num_ctas, int launch_cooperative_grid, int launch_pdl, int clusterDimX, int clusterDimY, int clusterDimZ, int shared_memory, CUstream stream, CUfunction function, CUdeviceptr global_scratchz) {
  void *params[] = { a,   };
  if (gridX*gridY*gridZ > 0) {
    // 4 attributes that we can currently pass maxmimum
    CUlaunchAttribute launchAttr[4];
    static cuLaunchKernelEx_t cuLaunchKernelExHandle = NULL;
    if (cuLaunchKernelExHandle == NULL) {
      cuLaunchKernelExHandle = getLaunchKernelExHandle();
    }
    CUlaunchConfig config;
    config.gridDimX = gridX;
    config.gridDimY = gridY;
    config.gridDimZ = gridZ;

    if (num_ctas != 1) {
      config.gridDimX *= clusterDimX;
      config.gridDimY *= clusterDimY;
      config.gridDimZ *= clusterDimZ;
    }

    config.blockDimX = 32 * num_warps;
    config.blockDimY = 1;
    config.blockDimZ = 1;
    config.sharedMemBytes = shared_memory;
    config.hStream = stream;
    config.attrs = launchAttr;
    int num_attrs = 0;

    if (launch_pdl != 0) {
      CUlaunchAttribute pdlAttr = { .id = CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION, .value = 1};
      launchAttr[num_attrs] = pdlAttr;
      ++num_attrs;
    }

    if (launch_cooperative_grid != 0) {
      CUlaunchAttribute coopAttr = { .id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE, .value = 1};
      launchAttr[num_attrs] = coopAttr;
      ++num_attrs;
    }

    if (num_ctas != 1) {
      CUlaunchAttribute clusterAttr = {};
      clusterAttr.id = CU_LAUNCH_ATTRIBUTE_CLUSTER_DIMENSION;
      clusterAttr.value.clusterDim.x = clusterDimX;
      clusterAttr.value.clusterDim.y = clusterDimY;
      clusterAttr.value.clusterDim.z = clusterDimZ;
      launchAttr[num_attrs] = clusterAttr;
      ++num_attrs;

      CUlaunchAttribute clusterSchedulingAttr = {};
      clusterSchedulingAttr.id = CU_LAUNCH_ATTRIBUTE_CLUSTER_SCHEDULING_POLICY_PREFERENCE;
      clusterSchedulingAttr.value.clusterSchedulingPolicyPreference = CU_CLUSTER_SCHEDULING_POLICY_SPREAD;
      launchAttr[num_attrs] = clusterSchedulingAttr;
      ++num_attrs;
    }

    config.numAttrs = num_attrs;

    CUDA_CHECK(cuLaunchKernelExHandle(&config, function, params, 0));
  }
}

typedef struct _DevicePtrInfo {
    CUdeviceptr dev_ptr;
    bool valid;
} DevicePtrInfo;

static inline DevicePtrInfo getPointer(PyObject *obj, int idx) {
  DevicePtrInfo ptr_info;
  ptr_info.dev_ptr = 0;
  ptr_info.valid = true;
  if (PyLong_Check(obj)) {
    ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(obj);
    return ptr_info;
  }
  if (obj == Py_None) {
    // valid nullptr
    return ptr_info;
  }
  PyObject *ptr = PyObject_GetAttrString(obj, "data_ptr");
  if(ptr){
    PyObject *empty_tuple = PyTuple_New(0);
    PyObject *ret = PyObject_Call(ptr, empty_tuple, NULL);
    Py_DECREF(empty_tuple);
    Py_DECREF(ptr);
    if (!PyLong_Check(ret)) {
      PyErr_SetString(PyExc_TypeError, "data_ptr method of Pointer object must return 64-bit int");
      ptr_info.valid = false;
      return ptr_info;
    }
    ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(ret);
    if(!ptr_info.dev_ptr)
      return ptr_info;
    uint64_t dev_ptr;
    int status = cuPointerGetAttribute(&dev_ptr, CU_POINTER_ATTRIBUTE_DEVICE_POINTER, ptr_info.dev_ptr);
    if (status == CUDA_ERROR_INVALID_VALUE) {
        PyErr_Format(PyExc_ValueError,
                     "Pointer argument (at %d) cannot be accessed from Triton (cpu tensor?)", idx);
        ptr_info.valid = false;
    } else if (status != CUDA_SUCCESS) {
        CUDA_CHECK(status);  // Catch any other cuda API errors
        ptr_info.valid = false;
    }
    ptr_info.dev_ptr = dev_ptr;
    Py_DECREF(ret);  // Thanks ChatGPT!
    return ptr_info;
  }
  PyErr_SetString(PyExc_TypeError, "Pointer argument must be either uint64 or have data_ptr method");
  ptr_info.valid = false;
  return ptr_info;
}

static inline CUtensorMap* getTmaDesc(PyObject *obj) {
  if (sizeof(CUtensorMap*) != 8) {
    PyErr_SetString(PyExc_SystemError, "getTmaDesc() requires 64-bit compilation");
    return NULL;
  }

  PyObject *method_handle = PyObject_GetAttrString(obj, "tma_desc_cpu_ptr");
  if (!method_handle) {
    PyErr_SetString(PyExc_TypeError, "tma_desc_cpu_ptr() method does not exist");
    return NULL;
  }

  PyObject *empty_tuple = PyTuple_New(0);
  if (!empty_tuple) {
    Py_DECREF(method_handle);
    PyErr_SetString(PyExc_SystemError, "Internal Python error!");
    return NULL;
  }
  PyObject *method_ret = PyObject_Call(method_handle, empty_tuple, NULL);
  Py_DECREF(empty_tuple);
  Py_DECREF(method_handle);
  if (!method_ret) {
    PyErr_SetString(PyExc_SystemError, "Internal Python error!");
    return NULL;
  }

  if (!PyLong_Check(method_ret)) {
    PyErr_SetString(PyExc_TypeError, "tma_desc_cpu_ptr() must return 64-bit int");
    Py_DECREF(method_ret);
    return NULL;
  }

  uint64_t ptr_as_uint = PyLong_AsUnsignedLongLong(method_ret);
  Py_DECREF(method_ret);
  if (!ptr_as_uint) {
    PyErr_SetString(PyExc_ValueError, "received NULL ptr from tma_desc_cpu_ptr()");
    return NULL;
  }
  if (ptr_as_uint % 64 != 0) {
    PyErr_SetString(PyExc_ValueError, "tma_desc_cpu_ptr() must be 64-byte aligned");
    return NULL;
  }

  return (CUtensorMap*)(ptr_as_uint);
}

static void ensureCudaContext() {
  CUcontext pctx;
  CUDA_CHECK(cuCtxGetCurrent(&pctx));
  if (!pctx) {
    // Ensure device context.
    CUdevice device;
    CUDA_CHECK(cuDeviceGet(&device, 0));
    CUDA_CHECK(cuDevicePrimaryCtxRetain(&pctx, device));
    CUDA_CHECK(cuCtxSetCurrent(pctx));
  }
}

static uint16_t pack_fp16(double f) {
    uint16_t result;
    // from https://github.com/python/pythoncapi-compat
#if 0x030600B1 <= PY_VERSION_HEX && PY_VERSION_HEX <= 0x030B00A1 && !defined(PYPY_VERSION)
    _PyFloat_Pack2(f, (unsigned char*)&result, 1);
#else
    PyFloat_Pack2(f, (unsigned char*)&result, 1);
#endif
    return result;
}

static uint16_t pack_bf16(double f) {
    float f32 = (float)f;
    uint32_t u32 = *(uint32_t*)&f32;
    return (uint16_t)(u32 >> 16);
}

static uint32_t pack_fp32(double f) {
    float f32 = (float)f;
    return *(uint32_t*)&f32;
}

static uint64_t pack_fp64(double f) {
    return *(uint64_t*)&f;
}

static PyObject* launch(PyObject* self, PyObject* args) {
  // ensure cuda context is valid before calling any CUDA APIs, e.g. before getPointer calls cuPointerGetAttributes
  ensureCudaContext();

  int gridX, gridY, gridZ;
  uint64_t _stream;
  uint64_t _function;
  int launch_cooperative_grid;
  int launch_pdl;
  PyObject *launch_enter_hook = NULL;
  PyObject *launch_exit_hook = NULL;
  PyObject *kernel_metadata = NULL;
  PyObject *launch_metadata = NULL;
  PyObject *global_scratch_obj = NULL;
  c                 ó8   •— g | ]\  }} ‰|¦  «        › d |› d�‘ŒS )rÁ   ú;r   )r   rª   rv   r    s      €r   r   z!make_launcher.<locals>.<listcomp>  s8   ø€ ÐRÐRÐR±u°q¸"�O�O BÑ'Ô'Ð2Ð2¨aÐ2Ð2Ð2ÐRÐRÐRr   z
  if(!PyArg_ParseTuple(args, "a7  ", &gridX, &gridY, &gridZ,
                                           &_stream, &_function, &launch_cooperative_grid, &launch_pdl, &global_scratch_obj,
                                           &kernel_metadata, &launch_metadata,
                                           &launch_enter_hook, &launch_exit_hookaˆ  )) {
    return NULL;
  }

  int num_warps, num_ctas, shared_memory, clusterDimX, clusterDimY, clusterDimZ;
  if (!PyArg_ParseTuple(kernel_metadata, "iiiiii", &num_warps, &num_ctas, &shared_memory, &clusterDimX, &clusterDimY, &clusterDimZ)) {
    PyErr_SetString(PyExc_TypeError, "kernel_metadata must be a tuple");
    return NULL;
  }

  // extract launch metadata
  if (launch_enter_hook != Py_None){
    PyObject* args = Py_BuildValue("(O)", launch_metadata);
    PyObject* ret = PyObject_CallObject(launch_enter_hook, args);
    Py_DECREF(args);
    if (!ret)
      return NULL;
    Py_DECREF(ret);
  }

  CUdeviceptr global_scratch = 0;
  if (global_scratch_obj != Py_None) {
    DevicePtrInfo global_scratch_info = getPointer(global_scratch_obj, -1);
    if (!global_scratch_info.valid) {
      return NULL;
    }
    global_scratch = global_scratch_info.dev_ptr;
  }

  // raise exception asap
  zâ
  Py_BEGIN_ALLOW_THREADS;
  _launch(gridX, gridY, gridZ, num_warps, num_ctas, launch_cooperative_grid, launch_pdl, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (CUstream)_stream, (CUfunction)_function, global_scratchaC  );
  Py_END_ALLOW_THREADS;
  if (PyErr_Occurred()) {
    return NULL;
  }

  if(launch_exit_hook != Py_None){
    PyObject* args = Py_BuildValue("(O)", launch_metadata);
    PyObject* ret = PyObject_CallObject(launch_exit_hook, args);
    Py_DECREF(args);
    if (!ret)
      return NULL;
    Py_DECREF(ret);
  }

  Py_RETURN_NONE;
}

static PyMethodDef ModuleMethods[] = {
  {"launch", launch, METH_VARARGS, "Entry point for all kernels with this signature"},
  {NULL, NULL, 0, NULL} // sentinel
};

static struct PyModuleDef ModuleDef = {
  PyModuleDef_HEAD_INIT,
  "__triton_launcher",
  NULL, //documentation
  -1, //size
  ModuleMethods
};

PyMODINIT_FUNC PyInit___triton_launcher(void) {
  PyObject *m = PyModule_Create(&ModuleDef);
  if(m == NULL) {
    return NULL;
  }
  PyModule_AddFunctions(m, ModuleMethods);
  return m;
}
)
ÚvaluesÚ	enumerater    Ú_BASE_ARGS_FORMATr‰   ÚitemsrÂ   r‡   rw   rˆ   )Ú	constantsrŠ   r“   r”   Úexpand_signatureÚargs_formatÚformatÚflat_signaturer�   Ú	args_listÚarg_decl_listrª   rv   Ú	arg_declsÚinternal_args_listÚparamsÚnewlineÚ	ptr_declsÚ	tma_declsÚfloat_storage_declsrE   r    r˜   r±   s     `                  @@@r   Úmake_launcherrÙ   |   s™  øøøø€ ð$ð $ð $ð $ð $ðLð ð ð ð ðð ð ð ð ðð ð ð ð ð. )Ð(¨×)9Ò)9Ñ);Ô);Ñ<Ô<ÐØ>Ð>¥)Ð,<Ñ"=Ô"=Ð>Ñ>Ô>€Ià—'’'ÐFÐFÐFÐF°9×3CÒ3CÑ3EÔ3EÐFÑFÔFÑGÔG€KÝ Ñ,€Fà€NØ×ÒÑ!Ô!ð 0ð 0ˆØÐ˜3 Ñ/Ô/Ð/Ð/Ø<Ð<¥)¨NÑ";Ô";Ð<Ñ<Ô<€IÝPSÐT]ÑP^ÔP^ÐabÒPbÐPb��t—y’yÐ LÐ L¸)¿/º/Ñ:KÔ:KÐ LÑ LÔ LÑLÔLÑLÐLÐhj€Ið €MØ—’Ñ"Ô"ð <ð <‰ˆˆ2Ø�ÒÐØØÕ#Ð#Ð#Ø× Ò Õ$6°rÔ$:Ð!CÐ!CÀÐ!CÐ!CÑDÔDÐDÐDà× Ò ¥I¨b¡M¤MÐ!:Ð!:°qÐ!:Ð!:Ñ;Ô;Ð;Ð;Ø—	’	˜-Ñ(Ô(€IØÐØ—’Ñ"Ô"ð 	2ð 	2‰ˆˆ2ØˆaŒ5�CŠ<ˆ<Ø×%Ò%Ð&<°Ð&<Ð&<Ð&<Ñ=Ô=Ð=Ð=ØÕ%Ð%Ð%Ø×%Ò%Ð&8¨QÐ&8Ð&8Ð&8Ñ9Ô9Ð9Ð9Ø�;ÒÐà×%Ò% n° n nÑ5Ô5Ð5Ð5Ø�;ÒÐØ×%Ò% j¨Q j jÑ1Ô1Ð1øÝ•3�y‘>”>Ñ"Ô"€Fð €Gðð à—_’_Ñ&Ô&ðñ ô €Ið
ð Øen×etÒetÑevÔevðñ ô €Iðð à—_’_Ñ&Ô&ðñ ô Ðð
 NÐM i§o¢oÑ&7Ô&7ÐMÑMÔM€FØ
‡M‚MÐ#Ñ$Ô$Ð$ðMõ\ [^ð  _hñ  [iô  [ið  lmò  [mð  [mð  GKð  NWñ  GWð  GWð  suð]Mð Mð^ —y’y Ñ(Ô(ð_Mð Mð@ ‡<‚<ÐRÐRÐRÐRÀ	ÇÂÑ@QÔ@QÐRÑRÔRÑSÔSðAMð MðB !'ðCMð MðH R[ðIMð MðD	 ‡<‚<�	ÑÔðE	Mð MðF	 ‡<‚<�	ÑÔðG	Mð MðH	 ‡<‚<Ð#Ñ$Ô$ðI	Mð MõL	 qtð  uGñ  qHô  qHð  KLò  qLð  qLð  IMð  PT÷  PYò  PYð  Zlñ  Pmô  Pmñ  Imð  Imð  RTðM	Mð Mð M€Cð\
 €Jr   c                   ó   — e Zd ZdZd„ Zd„ ZdS )ÚTmaDescKernelParamé€   c                 ó\   — dd l }|                     | j        |j        d¬¦  «        | _        d S )Nr   Úcpu©r�   Údevice)ÚtorchÚemptyÚTMA_DESC_SIZEÚuint8Údesc©rO   rá   s     r   rQ   zTmaDescKernelParam.__init__V  s-   € ØˆˆˆØ—K’K Ô 2¸%¼+Èe�KÑTÔTˆŒ	ˆ	ˆ	r   c                 ó4   — | j                              ¦   «         S r6   )rå   Údata_ptr)rO   s    r   Útma_desc_cpu_ptrz#TmaDescKernelParam.tma_desc_cpu_ptr[  s   € ØŒy×!Ò!Ñ#Ô#Ð#r   N)rR   rS   rT   rã   rQ   ré   r   r   r   rÛ   rÛ   S  s:   € € € € € Ø€MðUð Uð Uð
$ð $ð $ð $ð $r   rÛ   c              #   ó   K  — | ]}||fV — Œ	d S r6   r   )r   rª   s     r   r$   r$   `  s&   è è € Ð:Ð:¨1  A Ð:Ð:Ð:Ð:Ð:Ð:r   é   é
   é   é	   c           
      ó  — |€$| j         g| j        ¢| j        ¢| j        ¢| j        ¢S |d         }|d         }|d         }|d         }|d         }| j                              ¦   «         }| j        }| j        }	|	d         dk    sJ ‚t	          ¦   «         }
|
g|¢|	¢}|rt          |¦  «        }|dxx         dz  cc<   t          j        j        j	        j
                             |
                     ¦   «         |||t          |         |||	¦  «         |S )	NÚswizzleÚ	elem_sizeÚ	elem_typeÚ
block_sizeÚ
fp4_paddedr   r   r€   )Úbaser�   Ústridesrè   rÛ   ÚlistÚtritonÚruntimeÚdriverÚactiveÚutilsrN   ré   ÚTMA_DTYPE_DEVICE_TO_HOST)ÚargÚmetadatarð   rñ   rò   ró   rô   rè   r�   rö   rå   Úresults               r   Úmake_tensordesc_argr  f  s1  € ØÐð ”ÐM˜3œ9ÐM s¤{ÐM°S´YÐMÀÄÐMÐMà�yÔ!€GØ˜Ô%€IØ˜Ô%€IØ˜,Ô'€JØ˜,Ô'€JàŒx× Ò Ñ"Ô"€HØŒI€EØŒk€GØ�2Œ;˜!ÒÐÐÐåÑÔ€DØÐ%�UÐ%˜WÐ%€Fàð Ý�U‘”ˆØˆbˆ	ˆ	Œ	�Q‰ˆ	ˆ	‰	Ý
„NÔÔ Ô&×:Ò:Ø×ÒÑÔØØØÝ  Ô+ØØØñ	ô 	ð 	ð €Mr   c                 ó6   ‡ ‡‡‡— ddl mŠ ddlmŠ ˆˆˆ ˆfd„}|S )Nr   )ÚTensorDescriptorc                  ó   •— | d t          t          ¦  «        …         }| t          t          ¦  «        d …         }d}g }t          |¦  «        D ]a\  }}t          |‰	‰f¦  «        r5‰r‰|         nd }|dz  }|                     t          ||¦  «        ¦  «         ŒL|                     |¦  «         Œb‰r|t          ‰¦  «        k    sJ ‚ ‰
g |¢|¢R Ž S )Nr   r   )r‰   rÉ   rÈ   r‚   Úextendr  r‡   )ÚargsÚ	meta_argsÚraw_kernel_argsrŒ   Ú
final_argsrª   rþ   rŽ   ÚGluonTensorDescriptorr  Úlauncherr“   s           €€€€r   Úinnerz%wrap_handle_tensordesc.<locals>.inner’  s  ø€ ØÐ0�#Õ/Ñ0Ô0Ð0Ô1ˆ	Ø�sÕ#4Ñ5Ô5Ð6Ð6Ô7ˆØˆØˆ
Ý Ñ0Ô0ð 	'ð 	'‰FˆAˆsÝ˜#Ð 0Ð2GÐHÑIÔIð 'Ø:IÐS� ~Ô6Ð6Èt�Ø !Ñ#�Ø×!Ò!Õ"5°c¸4Ñ"@Ô"@ÑAÔAÐAÐAà×!Ò! #Ñ&Ô&Ð&Ð&Ø"ÐL n½¸OÑ8LÔ8LÒ&LÐ&LÐ&LÐ&LØˆxÐ0˜Ð0 ZÐ0Ð0Ð0Ð0r   )Útriton.tools.tensor_descriptorr  Ú'triton.experimental.gluon.nvidia.hopper)r  r“   r  r
  r  s   `` @@r   Úwrap_handle_tensordescr  Ž  sY   øøøø€ Ø?Ð?Ð?Ð?Ð?Ð?ØaÐaÐaÐaÐaÐað1ð 1ð 1ð 1ð 1ð 1ð 1ð 1ð €Lr   c                   ó   — e Zd Zd„ Zd„ ZdS )ÚCudaLauncherc                 óÌ  ‡‡— t          ‰d¦  «        r‰j        nt          ¦   «         }ˆfd„Šˆfd„|                     ¦   «         D ¦   «         }d„ ‰j                             ¦   «         D ¦   «         }t          |dd ¦  «        }t          |||¦  «        Št          ‰dt          ¦   «         t          t          ¬¦  «        }t          d„ |                     ¦   «         D ¦   «         ¦  «        }t          j        t          j        |j        d	¦  «        | _        |rt'          |j        |¦  «        n|j        | _        |j        | _        |j        | _        |j        | _        |j        | _        d S )
NrË   c                 ór   •— t          | t          ¦  «        r ‰j        j                             | ¦  «        fn| S r6   )r‚   r,   ÚfnÚ	arg_namesÚindex)r—   rE   s    €r   ú<lambda>z'CudaLauncher.__init__.<locals>.<lambda>¨  s2   ø€ ½ZÈÍ3Ñ=OÔ=OÐV˜SœVÔ-×3Ò3°AÑ6Ô6Ð9Ð9ÐUV€ r   c                 ó.   •— i | ]\  }} ‰|¦  «        |“ŒS r   r   )r   ÚidxÚvalueÚarg_idxs      €r   rµ   z)CudaLauncher.__init__.<locals>.<dictcomp>©  s'   ø€ ÐMÐMÐM©Z¨S°%�W�W˜S‘\”\ 5ÐMÐMÐMr   c                 ó   — i | ]\  }}||“Œ	S r   r   )r   r  r  s      r   rµ   z)CudaLauncher.__init__.<locals>.<dictcomp>ª  s   € ÐHÐHÐH¡J C¨�S˜%ÐHÐHÐHr   r“   Ú__triton_launcherrD   c              3   ój   K  — | ].}t          |t          ¦  «        o|                     d ¦  «        V — Œ/dS )rZ   N)r‚   r,   ru   )r   r�   s     r   r$   z(CudaLauncher.__init__.<locals>.<genexpr>´  s>   è è € Ð!vÐ!vÐ\_¥*¨Sµ#Ñ"6Ô"6Ð"W¸3¿>º>È,Ñ;WÔ;WÐ!vÐ!vÐ!vÐ!vÐ!vÐ!vr   r   )r=   rË   ÚdictrÊ   rŠ   ÚgetattrrÙ   r   r8   rG   rH   r-   rÇ   Ú	functoolsÚreduceÚoperatorÚmulÚcluster_dimsÚnum_ctasr  ÚlaunchÚglobal_scratch_sizeÚglobal_scratch_alignÚlaunch_cooperative_gridÚ
launch_pdl)	rO   rE   rÿ   rË   rŠ   r“   rP   Úhas_tensor_desc_argr  s	    `      @r   rQ   zCudaLauncher.__init__¦  sV  øø€ Ý%,¨S°+Ñ%>Ô%>ÐJ�C”M�MÅDÁFÄFˆ	ØVÐVÐVÐVˆØMÐMÐMÐM¸9¿?º?Ñ;LÔ;LÐMÑMÔMˆ	ØHÐH°#´-×2EÒ2EÑ2GÔ2GÐHÑHÔHˆ	Ý! (Ð,=¸tÑDÔDˆÝ˜I y°/ÑBÔBˆÝ%ØØ$Ý%™œÝ%Ýð
ñ 
ô 
ˆõ "Ð!vÐ!vÐcl×csÒcsÑcuÔcuÐ!vÑ!vÔ!vÑvÔvÐå!Ô(­¬°xÔ7LÈaÑPÔPˆŒØM`ÐpÕ,¨S¬Z¸ÑIÔIÐIÐfiÔfpˆŒØ#+Ô#?ˆÔ Ø$,Ô$AˆÔ!Ø'/Ô'GˆÔ$Ø"Ô-ˆŒˆˆr   c           
      óÊ   — | j         dk    r6||z  |z  }|| j        z  | j         z  }t          j        || j        |¦  «        }	nd }	 | j        |||||| j        | j        |	g|¢R Ž  d S ©Nr   )r(  r&  r   Ú
_allocatorr)  r'  r*  r+  )
rO   ÚgridXÚgridYÚgridZÚstreamÚfunctionr  Ú	grid_sizeÚ
alloc_sizeÚglobal_scratchs
             r   Ú__call__zCudaLauncher.__call__½  s–   € ØÔ# aÒ'Ð'Ø ™¨Ñ-ˆIØ" T¤]Ñ2°TÔ5MÑMˆJÝ(Ô3°JÀÔ@YÐ[aÑbÔbˆNˆNà!ˆNØˆŒ�E˜5 %¨°¸4Ô;WÐY]ÔYhØ"ð	+Ø%)ð	+ð 	+ð 	+ð 	+ð 	+ð 	+r   N)rR   rS   rT   rQ   r8  r   r   r   r  r  ¤  s2   € € € € € ð.ð .ð .ð.+ð +ð +ð +ð +r   r  c                   óX   ‡ — e Zd Zˆ fd„Zd„ Zd„ Zd„ Zed„ ¦   «         Zd„ Z	d„ Z
d„ Zˆ xZS )	Ú
CudaDriverc                 ó†   •— t          ¦   «         | _        t          | _        t	          ¦   «                              ¦   «          d S r6   )r:   rü   r  Úlauncher_clsr>   rQ   )rO   rA   s    €r   rQ   zCudaDriver.__init__Ê  s2   ø€ Ý‘[”[ˆŒ
Ý(ˆÔÝ‰Œ×ÒÑÔÐÐÐr   c                 ó¢   — |                       ¦   «         }|                      |¦  «        }|d         dz  |d         z   }d}t          d||¦  «        S )Nr   rì   r   é    r   )Úget_current_deviceÚget_device_capabilityr   )rO   rà   Ú
capabilityÚ	warp_sizes       r   Úget_current_targetzCudaDriver.get_current_targetÏ  sT   € Ø×(Ò(Ñ*Ô*ˆØ×/Ò/°Ñ7Ô7ˆ
Ø ”] RÑ'¨*°Q¬-Ñ7ˆ
Øˆ	Ý˜ ¨YÑ7Ô7Ð7r   c                 óZ   — dd l }|                     d|                      ¦   «         ¦  «        S )Nr   r   )rá   rà   r?  ræ   s     r   Úget_active_torch_devicez"CudaDriver.get_active_torch_deviceÖ  s+   € ØˆˆˆØ�|Š|˜F D×$;Ò$;Ñ$=Ô$=Ñ>Ô>Ð>r   c                 ó   — dd l }|j        S r.  )rá   r   ræ   s     r   Úget_device_interfacezCudaDriver.get_device_interfaceÚ  s   € ØˆˆˆØŒzÐr   c                  ó|   — 	 dd l } | j                             ¦   «         o| j        j        d u S # t
          $ r Y dS w xY w)Nr   F)rá   r   Úis_availableÚversionÚhipÚImportError)rá   s    r   Ú	is_activezCudaDriver.is_activeÞ  sU   € ð	ØˆLˆLˆLØ”:×*Ò*Ñ,Ô,ÐL°%´-Ô2CÀtÐ2KÐLøÝð 	ð 	ð 	Ø�5�5ð	øøøs   ‚*- ­
;º;c                 ó   — ddl m} |S )Nr   )Údo_bench)Útriton.testingrO  )rO   rO  s     r   Úget_benchmarkerzCudaDriver.get_benchmarkeræ  s   € Ø+Ð+Ð+Ð+Ð+Ð+Øˆr   c                 óh   — dd l }d}|                     t          |dz  ¦  «        |j        d¬¦  «        S )Nr   i   é   r   rß   )rá   râ   Úint)rO   rá   Ú
cache_sizes      r   Úget_empty_cache_for_benchmarkz(CudaDriver.get_empty_cache_for_benchmarkê  s8   € Øˆˆˆð
 'ˆ
Ø�{Š{�3˜z¨Q™Ñ/Ô/°u´yÈˆ{ÑPÔPÐPr   c                 ó.   — |                      ¦   «          d S r6   )Úzero_)rO   Úcaches     r   Úclear_cachezCudaDriver.clear_cacheó  s   € Ø�Š‰Œˆˆˆr   )rR   rS   rT   rQ   rC  rE  rG  ÚstaticmethodrM  rQ  rV  rZ  rU   rV   s   @r   r:  r:  È  s°   ø€ € € € € ðð ð ð ð ð
8ð 8ð 8ð?ð ?ð ?ðð ð ð ðð ñ „\ððð ð ðQð Qð Qðð ð ð ð ð ð r   r:  )+r!  r#  r   r'   rø   rƒ   Úpathlibr   r   Útriton.runtime.buildr   Útriton.runtimer   Útriton.backends.compilerr   Útriton.backends.driverr   r   r   ÚrealpathÚ__file__r    rG   r7   rH   Ú	lru_cacher4   r8   Úobjectr:   rw   rÂ   rÃ   rÉ   rÙ   rÛ   r  rˆ   rý   r  r  r  r:  r   r   r   ú<module>re     sÒ  ðØ Ð Ð Ð Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø 8Ð 8Ð 8Ð 8Ð 8Ð 8Ø &Ð &Ð &Ð &Ð &Ð &Ø .Ð .Ð .Ð .Ð .Ð .Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,à
Œ'�/Š/˜"œ'×*Ò*¨8Ñ4Ô4Ñ
5Ô
5€Ø”—’˜W iÑ0Ô0Ð1€Ø”—’˜W eÑ,Ô,€ØˆH€	ð €ÔÑÔðð ñ Ôðð. €ÔÑÔð,ð ,ñ Ôð,ð;ð ;ð ;ð ;ð ;�ñ ;ô ;ð ;ð6
ð 
ð 
ð4 ØØØØðð Ð ð ØØØØðð Ð ð #Ð ðTð Tð Tðn	$ð 	$ð 	$ð 	$ð 	$ñ 	$ô 	$ð 	$ð  ˜4Ð:Ð:°°°b±	´	Ð:Ñ:Ô:Ñ:Ô:Ð Ø Ð ˜Ñ ØÐ ˜Ñ Ø Ð ˜Ñ ð%ð %ð %ðPð ð ð,!+ð !+ð !+ð !+ð !+�6ñ !+ô !+ð !+ðH,ð ,ð ,ð ,ð ,�ñ ,ô ,ð ,ð ,ð ,r   