
    eߣ                    H   d dl m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mZmZmZmZmZ ddlmZmZmZmZ ej.                  r8dd	lmZ dd
lmZ ddl	mZ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% ddl&m'Z'  G d dejP                        Z) G d de      Z*y)    )annotationsN   )coredata)mlog)EnvironmentException
Popen_safe
is_windowsLibType	OptionKeyversion_compare   )Compilercuda_buildtype_argscuda_optimization_argscuda_debug_args)CompileCheckMode)BuildTarget)MutableKeyedOptionDictTypeKeyedOptionDictType)
Dependency)Environment)MachineInfo)DynamicLinker)MachineChoice)ExternalProgramc                      e Zd ZdZdZy)_PhasecompilerlinkerN)__name__
__module____qualname__COMPILERLINKER     ;/usr/lib/python3/dist-packages/mesonbuild/compilers/cuda.pyr   r   *   s    HFr&   r   c                     e Zd ZdZdZh dZ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#d$d%i d&d'd(d)d*d+d,d-d.d/d0d1d2d3d4d5d6d7d8d9d:d;d<d=d>d?d@dAdBdCdDdEdFdGi dHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd[d\d]d^d_d`dadbdcdddedfdgdhdidjdkdlZej                         D  ci c]  \  }}||
 c}}} ZdmZ		 	 d	 	 	 	 	 	 	 	 	 	 	 dfdoZ
edddq       Zeddr       Zej                  fddsZddtZdduZddvZdndndw	 	 	 	 	 	 	 ddxZdyZdzZd{Zdfd|Zdd}Zdd~ZddZ	 	 	 	 	 	 	 	 ddZddZddZddZddZ ddZ!ddZ"ddZ#ddZ$ddZ%ddZ&dddZ'ddZ(ddZ)ddZ*	 	 	 	 	 	 	 	 ddZ+ddZ,ddZ-	 	 	 	 ddZ.ddZ/ddZ0e1jd                  dpf	 	 	 	 	 ddZ3ddZ4ddZ5dfdZ6dfdZ7dfdZ8ddZ9ddZ:ddZ;ddZ<xZ=S c c}}} w )CudaCompilerz	-Xlinker=cuda>w   -E-G-M-V-g-h-v-MD-MM-MP-dc-dw-pg-MMD-ewp-lib-m32-m64-ptx--lib--m32--m64--ptx-cuda-dlto-keep-noeh--cuda--help--keep--link-clean-cubin-dlink--cubin--debug-dryrun-err-no-fatbin-noprof-shared--dryrun--fatbin--shared-arch-ls-code-ls-nodlink-objtemp	--compile	--profile	--verbose	--version	-Wreorder	-lineinfo	-restrict
--Wreorder
--device-c
--device-w
--restrict
-res-usage-save-temps-src-in-ptx--preprocess--save-temps--device-link-astoolspatch-nohdinitlist--device-debug--clean-targets--list-gpu-arch--list-gpu-code--no-exceptions--source-in-ptx--use_fast_math--dlink-time-opt--no-device-link--resource-usage-extended-lambda--extended-lambda--no-align-double--disable-warnings--dont-use-profile--objdir-as-tempdir--generate-line-info-expt-extended-lambda--display-error-number--expt-extended-lambda-keep-device-functions--generate-dependencies--keep-device-functions-Wdefault-stream-launch-expt-relaxed-constexpr--Wdefault-stream-launch--compile-as-tools-patch--expt-relaxed-constexpr--extensible-whole-program-Wext-lambda-captures-this--Wext-lambda-captures-this-Wno-deprecated-gpu-targets-allow-unsupported-compiler-extra-device-vectorization--Wno-deprecated-gpu-targets--allow-unsupported-compiler--extra-device-vectorization-Wno-deprecated-declarations--Wno-deprecated-declarations--generate-dependency-targets-forward-unknown-to-host-linker --forward-unknown-to-host-linker!--generate-nonsystem-dependencies!--no-host-device-initializer-list!-forward-unknown-to-host-compiler"--forward-unknown-to-host-compiler$--generate-dependencies-with-compile.--generate-nonsystem-dependencies-with-compile-c-w--run-use_fast_mathz--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-confignvccNc           
        t         |   ||||||	|
