o eߣ@sddlmZddlZddlZddlZddlZddlm Z ddlm Z ddl m Z m Z mZmZmZmZddlmZmZmZmZejrydd lmZdd lmZdd l mZmZdd lmZdd lm Z ddl!m"Z"ddl#m$Z$ddl m%Z%ddl&m'Z'Gdddej(Z)GdddeZ*dS)) annotationsN)coredata)mlog)EnvironmentException Popen_safe is_windowsLibType OptionKeyversion_compare)Compilercuda_buildtype_argscuda_optimization_argscuda_debug_args)CompileCheckMode) BuildTarget)MutableKeyedOptionDictTypeKeyedOptionDictType) Dependency) Environment) MachineInfo) DynamicLinker) MachineChoice)ExternalProgramc@seZdZdZdZdS)_PhasecompilerlinkerN)__name__ __module__ __qualname__COMPILERLINKERr#r#>/usr/lib/python3.10/site-packages/mesonbuild/compilers/cuda.pyr*srcseZdZdZdZhdZidddddd d d d d ddddddddddddddddddd d!d"d#d$d%id&d'd(d)d*d+d,d-d.d/d0d1d2d3d4d5d6d7d8d9d:d;dd?d@dAdBdCdDdEdFdGidHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd[d\d]d^d_d`dadbdcdddedfdgdhdidjdkdlZdmdneDZdoZ p pddfdd Z e ddddZ e dddZ ejfdddZd ddZd!ddZd"ddZdpdpdd#ddZdZdZdZd$fdd Zd%ddZd&ddZd&ddZd'ddZd(ddZd(ddZd)ddĄZd*ddDŽZ d*ddɄZ!d+dd̄Z"d(dd΄Z#d,ddфZ$d-ddԄZ%d.dd؄Z&d/d0dd݄Z'd1ddZ(d2ddZ)d-ddZ*d3ddZ+d4ddZ,d(ddZ-d5ddZ.d6ddZ/d(ddZ0e1j2dfd7ddZ3d8ddZ4d8ddZ5d9fd d Z6d:fd d Z7d:fdd Z8d&ddZ9d(ddZ:d(ddZ;d;ddZ<Z=S(< CudaCompilerz -Xlinker=Zcuda>wz-Wno-deprecated-gpu-targets-cz--disable-warningsz-dwz--resource-usagez--objdir-as-tempdirz --device-cz--dryrunz-nodlink-wz-cleanz-Wext-lambda-captures-thisz-keepz--compile-as-tools-patchz--ptxz--extra-device-vectorizationz-dltoz-ewpz --device-wz --device-link--runz--extended-lambdaz--source-in-ptxz-Wno-deprecated-declarationsz-dryrunz-Ez-forward-unknown-to-host-linkerz --Wreorderz-noehz-arch-lsz!-forward-unknown-to-host-compilerz$--generate-dependencies-with-compilez --verbosez -astoolspatchz--display-error-numberz-dcz-ptxz-expt-extended-lambdaz--device-debugz--Wno-deprecated-declarationsz-MDz-keep-device-functionsz -lineinfoz -nohdinitlistz--m64z-MMz--keepz-cudaz-MMDz-Vz--fatbinz--Wno-deprecated-gpu-targetsz--generate-line-infoz.--generate-nonsystem-dependencies-with-compilez--extensible-whole-programz-m64z -restrictz--allow-unsupported-compilerz-expt-relaxed-constexprz-code-lsz--cubinz --compilez --versionz --preprocessz-Mz--m32z--keep-device-functionsz--no-device-linkz --profilez"--forward-unknown-to-host-compilerz--use_fast_mathz-vz-m32z -save-tempsz--helpz -Wreorderz-cubinz-allow-unsupported-compilerz--linkz--libz--sharedz-objtempz--Wdefault-stream-launchz--Wext-lambda-captures-thisz--clean-targetsz-MPz-sharedz-hz-libz-dlink-use_fast_mathz-extra-device-vectorizationz--debugz -src-in-ptxz!--no-host-device-initializer-listz-fatbinz--generate-dependenciesz-err-noz--generate-dependency-targetsz-extended-lambdaz-Wdefault-stream-launchz -res-usagez-Gz--dlink-time-optz --restrictz--expt-extended-lambdaz--no-exceptionsz--dont-use-profilez-noprofz--cudaz-pgz--list-gpu-archz-gz --save-tempsz--no-align-doublez --forward-unknown-to-host-linkerz--list-gpu-codez!--generate-nonsystem-dependenciesz--expt-relaxed-constexprz --output-file-oz --pre-include-includez --library-lz--define-macro-Dz--undefine-macro-Uz--include-path-Iz--system-include-isystemz--library-path-Lz--output-directoryz-odirz--dependency-outputz-MFz--compiler-bindirz-ccbinz--archiver-binaryz-arbinz--cudart-cudartz --cudadevrtz -cudadevrtz--libdevice-directoryz-ldirz--target-directoryz -target-dirz--optimization-infoz -opt-infoz --optimize-Oz--ftemplate-backtrace-limitz-ftemplate-backtrace-limitz--ftemplate-depthz-ftemplate-depthz--xz-xz--std-stdz --machine-mz--compiler-options -Xcompilerz--linker-optionsz-Xlinkerz--archive-optionsz -Xarchivez--ptxas-optionsz-Xptxasz--nvlink-optionsz-Xnvlinkz --threads-tz --keep-dirz -keep-dirz --run-argsz -run-argsz--input-drive-prefixz-idpz--dependency-drive-prefixz-ddpz--drive-prefixz-dpz--dependency-target-namez-MTz--default-streamz-default-streamz--gpu-architecturez-archz --gpu-codez-codez--generate-codez-gencodez--relocatable-device-codez-rdcz --entriesz-ez--maxrregcountz -maxrregcountz--ftzz-ftzz --prec-divz -prec-divz --prec-sqrtz -prec-sqrtz--fmadz-fmadz--Werrorz-Werrorz --diag-errorz -diag-errorz--diag-suppressz-diag-suppressz --diag-warnz -diag-warnz--options-filez-optfz-timez -qpp-config)z--timez --qpp-configcCsi|]\}}||qSr#r#).0kvr#r#r$ zCudaCompiler.ZnvccNccache T.List[str]exelistversionstr for_machineris_crossbool exe_wrapperT.Optional['ExternalProgram'] host_compilerr info 'MachineInfo'rT.Optional['DynamicLinker'] full_versionT.Optional[str]c sNtj|||||| | |d|_|_|j_fdd|jD_dS)N)rrKrCcsi|] \}}||qSr#)_to_host_flags)r8levelflagsselfr#r$r;z)CudaCompiler.__init__..)super__init__rErGZ base_options warn_argsitems) rQr=r?r@rBrCrErGrHrrK __class__rPr$rTs zCudaCompiler.__init__Targlistmodereturnc sTd}d}d}d}|||ttjd}||vs|sO||vr.t||r,|||S|Sfdd||D}tfdd|Dd d g}d |Sd g}d } t|} | D]A} | |krh| sh|d qZ| |krx|d | 7<| } qZ| |krz |d t | 7<WqZt yYn w|d | 7<qZfd d|D}d|S)a  Shield an argument against both splitting by NVCC's list-argument parse logic, and interpretation by any shell. NVCC seems to consider every comma , that is neither escaped by \ nor inside a double-quoted string a split-point. Single-quotes do not provide protection against splitting; In fact, after splitting they are \-escaped. Unfortunately, double-quotes don't protect against shell expansion. What follows is a complex dance to accommodate everybody. '",\z"$`\csg|]}|qSr#_shield_nvcc_list_argr8sclsr#r$ r<z6CudaCompiler._shield_nvcc_list_arg..csg|]}|gqSr#r#rb)DQSQr#r$rfNFcsg|] }j|ddqS)F)rZr`rbrdr#r$rfsz\,) setstring whitespace intersectionsplitsumjoiniterappendnext StopIteration) rerYrZZSQZDQZCMZBSZquotablelinstringZargitcr#)rgrer$ras@          z"CudaCompiler._shield_nvcc_list_argrOcst|dkr|St|}g}dddddd dfd d }dfdd }d}|D]G}||s9d}||q+|rO|dd7<|d|||7<q+|rcd}||||||q+|rod}||q+td|S)z The flags to NVCC gets exceedingly verbose and unreadable when too many of them are shielded with -Xcompiler. Merge consecutive -Xcompiler-wrapped arguments into one. r flagrAr[rDcSs|dkS)Nr6r#ryr#r#r$is_xcompiler_flag_isolatedz=CudaCompiler._merge_flags..is_xcompiler_flag_isolatedcSs |dS)N -Xcompiler=) startswithrzr#r#r$is_xcompiler_flag_glued  z:CudaCompiler._merge_flags..is_xcompiler_flag_gluedcs|p|SNr#rzrr{r#r$is_xcompiler_flag z4CudaCompiler._merge_flags..is_xcompiler_flagflagitT.Iterator[str]cs8|r |tddSzt|WStyYdSw)Nr}rj)lenrtru)ryr)rr#r$get_xcompiler_vals  z4CudaCompiler._merge_flags..get_xcompiler_valFrir^Tz6-Xcompiler flag merging failed, unknown argument form!N)ryrAr[rD)ryrArrr[rA)rrrrs ValueError)rerOrxflagsrrZingroupryr#rr$ _merge_flagss2       zCudaCompiler._merge_flagsphaserc Csg}t|}|D]}||jvr||q|dddvr$||q|dddkrEd|vr2dnd}|d|jd |||qt|d kr|d d kr|dd vr|d ddkrozt|}WntynYnw|d dd kr~|dd}n|d d}|dd }n||jvs||jvrzt|}WntyYnw| d dd |jvs| d dd |jvr| d d\}}np| dr|dd }|dd}n\|dkr|d|d|nJ|dkr|d|d|d|d|n.|dkr|d|d|n|dkr)|d|d|n |d| |q|dus;J|j ||}|dvrut|d krY||| |q|dkrg||jvrgq|||| |q|dkr|dkr|d|d|d|||q|d vr|d|||q|||q|d!vr|||q|d"vr||d |q||||q||S)#z Translate generic "GCC-speak" plus particular "NVCC-speak" flags to NVCC flags. NVCC's "short" flags have broad similarities to the GCC standard, but have gratuitous, irritating differences. Nr z-//r^r]rjz-X=rr-Z IDULlmOxmter0z -ffast-mathr)r}z-fno-fast-mathz -ftz=falsez-prec-div=truez-prec-sqrt=truez-freciprocal-mathz-prec-div=falsez-fno-reciprocal-math>r+r0r/r1r,r3Zfastz-O3r6>rczg>r7r.r-r5>r4)rr_FLAG_PASSTHRU_NOARGSrsvaluerrtru_FLAG_LONG2SHORT_WITHARGS_FLAG_SHORT2LONG_WITHARGSror~stripragetrGZget_default_include_dirsr)rQrOrrrrywrapvalr#r#r$rM,s  3   $                                zCudaCompiler._to_host_flagscCdS)NFr#rPr#r#r$needs_static_linkerz CudaCompiler.needs_static_linker environment 'Environment'cC||j|tjSr)rMrGthread_link_flagsrr")rQrr#r#r$rzCudaCompiler.thread_link_flagswork_direnvNonecCsrtd|dd|jtdt|jd}d}|ddd }||jr,d nd 7}tj ||}tj ||d }t |d dd }| |Wdn1sVwYd |_ g} | ddd|g7} | | |jj7} |jr|jdur| |7} | ||7} |j| } tdd| t| |d\} } } tdt| tdt| td| jd krtd|d|jr|jdurdS|j|g} n |jdd|dg} tdd| t| |d\}} } tdt| tdt| td||jd kr&td|jd |d!| d kr0| |_ dStd"| dS)#NzSanity testing z compiler: zIs cross compiler: %s.zsanitycheckcuda.cua #include #include __global__ void kernel (void) {} int main(void){ struct cudaDeviceProp prop; int count, i; cudaError_t ret = cudaGetDeviceCount(&count); if(ret != cudaSuccess){ fprintf(stderr, "%d\n", (int)ret); }else{ for(i=0;i int main(void) {{ /* If it's not defined as a macro, try to use as a symbol */ #ifndef {symbol} {symbol}; #endif return 0; }}rTzw{prefix} #include <{header}> using {symbol}; int main(void) {{ return 0; }})Zcompiles format_map) rQrrrrrrZfargstfoundcachedr#r#r$has_header_symbolMs  zCudaCompiler.has_header_symbolz>=9.0z>=11.0z>=12.0'MutableKeyedOptionDictType'c st}td|j|jd}td|j|jd}gd}t|j|jr'|dg7}t|j|jr3|dg7}t|j|j r?|dg7}| |t d|d |t d d i|S) Nstdmachinelangccbindir)nonezc++03zc++11zc++14zc++17zc++20z&C++ language standard to use with CUDArz4CUDA non-default toolchain directory to use (-ccbin)rj)rS get_optionsr rBrr r@_CPP14_VERSION_CPP17_VERSION_CPP20_VERSIONupdaterZUserComboOptionZUserStringOption)rQoptsstd_keyZ ccbindir_keyZcpp_stdsrWr#r$rns&    zCudaCompiler.get_optionsr'KeyedOptionDictType'csFfdd|jD}td|j|jjd}|di}tj||dS)zM Convert an NVCC Option set to a host compiler's option set. csi|] \}}|||qSr#)r)r8keyoptrr#r$r;sz:CudaCompiler._to_host_compiler_options..rrr) overrides)rGrrVr rBrrZ OptionsView)rQrZ host_optionsrrr#rr$_to_host_compiler_optionssz&CudaCompiler._to_host_compiler_optionscCs`||}ts"td|j|jd}||}|jdkr"|d|j|||j | |S)Nrrrz--std=) rrr rBrrrsrMrGget_option_compile_argsr)rQrargsrrr#r#r$rs  z$CudaCompiler.get_option_compile_argscCs*||}|||j||tjSr)rrMrGget_option_link_argsrrr")rQrrr#r#r$rs  z!CudaCompiler.get_option_link_args shlib_namesuffix soversiondarwin_versionsT.Tuple[str, str]c Cs ||j||||||tjSr)rMrGget_soname_argsrr")rQrrrrrrr#r#r$rs  zCudaCompiler.get_soname_argscCdgS)Nr&r#rPr#r#r$rz"CudaCompiler.get_compile_only_argscCr)Nz-O0r#rPr#r#r$get_no_optimization_argsrz%CudaCompiler.get_no_optimization_argsoptimization_levelcCt|Sr)r)rQrr#r#r$get_optimization_argssz"CudaCompiler.get_optimization_argsrcC||j|Sr)rMrGsanitizer_compile_argsrQrr#r#r$rz#CudaCompiler.sanitizer_compile_argscCrr)rMrGsanitizer_link_argsrr#r#r$rrz CudaCompiler.sanitizer_link_argsis_debugcCrr)r)rQrr#r#r$get_debug_argsr|zCudaCompiler.get_debug_argscCr)NzB-Werror=cross-execution-space-call,deprecated-declarations,reorderr#rPr#r#r$get_werror_argsrzCudaCompiler.get_werror_argsrNcCs |j|Sr)rU)rQrNr#r#r$ get_warn_argsrzCudaCompiler.get_warn_args buildtypecCs(dd|j|D}t|||S)NcSsg|] }|dkr dn|qS)z/ZIz/Zir#r8rYr#r#r$rfrRz3CudaCompiler.get_buildtype_args..)rGget_buildtype_argsrrM)rQrZ host_argsr#r#r$rszCudaCompiler.get_buildtype_argsr is_systemcCs$|dkrd}|r d|gSd|gS)Nrjrz -isystem=r/r#)rQrrr#r#r$get_include_argsszCudaCompiler.get_include_argsFrel_objpchcC||j||Sr)rMrGget_compile_debugfile_args)rQrrr#r#r$rz'CudaCompiler.get_compile_debugfile_args targetfilecCrr)rMrGget_link_debugfile_argsrr")rQrr#r#r$rrz$CudaCompiler.get_link_debugfile_argscCr)Ndr#rPr#r#r$get_depfile_suffixrzCudaCompiler.get_depfile_suffixcCrr)rMrGget_buildtype_linker_argsrr")rQrr#r#r$rrz&CudaCompiler.get_buildtype_linker_args build_dirfrom_dir rpath_pathsT.Tuple[str, ...] build_rpath install_rpath"T.Tuple[T.List[str], T.Set[bytes]]c Cs,|j||||||\}}||tj|fSr)rGbuild_rpath_argsrMrr") rQrrrr r r Z rpath_argsZrpath_dirs_to_remover#r#r$rs zCudaCompiler.build_rpath_argsrcCs|Srr#)rQrr#r#r$linker_to_compiler_argsrz$CudaCompiler.linker_to_compiler_argscCs||jSr)rMrG get_pic_argsrPr#r#r$rrzCudaCompiler.get_pic_argsparameter_listcCsgSrr#)rQrrr#r#r$&compute_parameters_with_absolute_pathssz3CudaCompiler.compute_parameters_with_absolute_pathstargetcCsd|gS)Nr*r#rQrr#r#r$rr|zCudaCompiler.get_output_argscCs||jtjSr)rMrGget_std_exe_link_argsrr"rPr#r#r$rrz"CudaCompiler.get_std_exe_link_argslibname extra_dirslibtyper lib_prefix_warningT.Optional[T.List[str]]cCs d|gS)Nr,r#)rQrrrrrr#r#r$ find_librarys zCudaCompiler.find_librarycrt_valcCrr)rMrGget_crt_compile_args)rQrrr#r#r$rrz!CudaCompiler.get_crt_compile_argscCsJg}|j||}tdd|Dr|dg7}|||j||tjS)Ncss|]}|dvVqdS)>z/MDdz/MTdz/MDNr#rr#r#r$ sz1CudaCompiler.get_crt_link_args..z/NODEFAULTLIB:LIBCMT.lib)rGranyrMget_crt_link_argsrr")rQrrZhost_link_arg_overridesZhost_crt_compile_argsr#r#r$r s  zCudaCompiler.get_crt_link_args 'BuildTarget'c|t|tjSr)rMrSget_target_link_argsrr"rrWr#r$r#rz!CudaCompiler.get_target_link_argsdep 'Dependency'cs|t|Sr)rMrSget_dependency_compile_argsrQr$rWr#r$r&rz(CudaCompiler.get_dependency_compile_argscr"r)rMrSget_dependency_link_argsrr"r'rWr#r$r(rz%CudaCompiler.get_dependency_link_argscCsDtd|j|jd}||j}t|tr |dkr |d|dgSgS)Nrrrjz-ccbin=F)r rBrr isinstancerAra)rQrrrr#r#r$rs  zCudaCompiler.get_ccbin_argscCdd|jDS)NcSg|]}d|qSr}r#r8xr#r#r$rfrhz:CudaCompiler.get_profile_generate_args..)rGget_profile_generate_argsrPr#r#r$r/rz&CudaCompiler.get_profile_generate_argscCr*)NcSr+r,r#r-r#r#r$rfrhz5CudaCompiler.get_profile_use_args..)rGget_profile_use_argsrPr#r#r$r0rz!CudaCompiler.get_profile_use_argsdisablecCs |j|Sr)rGget_assert_args)rQr1r#r#r$r2s zCudaCompiler.get_assert_args)NN)r=r>r?r>r@rArBrrCrDrErFrGr rHrIrrJrKrL)T)rYrArZrDr[rA)rOr>r[r>)rOr>rrr[r>)r[rD)rrr[r>)rrArrr[r)rrArrArrArrrrrrr[r)r[r)rrr[r)rrr[r>)rrrrArrArrArrArrr[r>)r[r>)rrAr[r>)rrAr[r>)rrDr[r>)rNrAr[r>)rrAr[r>)rrArrDr[r>)F)rrArrDr[r>)rrAr[r>)r[rA)rrrrArrAr r r rAr rAr[r )rr>r[r>)rr>rrAr[r>)rrAr[r>) rrArrrr>rr rrDr[r)rrArrAr[r>)rr!r[r>)r$r%r[r>)r1rDr[r>)>rrr Z LINKER_PREFIXrrrrVridrT classmethodrarrr!rMrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr Z PREFER_SHAREDrrr r#r&r(rr/r0r2 __classcell__r#r#rWr$r%0sLC      !"#$%&'()*+,-./012348 :0 4  j                      r%)+ __future__renumos.pathrrltypingTrjrrZmesonlibrrrr r r compilersr rrr TYPE_CHECKINGrbuildrrrrrrrZ envconfigrZlinkers.linkersrrZprogramsrEnumrr%r#r#r#r$s*