ia) SSKJr SSKrSSKrSSKrSSKrSSKJrJ r SSK J r J r J r SSKJr SSKJrJr Sq\S 5rS r"S S 5r"S S\5r"SS\5r"SS\R65r"SS\5rg))contextmanagerN) FakeCUDAArrayFakeWithinKernelCUDAArray)Dim3FakeCUDAModuleswapped_cuda_module)normalize_kernel_dimensions)wrap_argArgHintc#H# [bS5eUqSv Sqg!Sqf=f7f)z" Push the current kernel context. Nz)concurrent simulated kernel not supported_kernel_context)mods V/var/www/html/trading/venv/lib/python3.13/site-packages/numba/cuda/simulator/kernel.py_push_kernel_contextrs0  "O$OO "O $s"""c[$)zL Get the current kernel context. This is usually done by a device function. rr_get_kernel_contextr$s  rc\rSrSrSrSrSrg) FakeOverload+z= Used only to provide the max_cooperative_grid_blocks method cg)Nrr)selfblockdims rmax_cooperative_grid_blocks(FakeOverload.max_cooperative_grid_blocks/srrN)__name__ __module__ __qualname____firstlineno____doc__r__static_attributes__rrrrr+s rrc\rSrSrSrSrg)FakeOverloadDict5c[5$N)r)rkeys r __getitem__FakeOverloadDict.__getitem__6s ~rrN)r r!r"r#r,r%rrrr'r'5srr'cj\rSrSrSrS/S4SjrSrSrSrSr SS jr \ S 5r \ S 5r S rg )FakeCUDAKernel<z Wraps a @cuda.jit-ed function. FcXlX lX0lXPl[ U5UlSUlSUlSUlSUl [R"X5 g)Nr) fn_device _fastmath_debuglist extensionsgrid_dim block_dimstreamdynshared_size functoolsupdate_wrapper)rr2devicefastmathr7debugs r__init__FakeCUDAKernel.__init__AsO ! z*    *rc ^^ TR(a7[TR[55 TR"U6sSSS5 $[ TR TR 5up#[X#TR5n[U5 /m U U4SjnUVs/sH oe"U5PM nn[TRU5 [R"U6H7n[TRX#TR5n U R"U/UQ76 M9 SSS5 T H n U "5 M SSS5 g!,(df  N=fs snf!,(df  N==f!,(df  g=f)Nc>[R"U4SjTRSU45up[U[R 5(a+UR S:a[U5RT5nOT[U[5(aURT5nO-[U[R5(a [U5nOUn[U[5(a [U5$U$)Nc*>UR"USTS.6$)Nr)r:retr) prepare_args)ty_val extensionrFs r;FakeCUDAKernel.__call__..fake_arg..bsi.D.D !/#rr) r<reducer7 isinstancenpndarrayndimr to_devicer voidrr)arg_retrFrs rfake_arg)FakeCUDAKernel.__call__..fake_arg_s"))#OO3K c2::..388a<"3-11$7CW-----CRWW--',CCc=114S99 r)r3r r2rr r8r9rr;rrNndindex BlockManagerr5run) rargsr8r9fake_cuda_modulerVrS fake_args grid_pointbmwbrFs ` @r__call__FakeCUDAKernel.__call__Os <<$TWW.A.CDww~ED:$--:>..J*(*.*=*=? !"2 3D .377$3#$I7$TWW.>?"$**h"7J%dggxDKKPBFF:2 2#8@ G4 3EDJ8??;4 3sBD5EE-EAE E5 EE E E E*ch[USS6uUlUl[U5S:Xa USUlU$)Nr )r r8r9lenr;)r configurations rr,FakeCUDAKernel.__getitem__s? 'r): ; & t~ }  ""/"2D  rcgr*rrs rbindFakeCUDAKernel.binds rcU$r*r)rr[s r specializeFakeCUDAKernel.specializes rc8US:a[SU-5eXSX44$)Nrz0Can't create ForAll with negative task count: %sr) ValueError)rntaskstpbr: sharedmems rforallFakeCUDAKernel.foralls1 A:O%&' 'Av011rc[5$r*)r'rjs r overloadsFakeCUDAKernel.overloadss !!rcUR$r*)r2rjs rpy_funcFakeCUDAKernel.py_funcs wwr) r5r3r4r9r;r7r2r8r:N)rrr)r r!r"r#r$rArar,rkrnrupropertyrxr{r%rrrr/r/<sW-2b +/b 2 ""rr/cV^\rSrSrSrU4SjrU4SjrSrSrSr Sr S r S r U=r $) BlockThreadz? Manages the execution of a function for a single CUDA thread. c>^U(a U4SjnUnOTn[[U] US9 [R"5UlSUlX l[U6Ul [U6Ul SUl SUl SUl XPl[URR6nURR UR URR"UR"URR$----Ulg)Nc>>[R"SS9 T"U0UD6 g)Nraise)divide)rNseterr)r[kwargsfs r debug_wrapper+BlockThread.__init__..debug_wrappers )4"6"r)targetFT)superrrA threadingEventsyncthreads_eventsyncthreads_blocked_managerrblockIdx threadIdx exceptiondaemonabortr@ _block_dimxyz thread_id) rrmanagerrrr@rrblockDim __class__s ` rrABlockThread.__init__s  ##FF k4))8!*!2#(  h y)   112))XZZ4>>;K;K;C::;?>>;K;K[[U] 5 g![anS[ UR 5-nS[ UR 5-n[U5S:Xa U<SU<3nO U<SU<SU<3n[R"5Sn[U5"U5U4Ul SnAgSnAff=f)Nztid=%szctaid=%s z: r ) rrrZ Exceptionr6rrstrsysexc_infotyper)retidctaidmsgtbrs rrZBlockThread.runs 0 +t ( * 0T$..11Cdmm!44E1v|!$e,%(%3"B#1gclB/DNN 0s B0BB++B0cUR(a [S5eSUlURR 5 URR 5 UR(a [S5eg)Nz"abort flag set on syncthreads callTz#abort flag set on syncthreads clear)r RuntimeErrorrrwaitclearrjs r syncthreadsBlockThread.syncthreadssY ::CD D#'  ##% $$& ::DE E rcLURRURRURR4nXRR U'UR 5 [R"URR 5nUR 5 U$r*) rrrrr block_staterrN count_nonzero)rvalueidxcounts rsyncthreads_countBlockThread.syncthreads_countsunn 0 0$..2B2BB). !!#&   !:!:;  rc^URRURRURR4nXRR U'UR 5 [R"URR 5nUR 5 U(aS$S$Nrr) rrrrrrrrNallrrrtests rsyncthreads_andBlockThread.syncthreads_and{nn 0 0$..2B2BB). !!#& vvdmm//0 qarc^URRURRURR4nXRR U'UR 5 [R"URR 5nUR 5 U(aS$S$r) rrrrrrrrNanyrs rsyncthreads_orBlockThread.syncthreads_orrrc@SUR<SUR<S3$)Nz Thread <<>>)rrrjs r__str__BlockThread.__str__s(, t~~FFr) rrrrr@rrrrr)r r!r"r#r$rArZrrrrrr% __classcell__)rs@rrrs4N0 0 F  GGrrc$\rSrSrSrSrSrSrg)rYa Manages the execution of a thread block. When run() is called, all threads are started. Each thread executes until it hits syncthreads(), at which point it sets its own syncthreads_blocked to True so that the BlockManager knows it is blocked. It then waits on its syncthreads_event. The BlockManager polls threads to determine if they are blocked in syncthreads(). If it finds a blocked thread, it adds it to the set of blocked threads. When all threads are blocked, it unblocks all the threads. The thread are unblocked by setting their syncthreads_blocked back to False and setting their syncthreads_event. The polling continues until no threads are alive, when execution is complete. cX lX0lXlX@l[R "U[R S9Ulg)N)dtype) _grid_dimr_fr5rNzerosbool_r)rrr8r9r@s rrABlockManager.__init__s-!# 88IRXX>rc^^[5n[5n[5n[R"TR6HTnUU4Sjn[ UTXTR 5nUR 5 URU5 URU5 MV U(GaUHnUR(aURU5 M'UR(dM:UH+n SU l SU lU RR5 M- URSRURS5e XE:Xa4UH$nSUlURR5 M& [5n[UVs/sHoR5(dMUPM sn5nU(aGMUH?nUR(dMURSRURS5e gs snf)Nc$>TR"T6 gr*)r)r[rsrr BlockManager.run..targetsrTFrr)setrNrXrrr5startaddrrrrwith_tracebackis_alive) rr^r[threads livethreadsblockedthreads block_pointrtt_others ` ` rrZBlockManager.run s~%e ::t7K FD*4;;OA GGI KKN OOA  8 (("&&q)[[[$+(, 6;311557$+ ++a.77 AGG!,'A,1A)''++-("%;H;a**,;HIK'k,A{{{kk!n33AKKNCCIs :G3G3)rr5rrrN)r r!r"r#r$rArZr%rrrrYrYs"?(DrrY) contextlibrr<rrnumpyrNcudadrv.devicearrayrr kernelapirrr errorsr r[r r rrrrdictr'objectr/ThreadrrYrrrrs% I@@0$   t_V_HPG)""PGfAD6ADr