|       || _        || _        |j                  | _        |j
                  j                         D ci c]  \  }}|| j                  |       c}}| _        y c c}}w )N)r   full_versionis_cross)super__init__exe_wrapperhost_compilerbase_options	warn_argsitems_to_host_flags)selfccacheexelistversionfor_machiner   r   r   infor   r   levelflags	__class__s                r'   r   zCudaCompiler.__init__   s    
 	';Vbn  zB  	C&*)66P]PgPgPmPmPopu%!4!4U!;;pps   A;Tc                   d}d}d}d}||z   |z   }t        t        j                  dz         }||vs|s||vr$t        |      j                  |      r||z   |z   S |S |j	                  |      D 	cg c]  }	| j                  |	       }
}	t        |
D 	cg c]  }	|	|g c}	dd g       }
dj                  |
      S dg}
d	}t        |      }|D ][  }||k(  r|s|
j                  d       ||k(  r|
dxx   |z  cc<   | }2||k(  r	 |
dxx   t        |      z  cc<   O|
dxx   |z  cc<   ] |
D 	cg c]  }	| j                  |	d	
       }
}	dj                  |
      S c c}	w c c}	w # t        $ r Y  Gw xY wc c}	w )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"$`\N F)listmodez\,)setstring
whitespaceintersectionsplit_shield_nvcc_list_argsumjoiniterappendnextStopIteration)clsargr   SQDQCMBSDQSQquotableslinstringargitcs                 r'   r   z"CudaCompiler._shield_nvcc_list_arg   s    "uRxv((01S=} s8((2c6"9$J <?99R=IaS..q1IIA.q!T.s3R8wwqz! AHIE 78HHRL"WbEQJE#+|H"W"e, bEQJE  HII!**1u*=IAI::a= 5 J." )  Js$   .EEE#,E3#	E0/E0c                  	 t        |      dk  r|S t        |      }g }dd	ddd	fd}dfd}d}|D ]  } ||      sd}|j                  |       |r"|dxx   dz  cc<   |dxx    |||      z  cc<   C 	|      r,d	}|j                  |       |j                   |||             w |      rd	}|j                  |       t        d
       |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   c                    | dk(  S )Nr   r%   flags    r'   is_xcompiler_flag_isolatedz=CudaCompiler._merge_flags.<locals>.is_xcompiler_flag_isolated  s    <''r&   c                $    | j                  d      S N-Xcompiler=)
startswithr   s    r'   is_xcompiler_flag_gluedz:CudaCompiler._merge_flags.<locals>.is_xcompiler_flag_glued
  s    ??=11r&   c                (     |       xs  |       S Nr%   )r   r   r   s    r'   is_xcompiler_flagz4CudaCompiler._merge_flags.<locals>.is_xcompiler_flag  s    -d3T7Nt7TTr&   c                f     |       r| t        d      d  S 	 t        |      S # t        $ r Y yw xY w)Nr   r   )lenr   r   )r   flagitr   s     r'   get_xcompiler_valz4CudaCompiler._merge_flags.<locals>.get_xcompiler_val  s?    &t,C./00<'$ s   
$ 	00Fr   r   Tz6-Xcompiler flag merging failed, unknown argument form!)r   strreturnbool)r   r   r   zT.Iterator[str]r   r   )r   r   r   
ValueError)
r   r   r   xflagsr   r   ingroupr   r   r   s
           @@r'   _merge_flagszCudaCompiler._merge_flags   s     u:?Le	(	2	U	  	[D$T*d#r
