o ڗi>@s<ddlmZmZddlmZmZmZmZddlm Z ddl Z ddl m Z m Z mZmZddlmZddlZddlZddlZddlZddlZddlZddlmZdefd d Ze d efd d Ze ddZe defddZ ddZ!e ddZ"e dddZ#e ddGdddZ$GdddeZ%dS)) BaseBackend GPUTarget)irpassesllvmnvidia) dataclassN)AnyDictTupleOptional) ModuleType)PathtargetcCsddS)NcSs|rdSdS)N) r)rrr)is_int8)lhsTyperhsTyper]/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/triton/backends/nvidia/compiler.pyszmin_dot_size..rrrrr min_dot_sizesrbinarycCstjd|ddtjtjtd|g}|D]5}tj|rPtj |rPt j |dgt j d}|durPt jd|dt jd }|durP||d fSqtd |) NTRITON__PATHbin --version)stderrz.*release (\d+\.\d+).*utf-8flagsz Cannot find )osenvirongetupperpathjoindirname__file__existsisfile subprocess check_outputSTDOUTresearchdecode MULTILINEgroup RuntimeError)rpathsrresultversionrrr_path_to_binarysr;cCs ttdddgd}|S)Nptxasrrr!)r/r0r;r4)r:rrrget_ptxas_version&sr=returncCsrt|tsJtt|d\}}|dkr#|dkrd|S|dkr#dS|dkr+d|S|dkr3d |Std |) zK Get the highest PTX version supported by the current CUDA driver. . PU F ?z?Triton only support CUDA 10.0 or higher, but got CUDA version: ) isinstancestrmapintsplitr7) cuda_versionmajorminorrrrptx_get_version,s rPcCs&|j}|durtd\}}t|}|S)Nr<) ptx_versionr;rP)optionsrQ_rMrrrget_ptx_version_from_options?s  rTcCs t|}td|}d|}|S)NSz+ptx)rTmin)rRrQllvm_ptx_versionfeaturesrrr get_featuresGs  rYcCs@t|d}t|WdS1swYdS)Nrb)openhashlibsha256read hexdigest)r)frrr file_hashUs $raT)frozenc@seZdZUdZeed<dZeed<dZeed<dZeed<dZ eed <dZ eed <dZ eed <d Z e eed <dZeed<d Zeed<dZeed<dZeeed<dZeeed<dZeed<dZeeed<d Zeed<d Zeed<dZeed<dZeed <dZeed!<d"d#Zd$d%Z d S)& CUDAOptions num_warpsr$num_ctas num_stagesrnum_buffers_warp_specnum_consumer_groupsreg_dec_producerreg_inc_consumerNmaxnreg)r$r$r$ cluster_dimsrQTenable_fp_fusion)fp8e5fp8e4b15supported_fp8_dtypesrdeprecated_fp8_dtypestf32default_dot_input_precision)rttf32x3ieeeallowed_dot_input_precisionsmax_num_imprecise_acc_default extern_libsFdebugcuda backend_namesanitize_overflowcCsttjd}|jdurint|j}|dds%tdt|d|d<t |dt | |j dkr?|j |j d@dksCJddS) Nlib libdeviceTRITON_LIBDEVICE_PATHzlibdevice.10.bcrzrr$znum_warps must be a power of 2)rr,parentrzdictr'r%getenvrIobject __setattr__tupleitemsre)selfdefault_libdirrzrrr __post_init__ts  zCUDAOptions.__post_init__cCsXt|j}tddt|dD|d<dddt|D}t|d S)Ncss |] \}}|t|fVqdSN)ra).0kvrrr sz#CUDAOptions.hash..rzrScSsg|] \}}|d|qS)-r)rnamevalrrr z$CUDAOptions.hash..r!) r__dict__rsortedr*rr\r]encoder_)r hash_dictkeyrrrhash}s zCUDAOptions.hash)!__name__ __module__ __qualname__rerK__annotations__rfrhrirjrkrlrmr rnrrQroboolrrr rIrsrurxryrzrr{r}r~rrrrrrrc[s.                 rccseZdZedefddZdeddffdd Zdefdd Zd d Z d d Z de e e ffddZddZeddZeddZeddZeddZeddZddZeddZZS) CUDABackendrcCs |jdkS)Nr|)backendrrrrsupports_targets zCUDABackend.supports_targetr>Ncs.t||j|_t|jtsJd|_dS)Ncubin)super__init__arch capabilityrHrK binary_ext)rr __class__rrrs  zCUDABackend.__init__csfddtjD}d|vr'ttj}|jdkr|dtt||d<d|vr4|jdkr4d|d<d |vrBt d d d k|d <|jdkrId nd |d<tdi|S)Ncsi|] }|vr||qSrr)rroptsrr rz-CUDABackend.parse_options..rrYfp8e4nvrsZ)rqroTRITON_DEFAULT_FP_FUSION1i@rryr) rc__dataclass_fields__keyssetrrraddrrr%r)rrargsrrrrr parse_optionss    zCUDABackend.parse_optionscCs(|j|j|j|jd|jd|jdfS)Nrr$)rerfsharedrn)rmetadatarrr pack_metadataszCUDABackend.pack_metadatacCs>ddlmmm}|jdkr|jn|jt|jd}|S)NrrB)convert_custom_typesr) triton.language.extra.cudalanguageextrar|rconvert_custom_float8_sm80convert_custom_float8_sm70rr)rr| codegen_fnsrrrget_codegen_implementations z&CUDABackend.get_codegen_implementationcCsddlm}d|iS)Nr)rztriton.language.extra.libdevice)rr)rrrrrget_module_maps zCUDABackend.get_module_mapcCst|dSr)r load_dialects)rctxrrrrszCUDABackend.load_dialectscCst|j}|tj|tj|tj |tj |tj |tj |tj |tj|tj||||Sr)r pass_managercontext enable_debugrcommon add_inlinerttiradd_rewrite_tensor_pointer add_combineadd_canonicalizeradd_reorder_broadcastadd_cseadd_licmadd_symbol_dceadd_loop_unrollrun)modroptpmrrr make_ttirs           zCUDABackend.make_ttircCs`t}|jdur|jd|_|jd|_|jd|_tjdddkr5t }t ||j }|j dt |j }|tj|d||jd |jtj||d d kr`tj|tjj||tj|tj|tj|tj|tj||d ktj ||d d krtj!|tj"|tj#||j$tj%||j$tj&||j$tj'||j(|j$|j)|j*tj+||j,tj-||j$tj.|tj||d ktj|tj/|tj0|tj |tj1||d d krtjj2|tjj3|tj4||5||j|j|jf|d<|S)Nrr$rMLIR_ENABLE_REMARK0rTzcuda:rrFrB rn)6r ClusterInforn clusterDimX clusterDimY clusterDimZr%r&r'r source_mgrrsource_mgr_diagrprintOpOnDiagnosticrrrradd_convert_to_ttgpuirrerfttgpuir add_coalesceadd_f32_dot_tc ttnvgpuir add_plan_ctaadd_remove_layout_conversionsadd_optimize_thread_localityadd_accelerate_matmuladd_optimize_dot_operandsrradd_optimize_accumulator_init add_combine_tensor_select_and_ifadd_ws_task_partitionrjadd_taskid_propagateadd_ws_data_partitionadd_ws_code_partitionrirkrl add_pipelinerhadd_ws_lowering add_prefetchadd_reduce_data_duplicationadd_reorder_instructionsradd_fence_insertionadd_tma_loweringrr)rrrr cluster_infosrcMgrdiagrrrr make_ttgirs^                         zCUDABackend.make_ttgircCst|}|d}|dur|d|9<|}t|j}|tjdddkr;t }t ||j} |j dt jj|tj|tj|tj|tj|t jj|||t jj|tj|tj|tj|tj|tjdddkrtj|| |t !t } t "|| } |dkrd nd |} t#|} d }t $| || | t %| |j&dur| 'D]}|(s|)r|*|j&q|j+rd d |j+D}t ,| |t -| t j.|d|d<t/| }~ ~ |S)Nz"triton_gpu.num-warp-groups-per-ctarerrrTTRITON_DISABLE_LINE_INFOrsm_90asm_nvptx64-nvidia-cudacSsg|]\}}|qSrr)rrr)rrrr1sz)CUDABackend.make_llir..ztriton_gpu.sharedr)0rT get_int_attrrrrrr%r&r'rrrrrrr%add_decompose_unsupported_conversionsrconvert add_scf_to_cfadd_index_to_llvmiradd_allocate_shared_memory add_to_llvmirradd_nvgpu_to_llvmadd_arith_to_llvmirrrrrllvmir add_di_scoper init_targets to_modulerYattach_datalayoutset_nvvm_reflect_ftzrm get_functionsis_declarationis_external_linkageset_nvvm_maxnregrzlink_extern_libsoptimize_module OPTIMIZE_O3rI)srcrrRrrQnum_warp_groupsrrrrrllvm_modprocrXtriplerr8retrrr make_llirs^                   zCUDABackend.make_llirc Cst|}d}|dkr dnd|}t|}t||||dg|jd}td|} t| dks0J| d |d <|d d |d }tjd d||tj d}tdd|}t j dddkrft dt ||S)Nrrrrznvptx-short-ptrFz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r$rrrFr?z\.version \d+\.\d+z .version r"z,\s*debug|debug,\s*rNVPTX_ENABLE_DUMPrrz // -----// NVPTX Dump //----- //)rTrYrtranslate_to_asmror2findalllensubr5r%r&r'print) rrrrrQr rrXr!namesrrrmake_ptx=s  zCUDABackend.make_ptxc Cs|td\}}tjdddd$}tjdddd}||||jd}tjd r0gnd g} |j r8gnd g} |d krAd nd} tjdddkrPddgng} |g| | d| d|| |jd|} z%t j | dd|dtj |jrt|jtj |jrt|jWn\t jy}zOt|j }|}Wdn1swYtj |jrt|j|jdkrd}n|jdtjkrd}nd|j}t|d|dd| d d}~wwt|d! }|}Wdn 1swYtj |rt|Wdn1swYWd|SWd|S1s7wY|S)"Nr<Fwz.ptx)deletemodesuffixrz.logz.orz -lineinfoz --fmad=falserarDISABLE_PTXAS_OPTrrz --opt-levelz-vz--gpu-name=sm_z-oT)check close_fdsr z!Internal Triton PTX codegen errorz`ptxas` raised SIGSEGVz`ptxas` failed with error code z `ptxas` stderr: z Repro command:   rZ)r;tempfileNamedTemporaryFilewriteflushrr%r&r'ror/rr)r-removeCalledProcessErrorr[r^ returncodesignalSIGSEGVr7r*)rrrrr<rSfsrcflogfbin line_infofmadr. opt_level ptxas_cmdelog_filelogerrorr`rrrr make_cubinSs               *( ((zCUDABackend.make_cubincs^fdd|d<fdd|d<fdd|d<fdd|d <fd d|d <dS) Ncs||Sr)rrrrRrrrrsz(CUDABackend.add_stages..rc||jSr)rrrMrNrrrttgircrOr)r"rrMrNrrrrPllircrOr)r*rrMrNrrrrPptxcrOr)rLrrMrNrrrrPrr)rstagesrRrrNr add_stagess zCUDABackend.add_stagescCst}|d|jS)Nr)r=r)rr:rrrrszCUDABackend.hash)rrr staticmethodrrrr rrrr rIr rrrrr"r*rLrU functools lru_cacher __classcell__rrrrrs,    4 <  ,r)&triton.backends.compilerrrtriton._C.libtritonrrrr dataclassesrrWtypingr r r r typesr r\r2r8r?r%r/pathlibrrrXrIr;r=rKrPrTrYrarcrrrrrs8      (