o YÛ·i›ã@sTdZddlZddlmZmZddlmZmZGdd„deƒZ e dkr(e  ¡dSdS)z Test basic language features éN)ÚcudaÚfloat64)ÚunittestÚ CUDATestCasec@s$eZdZdd„Zdd„Zdd„ZdS)ÚTestLangcsLd‰t d¡‡fdd„ƒ}t tˆƒ¡}|d|ƒ| t |ˆk¡¡dS)N)çð?g@g@úvoid(float64[:])cstˆƒD]\}}|||<qdS)N)Ú enumerate)ÚaÚiÚv©Útup©úW/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_lang.pyÚfoos ÿz$TestLang.test_enumerate..foo©ér)rÚjitÚnpÚzerosÚlenÚ assertTrueÚall)Úselfrr rr rÚtest_enumerate s  zTestLang.test_enumeratecsjd‰d‰t d¡‡‡fdd„ƒ}t d¡}|d|ƒt ˆ¡}t ˆ¡}| t ||| ¡k¡¡dS)N)réé)g@gffffff@gÍÌÌÌÌÌ@rcs0d}tˆˆƒD] \}}|||7}q||d<dS)Nr)Úzip)r Úcr Új©Út1Út2rrrs zTestLang.test_zip..foorr)rrrrÚarrayrrÚsum)rrr Úbrrr!rÚtest_zips     zTestLang.test_zipcCs,t d¡dd„ƒ}t d¡}|d|ƒdS)a? Ensure that typing and lowering of CUDA kernel API primitives works in more than one block. Was originally to ensure that macro expansion works for more than one block (issue #872), but macro expansion has been replaced by a "proper" implementation of all kernel API functions. zvoid(float64[:,:])cSs^tdƒD]}tjj}qtdƒD]}tjj}qtj dt¡}d|||f<|||f|||f<dS)Nrr©rrr)ÚrangerÚ threadIdxÚxÚyÚsharedr$r)Úaryr Útxr ÚtyÚsmrrrÚ"cuda_kernel_api_in_multiple_blocks1s     zCTestLang.test_issue_872..cuda_kernel_api_in_multiple_blocksr()rr(N)rrrr)rr2r rrrÚtest_issue_872)s  zTestLang.test_issue_872N)Ú__name__Ú __module__Ú __qualname__rr'r3rrrrr s rÚ__main__) Ú__doc__ÚnumpyrÚnumbarrÚnumba.cuda.testingrrrr4ÚmainrrrrÚs4 ÿ