c!
r
/f==
+D1d#/f=>(.d# !YZZ	[  r&   c           	     z   g }t        |      }|D ]  }|| j                  v r|j                  |       $|dd dvr|j                  |       =|dd dk(  r.d|v rdnd}|j                  d|j                   d	| | |        st	        |      d
k\  r>|d   dk(  r6|d   dv r/|d
d dk(  r	 t        |      }n|d
d d	k(  r|dd }n|d
d }|dd
 }n|| j                  v s|| j                  v r	 t        |      }nv|j                  d	d      d   | j                  v s!|j                  d	d      d   | j                  v r|j                  d	d      \  }}n|j                  d      r|dd j                         }|dd }n|dk(  r&|j                  d       |j                  d|z          n|dk(  rH|j                  d       |j                  d       |j                  d       |j                  d|z          ny|dk(  r&|j                  d       |j                  d|z          nN|dk(  r&|j                  d       |j                  d|z          n#|j                  d| j                  |      z          _J | j                  j                  ||      }|dv rt	        |      d
k(  r%|j                  || j                  |      z          |dk(  r|| j                  j                         v r|j                  |       |j                  | j                  |             |dk(  r|dk(  rI|j                  d       |j                  d       |j                  d       |j                  ||z          _|d v r'|j                  d       |j                  ||z          |j                  ||z          |d!v r|j                  ||z          |d"v r|j                  |d	z   |z          |j                  |       |j                  |        | j!                  |      S # t        $ r Y aw xY w# t        $ r Y w xY w)#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   r   z-X=r   r   -IDULlmOxmte   r      z-ffast-mathr   r   z-fno-fast-mathz
-ftz=falsez-prec-div=truez-prec-sqrt=truez-freciprocal-mathz-prec-div=falsez-fno-reciprocal-math>   r   r   r   r   r   r   fastz-O3r   >   gr   z>   r   r   r   r   >   r   )r   _FLAG_PASSTHRU_NOARGSr   valuer   r   r   _FLAG_LONG2SHORT_WITHARGS_FLAG_SHORT2LONG_WITHARGSr   r   stripr   getr   get_default_include_dirsr   )r   r   phaser   r   r   wrapvals           r'   r   zCudaCompiler._to_host_flags,  s-    e d	#Df t111d# BQxt#d#bqS "Tksr5;;-qtfTFCDTaDGsNtAw-7O !9?"6l !AY#%qr(Cqr(CBQx777D:::v,C C#A&$*H*HHJJsA&q)T-K-KK !JJsA.	c,12hnn&BQx =(MM"23MM-"45--MM,/MM"23MM"34MM-"4500MM"34MM-"4533MM"23MM-"45MM-0J0J40P"PQ?"? 1155dDADAA t9>MM$t'A'A#'F"FGZ'C43E3E3^3^3`,` MM$'MM$"<"<S"AB&=MM%(MM"23MM,/MM$s(+O+MM,/MM$s(+MM$s(+11d3h'!d3hsl+d#c"Id	#L   ((y )  % s$   )P,P-	P*)P*-	P:9P:c                     y)NFr%   r   s    r'   needs_static_linkerz CudaCompiler.needs_static_linker  s    r&   c                t    | j                  | j                  j                  |      t        j                        S r   )r   r   thread_link_flagsr   r$   )r   environments     r'   r  zCudaCompiler.thread_link_flags  s-    ""4#5#5#G#G#TV\VcVcddr&   c                   t        j                  d| j                         z   dz   dj                  | j                               t        j                  dt        | j                        z         d}d}|j                  dd      d	   }|| j                  rd
ndz  }t        j                  j                  ||      }t        j                  j                  ||dz         }t        |dd      5 }|j                  |       d d d        d| _        g }	|	ddd|gz  }	|	| j                  |j                  j                        z  }	| j                  r| j                   |	| j#                         z  }	|	| j%                  |      z  }	| j                  |	z   }
t        j                  ddj                  |
             t'        |
|      \  }}}t        j                  d       t        j                  |       t        j                  d       t        j                  |       t        j                  d       |j(                  d	k7  rt+        d| j-                          d      | j                  r,| j                   y | j                   j/                         |gz   }
n| j                  dd|z   dz   gz   }
t        j                  ddj                  |
             t'        |
