K i ddlZddlZddlZddlZddlZddlmZddlm Z ddl m Z m Z m Z ddlZddlmZddlmZmZmZgdZgdZeezZgdZed gzZeezZed gzZd d gZd gezd gzezd gzZeeehdz Z dZ!dZ"dZ#dZ$dZ%dZ&dZ'dZ(dZ)dZ*dZ+dZ,dZ-dZ.dZ/dZ0dZ1d2d e efd!Z2d3d"ejfd#e eejhffd$Z5d"e6d#ejnfd%Z8d#e6fd&Z9d'Z:d4d(Z;d4d)ZrrrcHtjjdddk(S)NTRITON_INTERPRET01)osenvironget^/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/triton/_internal_testing.pyis_interpreterr%s ::>>,c 2c 99r#c|trytjjjj SN)r%tritonruntimedriveractiveget_current_targetr"r#r$r,r,s* >> ' ' : : <!CCr#cbtxr$tjjddk(S)Nr r6r"r#r$ is_blackwellr=-s& 9 D99;A>"DDr#cbtxr$tjjddk\SNr r6r"r#r$is_hopper_or_newerrA1r:r#cbtxr$tjjddk(Sr?r6r"r#r$ is_hopperrC5r:r#c<t}|dS|jdk(S)NFhipr/r1s r$is_hiprF9"  !FN5?%(??r#cbt}|duxr |jdk(xr|jdk(S)NrEgfx90ar,r0archr1s r$ is_hip_cdna2rL>1  !F   U&..E"9 UfkkX>UUr#cbt}|duxr |jdk(xr|jdk(S)NrEgfx942rJr1s r$ is_hip_cdna3rPCrMr#cbt}|duxr |jdk(xr|jdk(S)NrEgfx950rJr1s r$ is_hip_cdna4rSHrMr#c`t}|duxr|jdk(xrd|jvS)NrEgfx11rJr1s r$ is_hip_gfx11rVM1  !F   T&..E"9 Tg>TTr#c`t}|duxr|jdk(xrd|jvS)NrEgfx12rJr1s r$ is_hip_gfx12rZRrWr#cFtxstxs tSr')rLrPrSr"r#r$ is_hip_cdnar\Ws > =\^ =|~=r#ctrdSdS)Nii)rSr"r#r$get_hip_lds_sizer^[s!^6..r#c<t}|dS|jdk(S)NFxpur/r1s r$is_xpura_rGr#cHt}|dSt|jS)N)r,strrKr1s r$get_archreds"  !F25S%55r#rscFt|tr|f}| td}|ttzvrt j tt |}| |jnt||j}| |jnt||j}tt |}|j||||}d||dk(<|S|r)d|vr%|jdd|t j}|S|tvr"|jdd|j|S|d k(rV|jdd|jd jd t j d zjd S|d vr|jdd|dkDSt#d|)zp Override `rs` if you're calling this function twice and don't want the same result for both calls. )seed)dtyperfloat8(rrrl)rint1bool_gzUnknown dtype ) isinstanceintr int_dtypes uint_dtypesnpiinfogetattrminmaxrandintr float_dtypesnormalastypeviewr RuntimeError)shape dtype_strrflowhighrvrjxs r$ numpy_randomris| %  z b !J,,Y/0;eiiCUYY,? Luyyc$ .BI& JJsD%uJ 5!q&  x9, JJr2uBGGJ 4 l "yyAu%,,Y77 j  !Q&--i8==hG"))T^J__eefopp / /yyAu%++^I;788r#rreturnc|jj}|tvr_|jd}|j t t |}ttj||t t|S|r3d|vr/ttj||t t|S|dk(r*|dk(r%tj||jStj||S)z Note: We need dst_type because the type of x can be different from dst_type. For example: x is of type `float32`, dst_type is `bfloat16`. If dst_type is None, we infer dst_type from x. u)devicerlrr) rjnamertlstripr}rwrur r7tensortlr)rrdst_typetsigned_type_namex_signeds r$ to_tritonrs  AK88C=88GB(89:5<<@'"a.QQ H,u||Af=wr8?TU U >h*4<<&1::< <||Af--r#c<tjt|dSr')r str_to_tyr rs r$str_to_triton_dtypers <<215t <.fresh_functions&__. 2MD' E4!5!5!7 800779 288rzz)&&txx&? ''1  2 2# r#cjD]\}}t||jD]&}|tjvstj|=(_yr')rrundorr r)rrkrrrrprev_propagate_envs r$reset_functionz)_fresh_knobs_impl..reset_functionsf&__. *MD' E4 ) *  "ABJJJJqM "1r#) r(rsetpytest MonkeyPatch__dict__rrq base_knobsr) rrrrrrrrrrs @@@@@r$_fresh_knobs_implrsu $$&K#^^113 D' gu// 0W@P@P5PUYamUm g IL,,  1 1 > ))Gs 6B&)NNNr')F)Frrrrur7r(triton.languagerrrtypingrrrr numpy.randomrtriton.runtime.jitr r r rsrtintegral_dtypesr{float_dtypes_with_bfloat16dtypesdtypes_with_bfloat16torch_float8_dtypes torch_dtypessortedr tma_dtypesr%r,r3r9r=rArCrFrLrPrSrVrZr\r^rarerndarrayrrrdrjrrrrrmarkskipif requires_tmarrrr)rrrr"r#r$rs  '' $UU 0 5 {*0 )ZL8 < ' ,& 6x*$y0<?:,N C,-0NN O := A DEDD@ V V V U U >/@ 6 9x '<9<..u]ELL=X7Y.&=3=288=GsGA a^{{!!ln"4\^!L >3>s> U5<<););)I)IIJ u|| +*HSX$6+*r#