|      \  }}}t        j                  d       t        j                  |       t        j                  d       t        j                  |       t        j                  d       |j1                          |j(                  d	k7  r*t+        d| j2                   d | j-                          d!      |dk(  r|| _        y t        j                  d"|z          y # 1 sw Y   xY w)#NzSanity testing z
 compiler: zIs cross compiler: %s.zsanitycheckcuda.cua  
        #include <cuda_runtime.h>
        #include <stdio.h>

        __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<count;i++){
                    if(cudaGetDeviceProperties(&prop, i) == cudaSuccess){
                        fprintf(stdout, "%d.%d\n", prop.major, prop.minor);
                    }
                }
            }
            fflush(stderr);
            fflush(stdout);
            return 0;
        }
        .r   r   _crossr   z.exewzutf-8)encodingr   r   staticz$Sanity check compiler command line: )cwdzSanity check compile stdout: z"-----
Sanity check compile stderr:z-----z	Compiler z cannot compile programs.r   r   zSanity check run command line: zSanity check run stdout: z-----
Sanity check run stderr:zExecutables created by z
 compiler z are not runnable.zcudaGetDeviceCount() returned )r   debugget_display_languager   r   r   r   rsplitospathopenwritedetected_ccget_ccbin_argsr   optionsr   get_compile_only_argsget_output_argsr   
returncoder   name_stringget_commandwaitlanguage)r   work_direnvsnamecodebinnamesource_namebinary_nameofiler   cmdlistpcstdostdepes                  r'   sanity_checkzCudaCompiler.sanity_check  s.   

$t'@'@'BB\QSVS[S[\`\h\hSij

+c$--.@@A$0 ,,sA&q)t}}8"4ggll8U3ggll8Wv-=>+sW5 	KK	  	$	8[99
 	$$S\\%9%9:: ==T--5 T//11E%%k22 ,,&

9388G;LM#G:D$

23

4

89

4

7==A&43C3C3E2FF_'`aa =='**668K=Hllgs[/@3/F%GGG

4chhw6GH#G:D$

./

4

45

4

7
	==A&)@zZ^ZjZjZlYmm  (A  B  B
 2:#DJJ7$>?K	 	s   ,N66O 
extra_argsdependenciesc                   |g }|||d}d}| j                  |j                  |      |||      \  }	}
|	rd|
fS d}| j                  |j                  |      |||      S )N)prefixheadersymbolz{prefix}
        #include <{header}>
        int main(void) {{
            /* If it's not defined as a macro, try to use as a symbol */
            #ifndef {symbol}
                {symbol};
            #endif
            return 0;
        }}r?  Tzw{prefix}
        #include <{header}>
        using {symbol};
        int main(void) {{
            return 0;
        }})compiles
format_map)r   hnamerE  rC  r2  r@  rA  fargstfoundcacheds              r'   has_header_symbolzCudaCompiler.has_header_symbolM  s     J!UfE all5&93:dpqv< }}Q\\%0#*[g}hhr&   z>=9.0z>=11.0z>=12.0c           	        t         |          }t        d| j                  | j                        }t        d| j                  | j                        }g d}t        | j                  | j                        r|dgz  }t        | j                  | j                        r|dgz  }t        | j                  | j                        r|dgz  }|j                  |t        j                  d|d	      |t        j                  d
d      i       |S )Nstdmachinelangccbindir)nonezc++03zc++11zc++14zc++17zc++20z&C++ language standard to use with CUDArT  z4CUDA non-default toolchain directory to use (-ccbin)r   )r   get_optionsr   r   r0  r   r   _CPP14_VERSION_CPP17_VERSION_CPP20_VERSIONupdater   UserComboOptionUserStringOption)r   optsstd_keyccbindir_keycpp_stdsr   s        r'   rU  zCudaCompiler.get_optionsn  s    w"$E4+;+;$--P T5E5EDMMZ-4<<)<)<=	!H4<<)<)<=	!H4<<)<)<=	!H(223[3;VE(334j468
 	 r&   c           	     2   | j                   j                         j                         D ci c]  \  }}||j                  ||       }}}t	        d| j
                  | j                   j                        }|di}t        j                  ||      S c c}}w )zM
        Convert an NVCC Option set to a host compiler's option set.
        rO  rP  rT  )	overrides)	r   rU  r   r  r   r   r0  r   OptionsView)r   r)  keyopthost_optionsr]  ra  s          r'   _to_host_compiler_optionsz&CudaCompiler._to_host_compiler_options  s     DHCUCUCaCaCcCiCiCklxsCW[[c22llE4+;+;$BTBTB]B]^f%	##LIFF ms   Bc                X   | j                  |      }t               sTt        d| j                  | j                        }||   }|j
                  dk7  r|j                  d|j
                  z          || j                  | j                  j                  | j                  |                  z   S )NrO  rP  rT  z--std=)r(  r	   r   r   r0  r	  r   r   r   get_option_compile_argsrf  )r   r)  argsrc  rO  s        r'   rh  z$CudaCompiler.get_option_compile_args  s    ""7+ |E4+;+;$--PC#,CyyF"Hsyy01d))$*<*<*T*TUYUsUst{U|*}~~~r&   c                    | j                  |      }|| j                  | j                  j                  | j	                  |            t
        j                        z   S r   )r(  r   r   get_option_link_argsrf  r   r$   )r   r)  ri  s      r'   rk  z!CudaCompiler.get_option_link_args  sd    ""7+d))$*<*<*Q*QRVRpRpqxRy*z  }C  }J  }J  K  K  	Kr&   c           
     ~    | j                  | j                  j                  ||||||      t        j                        S r   )r   r   get_soname_argsr   r$   )r   r2  rC  
shlib_namesuffix	soversiondarwin_versionss          r'   rm  zCudaCompiler.get_soname_args  sA     ""4#5#5#E#EVY$IJP--Y 	Yr&   c                    dgS )Nr   r%   r  s    r'   r*  z"CudaCompiler.get_compile_only_args  s	    vr&   c                    dgS )Nz-O0r%   r  s    r'   get_no_optimization_argsz%CudaCompiler.get_no_optimization_args  s	    wr&   c                    t         |   S r   )r   )r   optimization_levels     r'   get_optimization_argsz"CudaCompiler.get_optimization_args  s     &&899r&   c                V    | j                  | j                  j                  |            S r   )r   r   sanitizer_compile_argsr   r	  s     r'   ry  z#CudaCompiler.sanitizer_compile_args  s$    ""4#5#5#L#LU#STTr&   c                V    | j                  | j                  j                  |            S r   )r   r   sanitizer_link_argsrz  s     r'   r|  z CudaCompiler.sanitizer_link_args  s$    ""4#5#5#I#I%#PQQr&   c                    t         |   S r   )r   )r   is_debugs     r'   get_debug_argszCudaCompiler.get_debug_args  s    x((r&   c                    dgS )NzB-Werror=cross-execution-space-call,deprecated-declarations,reorderr%   r  s    r'   get_werror_argszCudaCompiler.get_werror_args  s    TUUr&   c                     | j                   |   S r   )r   )r   r   s     r'   get_warn_argszCudaCompiler.get_warn_args  s    ~~e$$r&   c                    | j                   j                  |      D cg c]  }|dk(  rdn| }}t        |   | j                  |      z   S c c}w )Nz/ZIz/Zi)r   get_buildtype_argsr   r   )r   	buildtyper   	host_argss       r'   r  zCudaCompiler.get_buildtype_args  sW     @D?Q?Q?d?den?opcUlU3p	p"9-0C0CI0NNN qs   Ac                ,    |dk(  rd}|rd|z   gS d|z   gS )Nr   r  z	-isystem=r   r%   )r   r$  	is_systems      r'   get_include_argszCudaCompiler.get_include_args  s)    2:D'0d"#Ctd{mCr&   c                X    | j                  | j                  j                  ||            S r   )r   r   get_compile_debugfile_args)r   rel_objpchs      r'   r  z'CudaCompiler.get_compile_debugfile_args  s(    ""4#5#5#P#PQXZ]#^__r&   c                t    | j                  | j                  j                  |      t        j                        S r   )r   r   get_link_debugfile_argsr   r$   )r   
targetfiles     r'   r  z$CudaCompiler.get_link_debugfile_args  s-    ""4#5#5#M#Mj#Y[a[h[hiir&   c                     y)Ndr%   r  s    r'   get_depfile_suffixzCudaCompiler.get_depfile_suffix  s    r&   c                t    | j                  | j                  j                  |      t        j                        S r   )r   r   get_buildtype_linker_argsr   r$   )r   r  s     r'   r  z&CudaCompiler.get_buildtype_linker_args  s.    ""4#5#5#O#OPY#Z\b\i\ijjr&   c                    | j                   j                  ||||||      \  }}| j                  |t        j                        |fS r   )r   build_rpath_argsr   r   r$   )	r   r2  	build_dirfrom_dirrpath_pathsbuild_rpathinstall_rpath
rpath_argsrpath_dirs_to_removes	            r'   r  zCudaCompiler.build_rpath_args  sK     .2-?-?-P-PHk;.O*)##J>@TUUr&   c                    |S r   r%   )r   ri  s     r'   linker_to_compiler_argsz$CudaCompiler.linker_to_compiler_args  s    r&   c                T    | j                  | j                  j                               S r   )r   r   get_pic_argsr  s    r'   r  zCudaCompiler.get_pic_args  s"    ""4#5#5#B#B#DEEr&   c                    g S r   r%   )r   parameter_listr  s      r'   &compute_parameters_with_absolute_pathsz3CudaCompiler.compute_parameters_with_absolute_paths  s    	r&   c                
    d|gS )Nr   r%   )r   targets     r'   r+  zCudaCompiler.get_output_args  s    f~r&   c                r    | j                  | j                  j                         t        j                        S r   )r   r   get_std_exe_link_argsr   r$   r  s    r'   r  z"CudaCompiler.get_std_exe_link_args  s(    ""4#5#5#K#K#Mv}}]]r&   c                    d|z   gS )Nr   r%   )r   libnamer2  
extra_dirslibtypelib_prefix_warnings         r'   find_libraryzCudaCompiler.find_library  s    wr&   c                X    | j                  | j                  j                  ||            S r   )r   r   get_crt_compile_args)r   crt_valr  s      r'   r  z!CudaCompiler.get_crt_compile_args  s'    ""4#5#5#J#J7T]#^__r&   c                    g }| j                   j                  ||      }t        d |D              r|dgz  }| j                  || j                   j	                  ||      z   t
        j                        S )Nc              3  $   K   | ]  }|d v  
 yw)>   /MD/MDd/MTdNr%   ).0r   s     r'   	<genexpr>z1CudaCompiler.get_crt_link_args.<locals>.<genexpr>  s     O#s--Os   z/NODEFAULTLIB:LIBCMT.lib)r   r  anyr   get_crt_link_argsr   r$   )r   r  r  host_link_arg_overrideshost_crt_compile_argss        r'   r  zCudaCompiler.get_crt_link_args  s     #% $ 2 2 G GQZ [O9NOO#(B'CC#""#:T=O=O=a=abikt=u#uw}  xE  xE  F  	Fr&   c                ^    | j                  t        | 	  |      t        j                        S r   )r   r   get_target_link_argsr   r$   )r   r  r   s     r'   r  z!CudaCompiler.get_target_link_args  s$    ""57#?#GWWr&   c                @    | j                  t        | 	  |            S r   )r   r   get_dependency_compile_argsr   depr   s     r'   r  z(CudaCompiler.get_dependency_compile_args  s    ""57#Fs#KLLr&   c                ^    | j                  t        | 	  |      t        j                        S r   )r   r   get_dependency_link_argsr   r$   r  s     r'   r  z%CudaCompiler.get_dependency_link_args  s$    ""57#CC#H&--XXr&   c                    t        d| j                  | j                        }||   j                  }t	        |t
              r|dk7  r| j                  d|z   d      gS g S )NrS  rP  r   z-ccbin=F)r   r   r0  r	  
isinstancer   r   )r   r)  rc  rS  s       r'   r(  zCudaCompiler.get_ccbin_args  sY    
D,<,<4==Q3<%%h$R..y/A5IJJIr&   c                b    | j                   j                         D cg c]  }d|z   	 c}S c c}w r   )r   get_profile_generate_argsr   xs     r'   r  z&CudaCompiler.get_profile_generate_args  s*    +/+=+=+W+W+YZa!ZZZ   ,c                b    | j                   j                         D cg c]  }d|z   	 c}S c c}w r   )r   get_profile_use_argsr  s     r'   r  z!CudaCompiler.get_profile_use_args  s*    +/+=+=+R+R+TUa!UUUr  c                8    | j                   j                  |      S r   )r   get_assert_args)r   disables     r'   r  zCudaCompiler.get_assert_args  s    !!11'::r&   )NN)r   T.List[str]r   r  r   r   r   r   r   r   r   zT.Optional['ExternalProgram']r   r   r   z'MachineInfo'r   zT.Optional['DynamicLinker']r   zT.Optional[str])T)r   r   r   r   r   r   )r   r  r   r  )r   r  r  r   r   r  )r   r   )r  'Environment'r   r  )r1  r   r2  r  r   None)rH  r   rE  r   rC  r   r2  r  r@  zGT.Union[None, T.List[str], T.Callable[[CompileCheckMode], T.List[str]]]rA  z T.Optional[T.List['Dependency']]r   zT.Tuple[bool, bool])r   z'MutableKeyedOptionDictType')r)  'KeyedOptionDictType'r   r  )r)  r  r   r  )r2  r  rC  r   rn  r   ro  r   rp  r   rq  zT.Tuple[str, str]r   r  )r   r  )rv  r   r   r  )r	  r   r   r  )r~  r   r   r  )r   r   r   r  )r  r   r   r  )r$  r   r  r   r   r  )F)r  r   r  r   r   r  )r  r   r   r  )r   r   )r2  r  r  r   r  r   r  zT.Tuple[str, ...]r  r   r  r   r   z"T.Tuple[T.List[str], T.Set[bytes]])ri  r  r   r  )r  r  r  r   r   r  )r  r   r   r  )r  r   r2  r  r  r  r  r
   r  r   r   zT.Optional[T.List[str]])r  r   r  r   r   r  )r  z'BuildTarget'r   r  )r  z'Dependency'r   r  )r  r   r   r  )>r    r!   r"   LINKER_PREFIXr0  r  r
  r   r  idr   classmethodr   r   r   r#   r   r  r  r>  rM  rV  rW  rX  rU  rf  rh  rk  rm  r*  rt  rw  ry  r|  r  r  r  r  r  r  r  r  r  r  r  r  r  r+  r  r
   PREFER_SHAREDr  r  r  r  r  r  r(  r  r  r  __classcell__)r  kvr   s   000@r'   r)   r)   0   s   MH@F6!6!
6! 	6! 		6!
 	6! 	6! 	
6! 	6! 	6! 	6! 	6! 	6! 		6! 	6! 	 6!  	!6!" 	#6!$ 	%6!& 	&0L'6!( 	0B)6!* 	+6!, 	-6!. 	/6!0 	16!2 	
36!4 	56!6 		76!8 	
96!: 	;6!< 	=6!> 	?6!@ 	A6!B 	$C6!D 	E6!F 	#G6!H 	0AI6!J 	K6!L 	M6!N 	
O6!P 	$Q6!R 	S6!T 	U6!V 	W6!X 	Y6!Z 	[6!\ 	]6!^ 		_6!` 	a6!b 	0@c6!d 	e6!f 	g6!h 180=k6!p 3L2Q2Q2S T T$!QA T	B
 8<15		q	q.K	q (	q0=	q 5	q  /		q 8! 8!t . .` BH q)fef@T quKOi,i&mi )Ii Uhi: NNN*
GKY #Y03Y):Y?JY:UR)V%OD
`jkV&7VFIV(+V0RVF:=BM^ )0(=(=Z^ % SW cz `FXMY[V;K !Us   G1r)   )+
__future__r   enumos.pathr#  r   typingTr   r   r   mesonlibr   r   r	   r
   r   r   	compilersr   r   r   r   TYPE_CHECKINGr   buildr   r   r   rA  r   r  r   	envconfigr   linkers.linkersr   r   programsr   Enumr   r)   r%   r&   r'   <module>r     sr    #       ) ) ??+#J))'/(*TYY g;8 g;r&   