| 1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950515253545556575859606162636465666768697071727374757677787980818283848586878889909192939495969798991001011021031041051061071081091101111121131141151161171181191201211221231241251261271281291301311321331341351361371381391401411421431441451461471481491501511521531541551561571581591601611621631641651661671681691701711721731741751761771781791801811821831841851861871881891901911921931941951961971981992002012022032042052062072082092102112122132142152162172182192202212222232242252262272282292302312322332342352362372382392402412422432442452462472482492502512522532542552562572582592602612622632642652662672682692702712722732742752762772782792802812822832842852862872882892902912922932942952962972982993003013023033043053063073083093103113123133143153163173183193203213223233243253263273283293303313323333343353363373383393403413423433443453463473483493503513523533543553563573583593603613623633643653663673683693703713723733743753763773783793803813823833843853863873883893903913923933943953963973983994004014024034044054064074084094104114124134144154164174184194204214224234244254264274284294304314324334344354364374384394404414424434444454464474484494504514524534544554564574584594604614624634644654664674684694704714724734744754764774784794804814824834844854864874884894904914924934944954964974984995005015025035045055065075085095105115125135145155165175185195205215225235245255265275285295305315325335345355365375385395405415425435445455465475485495505515525535545555565575585595605615625635645655665675685695705715725735745755765775785795805815825835845855865875885895905915925935945955965975985996006016026036046056066076086096106116126136146156166176186196206216226236246256266276286296306316326336346356366376386396406416426436446456466476486496506516526536546556566576586596606616626636646656666676686696706716726736746756766776786796806816826836846856866876886896906916926936946956966976986997007017027037047057067077087097107117127137147157167177187197207217227237247257267277287297307317327337347357367377387397407417427437447457467477487497507517527537547557567577587597607617627637647657667677687697707717727737747757767777787797807817827837847857867877887897907917927937947957967977987998008018028038048058068078088098108118128138148158168178188198208218228238248258268278288298308318328338348358368378388398408418428438448458468478488498508518528538548558568578588598608618628638648658668678688698708718728738748758768778788798808818828838848858868878888898908918928938948958968978988999009019029039049059069079089099109119129139149159169179189199209219229239249259269279289299309319329339349359369379389399409419429439449459469479489499509519529539549559569579589599609619629639649659669679689699709719729739749759769779789799809819829839849859869879889899909919929939949959969979989991000100110021003100410051006100710081009101010111012101310141015101610171018101910201021102210231024102510261027102810291030103110321033103410351036103710381039104010411042104310441045104610471048104910501051105210531054105510561057105810591060106110621063106410651066106710681069107010711072107310741075107610771078107910801081108210831084108510861087108810891090109110921093109410951096109710981099110011011102110311041105110611071108110911101111111211131114111511161117111811191120112111221123112411251126112711281129113011311132113311341135113611371138113911401141114211431144114511461147114811491150115111521153115411551156115711581159116011611162116311641165116611671168116911701171117211731174117511761177117811791180118111821183118411851186118711881189119011911192119311941195119611971198119912001201120212031204120512061207120812091210121112121213121412151216121712181219122012211222122312241225122612271228122912301231123212331234123512361237123812391240124112421243124412451246124712481249125012511252125312541255125612571258125912601261126212631264126512661267126812691270127112721273127412751276127712781279128012811282128312841285128612871288128912901291129212931294129512961297129812991300130113021303130413051306130713081309131013111312131313141315131613171318131913201321132213231324132513261327132813291330133113321333133413351336133713381339134013411342134313441345134613471348134913501351135213531354135513561357135813591360136113621363136413651366136713681369137013711372137313741375137613771378137913801381138213831384138513861387138813891390139113921393139413951396139713981399140014011402140314041405140614071408140914101411141214131414141514161417141814191420142114221423142414251426142714281429143014311432143314341435143614371438143914401441144214431444144514461447144814491450145114521453145414551456145714581459146014611462146314641465146614671468146914701471147214731474147514761477147814791480148114821483148414851486148714881489149014911492149314941495149614971498149915001501150215031504150515061507150815091510151115121513151415151516151715181519152015211522152315241525152615271528152915301531153215331534153515361537153815391540154115421543154415451546154715481549155015511552155315541555155615571558155915601561156215631564156515661567156815691570157115721573157415751576157715781579158015811582158315841585158615871588158915901591159215931594159515961597159815991600160116021603160416051606160716081609161016111612161316141615161616171618161916201621162216231624162516261627162816291630163116321633163416351636163716381639164016411642164316441645164616471648164916501651165216531654165516561657165816591660166116621663166416651666166716681669167016711672167316741675167616771678167916801681168216831684168516861687168816891690169116921693169416951696169716981699170017011702170317041705170617071708170917101711171217131714171517161717171817191720172117221723172417251726172717281729173017311732173317341735173617371738173917401741174217431744174517461747174817491750175117521753175417551756175717581759176017611762176317641765176617671768176917701771177217731774177517761777177817791780178117821783178417851786178717881789179017911792179317941795179617971798179918001801180218031804180518061807180818091810181118121813181418151816181718181819182018211822182318241825182618271828182918301831183218331834183518361837183818391840184118421843184418451846184718481849185018511852185318541855185618571858185918601861186218631864186518661867186818691870187118721873187418751876187718781879188018811882188318841885188618871888188918901891189218931894189518961897189818991900190119021903190419051906190719081909191019111912191319141915191619171918191919201921192219231924192519261927192819291930193119321933193419351936193719381939194019411942194319441945194619471948194919501951195219531954195519561957195819591960196119621963196419651966196719681969197019711972197319741975197619771978197919801981198219831984198519861987198819891990199119921993199419951996199719981999200020012002200320042005200620072008200920102011201220132014201520162017201820192020202120222023202420252026202720282029203020312032203320342035203620372038203920402041204220432044204520462047204820492050205120522053205420552056205720582059206020612062206320642065206620672068206920702071207220732074207520762077207820792080208120822083208420852086208720882089209020912092209320942095209620972098209921002101210221032104210521062107210821092110211121122113211421152116211721182119212021212122212321242125212621272128212921302131213221332134213521362137213821392140214121422143214421452146214721482149215021512152215321542155215621572158215921602161216221632164216521662167216821692170217121722173217421752176217721782179218021812182218321842185218621872188218921902191219221932194219521962197219821992200220122022203220422052206220722082209221022112212221322142215221622172218221922202221222222232224222522262227222822292230223122322233223422352236223722382239224022412242224322442245224622472248224922502251225222532254225522562257225822592260226122622263226422652266226722682269227022712272227322742275227622772278227922802281228222832284228522862287228822892290229122922293229422952296229722982299230023012302230323042305230623072308230923102311231223132314231523162317231823192320232123222323232423252326232723282329233023312332233323342335233623372338233923402341234223432344234523462347234823492350235123522353235423552356235723582359236023612362236323642365236623672368236923702371237223732374237523762377237823792380238123822383238423852386238723882389239023912392239323942395239623972398239924002401240224032404240524062407240824092410241124122413241424152416241724182419242024212422242324242425242624272428242924302431243224332434243524362437243824392440244124422443244424452446244724482449245024512452245324542455245624572458245924602461246224632464246524662467246824692470247124722473247424752476247724782479248024812482248324842485248624872488248924902491249224932494249524962497249824992500250125022503250425052506250725082509251025112512251325142515251625172518251925202521252225232524252525262527252825292530253125322533253425352536253725382539254025412542254325442545254625472548254925502551255225532554255525562557255825592560256125622563256425652566256725682569257025712572257325742575257625772578257925802581258225832584258525862587258825892590259125922593259425952596259725982599260026012602260326042605260626072608260926102611261226132614261526162617261826192620262126222623262426252626262726282629263026312632263326342635263626372638263926402641264226432644264526462647264826492650265126522653265426552656265726582659266026612662266326642665266626672668266926702671267226732674267526762677267826792680268126822683268426852686268726882689269026912692269326942695269626972698269927002701270227032704270527062707270827092710271127122713271427152716271727182719272027212722272327242725272627272728272927302731273227332734273527362737273827392740274127422743274427452746274727482749275027512752275327542755275627572758275927602761276227632764276527662767276827692770277127722773277427752776277727782779278027812782278327842785278627872788278927902791279227932794279527962797279827992800280128022803280428052806280728082809281028112812281328142815281628172818281928202821282228232824282528262827282828292830283128322833283428352836283728382839284028412842284328442845284628472848284928502851285228532854285528562857285828592860286128622863286428652866286728682869287028712872287328742875287628772878287928802881288228832884288528862887288828892890289128922893289428952896289728982899290029012902290329042905290629072908290929102911291229132914291529162917291829192920292129222923292429252926292729282929293029312932293329342935293629372938293929402941294229432944294529462947294829492950295129522953295429552956295729582959296029612962296329642965296629672968296929702971297229732974297529762977297829792980298129822983298429852986298729882989299029912992299329942995299629972998299930003001300230033004300530063007300830093010301130123013301430153016301730183019302030213022302330243025302630273028302930303031303230333034303530363037303830393040304130423043304430453046304730483049305030513052305330543055305630573058305930603061306230633064306530663067306830693070307130723073307430753076307730783079308030813082308330843085308630873088308930903091309230933094309530963097309830993100310131023103310431053106310731083109311031113112311331143115311631173118311931203121312231233124312531263127312831293130313131323133313431353136313731383139314031413142314331443145314631473148314931503151315231533154315531563157315831593160316131623163316431653166316731683169317031713172317331743175317631773178317931803181318231833184318531863187318831893190319131923193319431953196319731983199320032013202320332043205320632073208320932103211321232133214321532163217321832193220322132223223322432253226322732283229323032313232323332343235323632373238323932403241324232433244324532463247324832493250325132523253325432553256325732583259326032613262326332643265326632673268326932703271327232733274327532763277327832793280328132823283328432853286328732883289329032913292329332943295329632973298329933003301330233033304330533063307330833093310331133123313331433153316331733183319332033213322332333243325332633273328332933303331333233333334333533363337333833393340334133423343334433453346334733483349335033513352335333543355335633573358335933603361336233633364336533663367336833693370337133723373337433753376337733783379338033813382338333843385338633873388338933903391339233933394339533963397339833993400340134023403340434053406340734083409341034113412341334143415341634173418341934203421342234233424342534263427342834293430343134323433343434353436343734383439344034413442344334443445344634473448344934503451345234533454345534563457345834593460346134623463346434653466346734683469347034713472347334743475347634773478347934803481348234833484348534863487348834893490349134923493349434953496349734983499350035013502350335043505350635073508350935103511351235133514351535163517351835193520352135223523352435253526352735283529353035313532353335343535353635373538353935403541354235433544354535463547354835493550355135523553355435553556355735583559356035613562356335643565356635673568356935703571357235733574357535763577357835793580358135823583358435853586358735883589359035913592359335943595359635973598359936003601360236033604360536063607360836093610361136123613361436153616361736183619362036213622362336243625362636273628362936303631363236333634363536363637363836393640364136423643364436453646364736483649365036513652365336543655365636573658365936603661366236633664366536663667366836693670367136723673367436753676367736783679368036813682368336843685368636873688368936903691369236933694369536963697369836993700370137023703370437053706370737083709371037113712371337143715371637173718371937203721372237233724372537263727372837293730373137323733373437353736373737383739374037413742374337443745374637473748374937503751375237533754375537563757375837593760376137623763376437653766376737683769377037713772377337743775377637773778377937803781378237833784378537863787378837893790379137923793379437953796379737983799380038013802380338043805380638073808380938103811381238133814381538163817381838193820382138223823382438253826382738283829383038313832383338343835383638373838383938403841384238433844384538463847384838493850385138523853385438553856385738583859386038613862386338643865386638673868386938703871387238733874387538763877387838793880388138823883388438853886388738883889389038913892389338943895389638973898389939003901390239033904390539063907390839093910391139123913391439153916391739183919392039213922392339243925392639273928392939303931393239333934393539363937393839393940394139423943394439453946394739483949395039513952395339543955395639573958395939603961396239633964396539663967396839693970397139723973397439753976397739783979398039813982398339843985398639873988398939903991399239933994399539963997399839994000400140024003400440054006400740084009401040114012401340144015401640174018401940204021402240234024402540264027402840294030403140324033403440354036403740384039404040414042404340444045404640474048404940504051405240534054405540564057405840594060406140624063406440654066406740684069407040714072407340744075407640774078407940804081408240834084408540864087408840894090409140924093409440954096409740984099410041014102410341044105410641074108410941104111411241134114411541164117411841194120412141224123412441254126412741284129413041314132413341344135413641374138413941404141414241434144414541464147414841494150415141524153415441554156415741584159416041614162416341644165416641674168416941704171417241734174417541764177417841794180418141824183418441854186418741884189419041914192419341944195419641974198419942004201420242034204420542064207420842094210421142124213421442154216421742184219422042214222422342244225422642274228422942304231423242334234423542364237423842394240424142424243424442454246424742484249425042514252425342544255425642574258425942604261426242634264426542664267426842694270427142724273427442754276427742784279428042814282428342844285428642874288428942904291429242934294429542964297429842994300430143024303430443054306430743084309431043114312431343144315431643174318431943204321432243234324432543264327432843294330433143324333433443354336433743384339434043414342434343444345434643474348434943504351435243534354435543564357435843594360436143624363436443654366436743684369437043714372437343744375437643774378437943804381438243834384438543864387438843894390439143924393439443954396439743984399440044014402440344044405440644074408440944104411441244134414441544164417441844194420442144224423442444254426442744284429443044314432443344344435443644374438443944404441444244434444444544464447444844494450445144524453445444554456445744584459446044614462446344644465446644674468446944704471447244734474447544764477447844794480448144824483448444854486448744884489449044914492449344944495449644974498449945004501450245034504450545064507450845094510451145124513451445154516451745184519452045214522452345244525452645274528452945304531453245334534453545364537453845394540454145424543454445454546454745484549455045514552455345544555455645574558455945604561456245634564456545664567456845694570457145724573457445754576457745784579458045814582458345844585458645874588458945904591459245934594459545964597459845994600460146024603460446054606460746084609461046114612461346144615461646174618461946204621462246234624462546264627462846294630463146324633463446354636463746384639464046414642464346444645464646474648464946504651465246534654465546564657465846594660466146624663466446654666466746684669467046714672467346744675467646774678467946804681468246834684468546864687468846894690469146924693469446954696469746984699470047014702470347044705470647074708470947104711471247134714471547164717471847194720472147224723472447254726472747284729473047314732473347344735473647374738473947404741474247434744474547464747474847494750475147524753475447554756475747584759476047614762476347644765476647674768476947704771477247734774477547764777477847794780478147824783478447854786478747884789479047914792479347944795479647974798479948004801480248034804480548064807480848094810481148124813481448154816481748184819482048214822482348244825482648274828482948304831483248334834483548364837483848394840484148424843484448454846484748484849485048514852485348544855485648574858485948604861486248634864486548664867486848694870487148724873487448754876487748784879488048814882488348844885488648874888488948904891489248934894489548964897489848994900490149024903490449054906490749084909491049114912491349144915491649174918491949204921492249234924492549264927492849294930493149324933493449354936493749384939494049414942494349444945494649474948494949504951495249534954495549564957495849594960496149624963496449654966496749684969497049714972497349744975497649774978497949804981498249834984498549864987498849894990499149924993499449954996499749984999500050015002500350045005500650075008500950105011501250135014501550165017501850195020502150225023502450255026502750285029503050315032503350345035503650375038503950405041504250435044504550465047504850495050505150525053505450555056505750585059506050615062506350645065506650675068506950705071507250735074507550765077507850795080508150825083508450855086508750885089509050915092509350945095509650975098509951005101510251035104510551065107510851095110511151125113511451155116511751185119512051215122512351245125512651275128512951305131513251335134513551365137513851395140514151425143514451455146514751485149515051515152515351545155515651575158515951605161516251635164516551665167516851695170517151725173517451755176517751785179518051815182518351845185518651875188518951905191519251935194519551965197519851995200520152025203520452055206520752085209521052115212521352145215521652175218521952205221522252235224522552265227522852295230523152325233523452355236523752385239524052415242524352445245524652475248524952505251525252535254525552565257525852595260526152625263526452655266526752685269527052715272527352745275527652775278527952805281528252835284528552865287528852895290529152925293529452955296529752985299530053015302530353045305530653075308530953105311531253135314531553165317531853195320532153225323532453255326532753285329533053315332533353345335533653375338533953405341534253435344534553465347534853495350535153525353535453555356535753585359536053615362536353645365536653675368536953705371537253735374537553765377537853795380538153825383538453855386538753885389539053915392539353945395539653975398539954005401540254035404540554065407540854095410541154125413541454155416541754185419542054215422542354245425542654275428542954305431543254335434543554365437543854395440544154425443544454455446544754485449545054515452545354545455545654575458545954605461546254635464546554665467546854695470547154725473547454755476547754785479548054815482548354845485548654875488548954905491549254935494549554965497549854995500550155025503550455055506550755085509551055115512551355145515551655175518551955205521552255235524552555265527552855295530553155325533553455355536553755385539554055415542554355445545554655475548554955505551555255535554555555565557555855595560556155625563556455655566556755685569557055715572557355745575557655775578557955805581558255835584558555865587558855895590559155925593559455955596559755985599560056015602560356045605560656075608560956105611561256135614561556165617561856195620562156225623562456255626562756285629563056315632563356345635563656375638563956405641564256435644564556465647564856495650565156525653565456555656565756585659566056615662566356645665566656675668566956705671567256735674567556765677567856795680568156825683568456855686568756885689569056915692569356945695569656975698569957005701570257035704570557065707570857095710571157125713571457155716571757185719572057215722572357245725572657275728572957305731573257335734573557365737573857395740574157425743574457455746574757485749575057515752575357545755575657575758575957605761576257635764576557665767576857695770577157725773577457755776577757785779578057815782578357845785578657875788578957905791579257935794579557965797579857995800580158025803580458055806580758085809581058115812581358145815581658175818581958205821582258235824582558265827582858295830583158325833583458355836583758385839584058415842584358445845584658475848584958505851585258535854585558565857585858595860586158625863586458655866586758685869587058715872587358745875587658775878587958805881588258835884588558865887588858895890589158925893589458955896589758985899590059015902590359045905590659075908590959105911591259135914591559165917591859195920592159225923592459255926592759285929593059315932593359345935593659375938593959405941594259435944594559465947594859495950595159525953595459555956595759585959596059615962596359645965596659675968596959705971597259735974597559765977597859795980598159825983598459855986598759885989599059915992599359945995599659975998599960006001600260036004600560066007600860096010601160126013601460156016601760186019602060216022602360246025602660276028602960306031603260336034603560366037603860396040604160426043604460456046604760486049605060516052605360546055605660576058605960606061606260636064606560666067606860696070607160726073607460756076607760786079608060816082608360846085608660876088608960906091609260936094609560966097609860996100610161026103610461056106610761086109611061116112611361146115611661176118611961206121612261236124612561266127612861296130613161326133613461356136613761386139614061416142614361446145614661476148614961506151615261536154615561566157615861596160616161626163616461656166616761686169617061716172617361746175617661776178617961806181618261836184618561866187618861896190619161926193619461956196619761986199620062016202620362046205620662076208620962106211621262136214621562166217621862196220622162226223622462256226622762286229623062316232623362346235623662376238623962406241624262436244624562466247624862496250625162526253625462556256625762586259626062616262626362646265626662676268626962706271627262736274627562766277627862796280628162826283628462856286628762886289629062916292629362946295629662976298629963006301630263036304630563066307630863096310631163126313631463156316631763186319632063216322632363246325632663276328632963306331633263336334633563366337633863396340634163426343634463456346634763486349635063516352635363546355635663576358635963606361636263636364636563666367636863696370637163726373637463756376637763786379638063816382638363846385638663876388638963906391639263936394639563966397639863996400640164026403640464056406640764086409641064116412641364146415641664176418641964206421642264236424642564266427642864296430643164326433643464356436643764386439644064416442644364446445644664476448644964506451645264536454645564566457645864596460646164626463646464656466646764686469647064716472647364746475647664776478647964806481648264836484648564866487648864896490649164926493649464956496649764986499650065016502650365046505650665076508650965106511651265136514651565166517651865196520652165226523652465256526652765286529653065316532653365346535653665376538653965406541654265436544654565466547654865496550655165526553655465556556655765586559656065616562656365646565656665676568656965706571657265736574657565766577657865796580658165826583658465856586658765886589659065916592659365946595659665976598659966006601660266036604660566066607660866096610661166126613661466156616661766186619662066216622662366246625662666276628662966306631663266336634663566366637663866396640664166426643664466456646664766486649665066516652665366546655665666576658665966606661666266636664666566666667666866696670667166726673667466756676667766786679668066816682668366846685668666876688668966906691669266936694669566966697669866996700670167026703670467056706670767086709671067116712671367146715671667176718671967206721672267236724672567266727672867296730673167326733673467356736673767386739674067416742674367446745674667476748674967506751675267536754675567566757675867596760676167626763676467656766676767686769677067716772677367746775677667776778677967806781678267836784678567866787678867896790679167926793679467956796679767986799680068016802680368046805680668076808680968106811681268136814681568166817681868196820682168226823682468256826682768286829683068316832683368346835683668376838683968406841684268436844684568466847684868496850685168526853685468556856685768586859686068616862686368646865686668676868686968706871687268736874687568766877687868796880688168826883688468856886688768886889689068916892689368946895689668976898689969006901690269036904690569066907690869096910691169126913691469156916691769186919692069216922692369246925692669276928692969306931693269336934693569366937693869396940694169426943694469456946694769486949695069516952695369546955695669576958695969606961696269636964696569666967696869696970697169726973697469756976697769786979698069816982698369846985698669876988698969906991699269936994699569966997699869997000700170027003700470057006700770087009701070117012701370147015701670177018701970207021702270237024702570267027702870297030703170327033703470357036703770387039704070417042704370447045704670477048704970507051705270537054705570567057705870597060706170627063706470657066706770687069707070717072707370747075707670777078707970807081708270837084708570867087708870897090709170927093709470957096709770987099710071017102710371047105710671077108710971107111711271137114711571167117711871197120712171227123712471257126712771287129713071317132713371347135713671377138713971407141714271437144714571467147714871497150715171527153715471557156715771587159716071617162716371647165716671677168716971707171717271737174717571767177717871797180718171827183718471857186718771887189719071917192719371947195719671977198719972007201720272037204720572067207720872097210721172127213721472157216721772187219722072217222722372247225722672277228722972307231723272337234723572367237723872397240724172427243724472457246724772487249725072517252725372547255725672577258725972607261726272637264726572667267726872697270727172727273727472757276727772787279728072817282728372847285728672877288728972907291729272937294729572967297729872997300730173027303730473057306730773087309731073117312731373147315731673177318731973207321732273237324732573267327732873297330733173327333733473357336733773387339734073417342734373447345734673477348734973507351735273537354735573567357735873597360736173627363736473657366736773687369737073717372737373747375737673777378737973807381738273837384738573867387738873897390739173927393739473957396739773987399740074017402740374047405740674077408740974107411741274137414741574167417741874197420742174227423742474257426742774287429743074317432743374347435743674377438743974407441744274437444744574467447744874497450745174527453745474557456745774587459746074617462746374647465746674677468746974707471747274737474747574767477747874797480748174827483748474857486748774887489749074917492749374947495749674977498749975007501750275037504750575067507750875097510751175127513751475157516751775187519752075217522752375247525752675277528752975307531753275337534753575367537753875397540754175427543754475457546754775487549755075517552755375547555755675577558755975607561756275637564756575667567756875697570757175727573757475757576757775787579758075817582758375847585758675877588758975907591759275937594759575967597759875997600760176027603760476057606760776087609761076117612761376147615761676177618761976207621762276237624762576267627762876297630763176327633763476357636763776387639764076417642764376447645764676477648764976507651765276537654765576567657765876597660766176627663766476657666766776687669767076717672767376747675767676777678767976807681768276837684768576867687768876897690769176927693769476957696769776987699770077017702770377047705770677077708770977107711771277137714771577167717771877197720772177227723772477257726772777287729773077317732773377347735773677377738773977407741774277437744774577467747774877497750775177527753775477557756775777587759776077617762776377647765776677677768776977707771777277737774777577767777777877797780778177827783778477857786778777887789779077917792779377947795779677977798779978007801780278037804780578067807780878097810781178127813781478157816781778187819782078217822782378247825782678277828782978307831783278337834783578367837783878397840784178427843784478457846784778487849785078517852785378547855785678577858785978607861786278637864786578667867786878697870787178727873787478757876787778787879788078817882788378847885788678877888788978907891789278937894789578967897789878997900790179027903790479057906790779087909791079117912791379147915791679177918791979207921792279237924792579267927792879297930793179327933793479357936793779387939794079417942794379447945794679477948794979507951795279537954795579567957795879597960796179627963796479657966796779687969797079717972797379747975797679777978797979807981798279837984798579867987798879897990799179927993799479957996799779987999800080018002800380048005800680078008800980108011801280138014801580168017801880198020802180228023802480258026802780288029803080318032803380348035803680378038803980408041804280438044804580468047804880498050805180528053805480558056805780588059806080618062806380648065806680678068806980708071807280738074807580768077807880798080808180828083808480858086808780888089809080918092809380948095809680978098809981008101810281038104810581068107810881098110811181128113811481158116811781188119812081218122812381248125812681278128812981308131813281338134813581368137813881398140814181428143814481458146814781488149815081518152815381548155815681578158815981608161816281638164816581668167816881698170817181728173817481758176817781788179818081818182818381848185818681878188818981908191819281938194819581968197819881998200820182028203820482058206820782088209821082118212821382148215821682178218821982208221822282238224822582268227822882298230823182328233823482358236823782388239824082418242824382448245824682478248824982508251825282538254825582568257825882598260826182628263826482658266826782688269827082718272827382748275827682778278827982808281828282838284828582868287828882898290829182928293829482958296829782988299830083018302830383048305830683078308830983108311831283138314831583168317831883198320832183228323832483258326832783288329833083318332833383348335833683378338833983408341834283438344834583468347834883498350835183528353835483558356835783588359836083618362836383648365836683678368836983708371837283738374837583768377837883798380838183828383838483858386838783888389839083918392839383948395839683978398839984008401840284038404840584068407840884098410841184128413841484158416841784188419842084218422842384248425842684278428842984308431843284338434843584368437843884398440844184428443844484458446844784488449845084518452845384548455845684578458845984608461846284638464846584668467846884698470847184728473847484758476847784788479848084818482848384848485848684878488848984908491849284938494849584968497849884998500850185028503850485058506850785088509851085118512851385148515851685178518851985208521852285238524852585268527852885298530853185328533853485358536853785388539854085418542854385448545854685478548854985508551855285538554855585568557855885598560856185628563856485658566856785688569857085718572857385748575857685778578857985808581858285838584858585868587858885898590859185928593859485958596859785988599860086018602860386048605860686078608860986108611861286138614861586168617861886198620862186228623862486258626862786288629863086318632863386348635863686378638863986408641864286438644864586468647864886498650865186528653865486558656865786588659866086618662866386648665866686678668866986708671867286738674867586768677867886798680868186828683868486858686868786888689869086918692869386948695869686978698869987008701870287038704870587068707870887098710871187128713871487158716871787188719872087218722872387248725872687278728872987308731873287338734873587368737873887398740874187428743874487458746874787488749875087518752875387548755875687578758875987608761876287638764876587668767876887698770877187728773877487758776877787788779878087818782878387848785878687878788878987908791879287938794879587968797879887998800880188028803880488058806880788088809881088118812881388148815881688178818881988208821882288238824882588268827882888298830883188328833883488358836883788388839884088418842884388448845884688478848884988508851885288538854885588568857885888598860886188628863886488658866886788688869887088718872887388748875887688778878887988808881888288838884888588868887888888898890889188928893889488958896889788988899890089018902890389048905890689078908890989108911891289138914891589168917891889198920892189228923892489258926892789288929893089318932893389348935893689378938893989408941894289438944894589468947894889498950895189528953895489558956895789588959896089618962896389648965896689678968896989708971897289738974897589768977897889798980898189828983898489858986898789888989899089918992899389948995899689978998899990009001900290039004900590069007900890099010901190129013901490159016901790189019902090219022902390249025902690279028902990309031903290339034903590369037903890399040904190429043904490459046904790489049905090519052905390549055905690579058905990609061906290639064906590669067906890699070907190729073907490759076907790789079908090819082908390849085908690879088908990909091909290939094909590969097909890999100910191029103910491059106910791089109911091119112911391149115911691179118911991209121912291239124912591269127912891299130913191329133913491359136913791389139914091419142914391449145914691479148914991509151915291539154915591569157915891599160916191629163916491659166916791689169917091719172917391749175917691779178917991809181918291839184918591869187918891899190919191929193919491959196919791989199920092019202920392049205920692079208920992109211921292139214921592169217921892199220922192229223922492259226922792289229923092319232923392349235923692379238923992409241924292439244924592469247924892499250925192529253925492559256925792589259926092619262926392649265926692679268926992709271927292739274927592769277927892799280928192829283928492859286928792889289929092919292929392949295929692979298929993009301930293039304930593069307930893099310931193129313931493159316931793189319932093219322932393249325932693279328932993309331933293339334933593369337933893399340934193429343934493459346934793489349935093519352935393549355935693579358935993609361936293639364936593669367936893699370937193729373937493759376937793789379938093819382938393849385938693879388938993909391939293939394939593969397939893999400940194029403940494059406940794089409941094119412941394149415941694179418941994209421942294239424942594269427942894299430943194329433943494359436943794389439944094419442944394449445944694479448944994509451945294539454945594569457945894599460946194629463946494659466946794689469947094719472947394749475947694779478947994809481948294839484948594869487948894899490949194929493949494959496949794989499950095019502950395049505950695079508950995109511951295139514951595169517951895199520952195229523952495259526952795289529953095319532953395349535953695379538953995409541954295439544954595469547954895499550955195529553955495559556955795589559956095619562956395649565956695679568956995709571957295739574957595769577957895799580958195829583958495859586958795889589959095919592959395949595959695979598959996009601960296039604960596069607960896099610961196129613961496159616961796189619962096219622962396249625962696279628962996309631963296339634963596369637963896399640964196429643964496459646964796489649965096519652965396549655965696579658965996609661966296639664966596669667966896699670967196729673967496759676967796789679968096819682968396849685968696879688968996909691969296939694969596969697969896999700970197029703970497059706970797089709971097119712971397149715971697179718971997209721972297239724972597269727972897299730973197329733973497359736973797389739974097419742974397449745974697479748974997509751975297539754975597569757975897599760976197629763976497659766976797689769977097719772977397749775977697779778977997809781978297839784978597869787978897899790979197929793979497959796979797989799980098019802980398049805980698079808980998109811981298139814981598169817981898199820982198229823982498259826982798289829983098319832983398349835983698379838983998409841984298439844984598469847984898499850985198529853985498559856985798589859986098619862986398649865986698679868986998709871987298739874987598769877987898799880988198829883988498859886988798889889989098919892989398949895989698979898989999009901990299039904990599069907990899099910991199129913991499159916991799189919992099219922992399249925992699279928992999309931993299339934993599369937993899399940994199429943994499459946994799489949995099519952995399549955995699579958995999609961996299639964996599669967996899699970997199729973997499759976997799789979998099819982998399849985998699879988998999909991999299939994999599969997999899991000010001100021000310004100051000610007100081000910010100111001210013100141001510016100171001810019100201002110022100231002410025100261002710028100291003010031100321003310034100351003610037100381003910040100411004210043100441004510046100471004810049100501005110052100531005410055100561005710058100591006010061100621006310064100651006610067100681006910070100711007210073100741007510076100771007810079100801008110082100831008410085100861008710088100891009010091100921009310094100951009610097100981009910100101011010210103101041010510106101071010810109101101011110112101131011410115101161011710118101191012010121101221012310124101251012610127101281012910130101311013210133101341013510136101371013810139101401014110142101431014410145101461014710148101491015010151101521015310154101551015610157101581015910160101611016210163101641016510166101671016810169101701017110172101731017410175101761017710178101791018010181101821018310184101851018610187101881018910190101911019210193101941019510196101971019810199102001020110202102031020410205102061020710208102091021010211102121021310214102151021610217102181021910220102211022210223102241022510226102271022810229102301023110232102331023410235102361023710238102391024010241102421024310244102451024610247102481024910250102511025210253102541025510256102571025810259102601026110262102631026410265102661026710268102691027010271102721027310274102751027610277102781027910280102811028210283102841028510286102871028810289102901029110292102931029410295102961029710298102991030010301103021030310304103051030610307103081030910310103111031210313103141031510316103171031810319103201032110322103231032410325103261032710328103291033010331103321033310334103351033610337103381033910340103411034210343103441034510346103471034810349103501035110352103531035410355103561035710358103591036010361103621036310364103651036610367103681036910370103711037210373103741037510376103771037810379103801038110382103831038410385103861038710388103891039010391103921039310394103951039610397103981039910400104011040210403104041040510406104071040810409104101041110412104131041410415104161041710418104191042010421104221042310424104251042610427104281042910430104311043210433104341043510436104371043810439104401044110442104431044410445104461044710448104491045010451104521045310454104551045610457104581045910460104611046210463104641046510466104671046810469104701047110472104731047410475104761047710478104791048010481104821048310484104851048610487104881048910490104911049210493104941049510496104971049810499105001050110502105031050410505105061050710508105091051010511105121051310514105151051610517105181051910520105211052210523105241052510526105271052810529105301053110532105331053410535105361053710538105391054010541105421054310544105451054610547105481054910550105511055210553105541055510556105571055810559105601056110562105631056410565105661056710568105691057010571105721057310574105751057610577105781057910580105811058210583105841058510586105871058810589105901059110592105931059410595105961059710598105991060010601106021060310604106051060610607106081060910610106111061210613106141061510616106171061810619106201062110622106231062410625106261062710628106291063010631106321063310634106351063610637106381063910640106411064210643106441064510646106471064810649106501065110652106531065410655106561065710658106591066010661106621066310664106651066610667106681066910670106711067210673106741067510676106771067810679106801068110682106831068410685106861068710688106891069010691106921069310694106951069610697106981069910700107011070210703107041070510706107071070810709107101071110712107131071410715107161071710718107191072010721107221072310724107251072610727107281072910730107311073210733107341073510736107371073810739107401074110742107431074410745107461074710748107491075010751107521075310754107551075610757107581075910760107611076210763107641076510766107671076810769107701077110772107731077410775107761077710778107791078010781107821078310784107851078610787107881078910790107911079210793107941079510796107971079810799108001080110802108031080410805108061080710808108091081010811108121081310814108151081610817108181081910820108211082210823108241082510826108271082810829108301083110832108331083410835108361083710838108391084010841108421084310844108451084610847108481084910850108511085210853108541085510856108571085810859108601086110862108631086410865108661086710868108691087010871108721087310874108751087610877108781087910880108811088210883108841088510886108871088810889108901089110892108931089410895108961089710898108991090010901109021090310904109051090610907109081090910910109111091210913109141091510916109171091810919109201092110922109231092410925109261092710928109291093010931109321093310934109351093610937109381093910940109411094210943109441094510946109471094810949109501095110952109531095410955109561095710958109591096010961109621096310964109651096610967109681096910970109711097210973109741097510976109771097810979109801098110982109831098410985109861098710988109891099010991109921099310994109951099610997109981099911000110011100211003110041100511006110071100811009110101101111012110131101411015110161101711018110191102011021110221102311024110251102611027110281102911030110311103211033110341103511036110371103811039110401104111042110431104411045110461104711048110491105011051110521105311054110551105611057110581105911060110611106211063110641106511066110671106811069110701107111072110731107411075110761107711078110791108011081110821108311084110851108611087110881108911090110911109211093110941109511096110971109811099111001110111102111031110411105111061110711108111091111011111111121111311114111151111611117111181111911120111211112211123111241112511126111271112811129111301113111132111331113411135111361113711138111391114011141111421114311144111451114611147111481114911150111511115211153111541115511156111571115811159111601116111162111631116411165111661116711168111691117011171111721117311174111751117611177111781117911180111811118211183111841118511186111871118811189111901119111192111931119411195111961119711198111991120011201112021120311204112051120611207112081120911210112111121211213112141121511216112171121811219112201122111222112231122411225112261122711228112291123011231112321123311234112351123611237112381123911240112411124211243112441124511246112471124811249112501125111252112531125411255112561125711258112591126011261112621126311264112651126611267112681126911270112711127211273112741127511276112771127811279112801128111282112831128411285112861128711288112891129011291112921129311294112951129611297112981129911300113011130211303113041130511306113071130811309113101131111312113131131411315113161131711318113191132011321113221132311324113251132611327113281132911330113311133211333113341133511336113371133811339113401134111342113431134411345113461134711348113491135011351113521135311354113551135611357113581135911360113611136211363113641136511366113671136811369113701137111372113731137411375113761137711378113791138011381113821138311384113851138611387113881138911390113911139211393113941139511396113971139811399114001140111402114031140411405114061140711408114091141011411114121141311414114151141611417114181141911420114211142211423114241142511426114271142811429114301143111432114331143411435114361143711438114391144011441114421144311444114451144611447114481144911450114511145211453114541145511456114571145811459114601146111462114631146411465114661146711468114691147011471114721147311474114751147611477114781147911480114811148211483114841148511486114871148811489114901149111492114931149411495114961149711498114991150011501115021150311504115051150611507115081150911510115111151211513115141151511516115171151811519115201152111522115231152411525115261152711528115291153011531115321153311534115351153611537115381153911540115411154211543115441154511546115471154811549115501155111552115531155411555115561155711558115591156011561115621156311564115651156611567115681156911570115711157211573115741157511576115771157811579115801158111582115831158411585115861158711588115891159011591115921159311594115951159611597115981159911600116011160211603116041160511606116071160811609116101161111612116131161411615116161161711618116191162011621116221162311624116251162611627116281162911630116311163211633116341163511636116371163811639116401164111642116431164411645116461164711648116491165011651116521165311654116551165611657116581165911660116611166211663116641166511666116671166811669116701167111672116731167411675116761167711678116791168011681116821168311684116851168611687116881168911690116911169211693116941169511696116971169811699117001170111702117031170411705117061170711708117091171011711117121171311714117151171611717117181171911720117211172211723117241172511726117271172811729117301173111732117331173411735117361173711738117391174011741117421174311744117451174611747117481174911750117511175211753117541175511756117571175811759117601176111762117631176411765117661176711768117691177011771117721177311774117751177611777117781177911780117811178211783117841178511786117871178811789117901179111792117931179411795117961179711798117991180011801118021180311804118051180611807118081180911810118111181211813118141181511816118171181811819118201182111822118231182411825118261182711828118291183011831118321183311834118351183611837118381183911840118411184211843118441184511846118471184811849118501185111852118531185411855118561185711858118591186011861118621186311864118651186611867118681186911870118711187211873118741187511876118771187811879118801188111882118831188411885118861188711888118891189011891118921189311894118951189611897118981189911900119011190211903119041190511906119071190811909119101191111912119131191411915119161191711918119191192011921119221192311924119251192611927119281192911930119311193211933119341193511936119371193811939119401194111942119431194411945119461194711948119491195011951119521195311954119551195611957119581195911960119611196211963119641196511966119671196811969119701197111972119731197411975119761197711978119791198011981119821198311984119851198611987119881198911990119911199211993119941199511996119971199811999120001200112002120031200412005120061200712008120091201012011120121201312014120151201612017120181201912020120211202212023120241202512026120271202812029120301203112032120331203412035120361203712038120391204012041120421204312044120451204612047120481204912050120511205212053120541205512056120571205812059120601206112062120631206412065120661206712068120691207012071120721207312074120751207612077120781207912080120811208212083120841208512086120871208812089120901209112092120931209412095120961209712098120991210012101121021210312104121051210612107121081210912110121111211212113121141211512116121171211812119121201212112122121231212412125121261212712128121291213012131121321213312134121351213612137121381213912140121411214212143121441214512146121471214812149121501215112152121531215412155121561215712158121591216012161121621216312164121651216612167121681216912170121711217212173121741217512176121771217812179121801218112182121831218412185121861218712188121891219012191121921219312194121951219612197121981219912200122011220212203122041220512206122071220812209122101221112212122131221412215122161221712218122191222012221122221222312224122251222612227122281222912230122311223212233122341223512236122371223812239122401224112242122431224412245122461224712248122491225012251122521225312254122551225612257122581225912260122611226212263122641226512266122671226812269122701227112272122731227412275122761227712278122791228012281122821228312284122851228612287122881228912290122911229212293122941229512296122971229812299123001230112302123031230412305123061230712308123091231012311123121231312314123151231612317123181231912320123211232212323123241232512326123271232812329123301233112332123331233412335123361233712338123391234012341123421234312344123451234612347123481234912350123511235212353123541235512356123571235812359123601236112362123631236412365123661236712368123691237012371123721237312374123751237612377123781237912380123811238212383123841238512386123871238812389123901239112392123931239412395123961239712398123991240012401124021240312404124051240612407124081240912410124111241212413124141241512416124171241812419124201242112422124231242412425124261242712428124291243012431124321243312434124351243612437124381243912440124411244212443124441244512446124471244812449124501245112452124531245412455124561245712458124591246012461124621246312464124651246612467124681246912470124711247212473124741247512476124771247812479124801248112482124831248412485124861248712488124891249012491124921249312494124951249612497124981249912500125011250212503125041250512506125071250812509125101251112512125131251412515125161251712518125191252012521125221252312524125251252612527125281252912530125311253212533125341253512536125371253812539125401254112542125431254412545125461254712548125491255012551125521255312554125551255612557125581255912560125611256212563125641256512566125671256812569125701257112572125731257412575125761257712578125791258012581125821258312584125851258612587125881258912590125911259212593125941259512596125971259812599126001260112602126031260412605126061260712608126091261012611126121261312614126151261612617126181261912620126211262212623126241262512626126271262812629126301263112632126331263412635126361263712638126391264012641126421264312644126451264612647126481264912650126511265212653126541265512656126571265812659126601266112662126631266412665126661266712668126691267012671126721267312674126751267612677126781267912680126811268212683126841268512686126871268812689126901269112692126931269412695126961269712698126991270012701127021270312704127051270612707127081270912710127111271212713127141271512716127171271812719127201272112722127231272412725127261272712728127291273012731127321273312734127351273612737127381273912740127411274212743127441274512746127471274812749127501275112752127531275412755127561275712758127591276012761127621276312764127651276612767127681276912770127711277212773127741277512776127771277812779127801278112782127831278412785127861278712788127891279012791127921279312794127951279612797127981279912800128011280212803128041280512806128071280812809128101281112812128131281412815128161281712818128191282012821128221282312824128251282612827128281282912830128311283212833128341283512836128371283812839128401284112842128431284412845128461284712848128491285012851128521285312854128551285612857128581285912860128611286212863128641286512866128671286812869128701287112872128731287412875128761287712878128791288012881128821288312884128851288612887128881288912890128911289212893128941289512896128971289812899129001290112902129031290412905129061290712908129091291012911129121291312914129151291612917129181291912920129211292212923129241292512926129271292812929129301293112932129331293412935129361293712938129391294012941129421294312944129451294612947129481294912950129511295212953129541295512956129571295812959129601296112962129631296412965129661296712968129691297012971129721297312974129751297612977129781297912980129811298212983129841298512986129871298812989129901299112992129931299412995129961299712998129991300013001130021300313004130051300613007130081300913010130111301213013130141301513016130171301813019130201302113022130231302413025130261302713028130291303013031130321303313034130351303613037130381303913040130411304213043130441304513046130471304813049130501305113052130531305413055130561305713058130591306013061130621306313064130651306613067130681306913070130711307213073130741307513076130771307813079130801308113082130831308413085130861308713088130891309013091130921309313094130951309613097130981309913100131011310213103131041310513106131071310813109131101311113112131131311413115131161311713118131191312013121131221312313124131251312613127131281312913130131311313213133131341313513136131371313813139131401314113142131431314413145131461314713148131491315013151131521315313154131551315613157131581315913160131611316213163131641316513166131671316813169131701317113172131731317413175131761317713178131791318013181131821318313184131851318613187131881318913190131911319213193131941319513196131971319813199132001320113202132031320413205132061320713208132091321013211132121321313214132151321613217132181321913220132211322213223132241322513226132271322813229132301323113232132331323413235132361323713238132391324013241132421324313244132451324613247132481324913250132511325213253132541325513256132571325813259132601326113262132631326413265132661326713268132691327013271132721327313274132751327613277132781327913280132811328213283132841328513286132871328813289132901329113292132931329413295132961329713298132991330013301133021330313304133051330613307133081330913310133111331213313133141331513316133171331813319133201332113322133231332413325133261332713328133291333013331133321333313334133351333613337133381333913340133411334213343133441334513346133471334813349133501335113352133531335413355133561335713358133591336013361133621336313364133651336613367133681336913370133711337213373133741337513376133771337813379133801338113382133831338413385133861338713388133891339013391133921339313394133951339613397133981339913400134011340213403134041340513406134071340813409134101341113412134131341413415134161341713418134191342013421134221342313424134251342613427134281342913430134311343213433134341343513436134371343813439134401344113442134431344413445134461344713448134491345013451134521345313454134551345613457134581345913460134611346213463134641346513466134671346813469134701347113472134731347413475134761347713478134791348013481134821348313484134851348613487134881348913490134911349213493134941349513496134971349813499135001350113502135031350413505135061350713508135091351013511135121351313514135151351613517135181351913520135211352213523135241352513526135271352813529135301353113532135331353413535135361353713538135391354013541135421354313544135451354613547135481354913550135511355213553135541355513556135571355813559135601356113562135631356413565135661356713568135691357013571135721357313574135751357613577135781357913580135811358213583135841358513586135871358813589135901359113592135931359413595135961359713598135991360013601136021360313604136051360613607136081360913610136111361213613136141361513616136171361813619136201362113622136231362413625136261362713628136291363013631136321363313634136351363613637136381363913640136411364213643136441364513646136471364813649136501365113652136531365413655136561365713658136591366013661136621366313664136651366613667136681366913670136711367213673136741367513676136771367813679136801368113682136831368413685136861368713688136891369013691136921369313694136951369613697136981369913700137011370213703137041370513706137071370813709137101371113712137131371413715137161371713718137191372013721137221372313724137251372613727137281372913730137311373213733137341373513736137371373813739137401374113742137431374413745137461374713748137491375013751137521375313754137551375613757137581375913760137611376213763137641376513766137671376813769137701377113772137731377413775137761377713778137791378013781137821378313784137851378613787137881378913790137911379213793137941379513796137971379813799138001380113802138031380413805138061380713808138091381013811138121381313814138151381613817138181381913820138211382213823138241382513826138271382813829138301383113832138331383413835138361383713838138391384013841138421384313844138451384613847138481384913850138511385213853138541385513856138571385813859138601386113862138631386413865138661386713868138691387013871138721387313874138751387613877138781387913880138811388213883138841388513886138871388813889138901389113892138931389413895138961389713898138991390013901139021390313904139051390613907139081390913910139111391213913139141391513916139171391813919139201392113922139231392413925139261392713928139291393013931139321393313934139351393613937139381393913940139411394213943139441394513946139471394813949139501395113952139531395413955139561395713958139591396013961139621396313964139651396613967139681396913970139711397213973139741397513976139771397813979139801398113982139831398413985139861398713988139891399013991139921399313994139951399613997139981399914000140011400214003140041400514006140071400814009140101401114012140131401414015140161401714018140191402014021140221402314024140251402614027140281402914030140311403214033140341403514036140371403814039140401404114042140431404414045140461404714048140491405014051140521405314054140551405614057140581405914060140611406214063140641406514066140671406814069140701407114072140731407414075140761407714078140791408014081140821408314084140851408614087140881408914090140911409214093140941409514096140971409814099141001410114102141031410414105141061410714108141091411014111141121411314114141151411614117141181411914120141211412214123141241412514126141271412814129141301413114132141331413414135141361413714138141391414014141141421414314144141451414614147141481414914150141511415214153141541415514156141571415814159141601416114162141631416414165141661416714168141691417014171141721417314174141751417614177141781417914180141811418214183141841418514186141871418814189141901419114192141931419414195141961419714198141991420014201142021420314204142051420614207142081420914210142111421214213142141421514216142171421814219142201422114222142231422414225142261422714228142291423014231142321423314234142351423614237142381423914240142411424214243142441424514246142471424814249142501425114252142531425414255142561425714258142591426014261142621426314264142651426614267142681426914270142711427214273142741427514276142771427814279142801428114282142831428414285142861428714288142891429014291142921429314294142951429614297142981429914300143011430214303143041430514306143071430814309143101431114312143131431414315143161431714318143191432014321143221432314324143251432614327143281432914330143311433214333143341433514336143371433814339143401434114342143431434414345143461434714348143491435014351143521435314354143551435614357143581435914360143611436214363143641436514366143671436814369143701437114372143731437414375143761437714378143791438014381143821438314384143851438614387143881438914390143911439214393143941439514396143971439814399144001440114402144031440414405144061440714408144091441014411144121441314414144151441614417144181441914420144211442214423144241442514426144271442814429144301443114432144331443414435144361443714438144391444014441144421444314444144451444614447144481444914450144511445214453144541445514456144571445814459144601446114462144631446414465144661446714468144691447014471144721447314474144751447614477144781447914480144811448214483144841448514486144871448814489144901449114492144931449414495144961449714498144991450014501145021450314504145051450614507145081450914510145111451214513145141451514516145171451814519145201452114522145231452414525145261452714528145291453014531145321453314534145351453614537145381453914540145411454214543145441454514546145471454814549145501455114552145531455414555145561455714558145591456014561145621456314564145651456614567145681456914570145711457214573145741457514576145771457814579145801458114582145831458414585145861458714588145891459014591145921459314594145951459614597145981459914600146011460214603146041460514606146071460814609146101461114612146131461414615146161461714618146191462014621146221462314624146251462614627146281462914630146311463214633146341463514636146371463814639146401464114642146431464414645146461464714648146491465014651146521465314654146551465614657146581465914660146611466214663146641466514666146671466814669146701467114672146731467414675146761467714678146791468014681146821468314684146851468614687146881468914690146911469214693146941469514696146971469814699147001470114702147031470414705147061470714708147091471014711147121471314714147151471614717147181471914720147211472214723147241472514726147271472814729147301473114732147331473414735147361473714738147391474014741147421474314744147451474614747147481474914750147511475214753147541475514756147571475814759147601476114762147631476414765147661476714768147691477014771147721477314774147751477614777147781477914780147811478214783147841478514786147871478814789147901479114792147931479414795147961479714798147991480014801148021480314804148051480614807148081480914810148111481214813148141481514816148171481814819148201482114822148231482414825148261482714828148291483014831148321483314834148351483614837148381483914840148411484214843148441484514846148471484814849148501485114852148531485414855148561485714858148591486014861148621486314864148651486614867148681486914870148711487214873148741487514876148771487814879148801488114882148831488414885148861488714888148891489014891148921489314894148951489614897148981489914900149011490214903149041490514906149071490814909149101491114912149131491414915149161491714918149191492014921149221492314924149251492614927149281492914930149311493214933149341493514936149371493814939149401494114942149431494414945149461494714948149491495014951149521495314954149551495614957149581495914960149611496214963149641496514966149671496814969149701497114972149731497414975149761497714978149791498014981149821498314984149851498614987149881498914990149911499214993149941499514996149971499814999150001500115002150031500415005150061500715008150091501015011150121501315014150151501615017150181501915020150211502215023150241502515026150271502815029150301503115032150331503415035150361503715038150391504015041150421504315044150451504615047150481504915050150511505215053150541505515056150571505815059150601506115062150631506415065150661506715068150691507015071150721507315074150751507615077150781507915080150811508215083150841508515086150871508815089150901509115092150931509415095150961509715098150991510015101151021510315104151051510615107151081510915110151111511215113151141511515116151171511815119151201512115122151231512415125151261512715128151291513015131151321513315134151351513615137151381513915140151411514215143151441514515146151471514815149151501515115152151531515415155151561515715158151591516015161151621516315164151651516615167151681516915170151711517215173151741517515176151771517815179151801518115182151831518415185151861518715188151891519015191151921519315194151951519615197151981519915200152011520215203152041520515206152071520815209152101521115212152131521415215152161521715218152191522015221152221522315224152251522615227152281522915230152311523215233152341523515236152371523815239152401524115242152431524415245152461524715248152491525015251152521525315254152551525615257152581525915260152611526215263152641526515266152671526815269152701527115272152731527415275152761527715278152791528015281152821528315284152851528615287152881528915290152911529215293152941529515296152971529815299153001530115302153031530415305153061530715308153091531015311153121531315314153151531615317153181531915320153211532215323153241532515326153271532815329153301533115332153331533415335153361533715338153391534015341153421534315344153451534615347153481534915350153511535215353153541535515356153571535815359153601536115362153631536415365153661536715368153691537015371153721537315374153751537615377153781537915380153811538215383153841538515386153871538815389153901539115392153931539415395153961539715398153991540015401154021540315404154051540615407154081540915410154111541215413154141541515416154171541815419154201542115422154231542415425154261542715428154291543015431154321543315434154351543615437154381543915440154411544215443154441544515446154471544815449154501545115452154531545415455154561545715458154591546015461154621546315464154651546615467154681546915470154711547215473154741547515476154771547815479154801548115482154831548415485154861548715488154891549015491154921549315494154951549615497154981549915500155011550215503155041550515506155071550815509155101551115512155131551415515155161551715518155191552015521155221552315524155251552615527155281552915530155311553215533155341553515536155371553815539155401554115542155431554415545155461554715548155491555015551155521555315554155551555615557155581555915560155611556215563155641556515566155671556815569155701557115572155731557415575155761557715578155791558015581155821558315584155851558615587155881558915590155911559215593155941559515596155971559815599156001560115602156031560415605156061560715608156091561015611156121561315614156151561615617156181561915620156211562215623156241562515626156271562815629156301563115632156331563415635156361563715638156391564015641156421564315644156451564615647156481564915650156511565215653156541565515656156571565815659156601566115662156631566415665156661566715668156691567015671156721567315674156751567615677156781567915680156811568215683156841568515686156871568815689156901569115692156931569415695156961569715698156991570015701157021570315704157051570615707157081570915710157111571215713157141571515716157171571815719157201572115722157231572415725157261572715728157291573015731157321573315734157351573615737157381573915740157411574215743157441574515746157471574815749157501575115752157531575415755157561575715758157591576015761157621576315764157651576615767157681576915770157711577215773157741577515776157771577815779157801578115782157831578415785157861578715788157891579015791157921579315794157951579615797157981579915800158011580215803158041580515806158071580815809158101581115812158131581415815158161581715818158191582015821158221582315824158251582615827158281582915830158311583215833158341583515836158371583815839158401584115842158431584415845158461584715848158491585015851158521585315854158551585615857158581585915860158611586215863158641586515866158671586815869158701587115872158731587415875158761587715878158791588015881158821588315884158851588615887158881588915890158911589215893158941589515896158971589815899159001590115902159031590415905159061590715908159091591015911159121591315914159151591615917159181591915920159211592215923159241592515926159271592815929159301593115932159331593415935159361593715938159391594015941159421594315944159451594615947159481594915950159511595215953159541595515956159571595815959159601596115962159631596415965159661596715968159691597015971159721597315974159751597615977159781597915980159811598215983159841598515986159871598815989159901599115992159931599415995159961599715998159991600016001160021600316004160051600616007160081600916010160111601216013160141601516016160171601816019160201602116022160231602416025160261602716028160291603016031160321603316034160351603616037160381603916040160411604216043160441604516046160471604816049160501605116052160531605416055160561605716058160591606016061160621606316064160651606616067160681606916070160711607216073160741607516076160771607816079160801608116082160831608416085160861608716088160891609016091160921609316094160951609616097160981609916100161011610216103161041610516106161071610816109161101611116112161131611416115161161611716118161191612016121161221612316124161251612616127161281612916130161311613216133161341613516136161371613816139161401614116142161431614416145161461614716148161491615016151161521615316154161551615616157161581615916160161611616216163161641616516166161671616816169161701617116172161731617416175161761617716178161791618016181161821618316184161851618616187161881618916190161911619216193161941619516196161971619816199162001620116202162031620416205162061620716208162091621016211162121621316214162151621616217162181621916220162211622216223162241622516226162271622816229162301623116232162331623416235162361623716238162391624016241162421624316244162451624616247162481624916250162511625216253162541625516256162571625816259162601626116262162631626416265162661626716268162691627016271162721627316274162751627616277162781627916280162811628216283162841628516286162871628816289162901629116292162931629416295162961629716298162991630016301163021630316304163051630616307163081630916310163111631216313163141631516316163171631816319163201632116322163231632416325163261632716328163291633016331163321633316334163351633616337163381633916340163411634216343163441634516346163471634816349163501635116352163531635416355163561635716358163591636016361163621636316364163651636616367163681636916370163711637216373163741637516376163771637816379163801638116382163831638416385163861638716388163891639016391163921639316394163951639616397163981639916400164011640216403164041640516406164071640816409164101641116412164131641416415164161641716418164191642016421164221642316424164251642616427164281642916430164311643216433164341643516436164371643816439164401644116442164431644416445164461644716448164491645016451164521645316454164551645616457164581645916460164611646216463164641646516466164671646816469164701647116472164731647416475164761647716478164791648016481164821648316484164851648616487164881648916490164911649216493164941649516496164971649816499165001650116502165031650416505165061650716508165091651016511165121651316514165151651616517165181651916520165211652216523165241652516526165271652816529165301653116532165331653416535165361653716538165391654016541165421654316544165451654616547165481654916550165511655216553165541655516556165571655816559165601656116562165631656416565165661656716568165691657016571165721657316574165751657616577165781657916580165811658216583165841658516586165871658816589165901659116592165931659416595165961659716598165991660016601166021660316604166051660616607166081660916610166111661216613166141661516616166171661816619166201662116622166231662416625166261662716628166291663016631166321663316634166351663616637166381663916640166411664216643166441664516646166471664816649166501665116652166531665416655166561665716658166591666016661166621666316664166651666616667166681666916670166711667216673166741667516676166771667816679166801668116682166831668416685166861668716688166891669016691166921669316694166951669616697166981669916700167011670216703167041670516706167071670816709167101671116712167131671416715167161671716718167191672016721167221672316724167251672616727167281672916730167311673216733167341673516736167371673816739167401674116742167431674416745167461674716748167491675016751167521675316754167551675616757167581675916760167611676216763167641676516766167671676816769167701677116772167731677416775167761677716778167791678016781167821678316784167851678616787167881678916790167911679216793167941679516796167971679816799168001680116802168031680416805168061680716808168091681016811168121681316814168151681616817168181681916820168211682216823168241682516826168271682816829168301683116832168331683416835168361683716838168391684016841168421684316844168451684616847168481684916850168511685216853168541685516856168571685816859168601686116862168631686416865168661686716868168691687016871168721687316874168751687616877168781687916880168811688216883168841688516886168871688816889168901689116892168931689416895168961689716898168991690016901169021690316904169051690616907169081690916910169111691216913169141691516916169171691816919169201692116922169231692416925169261692716928169291693016931169321693316934169351693616937169381693916940169411694216943169441694516946169471694816949169501695116952169531695416955169561695716958169591696016961169621696316964169651696616967169681696916970169711697216973169741697516976169771697816979169801698116982169831698416985169861698716988169891699016991169921699316994169951699616997169981699917000170011700217003170041700517006170071700817009170101701117012170131701417015170161701717018170191702017021170221702317024170251702617027170281702917030170311703217033170341703517036170371703817039170401704117042170431704417045170461704717048170491705017051170521705317054170551705617057170581705917060170611706217063170641706517066170671706817069170701707117072170731707417075170761707717078170791708017081170821708317084170851708617087170881708917090170911709217093170941709517096170971709817099171001710117102171031710417105171061710717108171091711017111171121711317114171151711617117171181711917120171211712217123171241712517126171271712817129171301713117132171331713417135171361713717138171391714017141171421714317144171451714617147171481714917150171511715217153171541715517156171571715817159171601716117162171631716417165171661716717168171691717017171171721717317174171751717617177171781717917180171811718217183171841718517186171871718817189171901719117192171931719417195171961719717198171991720017201172021720317204172051720617207172081720917210172111721217213172141721517216172171721817219172201722117222172231722417225172261722717228172291723017231172321723317234172351723617237172381723917240172411724217243172441724517246172471724817249172501725117252172531725417255172561725717258172591726017261172621726317264172651726617267172681726917270172711727217273172741727517276172771727817279172801728117282172831728417285172861728717288172891729017291172921729317294172951729617297172981729917300173011730217303173041730517306173071730817309173101731117312173131731417315173161731717318173191732017321173221732317324173251732617327173281732917330173311733217333173341733517336173371733817339173401734117342173431734417345173461734717348173491735017351173521735317354173551735617357173581735917360173611736217363173641736517366173671736817369173701737117372173731737417375173761737717378173791738017381173821738317384173851738617387173881738917390173911739217393173941739517396173971739817399174001740117402174031740417405174061740717408174091741017411174121741317414174151741617417174181741917420174211742217423174241742517426174271742817429174301743117432174331743417435174361743717438174391744017441174421744317444174451744617447174481744917450174511745217453174541745517456174571745817459174601746117462174631746417465174661746717468174691747017471174721747317474174751747617477174781747917480174811748217483174841748517486174871748817489174901749117492174931749417495174961749717498174991750017501175021750317504175051750617507175081750917510175111751217513175141751517516175171751817519175201752117522175231752417525175261752717528175291753017531175321753317534175351753617537175381753917540175411754217543175441754517546175471754817549175501755117552175531755417555175561755717558175591756017561175621756317564175651756617567175681756917570175711757217573175741757517576175771757817579175801758117582175831758417585175861758717588175891759017591175921759317594175951759617597175981759917600176011760217603176041760517606176071760817609176101761117612176131761417615176161761717618176191762017621176221762317624176251762617627176281762917630176311763217633176341763517636176371763817639176401764117642176431764417645176461764717648176491765017651176521765317654176551765617657176581765917660176611766217663176641766517666176671766817669176701767117672176731767417675176761767717678176791768017681176821768317684176851768617687176881768917690176911769217693176941769517696176971769817699177001770117702177031770417705177061770717708177091771017711177121771317714177151771617717177181771917720177211772217723177241772517726177271772817729177301773117732177331773417735177361773717738177391774017741177421774317744177451774617747177481774917750177511775217753177541775517756177571775817759177601776117762177631776417765177661776717768177691777017771177721777317774177751777617777177781777917780177811778217783177841778517786177871778817789177901779117792177931779417795177961779717798177991780017801178021780317804178051780617807178081780917810178111781217813178141781517816178171781817819178201782117822178231782417825178261782717828178291783017831178321783317834178351783617837178381783917840178411784217843178441784517846178471784817849178501785117852178531785417855178561785717858178591786017861178621786317864178651786617867178681786917870178711787217873178741787517876178771787817879178801788117882178831788417885178861788717888178891789017891178921789317894178951789617897178981789917900179011790217903179041790517906179071790817909179101791117912179131791417915179161791717918179191792017921179221792317924179251792617927179281792917930179311793217933179341793517936179371793817939179401794117942179431794417945179461794717948179491795017951179521795317954179551795617957179581795917960179611796217963179641796517966179671796817969179701797117972179731797417975179761797717978179791798017981179821798317984179851798617987179881798917990179911799217993179941799517996179971799817999180001800118002180031800418005180061800718008180091801018011180121801318014180151801618017180181801918020180211802218023180241802518026180271802818029180301803118032180331803418035180361803718038180391804018041180421804318044180451804618047180481804918050180511805218053180541805518056180571805818059180601806118062180631806418065180661806718068180691807018071180721807318074180751807618077180781807918080180811808218083180841808518086180871808818089180901809118092180931809418095180961809718098180991810018101181021810318104181051810618107181081810918110181111811218113181141811518116181171811818119181201812118122181231812418125181261812718128181291813018131181321813318134181351813618137181381813918140181411814218143181441814518146181471814818149181501815118152181531815418155181561815718158181591816018161181621816318164181651816618167181681816918170181711817218173181741817518176181771817818179181801818118182181831818418185181861818718188181891819018191181921819318194181951819618197181981819918200182011820218203182041820518206182071820818209182101821118212182131821418215182161821718218182191822018221182221822318224182251822618227182281822918230182311823218233182341823518236182371823818239182401824118242182431824418245182461824718248182491825018251182521825318254182551825618257182581825918260182611826218263182641826518266182671826818269182701827118272182731827418275182761827718278182791828018281182821828318284182851828618287182881828918290182911829218293182941829518296182971829818299183001830118302183031830418305183061830718308183091831018311183121831318314183151831618317183181831918320183211832218323183241832518326183271832818329183301833118332183331833418335183361833718338183391834018341183421834318344183451834618347183481834918350183511835218353183541835518356183571835818359183601836118362183631836418365183661836718368183691837018371183721837318374183751837618377183781837918380183811838218383183841838518386183871838818389183901839118392183931839418395183961839718398183991840018401184021840318404184051840618407184081840918410184111841218413184141841518416184171841818419184201842118422184231842418425184261842718428184291843018431184321843318434184351843618437184381843918440184411844218443184441844518446184471844818449184501845118452184531845418455184561845718458184591846018461184621846318464184651846618467184681846918470184711847218473184741847518476184771847818479184801848118482184831848418485184861848718488184891849018491184921849318494184951849618497184981849918500185011850218503185041850518506185071850818509185101851118512185131851418515185161851718518185191852018521185221852318524185251852618527185281852918530185311853218533185341853518536185371853818539185401854118542185431854418545185461854718548185491855018551185521855318554185551855618557185581855918560185611856218563185641856518566185671856818569185701857118572185731857418575185761857718578185791858018581185821858318584185851858618587185881858918590185911859218593185941859518596185971859818599186001860118602186031860418605186061860718608186091861018611186121861318614186151861618617186181861918620186211862218623186241862518626186271862818629186301863118632186331863418635186361863718638186391864018641186421864318644186451864618647186481864918650186511865218653186541865518656186571865818659186601866118662186631866418665186661866718668186691867018671186721867318674186751867618677186781867918680186811868218683186841868518686186871868818689186901869118692186931869418695186961869718698186991870018701187021870318704187051870618707187081870918710187111871218713187141871518716187171871818719187201872118722187231872418725187261872718728187291873018731187321873318734187351873618737187381873918740187411874218743187441874518746187471874818749187501875118752187531875418755187561875718758187591876018761187621876318764187651876618767187681876918770187711877218773187741877518776187771877818779187801878118782187831878418785187861878718788187891879018791187921879318794187951879618797187981879918800188011880218803188041880518806188071880818809188101881118812188131881418815188161881718818188191882018821188221882318824188251882618827188281882918830188311883218833188341883518836188371883818839188401884118842188431884418845188461884718848188491885018851188521885318854188551885618857188581885918860188611886218863188641886518866188671886818869188701887118872188731887418875188761887718878188791888018881188821888318884188851888618887188881888918890188911889218893188941889518896188971889818899189001890118902189031890418905189061890718908189091891018911189121891318914189151891618917189181891918920189211892218923189241892518926189271892818929189301893118932189331893418935189361893718938189391894018941189421894318944189451894618947189481894918950189511895218953189541895518956189571895818959189601896118962189631896418965189661896718968189691897018971189721897318974189751897618977189781897918980189811898218983189841898518986189871898818989189901899118992189931899418995189961899718998189991900019001190021900319004190051900619007190081900919010190111901219013190141901519016190171901819019190201902119022190231902419025190261902719028190291903019031190321903319034190351903619037190381903919040190411904219043190441904519046190471904819049190501905119052190531905419055190561905719058190591906019061190621906319064190651906619067190681906919070190711907219073190741907519076190771907819079190801908119082190831908419085190861908719088190891909019091190921909319094190951909619097190981909919100191011910219103191041910519106191071910819109191101911119112191131911419115191161911719118191191912019121191221912319124191251912619127191281912919130191311913219133191341913519136191371913819139191401914119142191431914419145191461914719148191491915019151191521915319154191551915619157191581915919160191611916219163191641916519166191671916819169191701917119172191731917419175191761917719178191791918019181191821918319184191851918619187191881918919190191911919219193191941919519196191971919819199192001920119202192031920419205192061920719208192091921019211192121921319214192151921619217192181921919220192211922219223192241922519226192271922819229192301923119232192331923419235192361923719238192391924019241192421924319244192451924619247192481924919250192511925219253192541925519256192571925819259192601926119262192631926419265192661926719268192691927019271192721927319274192751927619277192781927919280192811928219283192841928519286192871928819289192901929119292192931929419295192961929719298192991930019301193021930319304193051930619307193081930919310193111931219313193141931519316193171931819319193201932119322193231932419325193261932719328193291933019331193321933319334193351933619337193381933919340193411934219343193441934519346193471934819349193501935119352193531935419355193561935719358193591936019361193621936319364193651936619367193681936919370193711937219373193741937519376193771937819379193801938119382193831938419385193861938719388193891939019391193921939319394193951939619397193981939919400194011940219403194041940519406194071940819409194101941119412194131941419415194161941719418194191942019421194221942319424194251942619427194281942919430194311943219433194341943519436194371943819439194401944119442194431944419445194461944719448194491945019451194521945319454194551945619457194581945919460194611946219463194641946519466194671946819469194701947119472194731947419475194761947719478194791948019481194821948319484194851948619487194881948919490194911949219493194941949519496194971949819499195001950119502195031950419505195061950719508195091951019511195121951319514195151951619517195181951919520195211952219523195241952519526195271952819529195301953119532195331953419535195361953719538195391954019541195421954319544195451954619547195481954919550195511955219553195541955519556195571955819559195601956119562195631956419565195661956719568195691957019571195721957319574195751957619577195781957919580195811958219583195841958519586195871958819589195901959119592195931959419595195961959719598195991960019601196021960319604196051960619607196081960919610196111961219613196141961519616196171961819619196201962119622196231962419625196261962719628196291963019631196321963319634196351963619637196381963919640196411964219643196441964519646196471964819649196501965119652196531965419655196561965719658196591966019661196621966319664196651966619667196681966919670196711967219673196741967519676196771967819679196801968119682196831968419685196861968719688196891969019691196921969319694196951969619697196981969919700197011970219703197041970519706197071970819709197101971119712197131971419715197161971719718197191972019721197221972319724197251972619727197281972919730197311973219733197341973519736197371973819739197401974119742197431974419745197461974719748197491975019751197521975319754197551975619757197581975919760197611976219763197641976519766197671976819769197701977119772197731977419775197761977719778197791978019781197821978319784197851978619787197881978919790197911979219793197941979519796197971979819799198001980119802198031980419805198061980719808198091981019811198121981319814198151981619817198181981919820198211982219823198241982519826198271982819829198301983119832198331983419835198361983719838198391984019841198421984319844198451984619847198481984919850198511985219853198541985519856198571985819859198601986119862198631986419865198661986719868198691987019871198721987319874198751987619877198781987919880198811988219883198841988519886198871988819889198901989119892198931989419895198961989719898198991990019901199021990319904199051990619907199081990919910199111991219913199141991519916199171991819919199201992119922199231992419925199261992719928199291993019931199321993319934199351993619937199381993919940199411994219943199441994519946199471994819949199501995119952199531995419955199561995719958199591996019961199621996319964199651996619967199681996919970199711997219973199741997519976199771997819979199801998119982199831998419985199861998719988199891999019991199921999319994199951999619997199981999920000200012000220003200042000520006200072000820009200102001120012200132001420015200162001720018200192002020021200222002320024200252002620027200282002920030200312003220033200342003520036200372003820039200402004120042200432004420045200462004720048200492005020051200522005320054200552005620057200582005920060200612006220063200642006520066200672006820069200702007120072200732007420075200762007720078200792008020081200822008320084200852008620087200882008920090200912009220093200942009520096200972009820099201002010120102201032010420105201062010720108201092011020111201122011320114201152011620117201182011920120201212012220123201242012520126201272012820129201302013120132201332013420135201362013720138201392014020141201422014320144201452014620147201482014920150201512015220153201542015520156201572015820159201602016120162201632016420165201662016720168201692017020171201722017320174201752017620177201782017920180201812018220183201842018520186201872018820189201902019120192201932019420195201962019720198201992020020201202022020320204202052020620207202082020920210202112021220213202142021520216202172021820219202202022120222202232022420225202262022720228202292023020231202322023320234202352023620237202382023920240202412024220243202442024520246202472024820249202502025120252202532025420255202562025720258202592026020261202622026320264202652026620267202682026920270202712027220273202742027520276202772027820279202802028120282202832028420285202862028720288202892029020291202922029320294202952029620297202982029920300203012030220303203042030520306203072030820309203102031120312203132031420315203162031720318203192032020321203222032320324203252032620327203282032920330203312033220333203342033520336203372033820339203402034120342203432034420345203462034720348203492035020351203522035320354203552035620357203582035920360203612036220363203642036520366203672036820369203702037120372203732037420375203762037720378203792038020381203822038320384203852038620387203882038920390203912039220393203942039520396203972039820399204002040120402204032040420405204062040720408204092041020411204122041320414204152041620417204182041920420204212042220423204242042520426204272042820429204302043120432204332043420435204362043720438204392044020441204422044320444204452044620447204482044920450204512045220453204542045520456204572045820459204602046120462204632046420465204662046720468204692047020471204722047320474204752047620477204782047920480204812048220483204842048520486204872048820489204902049120492204932049420495204962049720498204992050020501205022050320504205052050620507205082050920510205112051220513205142051520516205172051820519205202052120522205232052420525205262052720528205292053020531205322053320534205352053620537205382053920540205412054220543205442054520546205472054820549205502055120552205532055420555205562055720558205592056020561205622056320564205652056620567205682056920570205712057220573205742057520576205772057820579205802058120582205832058420585205862058720588205892059020591205922059320594205952059620597205982059920600206012060220603206042060520606206072060820609206102061120612206132061420615206162061720618206192062020621206222062320624206252062620627206282062920630206312063220633206342063520636206372063820639206402064120642206432064420645206462064720648206492065020651206522065320654206552065620657206582065920660206612066220663206642066520666206672066820669206702067120672206732067420675206762067720678206792068020681206822068320684206852068620687206882068920690206912069220693206942069520696206972069820699207002070120702207032070420705207062070720708207092071020711207122071320714207152071620717207182071920720207212072220723207242072520726207272072820729207302073120732207332073420735207362073720738207392074020741207422074320744207452074620747207482074920750207512075220753207542075520756207572075820759207602076120762207632076420765207662076720768207692077020771207722077320774207752077620777207782077920780207812078220783207842078520786207872078820789207902079120792207932079420795207962079720798207992080020801208022080320804208052080620807208082080920810208112081220813208142081520816208172081820819208202082120822208232082420825208262082720828208292083020831208322083320834208352083620837208382083920840208412084220843208442084520846208472084820849208502085120852208532085420855208562085720858208592086020861208622086320864208652086620867208682086920870208712087220873208742087520876208772087820879208802088120882208832088420885208862088720888208892089020891208922089320894208952089620897208982089920900209012090220903209042090520906209072090820909209102091120912209132091420915209162091720918209192092020921209222092320924209252092620927209282092920930209312093220933209342093520936209372093820939209402094120942209432094420945209462094720948209492095020951209522095320954209552095620957209582095920960209612096220963209642096520966209672096820969209702097120972209732097420975209762097720978209792098020981209822098320984209852098620987209882098920990209912099220993209942099520996209972099820999210002100121002210032100421005210062100721008210092101021011210122101321014210152101621017210182101921020210212102221023210242102521026210272102821029210302103121032210332103421035210362103721038210392104021041210422104321044210452104621047210482104921050210512105221053210542105521056210572105821059210602106121062210632106421065210662106721068210692107021071210722107321074210752107621077210782107921080210812108221083210842108521086210872108821089210902109121092210932109421095210962109721098210992110021101211022110321104211052110621107211082110921110211112111221113211142111521116211172111821119211202112121122211232112421125211262112721128211292113021131211322113321134211352113621137211382113921140211412114221143211442114521146211472114821149211502115121152211532115421155211562115721158211592116021161211622116321164211652116621167211682116921170211712117221173211742117521176211772117821179211802118121182211832118421185211862118721188211892119021191211922119321194211952119621197211982119921200212012120221203212042120521206212072120821209212102121121212212132121421215212162121721218212192122021221212222122321224212252122621227212282122921230212312123221233212342123521236212372123821239212402124121242212432124421245212462124721248212492125021251212522125321254212552125621257212582125921260212612126221263212642126521266212672126821269212702127121272212732127421275212762127721278212792128021281212822128321284212852128621287212882128921290212912129221293212942129521296212972129821299213002130121302213032130421305213062130721308213092131021311213122131321314213152131621317213182131921320213212132221323213242132521326213272132821329213302133121332213332133421335213362133721338213392134021341213422134321344213452134621347213482134921350213512135221353213542135521356213572135821359213602136121362213632136421365213662136721368213692137021371213722137321374213752137621377213782137921380213812138221383213842138521386213872138821389213902139121392213932139421395213962139721398213992140021401214022140321404214052140621407214082140921410214112141221413214142141521416214172141821419214202142121422214232142421425214262142721428214292143021431214322143321434214352143621437214382143921440214412144221443214442144521446214472144821449214502145121452214532145421455214562145721458214592146021461214622146321464214652146621467214682146921470214712147221473214742147521476214772147821479214802148121482214832148421485214862148721488214892149021491214922149321494214952149621497214982149921500215012150221503215042150521506215072150821509215102151121512215132151421515215162151721518215192152021521215222152321524215252152621527215282152921530215312153221533215342153521536215372153821539215402154121542215432154421545215462154721548215492155021551215522155321554215552155621557215582155921560215612156221563215642156521566215672156821569215702157121572215732157421575215762157721578215792158021581215822158321584215852158621587215882158921590215912159221593215942159521596215972159821599216002160121602216032160421605216062160721608216092161021611216122161321614216152161621617216182161921620216212162221623216242162521626216272162821629216302163121632216332163421635216362163721638216392164021641216422164321644216452164621647216482164921650216512165221653216542165521656216572165821659216602166121662216632166421665216662166721668216692167021671216722167321674216752167621677216782167921680216812168221683216842168521686216872168821689216902169121692216932169421695216962169721698216992170021701217022170321704217052170621707217082170921710217112171221713217142171521716217172171821719217202172121722217232172421725217262172721728217292173021731217322173321734217352173621737217382173921740217412174221743217442174521746217472174821749217502175121752217532175421755217562175721758217592176021761217622176321764217652176621767217682176921770217712177221773217742177521776217772177821779217802178121782217832178421785217862178721788217892179021791217922179321794217952179621797217982179921800218012180221803218042180521806218072180821809218102181121812218132181421815218162181721818218192182021821218222182321824218252182621827218282182921830218312183221833218342183521836218372183821839218402184121842218432184421845218462184721848218492185021851218522185321854218552185621857218582185921860218612186221863218642186521866218672186821869218702187121872218732187421875218762187721878218792188021881218822188321884218852188621887218882188921890218912189221893218942189521896218972189821899219002190121902219032190421905219062190721908219092191021911219122191321914219152191621917219182191921920219212192221923219242192521926219272192821929219302193121932219332193421935219362193721938219392194021941219422194321944219452194621947219482194921950219512195221953219542195521956219572195821959219602196121962219632196421965219662196721968219692197021971219722197321974219752197621977219782197921980219812198221983219842198521986219872198821989219902199121992219932199421995219962199721998219992200022001220022200322004220052200622007220082200922010220112201222013220142201522016220172201822019220202202122022220232202422025220262202722028220292203022031220322203322034220352203622037220382203922040220412204222043220442204522046220472204822049220502205122052220532205422055220562205722058220592206022061220622206322064220652206622067220682206922070220712207222073220742207522076220772207822079220802208122082220832208422085220862208722088220892209022091220922209322094220952209622097220982209922100221012210222103221042210522106221072210822109221102211122112221132211422115221162211722118221192212022121221222212322124221252212622127221282212922130221312213222133221342213522136221372213822139221402214122142221432214422145221462214722148221492215022151221522215322154221552215622157221582215922160221612216222163221642216522166221672216822169221702217122172221732217422175221762217722178221792218022181221822218322184221852218622187221882218922190221912219222193221942219522196221972219822199222002220122202222032220422205222062220722208222092221022211222122221322214222152221622217222182221922220222212222222223222242222522226222272222822229222302223122232222332223422235222362223722238222392224022241222422224322244222452224622247222482224922250222512225222253222542225522256222572225822259222602226122262222632226422265222662226722268222692227022271222722227322274222752227622277222782227922280222812228222283222842228522286222872228822289222902229122292222932229422295222962229722298222992230022301223022230322304223052230622307223082230922310223112231222313223142231522316223172231822319223202232122322223232232422325223262232722328223292233022331223322233322334223352233622337223382233922340223412234222343223442234522346223472234822349223502235122352223532235422355223562235722358223592236022361223622236322364223652236622367223682236922370223712237222373223742237522376223772237822379223802238122382223832238422385223862238722388223892239022391223922239322394223952239622397223982239922400224012240222403224042240522406224072240822409224102241122412224132241422415224162241722418224192242022421224222242322424224252242622427224282242922430224312243222433224342243522436224372243822439224402244122442224432244422445224462244722448224492245022451224522245322454224552245622457224582245922460224612246222463224642246522466224672246822469224702247122472224732247422475224762247722478224792248022481224822248322484224852248622487224882248922490224912249222493224942249522496224972249822499225002250122502225032250422505225062250722508225092251022511225122251322514225152251622517225182251922520225212252222523225242252522526225272252822529225302253122532225332253422535225362253722538225392254022541225422254322544225452254622547225482254922550225512255222553225542255522556225572255822559225602256122562225632256422565225662256722568225692257022571225722257322574225752257622577225782257922580225812258222583225842258522586225872258822589225902259122592225932259422595225962259722598225992260022601226022260322604226052260622607226082260922610226112261222613226142261522616226172261822619226202262122622226232262422625226262262722628226292263022631226322263322634226352263622637226382263922640226412264222643226442264522646226472264822649226502265122652226532265422655226562265722658226592266022661226622266322664226652266622667226682266922670226712267222673226742267522676226772267822679226802268122682226832268422685226862268722688226892269022691226922269322694226952269622697226982269922700227012270222703227042270522706227072270822709227102271122712227132271422715227162271722718227192272022721227222272322724227252272622727227282272922730227312273222733227342273522736227372273822739227402274122742227432274422745227462274722748227492275022751227522275322754227552275622757227582275922760227612276222763227642276522766227672276822769227702277122772227732277422775227762277722778227792278022781227822278322784227852278622787227882278922790227912279222793227942279522796227972279822799228002280122802228032280422805228062280722808228092281022811228122281322814228152281622817228182281922820228212282222823228242282522826228272282822829228302283122832228332283422835228362283722838228392284022841228422284322844228452284622847228482284922850228512285222853228542285522856228572285822859228602286122862228632286422865228662286722868228692287022871228722287322874228752287622877228782287922880228812288222883228842288522886228872288822889228902289122892228932289422895228962289722898228992290022901229022290322904229052290622907229082290922910229112291222913229142291522916229172291822919229202292122922229232292422925229262292722928229292293022931229322293322934229352293622937229382293922940229412294222943229442294522946229472294822949229502295122952229532295422955229562295722958229592296022961229622296322964229652296622967229682296922970229712297222973229742297522976229772297822979229802298122982229832298422985229862298722988229892299022991229922299322994229952299622997229982299923000230012300223003230042300523006230072300823009230102301123012230132301423015230162301723018230192302023021230222302323024230252302623027230282302923030230312303223033230342303523036230372303823039230402304123042230432304423045230462304723048230492305023051230522305323054230552305623057230582305923060230612306223063230642306523066230672306823069230702307123072230732307423075230762307723078230792308023081230822308323084230852308623087230882308923090230912309223093230942309523096230972309823099231002310123102231032310423105231062310723108231092311023111231122311323114231152311623117231182311923120231212312223123231242312523126231272312823129231302313123132231332313423135231362313723138231392314023141231422314323144231452314623147231482314923150231512315223153231542315523156231572315823159231602316123162231632316423165231662316723168231692317023171231722317323174231752317623177231782317923180231812318223183231842318523186231872318823189231902319123192231932319423195231962319723198231992320023201232022320323204232052320623207232082320923210232112321223213232142321523216232172321823219232202322123222232232322423225232262322723228232292323023231232322323323234232352323623237232382323923240232412324223243232442324523246232472324823249232502325123252232532325423255232562325723258232592326023261232622326323264232652326623267232682326923270232712327223273232742327523276232772327823279232802328123282232832328423285232862328723288232892329023291232922329323294232952329623297232982329923300233012330223303233042330523306233072330823309233102331123312233132331423315233162331723318233192332023321233222332323324233252332623327233282332923330233312333223333233342333523336233372333823339233402334123342233432334423345233462334723348233492335023351233522335323354233552335623357233582335923360233612336223363233642336523366233672336823369233702337123372233732337423375233762337723378233792338023381233822338323384233852338623387233882338923390233912339223393233942339523396233972339823399234002340123402234032340423405234062340723408234092341023411234122341323414234152341623417234182341923420234212342223423234242342523426234272342823429234302343123432234332343423435234362343723438234392344023441234422344323444234452344623447234482344923450234512345223453234542345523456234572345823459234602346123462234632346423465234662346723468234692347023471234722347323474234752347623477234782347923480234812348223483234842348523486234872348823489234902349123492234932349423495234962349723498234992350023501235022350323504235052350623507235082350923510235112351223513235142351523516235172351823519235202352123522235232352423525235262352723528235292353023531235322353323534235352353623537235382353923540235412354223543235442354523546235472354823549235502355123552235532355423555235562355723558235592356023561235622356323564235652356623567235682356923570235712357223573235742357523576235772357823579235802358123582235832358423585235862358723588235892359023591235922359323594235952359623597235982359923600236012360223603236042360523606236072360823609236102361123612236132361423615236162361723618236192362023621236222362323624236252362623627236282362923630236312363223633236342363523636236372363823639236402364123642236432364423645236462364723648236492365023651236522365323654236552365623657236582365923660236612366223663236642366523666236672366823669236702367123672236732367423675236762367723678236792368023681236822368323684236852368623687236882368923690236912369223693236942369523696236972369823699237002370123702237032370423705237062370723708237092371023711237122371323714237152371623717237182371923720237212372223723237242372523726237272372823729237302373123732237332373423735237362373723738237392374023741237422374323744237452374623747237482374923750237512375223753237542375523756237572375823759237602376123762237632376423765237662376723768237692377023771237722377323774237752377623777237782377923780237812378223783237842378523786237872378823789237902379123792237932379423795237962379723798237992380023801238022380323804238052380623807238082380923810238112381223813238142381523816238172381823819238202382123822238232382423825238262382723828238292383023831238322383323834238352383623837238382383923840238412384223843238442384523846238472384823849238502385123852238532385423855238562385723858238592386023861238622386323864238652386623867238682386923870238712387223873238742387523876238772387823879238802388123882238832388423885238862388723888238892389023891238922389323894238952389623897238982389923900239012390223903239042390523906239072390823909239102391123912239132391423915239162391723918239192392023921239222392323924239252392623927239282392923930239312393223933239342393523936239372393823939239402394123942239432394423945239462394723948239492395023951239522395323954239552395623957239582395923960239612396223963239642396523966239672396823969239702397123972239732397423975239762397723978239792398023981239822398323984239852398623987239882398923990239912399223993239942399523996239972399823999240002400124002240032400424005240062400724008240092401024011240122401324014240152401624017240182401924020240212402224023240242402524026240272402824029240302403124032240332403424035240362403724038240392404024041240422404324044240452404624047240482404924050240512405224053240542405524056240572405824059240602406124062240632406424065240662406724068240692407024071240722407324074240752407624077240782407924080240812408224083240842408524086240872408824089240902409124092240932409424095240962409724098240992410024101241022410324104241052410624107241082410924110241112411224113241142411524116241172411824119241202412124122241232412424125241262412724128241292413024131241322413324134241352413624137241382413924140241412414224143241442414524146241472414824149241502415124152241532415424155241562415724158241592416024161241622416324164241652416624167241682416924170241712417224173241742417524176241772417824179241802418124182241832418424185241862418724188241892419024191241922419324194241952419624197241982419924200242012420224203242042420524206242072420824209242102421124212242132421424215242162421724218242192422024221242222422324224242252422624227242282422924230242312423224233242342423524236242372423824239242402424124242242432424424245242462424724248242492425024251242522425324254242552425624257242582425924260242612426224263242642426524266242672426824269242702427124272242732427424275242762427724278242792428024281242822428324284242852428624287242882428924290242912429224293242942429524296242972429824299243002430124302243032430424305243062430724308243092431024311243122431324314243152431624317243182431924320243212432224323243242432524326243272432824329243302433124332243332433424335243362433724338243392434024341243422434324344243452434624347243482434924350243512435224353243542435524356243572435824359243602436124362243632436424365243662436724368243692437024371243722437324374243752437624377243782437924380243812438224383243842438524386243872438824389243902439124392243932439424395243962439724398243992440024401244022440324404244052440624407244082440924410244112441224413244142441524416244172441824419244202442124422244232442424425244262442724428244292443024431244322443324434244352443624437244382443924440244412444224443244442444524446244472444824449244502445124452244532445424455244562445724458244592446024461244622446324464244652446624467244682446924470244712447224473244742447524476244772447824479244802448124482244832448424485244862448724488244892449024491244922449324494244952449624497244982449924500245012450224503245042450524506245072450824509245102451124512245132451424515245162451724518245192452024521245222452324524245252452624527245282452924530245312453224533245342453524536245372453824539245402454124542245432454424545245462454724548245492455024551245522455324554245552455624557245582455924560245612456224563245642456524566245672456824569245702457124572245732457424575245762457724578245792458024581245822458324584245852458624587245882458924590245912459224593245942459524596245972459824599246002460124602246032460424605246062460724608246092461024611246122461324614246152461624617246182461924620246212462224623246242462524626246272462824629246302463124632246332463424635246362463724638246392464024641246422464324644246452464624647246482464924650246512465224653246542465524656246572465824659246602466124662246632466424665246662466724668246692467024671246722467324674246752467624677246782467924680246812468224683246842468524686246872468824689246902469124692246932469424695246962469724698246992470024701247022470324704247052470624707247082470924710247112471224713247142471524716247172471824719247202472124722247232472424725247262472724728247292473024731247322473324734247352473624737247382473924740247412474224743247442474524746247472474824749247502475124752247532475424755247562475724758247592476024761247622476324764247652476624767247682476924770247712477224773247742477524776247772477824779247802478124782247832478424785247862478724788247892479024791247922479324794247952479624797247982479924800248012480224803248042480524806248072480824809248102481124812248132481424815248162481724818248192482024821248222482324824248252482624827248282482924830248312483224833248342483524836248372483824839248402484124842248432484424845248462484724848248492485024851248522485324854248552485624857248582485924860248612486224863248642486524866248672486824869248702487124872248732487424875248762487724878248792488024881248822488324884248852488624887248882488924890248912489224893248942489524896248972489824899249002490124902249032490424905249062490724908249092491024911249122491324914249152491624917249182491924920249212492224923249242492524926249272492824929249302493124932249332493424935249362493724938249392494024941249422494324944249452494624947249482494924950249512495224953249542495524956249572495824959249602496124962249632496424965249662496724968249692497024971249722497324974249752497624977249782497924980249812498224983249842498524986249872498824989249902499124992249932499424995249962499724998249992500025001250022500325004250052500625007250082500925010250112501225013250142501525016250172501825019250202502125022250232502425025250262502725028250292503025031250322503325034250352503625037250382503925040250412504225043250442504525046250472504825049250502505125052250532505425055250562505725058250592506025061250622506325064250652506625067250682506925070250712507225073250742507525076250772507825079250802508125082250832508425085250862508725088250892509025091250922509325094250952509625097250982509925100251012510225103251042510525106251072510825109251102511125112251132511425115251162511725118251192512025121251222512325124251252512625127251282512925130251312513225133251342513525136251372513825139251402514125142251432514425145251462514725148251492515025151251522515325154251552515625157251582515925160251612516225163251642516525166251672516825169251702517125172251732517425175251762517725178251792518025181251822518325184251852518625187251882518925190251912519225193251942519525196251972519825199252002520125202252032520425205252062520725208252092521025211252122521325214252152521625217252182521925220252212522225223252242522525226252272522825229252302523125232252332523425235252362523725238252392524025241252422524325244252452524625247252482524925250252512525225253252542525525256252572525825259252602526125262252632526425265252662526725268252692527025271252722527325274252752527625277252782527925280252812528225283252842528525286252872528825289252902529125292252932529425295252962529725298252992530025301253022530325304253052530625307253082530925310253112531225313253142531525316253172531825319253202532125322253232532425325253262532725328253292533025331253322533325334253352533625337253382533925340253412534225343253442534525346253472534825349253502535125352253532535425355253562535725358253592536025361253622536325364253652536625367253682536925370253712537225373253742537525376253772537825379253802538125382253832538425385253862538725388253892539025391253922539325394253952539625397253982539925400254012540225403254042540525406254072540825409254102541125412254132541425415254162541725418254192542025421254222542325424254252542625427254282542925430254312543225433254342543525436254372543825439254402544125442254432544425445254462544725448254492545025451254522545325454254552545625457254582545925460254612546225463254642546525466254672546825469254702547125472254732547425475254762547725478254792548025481254822548325484254852548625487254882548925490254912549225493254942549525496254972549825499255002550125502255032550425505255062550725508255092551025511255122551325514255152551625517255182551925520255212552225523255242552525526255272552825529255302553125532255332553425535255362553725538255392554025541255422554325544255452554625547255482554925550255512555225553255542555525556255572555825559255602556125562255632556425565255662556725568255692557025571255722557325574255752557625577255782557925580255812558225583255842558525586255872558825589255902559125592255932559425595255962559725598255992560025601256022560325604256052560625607256082560925610256112561225613256142561525616256172561825619256202562125622256232562425625256262562725628256292563025631256322563325634256352563625637256382563925640256412564225643256442564525646256472564825649256502565125652256532565425655256562565725658256592566025661256622566325664256652566625667256682566925670256712567225673256742567525676256772567825679256802568125682256832568425685256862568725688256892569025691256922569325694256952569625697256982569925700257012570225703257042570525706257072570825709257102571125712257132571425715257162571725718257192572025721257222572325724257252572625727257282572925730257312573225733257342573525736257372573825739257402574125742257432574425745257462574725748257492575025751257522575325754257552575625757257582575925760257612576225763257642576525766257672576825769257702577125772257732577425775257762577725778257792578025781257822578325784257852578625787257882578925790257912579225793257942579525796257972579825799258002580125802258032580425805258062580725808258092581025811258122581325814258152581625817258182581925820258212582225823258242582525826258272582825829258302583125832258332583425835258362583725838258392584025841258422584325844258452584625847258482584925850258512585225853258542585525856258572585825859258602586125862258632586425865258662586725868258692587025871258722587325874258752587625877258782587925880258812588225883258842588525886258872588825889258902589125892258932589425895258962589725898258992590025901259022590325904259052590625907259082590925910259112591225913259142591525916259172591825919259202592125922259232592425925259262592725928259292593025931259322593325934259352593625937259382593925940259412594225943259442594525946259472594825949259502595125952259532595425955259562595725958259592596025961259622596325964259652596625967259682596925970259712597225973259742597525976259772597825979259802598125982259832598425985259862598725988259892599025991259922599325994259952599625997259982599926000260012600226003260042600526006260072600826009260102601126012260132601426015260162601726018260192602026021260222602326024260252602626027260282602926030260312603226033260342603526036260372603826039260402604126042260432604426045260462604726048260492605026051260522605326054260552605626057260582605926060260612606226063260642606526066260672606826069260702607126072260732607426075260762607726078260792608026081260822608326084260852608626087260882608926090260912609226093260942609526096260972609826099261002610126102261032610426105261062610726108261092611026111261122611326114261152611626117261182611926120261212612226123261242612526126261272612826129261302613126132261332613426135261362613726138261392614026141261422614326144261452614626147261482614926150261512615226153261542615526156261572615826159261602616126162261632616426165261662616726168261692617026171261722617326174261752617626177261782617926180261812618226183261842618526186261872618826189261902619126192261932619426195261962619726198261992620026201262022620326204262052620626207262082620926210262112621226213262142621526216262172621826219262202622126222262232622426225262262622726228262292623026231262322623326234262352623626237262382623926240262412624226243262442624526246262472624826249262502625126252262532625426255262562625726258262592626026261262622626326264262652626626267262682626926270262712627226273262742627526276262772627826279262802628126282262832628426285262862628726288262892629026291262922629326294262952629626297262982629926300263012630226303263042630526306263072630826309263102631126312263132631426315263162631726318263192632026321263222632326324263252632626327263282632926330263312633226333263342633526336263372633826339263402634126342263432634426345263462634726348263492635026351263522635326354263552635626357263582635926360263612636226363263642636526366263672636826369263702637126372263732637426375263762637726378263792638026381263822638326384263852638626387263882638926390263912639226393263942639526396263972639826399264002640126402264032640426405264062640726408264092641026411264122641326414264152641626417264182641926420264212642226423264242642526426264272642826429264302643126432264332643426435264362643726438264392644026441264422644326444264452644626447264482644926450264512645226453264542645526456264572645826459264602646126462264632646426465264662646726468264692647026471264722647326474264752647626477264782647926480264812648226483264842648526486264872648826489264902649126492264932649426495264962649726498264992650026501265022650326504265052650626507265082650926510265112651226513265142651526516265172651826519265202652126522265232652426525265262652726528265292653026531265322653326534265352653626537265382653926540265412654226543265442654526546265472654826549265502655126552265532655426555265562655726558265592656026561265622656326564265652656626567265682656926570265712657226573265742657526576265772657826579265802658126582265832658426585265862658726588265892659026591265922659326594265952659626597265982659926600266012660226603266042660526606266072660826609266102661126612266132661426615266162661726618266192662026621266222662326624266252662626627266282662926630266312663226633266342663526636266372663826639266402664126642266432664426645266462664726648266492665026651266522665326654266552665626657266582665926660266612666226663266642666526666266672666826669266702667126672266732667426675266762667726678266792668026681266822668326684266852668626687266882668926690266912669226693266942669526696266972669826699267002670126702267032670426705267062670726708267092671026711267122671326714267152671626717267182671926720267212672226723267242672526726267272672826729267302673126732267332673426735267362673726738267392674026741267422674326744267452674626747267482674926750267512675226753267542675526756267572675826759267602676126762267632676426765267662676726768267692677026771267722677326774267752677626777267782677926780267812678226783267842678526786267872678826789267902679126792267932679426795267962679726798267992680026801268022680326804268052680626807268082680926810268112681226813268142681526816 |
- // Experimental
- import * as base from './base.js';
- import * as text from './text.js';
- const mlir = {};
- const _ = {};
- mlir.ModelFactory = class {
- async match(context) {
- const stream = context.stream;
- const identifier = context.identifier;
- const extension = identifier.split('.').pop().toLowerCase();
- if (stream && stream.length > 4) {
- const buffer = stream.peek(4);
- const signature = String.fromCharCode.apply(null, buffer);
- if (signature === 'ML\xEFR') {
- return context.set('mlir.binary');
- }
- }
- try {
- const reader = await context.read('text', 0x10000);
- let whitespace = true;
- for (let line = reader.read('\n'); line !== undefined; line = reader.read('\n')) {
- if (/module\s+(@\w+|\w+|attributes|\{)/.test(line) ||
- /tensor<[\w\d]+>/.test(line) ||
- /func[.\s]*@\w+/.test(line) ||
- /%\w+\s*=\s*"[\w.]+/.test(line) ||
- /%\w+\s*=\s*\w+\./.test(line) ||
- /!\w+\s*=\s*![\w.]+</.test(line) ||
- /#\w+\s*=\s*#[\w.]+</.test(line) ||
- /#\w+\s*=\s*loc\s*\(/.test(line) ||
- /\w+\.\w+(?:\s+\w+)*\s+@\w+/.test(line) ||
- /\w+\.\w+\s+#[\w.]+</.test(line) ||
- /\w+\.\w+\s*<?\{/.test(line) ||
- /:\s*![\w.]+/.test(line) ||
- /(%\w+|\w{2,}|[)])\s*:\s*(\[|tensor<)/.test(line) ||
- /->\s*(![\w.]+|\(|tensor<)/.test(line)) {
- return context.set('mlir.text');
- }
- if (line && !line.trim().startsWith('//')) {
- whitespace = false;
- }
- }
- if (extension === 'mlir' && whitespace) {
- return context.set('mlir.text');
- }
- } catch {
- // continue
- }
- return null;
- }
- async open(context) {
- const metadata = await mlir.Metadata.open(context);
- switch (context.type) {
- case 'mlir.text': {
- const decoder = await context.read('text.decoder');
- const config = new _.ParserConfig(new _.DialectContext(metadata));
- const state = new _.ParserState(decoder, config);
- const parser = new _.TopLevelOperationParser(state);
- const block = new _.Block();
- parser.parse(block);
- const model = new mlir.Model(config, 'MLIR', '', block, state.attributeAliasDefinitions);
- return model;
- }
- case 'mlir.binary': {
- const binary = await context.read('binary');
- const config = new _.ParserConfig(new _.DialectContext(metadata));
- const reader = new _.BytecodeReader(binary, config);
- const block = reader.read();
- const format = `MLIR Bytecode v${reader.version}`;
- const producer = reader.producer;
- const model = new mlir.Model(config, format, producer, block, new Map());
- return model;
- }
- default: {
- throw new mlir.Error(`Unsupported MLIR format '${context.type}'.`);
- }
- }
- }
- };
- mlir.Model = class {
- constructor(config, format, producer, block, attributeAliasDefinitions) {
- this.format = format;
- this.producer = producer || '';
- this.modules = [];
- this.functions = [];
- this.metadata = [];
- const modules = [];
- const isFunc = (name) => name.endsWith('.func') || /\.func_v\d+$/.test(name);
- const isModule = (name) => name.endsWith('.module');
- const collectModules = (operations, path, attributes) => {
- let identifier = 0;
- const funcs = [];
- const ops = [];
- for (const op of operations) {
- if (isFunc(op.name.getStringRef())) {
- funcs.push(op);
- } else if (isModule(op.name.getStringRef())) {
- let name = op.getAttr('sym_name');
- name = name ? name.value : `$${identifier++}`;
- const modulePath = [...path, name];
- for (const region of op.regions || []) {
- for (const blk of region.blocks || []) {
- collectModules(blk.operations || [], modulePath, op.getAttrDictionary());
- }
- }
- } else {
- ops.push(op);
- }
- }
- if (funcs.length > 0 || ops.length > 0) {
- let name = null;
- if (attributes.get('sym_name')) {
- name = attributes.get('sym_name');
- name = `@${name.value}`;
- }
- modules.push({ path, symName: name, funcs, ops, attributes });
- }
- };
- collectModules(block.operations, [], new Map());
- const formatPrefix = (path, symName) => {
- if (symName) {
- return symName;
- }
- if (modules.length !== 1 && path.length > 0) {
- return path.map((path) => `${path}`).join('::');
- }
- return '';
- };
- const functions = new Map();
- let identifier = 0;
- for (const module of modules) {
- const prefix = formatPrefix(module.path, module.symName);
- for (const func of module.funcs) {
- const sym_name = func.getAttr('sym_name');
- const base = sym_name ? sym_name.value : `$${identifier}`;
- identifier++;
- const name = prefix ? `${prefix}::@${base}` : `@${base}`;
- functions.set(name, { func, prefix, base, module });
- }
- }
- const context = new mlir.Context(functions);
- for (const [name, info] of functions) {
- const graph = context.graph(info.func, name);
- this.functions.push(graph);
- }
- for (const module of modules) {
- if (module.ops.length > 0 || module.attributes.size > 0) {
- const name = formatPrefix(module.path, module.symName) || '';
- const opName = _.RegisteredOperationName.lookup('builtin.module', config.context);
- const state = new _.OperationState(null, opName);
- state.attributes = module.attributes;
- state.regions = [{ blocks: [{ operations: module.ops, arguments: [] }] }];
- const op = _.Operation.create(state);
- const graph = context.graph(op, name);
- this.modules.push(graph);
- }
- }
- for (const [name, attribute] of attributeAliasDefinitions) {
- let value = attribute.type;
- if (!value) {
- value = typeof attribute.value === 'string' ? attribute.value : JSON.stringify(attribute.value);
- }
- const metadata = new mlir.Argument(name, value, 'attribute');
- this.metadata.push(metadata);
- }
- }
- };
- mlir.Graph = class {
- constructor(func, context, name) {
- this.name = name || '';
- if (!name && func.attributes.has('sym_name')) {
- const sym_name = func.attributes.get('sym_name');
- this.name = sym_name.value;
- }
- this.type = 'graph';
- if (func.name === 'func' || func.name.getStringRef().endsWith('.func') || /\.func_v\d+$/.test(func.name.getStringRef())) {
- this.type = 'function';
- }
- this.inputs = [];
- this.outputs = [];
- this.nodes = [];
- this.metadata = [];
- const tensors = new Map();
- if (func.attributes.has('function_type')) {
- const function_type = func.attributes.get('function_type');
- const args = func.regions && func.regions[0] && func.regions[0].blocks && func.regions[0].blocks[0] && func.regions[0].blocks[0].arguments ? func.regions[0].blocks[0].arguments : [];
- const inputs = function_type.type.inputs;
- const results = function_type.type.results;
- for (let i = 0; i < inputs.length; i++) {
- const input = inputs[i];
- const name = args[i] && args[i].name ? args[i].name : `%arg${i}`;
- const type = mlir.Utility.valueType(input.type || input);
- const value = new mlir.Value(name, type, '', null);
- const argument = new mlir.Argument(name, [value]);
- this.inputs.push(argument);
- }
- for (let i = 0; i < results.length; i++) {
- const output = results[i];
- const name = output.value || i.toString();
- const type = mlir.Utility.valueType(output.type);
- const valueName = output.value || output.name || `%result${i}`;
- const value = new mlir.Value(valueName, type, '', null);
- const argument = new mlir.Argument(name, [value]);
- this.outputs.push(argument);
- }
- }
- const values = new Map();
- values.map = (name) => {
- if (!values.has(name)) {
- values.set(name, { name, to: [], from: [] });
- }
- return values.get(name);
- };
- const operations = [];
- for (const region of func.regions) {
- for (const block of region.blocks) {
- for (const op of block.operations) {
- const operation = {
- identifier: op.identifier,
- name: op.name,
- label: op.label,
- attributes: op.getAttrDictionary(),
- operands: [],
- results: [],
- inputs: [],
- outputs: [],
- regions: op.regions,
- delete: false,
- };
- const opMetadata = op.name.metadata;
- const operands = op.operands;
- let lastVariadicIndex = -1;
- let lastVariadicName = null;
- if (opMetadata && opMetadata.operands) {
- for (let j = opMetadata.operands.length - 1; j >= 0; j--) {
- const metaOp = opMetadata.operands[j];
- if (metaOp.type && metaOp.type.name === 'Variadic') {
- lastVariadicIndex = j;
- lastVariadicName = metaOp.name;
- break;
- }
- }
- }
- for (let i = 0; i < operands.length; i++) {
- const input = op.operands[i];
- let inputName = null;
- const isVariadicOverflow = lastVariadicIndex >= 0 && i >= lastVariadicIndex;
- if (opMetadata && opMetadata.operands && opMetadata.operands[i]) {
- inputName = opMetadata.operands[i].name;
- } else if (isVariadicOverflow) {
- inputName = lastVariadicName;
- } else {
- inputName = input.name || i.toString();
- }
- if (typeof input.name !== 'string' || !input.name) {
- throw new mlir.Error(`Invalid operand name '${JSON.stringify(input.name)}'.`);
- }
- const value = values.map(input.name);
- value.to.push(operation);
- const arg = { name: input.name, type: input.type };
- operation.operands.push(arg);
- if (isVariadicOverflow && operation.inputs.length > 0 && operation.inputs[operation.inputs.length - 1].name === inputName) {
- operation.inputs[operation.inputs.length - 1].value.push(arg);
- } else {
- operation.inputs.push({ name: inputName, value: [arg] });
- }
- }
- const results = op.results;
- let lastVariadicResultIndex = -1;
- let lastVariadicResultName = null;
- if (opMetadata && opMetadata.results) {
- for (let j = opMetadata.results.length - 1; j >= 0; j--) {
- const metaRes = opMetadata.results[j];
- if (metaRes.type && metaRes.type.name === 'Variadic') {
- lastVariadicResultIndex = j;
- lastVariadicResultName = metaRes.name;
- break;
- }
- }
- }
- for (let i = 0; i < results.length; i++) {
- const output = results[i];
- if (!output.name) {
- continue;
- }
- const value = values.map(output.name);
- value.type = mlir.Utility.valueType(output.type);
- value.from.push(operation);
- let outputName = null;
- const isVariadicOverflow = lastVariadicResultIndex >= 0 && i >= lastVariadicResultIndex;
- if (opMetadata && opMetadata.results && opMetadata.results[i]) {
- outputName = opMetadata.results[i].name;
- } else if (isVariadicOverflow) {
- outputName = lastVariadicResultName;
- } else {
- outputName = output.name;
- }
- operation.results.push(value);
- if (isVariadicOverflow && operation.outputs.length > 0 && operation.outputs[operation.outputs.length - 1].name === outputName) {
- operation.outputs[operation.outputs.length - 1].value.push(value);
- } else {
- operation.outputs.push({
- name: outputName,
- value: [value]
- });
- }
- }
- operations.push(operation);
- }
- }
- }
- const constantMap = new Map();
- const constantTypes = new Set([
- 'tosa.const', 'stablehlo.constant', 'arith.constant',
- 'mhlo.constant', 'torch.constant.tensor', 'onnx.Constant'
- ]);
- for (const op of operations) {
- if (constantTypes.has(op.name.getStringRef()) &&
- op.operands.length === 0 &&
- op.attributes.size === 1 &&
- op.results.length === 1) {
- const result = op.results[0];
- if (result.to && result.to.length === 1) {
- if (result.to[0].name.getStringRef().endsWith('.return')) {
- continue;
- }
- const valueAttr = op.attributes.get('value') || op.attributes.get('values');
- if ((valueAttr instanceof _.DenseElementsAttr || valueAttr instanceof _.DenseResourceElementsAttr) && valueAttr.value !== null && valueAttr.type && valueAttr.type.toString().startsWith('tensor<')) {
- const type = mlir.Utility.valueType(valueAttr.type);
- if (type instanceof mlir.TensorType) {
- constantMap.set(result.name, new mlir.Tensor(type, valueAttr.value));
- op.delete = true;
- }
- }
- }
- }
- }
- const torchConstantMap = new Map();
- for (const op of operations) {
- const opName = op.name.getStringRef();
- if (opName === 'torch.constant.int' ||
- opName === 'torch.constant.bool' ||
- opName === 'torch.constant.float' ||
- opName === 'torch.constant.str' ||
- opName === 'torch.constant.none') {
- if (op.operands.length === 0 &&
- op.results.length === 1) {
- const result = op.results[0];
- if (result.to && result.to.length === 1) {
- let value = null;
- let type = null;
- const attr = op.attributes.get('value');
- const attrValue = attr && typeof attr === 'object' ? attr.value : attr;
- if (opName === 'torch.constant.int') {
- value = attrValue === undefined ? 0 : attrValue;
- type = 'int64';
- } else if (opName === 'torch.constant.bool') {
- value = attrValue === undefined ? false : attrValue;
- type = 'boolean';
- } else if (opName === 'torch.constant.float') {
- value = attrValue === undefined ? 0.0 : attrValue;
- type = 'float64';
- } else if (opName === 'torch.constant.str') {
- value = attrValue === undefined ? '' : attrValue;
- type = 'string';
- } else if (opName === 'torch.constant.none') {
- value = null;
- type = 'none';
- }
- torchConstantMap.set(result.name, { value, type });
- op.delete = true;
- }
- }
- }
- }
- for (const op of operations) {
- if (op.name === 'torch.prim.ListConstruct' &&
- op.results.length === 1) {
- const result = op.results[0];
- if (result.to && result.to.length === 1) {
- const inputValues = [];
- let allConstant = true;
- for (const operand of op.operands) {
- if (torchConstantMap.has(operand.name)) {
- inputValues.push(torchConstantMap.get(operand.name).value);
- } else {
- allConstant = false;
- break;
- }
- }
- if (allConstant) {
- torchConstantMap.set(result.name, { value: inputValues, type: 'list' });
- op.delete = true;
- }
- }
- }
- }
- const tensor = (arg) => {
- if (!tensors.has(arg.name)) {
- const initializer = constantMap.get(arg.name) || null;
- let type = null;
- if (arg.type instanceof mlir.TensorType) {
- type = arg.type;
- } else if (arg.type) {
- type = mlir.Utility.valueType(arg.type);
- }
- tensors.set(arg.name, new mlir.Value(arg.name, type, null, initializer));
- }
- return tensors.get(arg.name);
- };
- for (const input of this.inputs) {
- for (const arg of input.value) {
- if (!tensors.has(arg.name)) {
- tensors.set(arg.name, arg);
- }
- }
- }
- const returnOp = operations.find((op) => op.name.getStringRef().endsWith('.return'));
- if (returnOp) {
- for (let i = 0; i < this.outputs.length && i < returnOp.operands.length; i++) {
- const returnValue = returnOp.operands[i];
- if (returnValue && typeof returnValue.name === 'string' && returnValue.name.startsWith('%')) {
- const output = this.outputs[i];
- const returnType = mlir.Utility.valueType(returnValue.type);
- const initializer = constantMap.get(returnValue.name) || null;
- output.value[0] = new mlir.Value(returnValue.name, returnType, '', initializer);
- }
- }
- returnOp.delete = true;
- }
- for (const output of this.outputs) {
- for (let i = 0; i < output.value.length; i++) {
- const arg = output.value[i];
- if (tensors.has(arg.name)) {
- output.value[i] = tensors.get(arg.name);
- } else {
- tensors.set(arg.name, arg);
- }
- }
- }
- for (const op of operations.filter((op) => !op.delete)) {
- const node = new mlir.Node(op, context, tensor, torchConstantMap);
- this.nodes.push(node);
- }
- for (const [name, value] of func.attributes) {
- if (name === 'sym_name' || name === 'function_type') {
- continue;
- }
- const metadata = new mlir.Argument(name, value, 'attribute');
- this.metadata.push(metadata);
- }
- }
- };
- mlir.Argument = class {
- constructor(name, value, type = null) {
- this.name = name;
- this.value = value;
- this.type = type;
- if (this.type) {
- const typeStr = this.type instanceof _.Type ? this.type.toString() : this.type;
- switch (typeStr) {
- case 'i64': case 'si64': this.type = 'int64'; break;
- case 'i48': case 'si48': this.type = 'int48'; break;
- case 'i32': case 'si32': this.type = 'int32'; break;
- case 'i16': case 'si16': this.type = 'int16'; break;
- case 'i8': case 'si8': this.type = 'int8'; break;
- case 'i1': this.type = 'int1'; break;
- case 'f32': case 'float32': this.type = 'float32'; break;
- case 'f64': case 'float64': this.type = 'float64'; break;
- case 'f16': this.type = 'float16'; break;
- case 'f80': this.type = 'float80'; break;
- case 'f128': this.type = 'float128'; break;
- case null:
- case 'attribute':
- case 'boolean':
- case 'string':
- case 'int64':
- case 'int32':
- case 'int16':
- case 'int8':
- case 'float16':
- case 'tensor':
- case 'type':
- case 'dense':
- case 'function':
- case 'symbol':
- case 'graph':
- case 'list':
- case 'none':
- break;
- default:
- if (/^[usi]i?[0-9]+$/.test(typeStr) || /^f[0-9]+$/.test(typeStr) ||
- /^f\d+E\d+M\d+/.test(typeStr) ||
- typeStr === 'bf16' || typeStr === 'tf32' || typeStr === 'index' || typeStr === 'none' ||
- typeStr === 'unit' || typeStr.startsWith('!') || typeStr.startsWith('tensor<') ||
- typeStr.startsWith('memref<') || typeStr.startsWith('vector<')) {
- this.type = typeStr;
- break;
- }
- throw new mlir.Error(`Unsupported argument type '${typeStr}'.`);
- }
- }
- }
- };
- mlir.Value = class {
- constructor(name, type, description, initializer) {
- if (typeof name !== 'string') {
- throw new mlir.Error(`Invalid value identifier '${JSON.stringify(name)}'.`);
- }
- this.name = name;
- this.type = !type && initializer ? initializer.type : type;
- this.description = description || null;
- this.initializer = initializer || null;
- }
- };
- mlir.Node = class {
- constructor(op, context, tensor, torchConstantMap) {
- if (!op.name) {
- throw new mlir.Error('Undefined node type.');
- }
- this.name = '';
- this.type = { ...op.name.getRegisteredInfo()?.metadata };
- this.type.name = op.label || op.identifier || '';
- this.type.identifier = op.name.getStringRef() || '';
- this.inputs = [];
- this.outputs = [];
- this.attributes = [];
- this.blocks = [];
- torchConstantMap = torchConstantMap || new Map();
- const segmentSizes = op.attributes && op.attributes.get('operandSegmentSizes');
- const operandMeta = this.type && this.type.operands;
- const operands = op.inputs || [];
- if (segmentSizes && operandMeta && Array.isArray(segmentSizes) && segmentSizes.length === operandMeta.length) {
- let offset = 0;
- for (let i = 0; i < segmentSizes.length; i++) {
- const size = segmentSizes[i];
- const name = operandMeta[i].name;
- for (let j = 0; j < size; j++) {
- if (offset + j < operands.length) {
- operands[offset + j] = { ...operands[offset + j], name };
- }
- }
- offset += size;
- }
- }
- const operandGroups = new Map();
- const operandOrder = [];
- for (const input of operands) {
- if (!operandGroups.has(input.name)) {
- operandGroups.set(input.name, []);
- operandOrder.push(input.name);
- }
- operandGroups.get(input.name).push(input);
- }
- for (const name of operandOrder) {
- const inputs = operandGroups.get(name);
- let argument = null;
- if (inputs.length === 1) {
- const [input] = inputs;
- if (Array.isArray(input.value) && input.value.length === 1) {
- const val = input.value[0];
- if (val && typeof val.name === 'string' && torchConstantMap.has(val.name)) {
- const constant = torchConstantMap.get(val.name);
- argument = new mlir.Argument(input.name, constant.value, constant.type);
- this.inputs.push(argument);
- continue;
- }
- }
- if (input.type) {
- const typeStr = input.type instanceof _.Type ? input.type.toString() : input.type;
- if (typeStr.startsWith('tensor<')) {
- const type = mlir.Utility.valueType(typeStr);
- const value = new mlir.Tensor(type, input.value);
- argument = new mlir.Argument(input.name, value, 'tensor');
- } else {
- argument = new mlir.Argument(input.name, input.value, input.type);
- }
- } else if (Array.isArray(input.value) && !input.value.every((value) => typeof value.name === 'string' && value.name.startsWith('%'))) {
- argument = new mlir.Argument(input.name, input.value, input.type || 'attribute');
- } else if (Array.isArray(input.value)) {
- argument = new mlir.Argument(input.name, input.value.map((arg) => tensor(arg)));
- } else {
- argument = new mlir.Argument(input.name, input.value, input.type || 'attribute');
- }
- } else {
- let allConstants = true;
- const constantValues = [];
- for (const input of inputs) {
- if (Array.isArray(input.value)) {
- for (const arg of input.value) {
- if (arg && typeof arg.name === 'string' && torchConstantMap.has(arg.name)) {
- constantValues.push(torchConstantMap.get(arg.name).value);
- } else {
- allConstants = false;
- break;
- }
- }
- } else {
- allConstants = false;
- }
- if (!allConstants) {
- break;
- }
- }
- if (allConstants && constantValues.length > 0) {
- argument = new mlir.Argument(name, constantValues, 'list');
- } else {
- const values = [];
- for (const input of inputs) {
- if (Array.isArray(input.value)) {
- values.push(...input.value.map((arg) => tensor(arg)));
- } else {
- values.push(tensor({ name: input.value, type: input.type }));
- }
- }
- argument = new mlir.Argument(name, values);
- }
- }
- this.inputs.push(argument);
- }
- for (const output of op.outputs || []) {
- const argument = new mlir.Argument(output.name, output.value.map((arg) => tensor(arg)));
- this.outputs.push(argument);
- }
- if (op.attributes) {
- for (const [name, attr] of op.attributes) {
- let value = attr;
- let type = null;
- if (attr instanceof _.SymbolRefAttr && context) {
- const graph = context.function(`${value.value}`);
- if (graph) {
- value = graph;
- type = 'function';
- }
- } else if (attr instanceof _.DenseElementsAttr) {
- value = new mlir.Tensor(mlir.Utility.valueType(attr.type), attr.value);
- type = 'tensor';
- } else if (attr instanceof _.DenseResourceElementsAttr) {
- value = new mlir.Tensor(mlir.Utility.valueType(attr.type), attr.value);
- type = 'tensor';
- } else if (attr instanceof _.SparseElementsAttr) {
- value = new mlir.Tensor(mlir.Utility.valueType(attr.type), attr.values);
- type = 'tensor';
- } else if (attr instanceof _.DenseArrayAttr) {
- value = attr.value;
- } else if (Array.isArray(attr)) {
- value = attr;
- } else if (attr) {
- value = attr.toString();
- }
- const attribute = new mlir.Argument(name, value, type || 'attribute');
- this.attributes.push(attribute);
- }
- }
- if (op.regions && op.regions.length > 0) {
- const opMetadata = this.type;
- for (let i = 0; i < op.regions.length; i++) {
- const region = op.regions[i];
- if (region.blocks && region.blocks.length > 0) {
- const name = (opMetadata.regions && opMetadata.regions[i] ? opMetadata.regions[i].name : null) || i.toString();
- const blockName = region.blocks[0].name || '';
- const func = { name: 'func', attributes: new Map(), regions: [region] };
- const graph = new mlir.Graph(func, context, blockName);
- const argument = new mlir.Argument(name, graph, 'graph');
- this.blocks.push(argument);
- }
- }
- }
- }
- };
- mlir.Tensor = class {
- constructor(type, data) {
- this.type = type;
- this.values = data;
- this.encoding = data instanceof Uint8Array ? '<' : '|';
- }
- };
- mlir.TensorType = class {
- constructor(dataType, shape) {
- this.dataType = mlir.Utility.dataType(dataType); // string
- this.shape = shape || new mlir.TensorShape([]); // mlir.TensorShape
- }
- toString() {
- return this.dataType + this.shape.toString();
- }
- };
- mlir.TensorShape = class {
- constructor(dimensions) {
- this.dimensions = dimensions;
- }
- toString() {
- if (!this.dimensions || this.dimensions.length === 0) {
- return '';
- }
- return `[${this.dimensions.map((dimension) => dimension.toString()).join(',')}]`;
- }
- };
- mlir.Context = class {
- constructor(functions) {
- this._functions = functions; // Map of fullName -> {func, prefix, base, module}
- this._graphs = new Map();
- this._constructing = new Set();
- }
- graph(module, name) {
- if (!this._graphs.has(name)) {
- this._constructing.add(name);
- const graph = new mlir.Graph(module, this, name);
- this._graphs.set(name, graph);
- this._constructing.delete(name);
- }
- return this._graphs.get(name);
- }
- function(name) {
- if (this._graphs.has(name)) {
- return this._graphs.get(name);
- }
- if (this._constructing.has(name)) {
- return { name, type: 'function', nodes: [], inputs: [], outputs: [] };
- }
- if (this._functions.has(name)) {
- const info = this._functions.get(name);
- return this.graph(info.func, name);
- }
- for (const [fullName, info] of this._functions) {
- if (info.base === name) {
- if (this._graphs.has(fullName)) {
- return this._graphs.get(fullName);
- }
- if (this._constructing.has(fullName)) {
- return { name: fullName, type: 'function', nodes: [], inputs: [], outputs: [] };
- }
- return this.graph(info.func, fullName);
- }
- }
- return null;
- }
- };
- mlir.Utility = class {
- static dataType(value) {
- if (value instanceof _.ComplexType) {
- const elementType = mlir.Utility.dataType(value.elementType);
- return `complex<${elementType}>`;
- }
- if (value instanceof _.Type) {
- value = value.toString();
- }
- switch (value) {
- case 'index': return 'int64';
- case 'f16': return 'float16';
- case 'f32': return 'float32';
- case 'f64': return 'float64';
- case 'f80': return 'float80';
- case 'f128': return 'float128';
- case 'bf16': return 'bfloat16';
- case 'fp8': return 'float8';
- case 'fp8e4m3': return 'float8e4m3';
- case 'fp8_e4m3': return 'float8e4m3';
- case 'fp8e4m3fn': return 'float8e4m3fn';
- case 'fp8e5m2': return 'float8e5m2';
- case 'fp8_e5m2': return 'float8e5m2';
- case 'f4E2M1FN': return 'float4e2m1fn';
- case 'f6E2M3FN': return 'float6e2m3fn';
- case 'f6E3M2FN': return 'float6e3m2fn';
- case 'f8E3M4': return 'float8e3m4';
- case 'f8E4M3': return 'float8e4m3';
- case 'f8E4M3B11FNUZ': return 'float8e4m3b11fnuz';
- case 'f8E4M3FN': return 'float8e4m3fn';
- case 'f8E4M3FNUZ': return 'float8e4m3fnuz';
- case 'f8E5M2': return 'float8e5m2';
- case 'f8E5M2FNUZ': return 'float8e5m2fnuz';
- case 'f8E8M0FNU': return 'float8e8m0fnu';
- case 'float8': return 'float8';
- case 'tf32': return 'tf32';
- case 'i1': return 'int1';
- case 'i2': return 'int2';
- case 'i4': return 'int4';
- case 'i8': return 'int8';
- case 'i16': return 'int16';
- case 'i32': return 'int32';
- case 'i48': return 'int48';
- case 'i64': return 'int64';
- case 'si8': return 'int8';
- case 'si16': return 'int16';
- case 'si32': return 'int32';
- case 'si64': return 'int64';
- case 'ui1': return 'uint1';
- case 'ui2': return 'uint2';
- case 'ui4': return 'uint4';
- case 'ui8': return 'uint8';
- case 'ui16': return 'uint16';
- case 'ui32': return 'uint32';
- case 'ui64': return 'uint64';
- case 'b8': return 'int8';
- case 'unk': return 'unk'; // torch dialect unknown dtype
- case '!tf_type.string': return 'string';
- case '!tosa.mxint8': return 'int8';
- case '!onnx.String': return 'string';
- default:
- if (value && value.startsWith('!')) {
- return value;
- }
- if (value && value.startsWith('vector<') && value.endsWith('>')) {
- return value;
- }
- if (value && value.startsWith('memref<') && value.endsWith('>')) {
- return value;
- }
- if (value && value.startsWith('tuple<') && value.endsWith('>')) {
- return value;
- }
- if (value && value.startsWith('complex<') && value.endsWith('>')) {
- const elementTypeStr = value.substring(8, value.length - 1);
- const convertedElementType = mlir.Utility.dataType(elementTypeStr);
- return `complex<${convertedElementType}>`;
- }
- if (value && /^[su]?i[0-9]+$/.test(value)) {
- const match = value.match(/^(s|u)?i([0-9]+)$/);
- if (match) {
- const [, signed, widthStr] = match;
- const width = parseInt(widthStr, 10);
- if (signed === 'u') {
- return `uint${width}`;
- } else if (signed === 's') {
- return `int${width}`;
- }
- return `int${width}`;
- }
- }
- throw new mlir.Error(`Unknown data type '${value}'.`);
- }
- }
- static valueType(type) {
- if (type === undefined) {
- return null;
- }
- const typeStr = type instanceof _.Type ? type.toString() : type;
- if (typeStr.startsWith('!') && !typeStr.startsWith('!torch.vtensor<')) {
- return typeStr;
- }
- if (typeStr.startsWith('tensor<') && typeStr.endsWith('>')) {
- const spec = typeStr.substring(7, typeStr.length - 1).trim();
- if (spec.startsWith('!')) {
- return mlir.Utility.valueType(spec);
- }
- let i = 0;
- const shape = [];
- while (i < spec.length) {
- if (spec[i] === '?' || spec[i] === '*') {
- shape.push('?');
- i++;
- } else if (/[0-9]/.test(spec[i])) {
- let numStr = '';
- while (i < spec.length && /[0-9]/.test(spec[i])) {
- numStr += spec[i];
- i++;
- }
- const dim = parseInt(numStr, 10);
- if (isNaN(dim)) {
- shape.push('?');
- } else {
- shape.push(dim);
- }
- } else {
- break;
- }
- if (i < spec.length && spec[i] === 'x') {
- i++;
- } else {
- break;
- }
- }
- let dataType = spec.substring(i);
- // Find encoding comma, but skip commas inside angle brackets
- let depth = 0;
- let encodingIndex = -1;
- for (let j = 0; j < dataType.length; j++) {
- if (dataType[j] === '<') {
- depth++;
- } else if (dataType[j] === '>') {
- depth--;
- } else if (dataType[j] === ',' && depth === 0) {
- encodingIndex = j;
- break;
- }
- }
- if (encodingIndex !== -1) {
- dataType = dataType.substring(0, encodingIndex).trim();
- }
- return new mlir.TensorType(dataType, new mlir.TensorShape(shape));
- }
- if (typeStr.startsWith('!torch.vtensor<') && typeStr.endsWith('>')) {
- const spec = typeStr.substring(15, typeStr.length - 1);
- let shape = null;
- let dataType = null;
- if (spec.startsWith('[')) {
- const bracketEnd = spec.indexOf(']');
- const shapeStr = spec.substring(0, bracketEnd + 1);
- const jsonStr = shapeStr.replace(/\?/g, '"?"');
- shape = JSON.parse(jsonStr);
- const rest = spec.substring(bracketEnd + 1);
- if (rest.startsWith(',')) {
- const parts = rest.substring(1).split(',');
- dataType = parts[0].trim();
- }
- } else if (spec.startsWith('*')) {
- if (spec.includes(',')) {
- const parts = spec.split(',');
- dataType = parts[1].trim();
- }
- } else {
- const parts = spec.split(',');
- dataType = parts[0].trim();
- }
- return new mlir.TensorType(dataType, shape ? new mlir.TensorShape(shape) : null);
- }
- if (typeStr.startsWith('tuple<') && typeStr.endsWith('>')) {
- return typeStr;
- }
- return typeStr;
- }
- };
- _.Block = class {
- constructor() {
- this.operations = [];
- }
- };
- _.OperationState = class {
- constructor(location, name) {
- this.location = location;
- this.name = name;
- if (this.name instanceof _.OperationName === false) {
- throw new mlir.Error('Invalid operation name.');
- }
- this.identifier = name.identifier || name.getStringRef();
- delete name.identifier;
- this.attributes = new Map();
- this.operands = [];
- this.types = [];
- this.regions = [];
- this.propertiesAttr = null;
- }
- get op() {
- return this.name.getStringRef(); // Workaround
- }
- addRegion() {
- const region = {};
- this.regions.push(region);
- return region;
- }
- addTypes(newTypes) {
- if (!Array.isArray(newTypes) || !newTypes.every((type) => type instanceof _.Type)) {
- throw new mlir.Error(`Invalid types '${JSON.stringify(newTypes.filter((type) => type instanceof _.Type === false))}'.`);
- }
- for (const type of newTypes) {
- this.types.push(type);
- }
- }
- addAttribute(name, value) {
- if (typeof name !== 'string' || name.length === 0) {
- throw new mlir.Error(`Invalid attribute name '${JSON.stringify(name)}'.`);
- }
- this.attributes.set(name, value);
- }
- getAttr(name) {
- if (this.propertiesAttr instanceof _.DictionaryAttr) {
- const value = this.propertiesAttr.get(name);
- if (value !== undefined) {
- return value;
- }
- }
- return this.attributes.get(name);
- }
- getAttrDictionary() {
- if (this.propertiesAttr instanceof _.DictionaryAttr) {
- const result = new Map(this.attributes);
- for (const [name, value] of this.propertiesAttr.value) {
- result.set(name, value);
- }
- return result;
- }
- return this.attributes;
- }
- };
- _.OperationName = class {
- constructor(dialect, name) {
- this.dialect = dialect;
- this.name = name;
- }
- getStringRef() {
- return this.name;
- }
- getRegisteredInfo() {
- return null;
- }
- };
- _.RegisteredOperationName = class extends _.OperationName {
- constructor(dialect, name, metadata) {
- super(dialect, name);
- this.metadata = metadata;
- if (metadata.assemblyFormat) {
- const parser = new _.AssemblyFormatParser(metadata);
- this.directives = parser.parse();
- }
- }
- static lookup(name, context) {
- const index = name.indexOf('.');
- if (index !== -1) {
- const dialectName = name.substring(0, index);
- const dialect = context.getOrLoadDialect(dialectName);
- if (dialect) {
- return dialect.getOperation(name);
- }
- }
- return null;
- }
- getRegisteredInfo() {
- return this;
- }
- hasTrait(type) {
- if (this.metadata && Array.isArray(this.metadata.traits)) {
- for (const trait of this.metadata.traits) {
- if (trait.type && trait.type.name === type) {
- return true;
- }
- }
- }
- return false;
- }
- };
- _.Operation = class {
- static create(state) {
- return new _.Operation(state);
- }
- constructor(state) {
- this.name = state.name; // registered operation name
- this.identifier = state.identifier; // original parsed identifier
- this.label = state.label; // human-readable name
- this.attributes = state.attributes;
- this.operands = state.operands;
- this.regions = state.regions;
- this.propertiesAttr = state.propertiesAttr;
- this.loc = state.loc;
- this.results = [];
- if (Array.isArray(state.types)) {
- for (let i = 0; i < state.types.length; i++) {
- const result = new _.OpResult(this, i, state.types[i]);
- this.results.push(result);
- }
- }
- }
- getAttr(name) {
- if (this.propertiesAttr instanceof _.DictionaryAttr) {
- const value = this.propertiesAttr.get(name);
- if (value !== undefined) {
- return value;
- }
- }
- return this.attributes.get(name);
- }
- getAttrDictionary() {
- if (this.propertiesAttr instanceof _.DictionaryAttr) {
- const result = new Map(this.attributes);
- for (const [key, value] of this.propertiesAttr.value) {
- result.set(key, value);
- }
- return result;
- }
- return this.attributes;
- }
- };
- _.UnresolvedOperand = class {
- constructor(location, name, number) {
- this.location = location;
- this.name = name;
- this.number = number;
- }
- toString() {
- return this.number > 0 ? `${this.name}#${this.number}` : this.name;
- }
- };
- _.Value = class {
- constructor(name, type) {
- this.name = name;
- this.type = type;
- }
- toString() {
- return this.name;
- }
- };
- _.OpResult = class extends _.Value {
- constructor(owner, resultNumber, type) {
- super(null, type);
- this.owner = owner;
- this.resultNumber = resultNumber;
- }
- };
- _.Attribute = class {
- };
- _.TypedAttr = class extends _.Attribute {
- constructor(value, type) {
- super();
- this.value = value;
- this.type = type;
- }
- toString() {
- return this.value;
- }
- };
- _.StringAttr = class extends _.TypedAttr {
- constructor(value, type) {
- super(value, type || new _.PrimitiveType('string'));
- }
- toString() {
- return this.value;
- }
- };
- _.UnitAttr = class extends _.Attribute {
- toString() {
- return '';
- }
- };
- _.IntegerAttr = class extends _.Attribute {
- constructor(type, value) {
- super();
- this.value = value;
- this.type = type;
- }
- toString() {
- return this.value.toString();
- }
- };
- _.BoolAttr = class extends _.Attribute {
- constructor(value) {
- super();
- this.value = value;
- this.type = new _.IntegerType('i1');
- }
- toString() {
- return this.value ? 'true' : 'false';
- }
- };
- _.FloatAttr = class extends _.Attribute {
- constructor(type, value) {
- super();
- this.value = value;
- this.type = type;
- }
- toString() {
- return String(this.value);
- }
- };
- _.AffineMapAttr = class extends _.Attribute {
- constructor(map) {
- super();
- this._map = map;
- }
- get value() {
- return this._map;
- }
- toString() {
- return `affine_map<${this._map.toString()}>`;
- }
- };
- _.IntegerSetAttr = class extends _.Attribute {
- constructor(set) {
- super();
- this._set = set;
- }
- get value() {
- return this._set;
- }
- toString() {
- return `affine_set<${this._set.toString()}>`;
- }
- };
- _.SymbolRefAttr = class extends _.Attribute {
- constructor(rootReference, nestedReferences) {
- super();
- this.rootReference = rootReference;
- this.nestedReferences = nestedReferences;
- }
- get value() {
- if (this.nestedReferences && this.nestedReferences.length > 0) {
- return `${this.rootReference}${this.nestedReferences.map((ref) => `::${ref}`).join('')}`;
- }
- return this.rootReference;
- }
- toString() {
- return this.value;
- }
- };
- _.DenseElementsAttr = class extends _.Attribute {
- constructor(value, type) {
- super();
- this.value = value;
- this.type = type;
- }
- };
- _.SparseElementsAttr = class extends _.Attribute {
- constructor(type, indices, values) {
- super();
- this.type = type;
- this.indices = indices;
- this.values = values;
- }
- };
- _.DenseResourceElementsHandle = class {
- constructor(key, blob = null) {
- this.key = key;
- this.blob = blob;
- }
- };
- _.DenseResourceElementsAttr = class extends _.Attribute {
- constructor(type, handle) {
- super();
- this.type = type;
- this.rawHandle = handle;
- }
- get value() {
- return this.rawHandle ? this.rawHandle.blob : null;
- }
- toString() {
- const key = this.rawHandle ? this.rawHandle.key : 'unknown';
- return `dense_resource<${key}>`;
- }
- };
- _.OpaqueAttr = class extends _.Attribute {
- constructor(dialectName, symbolData, type) {
- super();
- this.dialectName = dialectName;
- this.symbolData = symbolData;
- this.type = type;
- }
- get value() {
- return this.toString();
- }
- toString() {
- if (this.symbolData) {
- return `${this.dialectName}.${this.symbolData}`;
- }
- return this.dialectName;
- }
- };
- _.AffineLowPrecOp = {
- LNoOp: '',
- Add: 'Add',
- Sub: 'Sub'
- };
- _.AffineHighPrecOp = {
- HNoOp: '',
- Mul: 'Mul',
- FloorDiv: 'FloorDiv',
- CeilDiv: 'CeilDiv',
- Mod: 'Mod'
- };
- _.AffineExprKind = {
- Add: 'Add',
- Mul: 'Mul',
- Mod: 'Mod',
- FloorDiv: 'FloorDiv',
- CeilDiv: 'CeilDiv',
- Constant: 'Constant',
- DimId: 'DimId',
- SymbolId: 'SymbolId'
- };
- _.AffineExpr = class {
- constructor(kind) {
- this.kind = kind;
- }
- isSymbolicOrConstant() {
- return this.kind === _.AffineExprKind.Constant || this.kind === _.AffineExprKind.SymbolId;
- }
- toString() {
- return '';
- }
- };
- _.AffineDimExpr = class extends _.AffineExpr {
- constructor(position) {
- super(_.AffineExprKind.DimId);
- this._position = position;
- }
- getPosition() {
- return this._position;
- }
- toString() {
- return `d${this._position}`;
- }
- };
- _.AffineSymbolExpr = class extends _.AffineExpr {
- constructor(position) {
- super(_.AffineExprKind.SymbolId);
- this._position = position;
- }
- getPosition() {
- return this._position;
- }
- toString() {
- return `s${this._position}`;
- }
- };
- _.AffineConstantExpr = class extends _.AffineExpr {
- constructor(value) {
- super(_.AffineExprKind.Constant);
- this.value = value;
- }
- toString() {
- return String(this.value);
- }
- };
- _.AffineBinaryOpExpr = class extends _.AffineExpr {
- constructor(kind, lhs, rhs) {
- super(kind);
- this.lhs = lhs;
- this.rhs = rhs;
- }
- isSymbolicOrConstant() {
- return this.lhs.isSymbolicOrConstant() && this.rhs.isSymbolicOrConstant();
- }
- toString() {
- let op = '';
- switch (this.kind) {
- case _.AffineExprKind.Add: op = ' + '; break;
- case _.AffineExprKind.Mul: op = ' * '; break;
- case _.AffineExprKind.Mod: op = ' mod '; break;
- case _.AffineExprKind.FloorDiv: op = ' floordiv '; break;
- case _.AffineExprKind.CeilDiv: op = ' ceildiv '; break;
- default: throw new Error(`Unexpected affine expression kind '${this.kind}'.`);
- }
- return `${this.lhs.toString()}${op}${this.rhs.toString()}`;
- }
- };
- _.AffineMap = class {
- constructor(numDims, numSymbols, results) {
- this._numDims = numDims;
- this._numSymbols = numSymbols;
- this._results = results; // Array of AffineExpr or string for backward compat
- }
- getNumResults() {
- return this._results.length;
- }
- toString() {
- const dims = [];
- for (let i = 0; i < this._numDims; i++) {
- dims.push(`d${i}`);
- }
- const symbols = [];
- for (let i = 0; i < this._numSymbols; i++) {
- symbols.push(`s${i}`);
- }
- const dimsStr = `(${dims.join(', ')})`;
- const symbolsStr = symbols.length > 0 ? `[${symbols.join(', ')}]` : '';
- const resultsStr = this._results.map((r) => r.toString ? r.toString() : String(r)).join(', ');
- return `${dimsStr}${symbolsStr} -> (${resultsStr})`;
- }
- static get(numDims, numSymbols, results) {
- return new _.AffineMap(numDims, numSymbols, results);
- }
- // Backward compatible string-based parse
- static parse(str) {
- const arrowIdx = str.indexOf('->');
- if (arrowIdx === -1) {
- return new _.AffineMap(0, 0, [str]);
- }
- let resultPart = str.substring(arrowIdx + 2).trim();
- const domainIdx = resultPart.indexOf(', domain:');
- if (domainIdx !== -1) {
- resultPart = resultPart.substring(0, domainIdx).trim();
- }
- // Parse result expressions from the result part (e.g., "(d0, d1 + s0)")
- const results = [];
- if (resultPart.startsWith('(')) {
- let depth = 0;
- let start = 1; // Skip opening paren
- for (let i = 0; i < resultPart.length; i++) {
- const ch = resultPart[i];
- if (ch === '(' || ch === '[') {
- depth++;
- } else if (ch === ')' || ch === ']') {
- depth--;
- if (depth === 0) {
- const expr = resultPart.substring(start, i).trim();
- if (expr) {
- results.push(expr);
- }
- break;
- }
- } else if (ch === ',' && depth === 1) {
- const expr = resultPart.substring(start, i).trim();
- if (expr) {
- results.push(expr);
- }
- start = i + 1;
- }
- }
- } else {
- results.push(resultPart);
- }
- // For backward compat, store string representation of results
- return new _.AffineMap(0, 0, results.length > 0 ? results : [str]);
- }
- };
- // Integer set for affine_set constraints
- _.IntegerSet = class {
- constructor(numDims, numSymbols, constraints, eqFlags) {
- this._numDims = numDims;
- this._numSymbols = numSymbols;
- this._constraints = constraints; // Array of AffineExpr
- this._eqFlags = eqFlags; // Array of bool (true = equality, false = inequality)
- }
- getNumConstraints() {
- return this._constraints.length;
- }
- getConstraints() {
- return this._constraints;
- }
- getConstraint(idx) {
- return this._constraints[idx];
- }
- getEqFlags() {
- return this._eqFlags;
- }
- isEq(idx) {
- return this._eqFlags[idx];
- }
- toString() {
- const dims = [];
- for (let i = 0; i < this._numDims; i++) {
- dims.push(`d${i}`);
- }
- const symbols = [];
- for (let i = 0; i < this._numSymbols; i++) {
- symbols.push(`s${i}`);
- }
- const dimsStr = `(${dims.join(', ')})`;
- const symbolsStr = symbols.length > 0 ? `[${symbols.join(', ')}]` : '';
- const constraintsStr = this._constraints.map((c, i) => {
- const exprStr = c.toString ? c.toString() : String(c);
- return this._eqFlags[i] ? `${exprStr} == 0` : `${exprStr} >= 0`;
- }).join(', ');
- return `${dimsStr}${symbolsStr} : (${constraintsStr})`;
- }
- static get(numDims, numSymbols, constraints, eqFlags) {
- return new _.IntegerSet(numDims, numSymbols, constraints, eqFlags);
- }
- };
- _.IndexingMap = class {
- constructor(affineMap) {
- this._affineMap = affineMap;
- }
- GetAffineMap() {
- return this._affineMap;
- }
- };
- _.IndexingMapAttr = class extends _.Attribute {
- constructor(indexingMap) {
- super();
- this._indexingMap = indexingMap;
- }
- get value() {
- return this.toString();
- }
- getIndexingMap() {
- return this._indexingMap;
- }
- static parse(parser /*, type */) {
- parser.parseLess();
- const indexingMap = _.IndexingMapAttr.parseChainOfStringsAsIndexingMap(parser);
- parser.parseGreater();
- return new _.IndexingMapAttr(indexingMap);
- }
- static parseChainOfStringsAsIndexingMap(parser) {
- let indexingMapStr = '';
- let str = parser.parseOptionalString();
- while (str !== null) {
- indexingMapStr += str;
- str = parser.parseOptionalString();
- }
- return _.IndexingMapAttr.parseIndexingMap(indexingMapStr);
- }
- static parseIndexingMap(str) {
- const domainIdx = str.indexOf(', domain:');
- const mapStr = domainIdx >= 0 ? str.substring(0, domainIdx) : str.replace(/,$/, '');
- const affineMap = _.AffineMap.parse(mapStr);
- return new _.IndexingMap(affineMap);
- }
- toString() {
- return `#xla.indexing_map<"${this._indexingMap.GetAffineMap().toString()}">`;
- }
- };
- _.ArrayAttr = class extends _.Attribute {
- constructor(elements) {
- super();
- this.elements = elements; // Array of Attribute objects
- }
- get value() {
- return this.elements.map((e) => e && e.value !== undefined ? e.value : e);
- }
- toString() {
- return `${this.elements.map((e) => e && e.toString ? e.toString() : String(e)).join(', ')}`;
- }
- };
- _.DenseI64ArrayAttr = class extends _.Attribute {
- constructor(values) {
- super();
- this._values = values; // Array of i64 integers
- }
- get value() {
- return this._values;
- }
- toString() {
- return `[${this._values.join(', ')}]`;
- }
- static parse(parser) {
- if (parser.getToken().isNot(_.Token.l_square)) {
- return null;
- }
- parser.consumeToken(_.Token.l_square);
- const values = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- const value = parser.parseOptionalInteger();
- if (value === null) {
- break;
- }
- values.push(value);
- parser.parseOptionalComma();
- }
- parser.consumeToken(_.Token.r_square);
- return new _.DenseI64ArrayAttr(values);
- }
- };
- _.DictionaryAttr = class extends _.Attribute {
- constructor(value) {
- super();
- this._value = value; // Map of name -> Attribute
- }
- get value() {
- return this._value;
- }
- get(name) {
- return this._value.get(name);
- }
- toString() {
- const entries = Array.from(this._value.entries())
- .map(([k, v]) => `${k} = ${v && v.toString ? v.toString() : String(v)}`);
- return `{${entries.join(', ')}}`;
- }
- };
- _.DenseArrayAttr = class extends _.Attribute {
- constructor(type, size, data) {
- super();
- this.type = type;
- this.size = size;
- if (data instanceof Uint8Array) {
- this.values = null;
- this.data = data;
- } else {
- this.values = data;
- }
- }
- get value() {
- if (this.values === null) {
- const blob = this.data;
- const typeStr = this.type.toString();
- const view = new DataView(blob.buffer, blob.byteOffset, blob.length);
- const values = [];
- if (typeStr.startsWith('i') || typeStr.startsWith('si') || typeStr.startsWith('ui') || typeStr === 'index') {
- const match = typeStr.match(/[su]?i(\d+)/);
- const bitWidth = match ? parseInt(match[1], 10) : 64;
- const unsigned = typeStr.startsWith('ui');
- const byteWidth = Math.ceil(bitWidth / 8);
- const size = blob.length / byteWidth;
- for (let i = 0; i < size; i++) {
- if (bitWidth <= 8) {
- values.push(unsigned ? view.getUint8(i * byteWidth) : view.getInt8(i * byteWidth));
- } else if (bitWidth <= 16) {
- values.push(unsigned ? view.getUint16(i * byteWidth, true) : view.getInt16(i * byteWidth, true));
- } else if (bitWidth <= 32) {
- values.push(unsigned ? view.getUint32(i * byteWidth, true) : view.getInt32(i * byteWidth, true));
- } else {
- values.push(unsigned ? view.getBigUint64(i * byteWidth, true) : view.getBigInt64(i * byteWidth, true));
- }
- }
- } else if (typeStr === 'f32') {
- const size = blob.length / 4;
- for (let i = 0; i < size; i++) {
- values.push(view.getFloat32(i * 4, true));
- }
- } else if (typeStr === 'f64') {
- const size = blob.length / 8;
- for (let i = 0; i < size; i++) {
- values.push(view.getFloat64(i * 8, true));
- }
- } else if (typeStr === 'f16') {
- const size = blob.length / 2;
- for (let i = 0; i < size; i++) {
- values.push(view.getFloat16(i * 2, true));
- }
- }
- this.values = values;
- this.data = null;
- }
- return this.values;
- }
- toString() {
- const typeStr = this.type ? this.type.toString() : '';
- return `array<${typeStr}: ${this.value.join(', ')}>`;
- }
- };
- _.TypeAttrOf = class extends _.Attribute {
- constructor(type) {
- super();
- this.type = type; // the type IS the value
- }
- toString() {
- return this.type.toString();
- }
- };
- _.ConvDimensionNumbersAttr = class extends _.Attribute {
- constructor(input, kernel, output) {
- super();
- this.input = input;
- this.kernel = kernel;
- this.output = output;
- }
- toString() {
- const formatDim = (dims) => `[${dims.join(', ')}]`;
- return `${formatDim(this.input)}x${formatDim(this.kernel)}->${formatDim(this.output)}`;
- }
- };
- _.StringRef = class {
- constructor(value, position) {
- this.value = value;
- this.position = position;
- }
- data() {
- return this.position;
- }
- str() {
- return this.value;
- }
- drop_front(n = 1) {
- return this.value.substring(n);
- }
- getAsInteger(radix) {
- const str = this.value;
- if (str.length === 0) {
- return null;
- }
- const value = radix === 0 ? Number(str) : parseInt(str, radix);
- if (!Number.isInteger(value)) {
- return null;
- }
- return value;
- }
- toString() {
- return this.value;
- }
- };
- _.Type = class {
- constructor(value) {
- if (value && value instanceof _.RankedTensorType === false && value.startsWith('tensor<')) {
- // Do not remove. Investigate why RankedTensorType is not getting parsed.
- throw new mlir.Error(`Invalid type '${value}'.`);
- }
- this._value = value;
- }
- get name() {
- return this._value;
- }
- toString() {
- return this._value;
- }
- };
- _.OpaqueType = class extends _.Type {
- constructor(dialect, symbolData) {
- super();
- this.dialect = dialect;
- this.symbolData = symbolData;
- }
- toString() {
- return `!${this.dialect}${this.symbolData}`;
- }
- };
- _.util = {};
- _.util.VariantType = class extends _.Type {
- toString() {
- return '?';
- }
- };
- _.util.ListType = class extends _.Type {
- constructor(elementType) {
- super();
- this.elementType = elementType;
- }
- toString() {
- return `!util.list<${this.elementType}>`;
- }
- };
- _.PrimitiveType = class extends _.Type {
- };
- _.IntegerType = class extends _.Type {
- };
- _.IntegerType.kMaxWidth = (1 << 24) - 1;
- _.FloatType = class extends _.Type {
- };
- _.NoneType = class extends _.Type {
- constructor() {
- super('none');
- }
- };
- _.IndexType = class extends _.Type {
- constructor() {
- super('index');
- }
- };
- _.FunctionType = class extends _.Type {
- constructor(inputs, results) {
- super(null);
- this.inputs = inputs || [];
- this.results = results || [];
- }
- toString() {
- const inputs = this.inputs.map((t) => t.toString());
- const results = this.results.map((t) => t.toString());
- const result = results.length === 1 ? results[0] : `(${results.join(', ')})`;
- return `(${inputs.join(', ')}) -> ${result}`;
- }
- };
- _.ComplexType = class extends _.Type {
- constructor(elementType) {
- super(null);
- this.elementType = elementType;
- }
- toString() {
- const elementTypeStr = this.elementType?.toString ? this.elementType.toString() : this.elementType;
- return `complex<${elementTypeStr}>`;
- }
- };
- _.ShapedType = class extends _.Type {
- getNumElements() {
- if (this.shape.some((d) => d < 0 || d === _.ShapedType.kDynamic)) {
- return 0;
- }
- return this.shape.length === 0 ? 1 : this.shape.reduce((a, b) => a * b, 1);
- }
- };
- _.ShapedType.kDynamic = Number.MIN_SAFE_INTEGER;
- _.RankedTensorType = class extends _.ShapedType {
- constructor(shape, elementType, encoding) {
- super(null);
- this.shape = shape || [];
- this.elementType = elementType;
- this.encoding = encoding;
- }
- toString() {
- const shapeStr = this.shape.map((d) => d < 0 ? '?' : d).join('x');
- const elementTypeStr = this.elementType?.toString ? this.elementType.toString() : this.elementType;
- const prefix = shapeStr ? `${shapeStr}x` : '';
- if (this.encoding) {
- return `tensor<${prefix}${elementTypeStr}, ${this.encoding}>`;
- }
- return `tensor<${prefix}${elementTypeStr}>`;
- }
- };
- _.UnrankedTensorType = class extends _.Type {
- constructor(elementType) {
- super(null);
- this.elementType = elementType;
- }
- toString() {
- const elementTypeStr = this.elementType.toString();
- return `tensor<*x${elementTypeStr}>`;
- }
- };
- _.VectorType = class extends _.ShapedType {
- constructor(shape, elementType, scalableDims) {
- super(null);
- this.shape = shape || [];
- this.elementType = elementType;
- this.scalableDims = scalableDims || [];
- }
- toString() {
- const parts = this.shape.map((d, i) => {
- const isScalable = this.scalableDims[i];
- return isScalable ? `[${d}]` : String(d);
- });
- const shapeStr = parts.join('x');
- const elementTypeStr = this.elementType?.toString ? this.elementType.toString() : this.elementType;
- const prefix = shapeStr ? `${shapeStr}x` : '';
- return `vector<${prefix}${elementTypeStr}>`;
- }
- };
- _.MemRefType = class extends _.ShapedType {
- constructor(shape, elementType, layout, memorySpace) {
- super(null);
- this.shape = shape || [];
- this.elementType = elementType;
- this.layout = layout || null;
- this.memorySpace = memorySpace || null;
- }
- toString() {
- const shapeStr = this.shape.map((d) => d < 0 ? '?' : d).join('x');
- const elementTypeStr = this.elementType?.toString ? this.elementType.toString() : this.elementType;
- const prefix = shapeStr ? `${shapeStr}x` : '';
- let result = `memref<${prefix}${elementTypeStr}`;
- if (this.layout) {
- const layoutStr = typeof this.layout === 'object' ? JSON.stringify(this.layout) : this.layout;
- result += `, ${layoutStr}`;
- }
- if (this.memorySpace) {
- const memorySpaceStr = typeof this.memorySpace === 'object' ? JSON.stringify(this.memorySpace) : this.memorySpace;
- result += `, ${memorySpaceStr}`;
- }
- result += '>';
- return result;
- }
- };
- _.UnrankedMemRefType = class extends _.Type {
- constructor(elementType, memorySpace) {
- super(null);
- this.elementType = elementType;
- this.memorySpace = memorySpace || null;
- }
- toString() {
- const elementTypeStr = this.elementType?.toString ? this.elementType.toString() : this.elementType;
- let result = `memref<*x${elementTypeStr}`;
- if (this.memorySpace) {
- const memorySpaceStr = typeof this.memorySpace === 'object' ? JSON.stringify(this.memorySpace) : this.memorySpace;
- result += `, ${memorySpaceStr}`;
- }
- result += '>';
- return result;
- }
- };
- _.TupleType = class extends _.Type {
- constructor(types) {
- super(null);
- this.types = types || [];
- }
- getTypes() {
- return this.types;
- }
- getNumTypes() {
- return this.types.length;
- }
- getType(index) {
- return this.types[index];
- }
- toString() {
- const typeStrs = this.types.map((t) => t?.toString ? t.toString() : t);
- return `tuple<${typeStrs.join(', ')}>`;
- }
- };
- _.SMLoc = class {
- constructor(decoder, position) {
- this.decoder = decoder;
- this.position = position || 0;
- }
- copy() {
- return new _.SMLoc(this.decoder, this.position);
- }
- toString() {
- let line = 1;
- let column = 1;
- const position = this.decoder.position;
- this.decoder.position = 0;
- let c = '';
- do {
- if (this.decoder.position === this.position) {
- this.decoder.position = position;
- return `at ${line}:${column}.`;
- }
- c = this.decoder.decode();
- if (c === '\n') {
- line++;
- column = 1;
- } else {
- column++;
- }
- }
- while (c !== undefined);
- this.decoder.position = position;
- return `at ${line}:${column}.`;
- }
- };
- _.Token = class {
- constructor(decoder) {
- this.loc = new _.SMLoc(decoder);
- this.kind = null;
- this.spelling = new _.StringRef('', decoder.position);
- }
- is(kind) {
- return this.kind === kind;
- }
- isAny(...args) {
- return args.some((kind) => this.kind === kind);
- }
- isNot(kind) {
- return this.kind !== kind;
- }
- isKeyword() {
- return this.kind.startsWith('kw_');
- }
- getSpelling() {
- return this.spelling;
- }
- getUnsignedIntegerValue() {
- const isHex = this.spelling.str().length > 1 && this.spelling.str()[1] === 'x';
- const value = this.spelling.getAsInteger(isHex ? 0 : 10);
- if (value === null || value < 0 || value > 0xFFFFFFFF) {
- return null;
- }
- return value;
- }
- getUInt64IntegerValue() {
- const spelling = this.spelling.str();
- if (/^-?(\d+|0x[\da-f]+|0o[0-7]+|0b[01]+)$/i.test(spelling)) {
- return BigInt(spelling);
- }
- return null;
- }
- getFloatingPointValue() {
- const result = parseFloat(this.spelling.str());
- if (isNaN(result)) {
- return null;
- }
- return result;
- }
- getIntTypeBitwidth() {
- const spelling = this.spelling.str();
- const start = spelling[0] === 'i' ? 1 : 2;
- const result = parseInt(spelling.slice(start), 10);
- if (isNaN(result)) {
- return null;
- }
- return result;
- }
- getIntTypeSignedness() {
- const spelling = this.spelling.str();
- if (spelling[0] === 'i') {
- return null;
- }
- if (spelling[0] === 's') {
- return true;
- }
- return false;
- }
- getHexStringValue() {
- const spelling = this.spelling.str();
- const bytes = spelling.slice(1, -1);
- if (!bytes.startsWith('0x') || (bytes.length - 2) % 2 !== 0) {
- return null;
- }
- const hex = bytes.slice(2);
- const result = [];
- for (let i = 0; i < hex.length; i += 2) {
- const hi = parseInt(hex[i], 16);
- const lo = parseInt(hex[i + 1], 16);
- if (isNaN(hi) || isNaN(lo)) {
- return null;
- }
- result.push((hi << 4) | lo);
- }
- return result;
- }
- getSymbolReference() {
- const nameStr = this.spelling.str().slice(1);
- if (nameStr[0] === '"') {
- return this.getStringValue();
- }
- return nameStr;
- }
- getHashIdentifierNumber() {
- const str = this.spelling.str().slice(1);
- const result = parseInt(str, 10);
- if (isNaN(result)) {
- return null;
- }
- return result;
- }
- getStringValue() {
- let bytes = this.spelling.str();
- bytes = bytes.slice(1);
- bytes = bytes.slice(0, -1);
- if (this.kind === _.Token.at_identifier) {
- bytes = bytes.slice(1);
- }
- if (!bytes.includes('\\')) {
- return bytes;
- }
- const parts = [];
- for (let i = 0; i < bytes.length;) {
- const c = bytes[i++];
- if (c !== '\\') {
- parts.push(c);
- continue;
- }
- const c1 = bytes[i++];
- switch (c1) {
- case '"':
- case '\\':
- parts.push(c1);
- continue;
- case 'n':
- parts.push('\n');
- continue;
- case 't':
- parts.push('\t');
- continue;
- default:
- break;
- }
- const c2 = bytes[i++];
- const hexValue = (parseInt(c1, 16) << 4) | parseInt(c2, 16);
- parts.push(String.fromCharCode(hexValue));
- }
- return parts.join('');
- }
- };
- _.Token.arrow = '->';
- _.Token.at_identifier = '@';
- _.Token.bare_identifier = 'id';
- _.Token.caret_identifier = '^';
- _.Token.colon = ':';
- _.Token.comma = ',';
- _.Token.ellipsis = 'ellipsis';
- _.Token.eof = 'eof';
- _.Token.equal = '=';
- _.Token.exclamation_identifier = '!';
- _.Token.file_metadata_begin = '{-#';
- _.Token.file_metadata_end = '#-}';
- _.Token.floatliteral = 'float';
- _.Token.greater = '>';
- _.Token.hash_identifier = '#';
- _.Token.integer = 'int';
- _.Token.inttype = 'inttype';
- _.Token.kw_ceildiv = 'kw_ceildiv';
- _.Token.kw_floordiv = 'kw_floordiv';
- _.Token.kw_mod = 'kw_mod';
- _.Token.l_brace = '{';
- _.Token.l_paren = '(';
- _.Token.l_square = '[';
- _.Token.less = '<';
- _.Token.minus = 'minus';
- _.Token.percent_identifier = '%';
- _.Token.plus = 'plus';
- _.Token.question = '?';
- _.Token.r_brace = '}';
- _.Token.r_paren = ')';
- _.Token.r_square = ']';
- _.Token.slash = '/';
- _.Token.star = '*';
- _.Token.string = 'string';
- _.Token.vertical_bar = '|';
- _.Token.kw_affine_map = 'kw_affine_map';
- _.Token.kw_affine_set = 'kw_affine_set';
- _.Token.kw_array = 'kw_array';
- _.Token.kw_dense = 'kw_dense';
- _.Token.kw_dense_resource = 'kw_dense_resource';
- _.Token.kw_distinct = 'kw_distinct';
- _.Token.kw_loc = 'kw_loc';
- _.Token.kw_sparse = 'kw_sparse';
- _.Token.kw_strided = 'kw_strided';
- _.Token.kw_unit = 'kw_unit';
- _.Token.kw_offset = 'kw_offset';
- _.Token.kw_symbol = 'kw_symbol';
- _.Token.kw_true = 'kw_true';
- _.Token.kw_false = 'kw_false';
- _.Lexer = class {
- constructor(decoder) {
- this._decoder = decoder;
- this._currentPosition = this._decoder.position;
- this._current = this._decoder.decode();
- this._nextPosition = this._decoder.position;
- this._next = this._decoder.decode();
- this._tokens = [new _.Token(decoder), new _.Token(decoder), new _.Token(decoder), new _.Token(decoder)];
- this._index = 0;
- this._errorLoc = new _.SMLoc(decoder);
- }
- location() {
- const loc = new _.SMLoc(this._decoder, this._position);
- return loc.toString();
- }
- lexToken() {
- this._position = this._currentPosition;
- while (this._current) {
- switch (this._current) {
- case ' ':
- case '\t':
- case '\n':
- case '\r':
- this._skipWhitespace();
- this._position = this._currentPosition;
- continue;
- case '/':
- if (this._peek() !== '/') {
- this._read();
- return this.formToken('/', '/');
- }
- this.lexComment();
- this._position = this._currentPosition;
- continue;
- case '.':
- this._read();
- if (this._current === '.' && this._next === '.') {
- this._read();
- this._read();
- return this.formToken('ellipsis', '...');
- }
- throw new mlir.Error(`Expected three consecutive dots for an ellipsis ${this.location()}`);
- case '-':
- if (this._peek() === '>') {
- this._read();
- this._read();
- return this.formToken('->', '->');
- }
- this._read();
- return this.formToken('minus', '-');
- case '+':
- this._read();
- return this.formToken('plus', '+');
- case '"':
- return this.lexString();
- case '@':
- return this.lexPrefixedIdentifier('@');
- case '%':
- return this.lexPrefixedIdentifier('%');
- case '#':
- if (this._peek() === '-') {
- const position = this._decoder.position;
- const next = this._decoder.decode();
- this._decoder.position = position;
- if (next === '}') {
- this._read();
- this._read();
- this._read();
- return this.formToken('#-}', '#-}');
- }
- }
- return this.lexPrefixedIdentifier('#');
- case '!':
- return this.lexPrefixedIdentifier('!');
- case '^':
- return this.lexPrefixedIdentifier('^');
- case '=':
- this._read();
- return this.formToken('=', '=');
- case ':':
- this._read();
- return this.formToken(':', ':');
- case ',':
- case '(':
- case ')':
- case '}':
- case '[':
- case ']':
- case '<':
- case '>':
- case '?':
- case '*':
- case '|': {
- const value = this._read();
- return this.formToken(value, value);
- }
- case '{':
- if (this._peek() === '-') {
- const position = this._decoder.position;
- const next = this._decoder.decode();
- this._decoder.position = position;
- if (next === '#') {
- this._read();
- this._read();
- this._read();
- return this.formToken('{-#', '{-#');
- }
- }
- this._read();
- return this.formToken('{', '{');
- default:
- if (/[a-zA-Z_]/.test(this._current)) {
- return this.lexBareIdentifierOrKeyword();
- }
- if (/[0-9]/.test(this._current)) {
- return this.lexNumber();
- }
- throw new mlir.Error(`Unexpected character '${this._current}' ${this.location()}`);
- }
- }
- return this.formToken('eof', null);
- }
- resetPointer(newPointer) {
- if (newPointer < 0) {
- throw new mlir.Error('Invalid negative offset.');
- }
- this._decoder.position = newPointer;
- this._nextPosition = this._decoder.position;
- this._next = this._decoder.decode();
- this._read();
- }
- _read() {
- const current = this._current;
- this._current = this._next;
- this._currentPosition = this._nextPosition;
- this._nextPosition = this._decoder.position;
- this._next = this._decoder.decode();
- return current;
- }
- _peek() {
- return this._next;
- }
- _eat(value) {
- if (this._current === value) {
- this._read();
- return true;
- }
- return false;
- }
- _skipWhitespace() {
- while (this._current !== undefined && (this._current === ' ' || this._current === '\t' || this._current === '\n' || this._current === '\r')) {
- this._read();
- }
- }
- lexComment() {
- this._read('/');
- if (this._current !== '/') {
- throw new mlir.Error(`Invalid comment.`);
- }
- while (this._current && this._current !== '\n' && this._current !== '\r') {
- this._read();
- }
- }
- lexNumber() {
- let v = '';
- let type = 'int';
- while (this._current && /[0-9]/.test(this._current)) {
- v += this._read();
- }
- if (v === '0' && this._current === 'x' && /[0-9a-fA-F]/.test(this._peek())) {
- v += this._read();
- while (this._current && /[0-9a-fA-F]/.test(this._current)) {
- v += this._read();
- }
- return this.formToken(type, v);
- }
- if (this._current === '.') {
- v += this._read();
- type = 'float';
- while (this._current && /[0-9]/.test(this._current)) {
- v += this._read();
- }
- if (this._current === 'e' || this._current === 'E') {
- const next1 = this._peek();
- let next2 = '';
- if (next1 === '+' || next1 === '-') {
- const position = this._decoder.position;
- next2 = this._decoder.decode();
- this._decoder.position = position;
- }
- if (/[0-9]/.test(next1) || ((next1 === '+' || next1 === '-') && /[0-9]/.test(next2))) {
- v += this._read();
- if (this._current === '+' || this._current === '-') {
- v += this._read();
- }
- while (this._current && /[0-9]/.test(this._current)) {
- v += this._read();
- }
- }
- }
- return this.formToken(type, v);
- }
- return this.formToken(type, v);
- }
- lexString() {
- const parts = ['"'];
- this._read();
- let chunk = '';
- while (this._current && this._current !== '"') {
- if (this._current === '\n' || this._current === '\v' ||
- this._current === '\f' || this._current === '\r') {
- throw new mlir.Error(`Expected '"' in string literal ${this.location()}`);
- }
- if (this._current === '\\') {
- if (chunk) {
- parts.push(chunk);
- chunk = '';
- }
- parts.push(this._current);
- this._read();
- if (this._current === '"' || this._current === '\\' || this._current === 'n' || this._current === 't') {
- parts.push(this._current);
- this._read();
- } else if (this._current && /[0-9a-fA-F]/.test(this._current) && this._next && /[0-9a-fA-F]/.test(this._next)) {
- parts.push(this._current);
- this._read();
- parts.push(this._current);
- this._read();
- } else {
- throw new mlir.Error(`Unknown escape in string literal ${this.location()}`);
- }
- } else {
- chunk += this._current;
- if (chunk.length >= 4096) {
- parts.push(chunk);
- chunk = '';
- }
- this._read();
- }
- }
- if (chunk) {
- parts.push(chunk);
- }
- if (this._eat('"')) {
- parts.push('"');
- return this.formToken(_.Token.string, parts.join(''));
- }
- throw new mlir.Error('Unterminated string literal');
- }
- lexBareIdentifierOrKeyword() {
- let result = '';
- while (this._current && (/[a-zA-Z_$.]/.test(this._current) || /[0-9]/.test(this._current))) {
- result += this._read();
- }
- const isAllDigit = (str) => /^[0-9]+$/.test(str);
- if ((result.length > 1 && result[0] === 'i' && isAllDigit(result.slice(1))) ||
- (result.length > 2 && result[1] === 'i' && (result[0] === 's' || result[0] === 'u') && isAllDigit(result.slice(2)))) {
- return this.formToken('inttype', result);
- }
- switch (result) {
- case 'loc':
- case 'affine_map':
- case 'affine_set':
- case 'dense':
- case 'dense_resource':
- case 'sparse':
- case 'array':
- case 'strided':
- case 'distinct':
- case 'unit':
- case 'floordiv':
- case 'ceildiv':
- case 'mod':
- case 'offset':
- case 'symbol':
- return this.formToken(`kw_${result}`, result);
- case 'true':
- return this.formToken(_.Token.kw_true, result);
- case 'false':
- return this.formToken(_.Token.kw_false, result);
- case 'unknown':
- return this.formToken(_.Token.bare_identifier, result);
- default:
- return this.formToken(_.Token.bare_identifier, result);
- }
- }
- lexPrefixedIdentifier(prefix) {
- let result = prefix;
- this._read();
- if (prefix === '@' && this._current === '"') {
- result += this.lexString().getSpelling().str();
- return this.formToken(prefix, result);
- }
- if (prefix === '@') {
- // symbol-ref-id ::= `@` (bare-id | string-literal)
- // bare-id ::= (letter|[_]) (letter|digit|[_$.])*
- if (this._current && /[a-zA-Z_]/.test(this._current)) {
- while (this._current && /[a-zA-Z0-9_$.]/.test(this._current)) {
- result += this._read();
- }
- } else if (result.length === 1) {
- throw new mlir.Error(`@ identifier expected to start with letter or '_' ${this.location()}`);
- }
- } else if (this._current && /[0-9]/.test(this._current)) {
- // suffix-id ::= digit+ | ...
- while (this._current && /[0-9]/.test(this._current)) {
- result += this._read();
- }
- } else if (this._current && /[a-zA-Z_$.-]/.test(this._current)) {
- // suffix-id ::= ... | (letter|id-punct) (letter|id-punct|digit)*
- while (this._current && /[a-zA-Z_$0-9.-]/.test(this._current)) {
- if (this._current === '-' && this._peek() === '>') {
- break;
- }
- result += this._read();
- }
- } else if (result.length === 1) {
- const errorKinds = { '#': 'Invalid attribute name', '%': 'Invalid SSA name', '^': 'Invalid block name', '!': 'Invalid type identifier' };
- throw new mlir.Error(`${errorKinds[prefix] || 'Invalid identifier'} ${this.location()}`);
- }
- if (prefix === '@' && this._current === ':' && this._peek() === ':') {
- result += this._read();
- result += this._read();
- result += this.lexPrefixedIdentifier('@').getSpelling().str();
- }
- const kind = prefix === '$' ? '%' : prefix;
- return this.formToken(kind, result);
- }
- formToken(kind, value) {
- const token = this._tokens[this._index];
- this._index = (this._index + 1) % this._tokens.length;
- token.loc.position = this._position;
- token.kind = kind;
- token.spelling.value = value;
- token.spelling.position = this._position;
- return token;
- }
- };
- _.AsmResourceParser = class {
- constructor(name) {
- this.name = name;
- }
- };
- _.ParsedResourceEntry = {};
- _.ParsedResourceEntry.Text = class {
- constructor(key, keyLoc, value, parser) {
- this.key = key;
- this.keyLoc = keyLoc;
- this.value = value;
- this.parser = parser;
- }
- parseAsBool() {
- if (this.value.kind === 'boolean') {
- return this.value.value;
- }
- throw new mlir.Error(`Expected boolean value for resource entry '${this.key}'`);
- }
- parseAsString() {
- if (this.value.kind === _.Token.string) {
- return this.value.getStringValue();
- }
- throw new mlir.Error(`Expected string value for resource entry '${this.key}'`);
- }
- parseAsBlob() {
- if (this.value.kind === _.Token.string) {
- const hexStr = this.value.getStringValue();
- if (hexStr.startsWith('0x')) {
- const hex = hexStr.slice(2);
- const bytes = new Uint8Array(hex.length / 2);
- for (let i = 0; i < bytes.length; i++) {
- bytes[i] = parseInt(hex.slice(i * 2, i * 2 + 2), 16);
- }
- if (bytes.length >= 4) {
- return bytes.slice(4);
- }
- return bytes;
- }
- }
- throw new mlir.Error(`Expected hex string blob value for resource entry '${this.key}'`);
- }
- };
- _.ParserConfig = class {
- constructor(context) {
- this.context = context;
- this.resourceParsers = new Map();
- }
- getResourceParser(name) {
- return this.resourceParsers.get(name);
- }
- attachResourceParser(parser) {
- this.resourceParsers.set(parser.name, parser);
- }
- };
- _.ParserState = class {
- constructor(decoder, config) {
- this.config = config;
- this.defaultDialectStack = ['builtin'];
- this.attributeAliasDefinitions = new Map();
- this.typeAliasDefinitions = new Map();
- this.dialectResources = new Map();
- this.deferredLocsReferences = [];
- this.lex = new _.Lexer(decoder);
- this.curToken = this.lex.lexToken();
- }
- };
- _.IsolatedSSANameScope = class {
- constructor() {
- this.values = new Map();
- this.definitionsPerScope = [];
- }
- recordDefinition(def) {
- this.definitionsPerScope[this.definitionsPerScope.length - 1].add(def);
- }
- pushSSANameScope() {
- this.definitionsPerScope.push(new Set());
- }
- popSSANameScope() {
- for (const def of this.definitionsPerScope.pop()) {
- this.values.delete(def);
- }
- }
- };
- _.Parser = class {
- constructor(state) {
- this.state = state;
- }
- get context() {
- return this.state.config.context;
- }
- parseCommaSeparatedListUntil(rightToken, parseElement, allowEmptyList = true) {
- if (this.getToken().is(rightToken)) {
- if (!allowEmptyList) {
- throw new mlir.Error(`Expected list element ${this.location()}`);
- }
- this.consumeToken(rightToken);
- return;
- }
- this.parseCommaSeparatedList('none', parseElement);
- this.parseToken(rightToken, `Expected '${rightToken}'`);
- }
- parseResourceFileMetadata(parseBody) {
- this.parseToken(_.Token.l_brace, "Expected '{'");
- this.parseCommaSeparatedListUntil(_.Token.r_brace, () => {
- const nameLoc = this.getToken().loc;
- const name = this.parseOptionalKeyword();
- if (!name) {
- throw new mlir.Error(`Expected identifier key for 'resource' entry ${this.location()}`);
- }
- this.parseToken(_.Token.colon);
- this.parseToken(_.Token.l_brace);
- parseBody(name, nameLoc);
- });
- }
- parseDialectResourceFileMetadata() {
- this.parseResourceFileMetadata((name /*, nameLoc */) => {
- const dialect = this.context.getOrLoadDialect(name);
- if (!dialect) {
- throw new mlir.Error(`Dialect '${name}' is unknown ${this.location()}`);
- }
- this.parseCommaSeparatedListUntil(_.Token.r_brace, () => {
- const keyLoc = this.getToken().loc.copy();
- const handle = this.parseResourceHandle(dialect);
- const key = dialect.getResourceKey(handle);
- this.parseToken(_.Token.colon, "Expected ':'");
- const valueTok = this.getToken();
- this.consumeToken();
- const entry = new _.ParsedResourceEntry.Text(key, keyLoc, valueTok, this);
- dialect.parseResource(entry);
- });
- });
- }
- parseExternalResourceFileMetadata() {
- this.parseResourceFileMetadata((name /*, nameLoc */) => {
- const handler = this.state.config.getResourceParser(name);
- this.parseCommaSeparatedListUntil(_.Token.r_brace, () => {
- const keyLoc = this.getToken().loc;
- const key = this.parseOptionalKeywordOrString();
- if (!key) {
- throw new mlir.Error(`Expected identifier key for 'external_resources' entry ${this.location()}`);
- }
- this.parseToken(_.Token.colon, "Expected ':'");
- const valueTok = this.getToken();
- this.consumeToken();
- if (handler && handler.parseResource) {
- const entry = new _.ParsedResourceEntry.Text(key, keyLoc, valueTok, this);
- handler.parseResource(entry);
- }
- });
- });
- }
- parseOptionalKeywordOrString() {
- const keyword = this.parseOptionalKeyword();
- if (keyword) {
- return keyword;
- }
- return this.parseOptionalString();
- }
- finalize(/* block */) {
- }
- isCurrentTokenAKeyword() {
- return this.getToken().isAny(_.Token.bare_identifier, _.Token.inttype) || this.getToken().isKeyword();
- }
- parseOptionalKeyword() {
- if (!this.isCurrentTokenAKeyword()) {
- return null;
- }
- const keyword = this.getTokenSpelling().str();
- this.consumeToken();
- return keyword;
- }
- parseTypeListNoParens() {
- return this.parseCommaSeparatedList('none', () => this.parseType());
- }
- parseTypeListParens() {
- this.parseToken(_.Token.l_paren, "Expected '('");
- if (this.consumeIf(_.Token.r_paren)) {
- return [];
- }
- const types = this.parseTypeListNoParens();
- this.parseToken(_.Token.r_paren, "Expected ')'");
- return types;
- }
- skip(open) {
- const closingFor = { '<': '>', '[': ']', '(': ')', '{': '}' };
- const openingFor = { '>': '<', ']': '[', ')': '(', '}': '{' };
- const delimiters = new Set(['<', '>', '[', ']', '(', ')', '{', '}', ',', ':', '=']);
- let value = '';
- let prevToken = '';
- if (this.getToken().is(open)) {
- const stack = [open];
- prevToken = this.getToken().getSpelling().str();
- this.consumeToken();
- value += prevToken;
- while (stack.length > 0) {
- if (this.getToken().is(_.Token.eof)) {
- throw new mlir.Error(`Unbalanced '${stack[stack.length - 1]}' ${this.location()}`);
- }
- const token = this.getToken().getSpelling().str();
- if (closingFor[token]) {
- stack.push(token);
- } else if (openingFor[token]) {
- if (stack[stack.length - 1] === openingFor[token]) {
- stack.pop();
- } else if (token !== '>') {
- throw new mlir.Error(`Unbalanced '${stack[stack.length - 1]}' ${this.location()}`);
- }
- }
- const curToken = this.getToken().getSpelling().str();
- this.consumeToken();
- if (!delimiters.has(prevToken) && !delimiters.has(curToken)) {
- value += ' ';
- }
- value += curToken;
- prevToken = curToken;
- }
- }
- return value;
- }
- parseSuccessors(successors) {
- const parsed = this.parseCommaSeparatedList('square', () => {
- const label = this.getTokenSpelling().str();
- this.consumeToken(_.Token.caret_identifier);
- return { label };
- });
- if (parsed.length === 0) {
- throw new mlir.Error(`Expected at least one successor ${this.location()}`);
- }
- for (const s of parsed) {
- successors.push(s);
- }
- }
- parseAttributeDict(attributes) {
- const seenKeys = new Set();
- this.parseCommaSeparatedList('brace', () => {
- // The name of an attribute can either be a bare identifier, or a string.
- let name = null;
- if (this.getToken().is(_.Token.string)) {
- name = this.getToken().getStringValue();
- } else if (this.getToken().isAny(_.Token.bare_identifier, _.Token.inttype) || this.getToken().isKeyword()) {
- name = this.getTokenSpelling().str();
- } else {
- throw new mlir.Error(`Expected attribute name ${this.location()}`);
- }
- if (!name) {
- throw new mlir.Error(`Expected attribute name ${this.location()}`);
- }
- this.consumeToken();
- if (seenKeys.has(name)) {
- throw new mlir.Error(`Duplicate key '${name}' in dictionary attribute ${this.location()}`);
- }
- seenKeys.add(name);
- if (!this.consumeIf(_.Token.equal)) {
- attributes.set(name, new _.UnitAttr());
- return;
- }
- const attr = this.parseAttribute();
- attributes.set(name, attr);
- });
- }
- parseLocationInstance() {
- if (this.getToken().is(_.Token.hash_identifier)) {
- const locAttr = this.parseExtendedAttr();
- if (locAttr instanceof _.LocationAttr === false && locAttr instanceof _.OpaqueAttr === false) {
- throw new mlir.Error(`Expected location attribute, but got '${locAttr}' ${this.location()}`);
- }
- return locAttr;
- }
- if (this.getToken().is(_.Token.string)) {
- return this.parseNameOrFileLineColRange();
- }
- if (!this.getToken().is(_.Token.bare_identifier)) {
- throw new mlir.Error(`Expected location instance, got '${this.getTokenSpelling().str()}' ${this.location()}`);
- }
- if (this.getToken().spelling.str() === 'callsite') {
- return this.parseCallSiteLocation();
- }
- if (this.getToken().spelling.str() === 'fused') {
- return this.parseFusedLocation();
- }
- if (this.getToken().spelling.str() === 'unknown') {
- this.consumeToken(_.Token.bare_identifier);
- return new _.UnknownLoc();
- }
- throw new mlir.Error(`Expected location instance, got '${this.getTokenSpelling().str()}' ${this.location()}`);
- }
- parseCallSiteLocation() {
- this.consumeToken(_.Token.bare_identifier);
- this.parseToken(_.Token.l_paren, "Expected '(' in callsite location");
- const callee = this.parseLocationInstance();
- if (this.getToken().isNot(_.Token.bare_identifier) || this.getToken().getSpelling().str() !== 'at') {
- throw new mlir.Error(`Expected 'at' in callsite location ${this.location()}`);
- }
- this.consumeToken(_.Token.bare_identifier);
- const caller = this.parseLocationInstance();
- this.parseToken(_.Token.r_paren, "Expected ')' in callsite location");
- return new _.CallSiteLoc(callee, caller);
- }
- parseFusedLocation() {
- this.consumeToken(_.Token.bare_identifier);
- let metadata = null;
- if (this.consumeIf(_.Token.less)) {
- metadata = this.parseAttribute();
- if (!metadata) {
- throw new mlir.Error(`Expected attribute in fused location metadata ${this.location()}`);
- }
- this.parseToken(_.Token.greater, "Expected '>' after fused location metadata");
- }
- const locations = this.parseCommaSeparatedList('square', () => this.parseLocationInstance());
- return new _.FusedLoc(locations, metadata);
- }
- parseNameOrFileLineColRange() {
- const str = this.getToken().getStringValue();
- this.consumeToken(_.Token.string);
- if (this.consumeIf(_.Token.colon)) {
- if (this.getToken().isNot(_.Token.integer)) {
- throw new mlir.Error(`Expected integer line number in FileLineColRange ${this.location()}`);
- }
- const startLine = this.getToken().getUnsignedIntegerValue();
- if (startLine === null) {
- throw new mlir.Error(`Expected integer line number in FileLineColRange ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- if (this.getToken().isNot(_.Token.colon)) {
- return new _.FileLineColRange(str, startLine);
- }
- this.consumeToken(_.Token.colon);
- if (this.getToken().isNot(_.Token.integer)) {
- throw new mlir.Error(`Expected integer column number in FileLineColRange ${this.location()}`);
- }
- const startCol = this.getToken().getUnsignedIntegerValue();
- if (startCol === null) {
- throw new mlir.Error(`Expected integer column number in FileLineColRange ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- if (!this.isCurrentTokenAKeyword() || this.getTokenSpelling().str() !== 'to') {
- return new _.FileLineColLoc(str, startLine, startCol);
- }
- this.consumeToken();
- let endLine = null;
- if (this.getToken().is(_.Token.integer)) {
- endLine = this.getToken().getUnsignedIntegerValue();
- if (endLine === null) {
- throw new mlir.Error(`Expected integer line number in FileLineColRange ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- }
- if (this.getToken().isNot(_.Token.colon)) {
- throw new mlir.Error(`Expected either integer or ':' post 'to' in FileLineColRange ${this.location()}`);
- }
- this.consumeToken(_.Token.colon);
- if (this.getToken().isNot(_.Token.integer)) {
- throw new mlir.Error(`Expected integer column number in FileLineColRange ${this.location()}`);
- }
- const endCol = this.getToken().getUnsignedIntegerValue();
- if (endCol === null) {
- throw new mlir.Error(`Expected integer column number in FileLineColRange ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- if (endLine !== null) {
- return new _.FileLineColRange(str, startLine, startCol, endLine, endCol);
- }
- return new _.FileLineColRange(str, startLine, startCol, undefined, endCol);
- }
- if (this.consumeIf(_.Token.l_paren)) {
- const childLoc = this.parseLocationInstance();
- this.parseToken(_.Token.r_paren, "Expected ')' after child location of NameLoc");
- return new _.NameLoc(str, childLoc);
- }
- return new _.NameLoc(str);
- }
- parseLocationAlias() {
- const tok = this.getToken();
- this.consumeToken(_.Token.hash_identifier);
- const identifier = tok.getSpelling().str().substring(1); // drop_front - remove '#' prefix
- const attr = this.state.attributeAliasDefinitions.get(tok.getSpelling().str());
- if (attr) {
- if (attr instanceof _.LocationAttr) {
- return attr;
- }
- return attr;
- }
- const index = this.state.deferredLocsReferences.length;
- this.state.deferredLocsReferences.push({ loc: tok.loc, identifier });
- return new _.OpaqueLoc(index, identifier, new _.UnknownLoc());
- }
- parseOptionalLocationSpecifier() {
- if (!this.consumeIf(_.Token.kw_loc)) {
- return null;
- }
- this.parseToken(_.Token.l_paren, "expected '(' in location");
- const tok = this.getToken();
- let directLoc = null;
- if (tok.is(_.Token.hash_identifier) && !tok.getSpelling().str().includes('.')) {
- directLoc = this.parseLocationAlias();
- } else {
- directLoc = this.parseLocationInstance();
- }
- this.parseToken(_.Token.r_paren, "expected ')' in location");
- return directLoc;
- }
- parseXInDimensionList() {
- // Matches ref impl: if token is 'xf32', reset lexer to after 'x'
- if (this.getToken().isNot(_.Token.bare_identifier)) {
- return false;
- }
- const token = this.getToken().getSpelling().str();
- if (!token.startsWith('x')) {
- return false;
- }
- // If we had a prefix of 'x' (e.g., 'xf32'), reset lexer to after the 'x'
- if (token.length > 1) {
- this.state.lex.resetPointer(this.getToken().loc.position + 1);
- }
- this.state.curToken = this.state.lex.lexToken();
- return true;
- }
- parseIntegerInDimensionList() {
- const spelling = this.getTokenSpelling().str();
- if (spelling[0] === '0' && spelling.length > 1 && spelling[1] === 'x') {
- this.state.lex.resetPointer(this.getToken().loc.position + 1);
- this.state.curToken = this.state.lex.lexToken();
- return 0;
- }
- const dimension = this.getToken().getUInt64IntegerValue();
- if (dimension === null) {
- throw new mlir.Error(`Invalid dimension ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- return dimension.toNumber();
- }
- parseVectorDimensionList() {
- const dimensions = [];
- const scalableDims = [];
- while (this.getToken().is(_.Token.integer) || this.getToken().is(_.Token.l_square)) {
- const scalable = this.consumeIf(_.Token.l_square);
- dimensions.push(this.parseIntegerInDimensionList());
- if (scalable) {
- if (!this.consumeIf(_.Token.r_square)) {
- throw new mlir.Error(`Missing ']' closing scalable dimension ${this.location()}`);
- }
- }
- scalableDims.push(scalable);
- if (!this.parseXInDimensionList()) {
- break;
- }
- }
- return { dimensions, scalableDims };
- }
- parseDimensionListRanked(allowDynamic = true, withTrailingX = true) {
- const dimensions = [];
- const parseDim = () => {
- if (allowDynamic && this.consumeIf(_.Token.question)) {
- dimensions.push(_.ShapedType.kDynamic);
- return true;
- }
- if (this.getToken().is(_.Token.integer)) {
- dimensions.push(this.parseIntegerInDimensionList());
- return true;
- }
- return false;
- };
- const hasDimToken = () => this.getToken().isAny(_.Token.integer, _.Token.question);
- if (withTrailingX) {
- while (hasDimToken()) {
- if (!parseDim() || !this.parseXInDimensionList()) {
- break;
- }
- }
- } else if (parseDim()) {
- while (this.getToken().is(_.Token.bare_identifier) && this.getTokenSpelling().str().startsWith('x')) {
- if (!this.parseXInDimensionList() || !parseDim()) {
- break;
- }
- }
- }
- return { dimensions };
- }
- parseTensorType() {
- this.parseToken(_.Token.less, "Expected '<' in tensor type");
- let isUnranked = false;
- let dimensions = [];
- if (this.consumeIf(_.Token.star)) {
- isUnranked = true;
- this.parseXInDimensionList();
- } else {
- const dimInfo = this.parseDimensionListRanked();
- dimensions = dimInfo.dimensions;
- }
- const elementType = this.parseType();
- let encoding = null;
- if (this.consumeIf(_.Token.comma)) {
- encoding = this.parseAttribute();
- }
- this.parseToken(_.Token.greater, "Expected '>' in tensor type");
- if (elementType instanceof _.NoneType || elementType instanceof _.FunctionType ||
- elementType instanceof _.TupleType || elementType instanceof _.RankedTensorType ||
- elementType instanceof _.UnrankedTensorType || elementType instanceof _.MemRefType ||
- elementType instanceof _.UnrankedMemRefType) {
- throw new mlir.Error(`Invalid tensor element type ${this.location()}`);
- }
- if (isUnranked) {
- if (encoding) {
- throw new mlir.Error(`Cannot apply encoding to unranked tensor ${this.location()}`);
- }
- return new _.UnrankedTensorType(elementType);
- }
- return new _.RankedTensorType(dimensions, elementType, encoding);
- }
- parseMemRefType() {
- this.parseToken(_.Token.less, "Expected '<' in memref type");
- let isUnranked = false;
- let dimensions = [];
- if (this.consumeIf(_.Token.star)) {
- isUnranked = true;
- this.parseXInDimensionList();
- } else {
- const dimInfo = this.parseDimensionListRanked();
- dimensions = dimInfo.dimensions;
- }
- const elementType = this.parseType();
- if (elementType instanceof _.NoneType || elementType instanceof _.FunctionType ||
- elementType instanceof _.TupleType || elementType instanceof _.RankedTensorType ||
- elementType instanceof _.UnrankedTensorType) {
- throw new mlir.Error(`Invalid memref element type ${this.location()}`);
- }
- let layout = null;
- let memorySpace = null;
- while (this.consumeIf(_.Token.comma)) {
- const attr = this.parseAttribute();
- const isLayout = attr instanceof _.AffineMapAttr || (attr && attr.type === 'strided') || (attr instanceof _.OpaqueAttr && attr.toString().includes('layout'));
- if (isLayout) {
- layout = attr;
- if (isUnranked) {
- throw new mlir.Error(`Cannot have affine map for unranked memref type ${this.location()}`);
- }
- if (memorySpace) {
- throw new mlir.Error(`Expected memory space to be last in memref type ${this.location()}`);
- }
- } else {
- if (memorySpace) {
- throw new mlir.Error(`Multiple memory spaces specified in memref type ${this.location()}`);
- }
- memorySpace = attr;
- }
- }
- this.parseToken(_.Token.greater, "Expected '>' in memref type");
- if (isUnranked) {
- return new _.UnrankedMemRefType(elementType, memorySpace);
- }
- return new _.MemRefType(dimensions, elementType, layout, memorySpace);
- }
- parseVectorType() {
- this.parseToken(_.Token.less, "Expected '<' in vector type");
- const dimInfo = this.parseVectorDimensionList();
- const elementType = this.parseType();
- this.parseToken(_.Token.greater, "Expected '>' in vector type");
- return new _.VectorType(dimInfo.dimensions, elementType, dimInfo.scalableDims);
- }
- parseComplexType() {
- this.parseToken(_.Token.less, "Expected '<' in complex type");
- const elementType = this.parseType();
- if (!(elementType instanceof _.FloatType) && !(elementType instanceof _.IntegerType)) {
- throw new mlir.Error(`Invalid element type for complex ${this.location()}`);
- }
- this.parseToken(_.Token.greater, "Expected '>' in complex type");
- return new _.ComplexType(elementType);
- }
- parseTupleType() {
- this.parseToken(_.Token.less, "Expected '<' in tuple type");
- if (this.consumeIf(_.Token.greater)) {
- return new _.TupleType([]);
- }
- const types = this.parseCommaSeparatedList('none', () => this.parseType());
- this.parseToken(_.Token.greater, "Expected '>' in tuple type");
- return new _.TupleType(types);
- }
- parseCustomTypeWithFallback(typeT) {
- if (typeT && this.getToken().isNot(_.Token.exclamation_identifier)) {
- if (typeof typeT === 'function') {
- return typeT(this);
- }
- const index = typeT.name.indexOf('.');
- if (index === -1) {
- throw new mlir.Error(`Invalid type name '${typeT.name}.`);
- }
- const dialectName = typeT.name.substring(0, index);
- const dialect = this.context.getOrLoadDialect(dialectName);
- this.context.checkDialect(dialect, dialectName, 'custom type');
- return dialect.parseCustomTypeWithFallback(this, typeT.type);
- }
- return this.parseType();
- }
- parseCustomAttributeWithFallback(attrT, type) {
- if (attrT) {
- return attrT(this, type);
- }
- return this.parseAttribute(type);
- }
- parseType() {
- if (this.getToken().is(_.Token.l_paren)) {
- return this.parseFunctionType();
- }
- return this.parseNonFunctionType();
- }
- parseOptionalType() {
- switch (this.getToken().kind) {
- case _.Token.l_paren:
- case _.Token.exclamation_identifier:
- case _.Token.inttype:
- return this.parseType();
- case _.Token.bare_identifier: {
- switch (this.getToken().getSpelling().str()) {
- case 'memref':
- case 'tensor':
- case 'complex':
- case 'tuple':
- case 'vector':
- case 'f4E2M1FN':
- case 'f6E2M3FN':
- case 'f6E3M2FN':
- case 'f8E5M2':
- case 'f8E4M3':
- case 'f8E4M3FN':
- case 'f8E5M2FNUZ':
- case 'f8E4M3FNUZ':
- case 'f8E4M3B11FNUZ':
- case 'f8E3M4':
- case 'f8E8M0FNU':
- case 'bf16':
- case 'f16':
- case 'tf32':
- case 'f32':
- case 'f64':
- case 'f80':
- case 'f128':
- case 'index':
- case 'none':
- return this.parseType();
- default:
- break;
- }
- break;
- }
- default: {
- break;
- }
- }
- return null;
- }
- parseNonFunctionType() {
- switch (this.getToken().kind) {
- case _.Token.inttype: {
- const width = this.getToken().getIntTypeBitwidth();
- if (width === null) {
- throw new mlir.Error(`Invalid integer width ${this.location()}`);
- }
- if (width > _.IntegerType.kMaxWidth) {
- throw new mlir.Error(`Integer bitwidth is limited to ${_.IntegerType.kMaxWidth} bits ${this.location()}`);
- }
- const signedness = this.getToken().getIntTypeSignedness();
- this.consumeToken(_.Token.inttype);
- let prefix = 'i';
- if (signedness === true) {
- prefix = 'si';
- } else if (signedness === false) {
- prefix = 'ui';
- }
- return new _.IntegerType(`${prefix}${width}`);
- }
- case _.Token.exclamation_identifier: {
- return this.parseExtendedType();
- }
- case _.Token.bare_identifier: {
- const value = this.getTokenSpelling().str();
- this.consumeToken(_.Token.bare_identifier);
- switch (value) {
- case 'tensor':
- return this.parseTensorType();
- case 'vector':
- return this.parseVectorType();
- case 'memref':
- return this.parseMemRefType();
- case 'complex':
- return this.parseComplexType();
- case 'tuple':
- return this.parseTupleType();
- case 'none':
- return new _.NoneType();
- case 'index':
- return new _.IndexType();
- case 'bf16':
- case 'f16':
- case 'f32':
- case 'f64':
- case 'f80':
- case 'f128':
- case 'tf32':
- case 'f8E5M2':
- case 'f8E4M3':
- case 'f8E4M3FN':
- case 'f8E5M2FNUZ':
- case 'f8E4M3FNUZ':
- case 'f8E4M3B11FNUZ':
- case 'f8E3M4':
- case 'f8E8M0FNU':
- case 'f4E2M1FN':
- case 'f6E2M3FN':
- case 'f6E3M2FN':
- return new _.FloatType(value);
- default:
- throw new mlir.Error(`Invalid type '${value}' ${this.location()}`);
- }
- }
- default: {
- break;
- }
- }
- throw new mlir.Error(`Invalid type '${this.getTokenSpelling().str()}' ${this.location()}`);
- }
- parseExtendedType() {
- return this.parseExtendedSymbol(this.state.asmState, this.state.typeAliasDefinitions, (dialectName, symbolData) => {
- const dialect = this.context.getOrLoadDialect(dialectName);
- if (dialect) {
- const curLexerPos = this.getToken().loc.position;
- this.resetToken(symbolData.data());
- const customParser = new _.CustomDialectAsmParser(symbolData, this);
- const type = dialect.parseType(customParser, dialectName);
- this.resetToken(curLexerPos);
- return type;
- }
- return new _.OpaqueType(dialectName, symbolData);
- });
- }
- parseExtendedSymbol(asmState, aliases, createSymbol) {
- const tok = this.getToken();
- const startPos = tok.loc.position;
- const tokSpelling = tok.getSpelling().str();
- const isType = tokSpelling.startsWith('!');
- const identifier = tok.getSpelling().drop_front();
- this.consumeToken();
- const dotIndex = identifier.indexOf('.');
- let dialectName = identifier;
- let symbolData = '';
- if (dotIndex !== -1) {
- dialectName = identifier.substring(0, dotIndex);
- symbolData = identifier.substring(dotIndex + 1);
- }
- const isPrettyName = symbolData.length > 0;
- const hasTrailingData = this.getToken().is(_.Token.less);
- if (!hasTrailingData && !isPrettyName) {
- if (!aliases.has(tokSpelling)) {
- throw new mlir.Error(`Undefined symbol alias '${identifier}' ${this.location()}`);
- }
- if (asmState) {
- if (isType) {
- asmState.addTypeAliasUses(identifier);
- } else {
- asmState.addAttrAliasUses(identifier);
- }
- }
- return aliases.get(tokSpelling);
- }
- if (hasTrailingData) {
- if (this.getToken().is(_.Token.less)) {
- symbolData += this.skip(_.Token.less);
- } else if (this.getToken().is(_.Token.l_paren)) {
- symbolData += this.skip(_.Token.l_paren);
- }
- if (!isPrettyName) {
- symbolData = symbolData.slice(1, -1);
- }
- }
- const posOffset = dotIndex === -1 ? dialectName.length : dotIndex;
- symbolData = new _.StringRef(symbolData, startPos + 1 + posOffset + 1);
- return createSymbol(dialectName, symbolData);
- }
- parseFunctionType() {
- const inputs = this.parseTypeListParens();
- this.parseToken(_.Token.arrow, "Expected '->' in function type");
- const results = this.parseFunctionResultTypes();
- return new _.FunctionType(inputs, results);
- }
- parseFunctionResultTypes() {
- if (this.getToken().is(_.Token.l_paren)) {
- return this.parseTypeListParens();
- }
- const type = this.parseNonFunctionType();
- if (type) {
- return [type];
- }
- return [];
- }
- parseCommaSeparatedList(delimiter, parseElement) {
- const results = [];
- const delimiters = {
- none: [null, null],
- paren: [_.Token.l_paren, _.Token.r_paren],
- square: [_.Token.l_square, _.Token.r_square],
- angle: [_.Token.less, _.Token.greater],
- brace: [_.Token.l_brace, _.Token.r_brace],
- optionalParen: [_.Token.l_paren, _.Token.r_paren],
- optionalSquare: [_.Token.l_square, _.Token.r_square],
- optionalAngle: [_.Token.less, _.Token.greater],
- optionalBrace: [_.Token.l_brace, _.Token.r_brace]
- };
- const [open, close] = delimiters[delimiter] || [null, null];
- const isOptional = delimiter && delimiter.startsWith('optional');
- if (open) {
- if (isOptional) {
- if (!this.consumeIf(open)) {
- return results;
- }
- } else {
- this.parseToken(open, `Expected '${open}'`);
- }
- if (close && this.consumeIf(close)) {
- return results;
- }
- }
- const first = parseElement();
- if (first !== null && first !== undefined) {
- results.push(first);
- }
- while (this.consumeIf(_.Token.comma)) {
- const elem = parseElement();
- if (elem !== null && elem !== undefined) {
- results.push(elem);
- }
- }
- if (close) {
- this.parseToken(close, `Expected '${close}'`);
- }
- return results;
- }
- parseFloatAttr(type = null, isNegative = false) {
- const val = this.getToken().getFloatingPointValue();
- if (val === null) {
- throw new mlir.Error(`Floating point value too large for attribute ${this.location()}`);
- }
- this.consumeToken(_.Token.floatliteral);
- if (!type) {
- if (this.consumeIf(_.Token.colon)) {
- type = this.parseType();
- } else {
- type = new _.FloatType('f64');
- }
- }
- if (type instanceof _.FloatType === false) {
- throw new mlir.Error(`Floating point value not valid for specified type ${this.location()}`);
- }
- return new _.FloatAttr(type, isNegative ? -val : val);
- }
- parseDecOrHexAttr(type = null, isNegative = false) {
- const tok = this.getToken();
- const spelling = tok.getSpelling().str();
- this.consumeToken(_.Token.integer);
- if (!type) {
- if (this.consumeIf(_.Token.colon)) {
- type = this.parseType();
- } else {
- type = new _.IntegerType('i64');
- }
- }
- if (type instanceof _.FloatType) {
- return new _.FloatAttr(type, isNegative ? -spelling : spelling);
- }
- if (!(type instanceof _.IntegerType) && !(type instanceof _.IndexType)) {
- throw new mlir.Error(`Integer literal not valid for specified type ${this.location()}`);
- }
- if (isNegative && type instanceof _.IntegerType && type.name.startsWith('ui')) {
- throw new mlir.Error(`Negative integer literal not valid for unsigned integer type ${this.location()}`);
- }
- const val = tok.spelling.getAsInteger(spelling.length > 1 && spelling[1] === 'x' ? 0 : 10);
- if (val === null) {
- throw new mlir.Error(`Integer constant out of range for attribute ${this.location()}`);
- }
- return new _.IntegerAttr(type, isNegative ? -val : val);
- }
- parseAttribute(type = null) {
- switch (this.getToken().kind) {
- case _.Token.kw_affine_map: {
- this.consumeToken(_.Token.kw_affine_map);
- this.parseToken(_.Token.less, "Expected '<' in affine map");
- const map = this.parseAffineMapReference();
- this.parseToken(_.Token.greater, "Expected '>' in affine map");
- return new _.AffineMapAttr(map);
- }
- case _.Token.kw_affine_set: {
- this.consumeToken(_.Token.kw_affine_set);
- this.parseToken(_.Token.less, "Expected '<' in integer set");
- const set = this.parseIntegerSetReference();
- this.parseToken(_.Token.greater, "Expected '>' in integer set");
- return new _.IntegerSetAttr(set);
- }
- case _.Token.l_square: {
- this.consumeToken(_.Token.l_square);
- const elements = [];
- this.parseCommaSeparatedListUntil(_.Token.r_square, () => {
- elements.push(this.parseAttribute());
- });
- return new _.ArrayAttr(elements);
- }
- case _.Token.kw_false:
- this.consumeToken(_.Token.kw_false);
- return new _.BoolAttr(false);
- case _.Token.kw_true:
- this.consumeToken(_.Token.kw_true);
- return new _.BoolAttr(true);
- case _.Token.kw_dense:
- return this.parseDenseElementsAttr(type);
- case _.Token.kw_dense_resource:
- return this.parseDenseResourceElementsAttr(type);
- case _.Token.kw_array:
- return this.parseDenseArrayAttr(type);
- case _.Token.l_brace: {
- const attributes = new Map();
- this.parseAttributeDict(attributes);
- return new _.DictionaryAttr(attributes);
- }
- case _.Token.hash_identifier:
- return this.parseExtendedAttr(type);
- case _.Token.floatliteral:
- return this.parseFloatAttr(type, false);
- case _.Token.integer:
- return this.parseDecOrHexAttr(type, false);
- case _.Token.minus: {
- this.consumeToken(_.Token.minus);
- if (this.getToken().is(_.Token.integer)) {
- return this.parseDecOrHexAttr(type, true);
- }
- if (this.getToken().is(_.Token.floatliteral)) {
- return this.parseFloatAttr(type, true);
- }
- throw new mlir.Error(`Expected constant integer or floating point value ${this.location()}`);
- }
- case _.Token.kw_loc:
- return this.parseOptionalLocationSpecifier();
- case _.Token.kw_sparse:
- return this.parseSparseElementsAttr(type);
- case _.Token.kw_strided:
- return this.parseStridedLayoutAttr();
- case _.Token.kw_distinct:
- return this.parseDistinctAttr(type);
- case _.Token.string: {
- const value = this.getToken().getStringValue();
- this.consumeToken(_.Token.string);
- if (!type && this.consumeIf(_.Token.colon)) {
- type = this.parseType();
- }
- return new _.StringAttr(value, type);
- }
- case _.Token.at_identifier: {
- const nameStr = this.getToken().getSymbolReference();
- this.consumeToken(_.Token.at_identifier);
- const nestedRefs = [];
- while (this.getToken().is(_.Token.colon)) {
- const curPointer = this.getToken().loc.position;
- this.consumeToken(_.Token.colon);
- if (!this.consumeIf(_.Token.colon)) {
- this.resetToken(curPointer);
- break;
- }
- if (this.getToken().isNot(_.Token.at_identifier)) {
- throw new mlir.Error(`Expected nested symbol reference identifier ${this.location()}`);
- }
- nestedRefs.push(this.getToken().getSymbolReference());
- this.consumeToken(_.Token.at_identifier);
- }
- return new _.SymbolRefAttr(nameStr, nestedRefs);
- }
- case _.Token.kw_unit:
- this.consumeToken(_.Token.kw_unit);
- return new _.UnitAttr();
- case _.Token.bare_identifier: {
- const tokenValue = this.getTokenSpelling().str();
- if (tokenValue === 'tensor' || tokenValue === 'vector' || tokenValue === 'memref' ||
- tokenValue === 'none' || tokenValue === 'index' || /^[su]?i[0-9]+$/.test(tokenValue) ||
- /^f[0-9]+$/.test(tokenValue) || tokenValue === 'bf16' || tokenValue === 'tf32' ||
- /^f\d+E\d+M\d+/.test(tokenValue) || tokenValue === 'complex' || tokenValue === 'tuple') {
- const parsedType = this.parseType();
- return { value: parsedType, type: new _.PrimitiveType('type') };
- }
- if (tokenValue === 'DEFAULT') {
- this.consumeToken(_.Token.bare_identifier);
- return { value: tokenValue };
- }
- this.consumeToken(_.Token.bare_identifier);
- if (this.getToken().is(_.Token.less)) {
- return { value: tokenValue + this.skip('<') };
- }
- return { value: tokenValue };
- }
- case _.Token.exclamation_identifier: {
- const parsedType = this.parseType();
- return { value: parsedType, type: new _.PrimitiveType('type') };
- }
- case _.Token.percent_identifier: {
- const value = this.getTokenSpelling().str();
- this.consumeToken(_.Token.percent_identifier);
- return { value };
- }
- case _.Token.less: {
- const value = this.skip('<');
- return { value };
- }
- default: {
- const parsedType = this.parseOptionalType();
- if (parsedType) {
- return new _.TypeAttrOf(parsedType);
- }
- throw new mlir.Error(`Unexpected attribute token '${this.getTokenSpelling().str()}' ${this.location()}`);
- }
- }
- }
- parseExtendedAttr(type = null) {
- const attr = this.parseExtendedSymbol(this.state.asmState, this.state.attributeAliasDefinitions, (dialectName, symbolData) => {
- let attrType = type;
- if (this.consumeIf(_.Token.colon)) {
- attrType = this.parseType();
- if (!attrType) {
- return new _.Attribute();
- }
- }
- const dialect = this.context.getOrLoadDialect(dialectName);
- if (dialect) {
- const curLexerPos = this.getToken().loc.position;
- this.resetToken(symbolData.data());
- const customParser = new _.CustomDialectAsmParser(symbolData, this);
- const attr = dialect.parseAttribute(customParser, attrType);
- this.resetToken(curLexerPos);
- if (attr) {
- return attr;
- }
- }
- return new _.OpaqueAttr(`#${dialectName}`, symbolData, attrType);
- });
- return attr;
- }
- parseResourceHandle(dialect) {
- const name = this.parseOptionalKeywordOrString();
- if (!name) {
- throw new mlir.Error(`Expected identifier key for 'resource' entry ${this.location()}`);
- }
- const resources = this.state.dialectResources;
- if (!resources.has(dialect)) {
- resources.set(dialect, new Map());
- }
- const dialectEntries = resources.get(dialect);
- if (!dialectEntries.has(name)) {
- const handle = dialect.declareResource(name);
- const key = dialect.getResourceKey(handle);
- dialectEntries.set(name, { key, handle });
- }
- const entry = dialectEntries.get(name);
- return entry.handle;
- }
- parseDenseElementsAttr(attrType) {
- this.consumeToken(_.Token.kw_dense);
- this.parseToken(_.Token.less, "Expected '<' after 'dense'");
- let literalParser = null;
- if (!this.consumeIf(_.Token.greater)) {
- literalParser = new _.TensorLiteralParser(this);
- literalParser.parse(/* allowHex */ true);
- this.parseToken(_.Token.greater, "Expected '>'");
- }
- const type = this.parseElementsLiteralType(attrType);
- const value = literalParser ? literalParser.getAttr(type) : null;
- return new _.DenseElementsAttr(value, type);
- }
- parseDenseResourceElementsAttr(attrType) {
- this.consumeToken(_.Token.kw_dense_resource);
- this.parseToken(_.Token.less, "Expected '<' after 'dense_resource'");
- const rawHandle = this.parseResourceHandle(this.context.getOrLoadDialect('builtin'));
- this.parseToken(_.Token.greater, "Expected '>'");
- let type = attrType;
- if (!type) {
- this.parseToken(_.Token.colon, "Expected ':'");
- type = this.parseType();
- }
- return new _.DenseResourceElementsAttr(type, rawHandle);
- }
- parseDenseArrayAttr(/* attrType */) {
- this.consumeToken(_.Token.kw_array);
- this.parseToken(_.Token.less, "Expected '<' after 'array'");
- const arrayType = this.parseType();
- if (!(arrayType instanceof _.IntegerType) && !(arrayType instanceof _.FloatType)) {
- throw new mlir.Error(`Expected integer or float type, got '${arrayType}' ${this.location()}`);
- }
- if (this.consumeIf(_.Token.greater)) {
- return new _.DenseArrayAttr(arrayType, 0, []);
- }
- const arrayValues = [];
- if (this.consumeIf(_.Token.colon)) {
- while (this.getToken().isNot(_.Token.greater)) {
- const val = this.parseAttribute();
- arrayValues.push(val && val.value !== undefined ? val.value : val);
- this.consumeIf(_.Token.comma);
- }
- }
- this.parseToken(_.Token.greater, "Expected '>' to close an array attribute");
- return new _.DenseArrayAttr(arrayType, arrayValues.length, arrayValues);
- }
- parseSparseElementsAttr(attrType) {
- this.consumeToken(_.Token.kw_sparse);
- this.parseToken(_.Token.less, "Expected '<' after 'sparse'");
- let indices = null;
- let values = null;
- if (!this.consumeIf(_.Token.greater)) {
- const indiceParser = new _.TensorLiteralParser(this);
- indiceParser.parse(/* allowHex */ false);
- indices = indiceParser._storage.map((elem) => {
- const val = elem.isNegative ? -elem.value : elem.value;
- return typeof val === 'string' ? parseInt(val, 10) : val;
- });
- this.parseToken(_.Token.comma, "Expected ','");
- const valuesParser = new _.TensorLiteralParser(this);
- valuesParser.parse(/* allowHex */ true);
- if (valuesParser._hexStorage) {
- values = valuesParser._hexStorage;
- } else {
- values = valuesParser._storage.map((elem) => {
- if (elem.kind === 'float') {
- const val = parseFloat(elem.value);
- return elem.isNegative ? -val : val;
- }
- const val = elem.isNegative ? -elem.value : elem.value;
- return typeof val === 'string' ? parseFloat(val) : val;
- });
- }
- this.parseToken(_.Token.greater, "Expected '>'");
- }
- const type = this.parseElementsLiteralType(attrType);
- return new _.SparseElementsAttr(type, indices, values);
- }
- parseStridedLayoutAttr() {
- this.consumeToken(_.Token.kw_strided);
- this.parseToken(_.Token.less, "Expected '<' after 'strided'");
- this.parseToken(_.Token.l_square, "Expected '['");
- // Parse dimension list: integer or '?' separated by commas
- const strides = [];
- while (this.getToken().isNot(_.Token.r_square)) {
- const value = this.parseOptionalInteger();
- if (value !== null) {
- strides.push(value);
- } else if (this.consumeIf(_.Token.question)) {
- strides.push('?');
- } else {
- throw new mlir.Error(`Expected a 64-bit signed integer or '?' in strided layout ${this.location()}`);
- }
- if (this.getToken().isNot(_.Token.r_square)) {
- this.consumeIf(_.Token.comma);
- }
- }
- this.parseToken(_.Token.r_square, "Expected ']'");
- let offset = null;
- if (this.consumeIf(_.Token.comma)) {
- this.parseToken(_.Token.kw_offset, "Expected 'offset' after comma");
- this.parseToken(_.Token.colon, "Expected ':' after 'offset'");
- offset = this.parseOptionalInteger();
- if (offset === null) {
- if (this.consumeIf(_.Token.question)) {
- offset = '?';
- } else {
- throw new mlir.Error(`Expected a 64-bit signed integer or '?' for offset in strided layout ${this.location()}`);
- }
- }
- }
- this.parseToken(_.Token.greater, "Expected '>'");
- const stridesStr = `[${strides.join(', ')}]`;
- const offsetStr = offset === null ? '' : `, offset: ${offset}`;
- return { value: `strided<${stridesStr}${offsetStr}>`, type: 'strided', strides, offset };
- }
- // Parse an affine map reference using AffineParser
- // Following reference: Parser::parseAffineMapReference
- parseAffineMapReference() {
- const affineParser = new _.AffineParser(this.state);
- const result = affineParser.parseAffineMapOrIntegerSetInline();
- if (result.set) {
- throw new mlir.Error(`Expected AffineMap, but got IntegerSet ${this.location()}`);
- }
- return result.map;
- }
- // Parse an integer set reference using AffineParser
- // Following reference: Parser::parseIntegerSetReference
- parseIntegerSetReference() {
- const affineParser = new _.AffineParser(this.state);
- const result = affineParser.parseAffineMapOrIntegerSetInline();
- if (result.map) {
- throw new mlir.Error(`Expected IntegerSet, but got AffineMap ${this.location()}`);
- }
- return result.set;
- }
- // Parse an AffineMap where dims/symbols are SSA ids
- // Following reference: Parser::parseAffineMapOfSSAIds
- parseAffineMapOfSSAIds(parseElement, delimiter = 'Paren') {
- const affineParser = new _.AffineParser(this.state, true, parseElement);
- return affineParser.parseAffineMapOfSSAIds(delimiter);
- }
- // Parse an AffineExpr where dims/symbols are SSA ids
- // Following reference: Parser::parseAffineExprOfSSAIds
- parseAffineExprOfSSAIds(parseElement) {
- const affineParser = new _.AffineParser(this.state, true, parseElement);
- return affineParser.parseAffineExprOfSSAIds();
- }
- parseDistinctAttr(type) {
- this.consumeToken(_.Token.kw_distinct);
- this.parseToken(_.Token.l_square, "Expected '[' after 'distinct'");
- if (this.getToken().isNot(_.Token.integer)) {
- throw new mlir.Error(`Expected distinct ID ${this.location()}`);
- }
- const token = this.getToken();
- const id = token.getUInt64IntegerValue();
- if (id === null) {
- throw new mlir.Error(`Expected an unsigned 64-bit integer ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- this.parseToken(_.Token.r_square, "Expected ']' to close distinct ID");
- this.parseToken(_.Token.less, "Expected '<' after distinct ID");
- let referencedAttr = null;
- if (this.consumeIf(_.Token.greater)) {
- referencedAttr = new _.UnitAttr();
- } else {
- referencedAttr = this.parseAttribute(type);
- this.parseToken(_.Token.greater, "Expected '>' to close distinct attribute");
- }
- return { value: `distinct[${id.toString()}]`, referencedAttr, type: 'distinct' };
- }
- parseElementsLiteralType(type) {
- if (!type) {
- this.parseToken(_.Token.colon, "Expected ':'");
- return this.parseType();
- }
- // Type is a concrete type object - use it directly
- return type;
- }
- parseOptionalAttribute(type) {
- switch (this.getToken().kind) {
- case _.Token.at_identifier:
- case _.Token.percent_identifier:
- case _.Token.integer:
- case _.Token.floatliteral:
- case _.Token.hash_identifier:
- case _.Token.l_square:
- case _.Token.l_brace:
- case _.Token.less:
- case _.Token.string:
- return this.parseAttribute(type);
- case _.Token.minus:
- case _.Token.kw_loc:
- case _.Token.kw_affine_map:
- case _.Token.kw_affine_set:
- case _.Token.kw_dense:
- case _.Token.kw_dense_resource:
- case _.Token.kw_array:
- case _.Token.kw_sparse:
- case _.Token.kw_strided:
- case _.Token.kw_distinct:
- case _.Token.kw_unit:
- case _.Token.kw_true:
- case _.Token.kw_false:
- return this.parseAttribute(type);
- default: {
- const value = this.parseOptionalType(type);
- if (value) {
- return { value, type: 'type' };
- }
- return null;
- }
- }
- }
- parseOptionalInteger() {
- // Parse `false` and `true` keywords as 0 and 1 respectively.
- if (this.consumeIf(_.Token.kw_false)) {
- return 0;
- }
- if (this.consumeIf(_.Token.kw_true)) {
- return 1;
- }
- if (this.getToken().isNot(_.Token.integer) && this.getToken().isNot(_.Token.minus)) {
- return null;
- }
- const negative = this.consumeIf(_.Token.minus);
- const curTok = this.getToken();
- this.parseToken(_.Token.integer, 'Expected integer value');
- const spelling = curTok.spelling;
- const isHex = spelling.length > 1 && spelling[1] === 'x';
- const result = spelling.getAsInteger(isHex ? 0 : 10);
- if (result === null) {
- throw new mlir.Error(`Integer value too large ${this.location()}`);
- }
- return negative ? -result : result;
- }
- parseInteger() {
- const result = this.parseOptionalInteger();
- if (result === null) {
- throw new mlir.Error(`Expected integer value ${this.location()}`);
- }
- return result;
- }
- parseString() {
- const value = this.getToken().getSpelling().str();
- this.consumeToken(_.Token.string);
- return value;
- }
- parseOptionalVerticalBar() {
- return this.consumeIf(_.Token.vertical_bar);
- }
- parseOptionalString() {
- if (!this.getToken().is(_.Token.string)) {
- return null;
- }
- const string = this.getToken().getStringValue();
- this.consumeToken();
- return string;
- }
- getToken() {
- return this.state.curToken;
- }
- getTokenSpelling() {
- return this.state.curToken.spelling;
- }
- resetToken(tokPos) {
- this.state.lex.resetPointer(tokPos);
- this.state.curToken = this.state.lex.lexToken();
- }
- consumeToken(kind) {
- if (kind !== undefined && this.state.curToken.kind !== kind) {
- throw new mlir.Error(`consumeToken: Expected '${kind}' ${this.location()}`);
- }
- this.state.curToken = this.state.lex.lexToken();
- }
- parseToken(kind, message) {
- if (this.consumeIf(kind)) {
- return;
- }
- throw new mlir.Error(`${message || `Expected '${kind}'`} ${this.location()}`);
- }
- consumeIf(kind) {
- if (this.state.curToken.isNot(kind)) {
- return false;
- }
- this.consumeToken(kind);
- return true;
- }
- location() {
- return this.getToken().loc.toString();
- }
- static parseSymbol(inputStr, context, parserFn) {
- const decoder = text.Decoder.open(inputStr);
- const config = new _.ParserConfig(context);
- const state = new _.ParserState(decoder, config);
- const parser = new _.Parser(state, context);
- const startPos = state.curToken.position || 0;
- const symbol = parserFn(parser);
- if (!symbol) {
- return { symbol: null, numRead: 0 };
- }
- const endPos = parser.getToken().loc.position;
- const numRead = endPos - startPos;
- return { symbol, numRead };
- }
- static parseAttribute(attrStr, context, type = null) {
- const result = _.Parser.parseSymbol(attrStr, context, (parser) => {
- return parser.parseAttribute(type);
- });
- if (!result.symbol) {
- return { attribute: null, numRead: 0 };
- }
- return { attribute: result.symbol, numRead: result.numRead };
- }
- static parseType(typeStr, context) {
- const result = _.Parser.parseSymbol(typeStr, context, (parser) => {
- return parser.parseType();
- });
- if (!result.symbol) {
- return { type: null, numRead: 0 };
- }
- return { type: result.symbol, numRead: result.numRead };
- }
- };
- _.TopLevelOperationParser = class extends _.Parser {
- parse(block) {
- const opParser = new _.OperationParser(this.state);
- while (true) {
- switch (this.getToken().kind) {
- case _.Token.eof: {
- opParser.finalize(block);
- return block;
- }
- case _.Token.hash_identifier: {
- this.parseAttributeAliasDef();
- break;
- }
- case _.Token.exclamation_identifier: {
- this.parseTypeAliasDef();
- break;
- }
- case _.Token.file_metadata_begin: {
- this.parseFileMetadataDictionary();
- break;
- }
- default: {
- const op = opParser.parseOperation();
- block.operations.push(op);
- }
- }
- }
- }
- parseAttributeAliasDef() {
- const aliasName = this.getToken().getSpelling().str();
- this.consumeToken(_.Token.hash_identifier);
- this.parseToken(_.Token.equal, "Expected '=' in attribute alias definition");
- let attr = null;
- if (this.getToken().is(_.Token.l_paren)) {
- const dims = this.skip(_.Token.l_paren);
- const symbols = this.getToken().is(_.Token.l_square) ? this.skip(_.Token.l_square) : '';
- this.parseToken(_.Token.arrow, "Expected '->'");
- const results = this.getToken().is(_.Token.l_paren) ? this.skip(_.Token.l_paren) : '';
- attr = { value: `affine_map<${dims}${symbols} -> ${results}>`, name: 'affine_map' };
- } else {
- attr = this.parseAttribute();
- }
- this.state.attributeAliasDefinitions.set(aliasName, attr);
- }
- parseTypeAliasDef() {
- const aliasName = this.getTokenSpelling().str();
- this.consumeToken(_.Token.exclamation_identifier);
- this.parseToken(_.Token.equal, "Expected '=' in type alias definition");
- if (this.getToken().is(_.Token.bare_identifier) && this.getTokenSpelling().str() === 'type') {
- this.consumeToken(_.Token.bare_identifier);
- }
- const type = this.parseType();
- this.state.typeAliasDefinitions.set(aliasName, type);
- }
- parseFileMetadataDictionary() {
- this.consumeToken(_.Token.file_metadata_begin);
- this.parseCommaSeparatedListUntil(_.Token.file_metadata_end, () => {
- const keyLoc = this.getToken().loc;
- const key = this.parseOptionalKeyword();
- if (!key) {
- throw new mlir.Error(`Expected identifier key in file metadata dictionary ${this.location()}`);
- }
- this.parseToken(_.Token.colon, "Expected ':'");
- if (key === 'dialect_resources') {
- this.parseDialectResourceFileMetadata();
- } else if (key === 'external_resources') {
- this.parseExternalResourceFileMetadata();
- } else {
- throw new mlir.Error(`Unknown key '${key}' in file metadata dictionary ${keyLoc.toString()}`);
- }
- });
- }
- };
- // Specialized parser for affine structures (affine maps, affine expressions, integer sets)
- // Following the reference implementation in AffineParser.cpp
- _.AffineParser = class extends _.Parser {
- constructor(state, allowParsingSSAIds = false, parseElement = null) {
- super(state);
- this.allowParsingSSAIds = allowParsingSSAIds;
- this.parseElement = parseElement;
- this.dimsAndSymbols = []; // Array of {name, expr} pairs
- this.numDimOperands = 0;
- this.numSymbolOperands = 0;
- }
- // Parse an ambiguous affine map or integer set inline
- // affine-map-or-integer-set ::= dim-and-symbol-id-lists (`->` | `:`) ...
- parseAffineMapOrIntegerSetInline() {
- const { numDims, numSymbols } = this.parseDimAndOptionalSymbolIdList();
- if (this.consumeIf(_.Token.arrow)) {
- return { map: this.parseAffineMapRange(numDims, numSymbols), set: null };
- }
- if (!this.consumeIf(_.Token.colon)) {
- throw new mlir.Error(`Expected '->' or ':' ${this.location()}`);
- }
- return { map: null, set: this.parseIntegerSetConstraints(numDims, numSymbols) };
- }
- // Parse dimension and optional symbol identifier lists: (d0, d1, ...)[s0, s1, ...]
- parseDimAndOptionalSymbolIdList() {
- const numDims = this.parseDimIdList();
- let numSymbols = 0;
- if (this.getToken().is(_.Token.l_square)) {
- numSymbols = this.parseSymbolIdList();
- }
- return { numDims, numSymbols };
- }
- // Parse dimension identifier list: (d0, d1, ...)
- parseDimIdList() {
- let numDims = 0;
- this.parseToken(_.Token.l_paren, "Expected '(' in dimensional identifier list");
- if (this.getToken().isNot(_.Token.r_paren)) {
- do {
- const dimExpr = this.getAffineDimExpr(numDims);
- this.parseIdentifierDefinition(dimExpr);
- numDims++;
- } while (this.consumeIf(_.Token.comma));
- }
- this.parseToken(_.Token.r_paren, "Expected ')' in dimensional identifier list");
- return numDims;
- }
- // Parse symbol identifier list: [s0, s1, ...]
- parseSymbolIdList() {
- let numSymbols = 0;
- this.parseToken(_.Token.l_square, "Expected '[' in symbol list");
- if (this.getToken().isNot(_.Token.r_square)) {
- do {
- const symbolExpr = this.getAffineSymbolExpr(numSymbols);
- this.parseIdentifierDefinition(symbolExpr);
- numSymbols++;
- } while (this.consumeIf(_.Token.comma));
- }
- this.parseToken(_.Token.r_square, "Expected ']' in symbol list");
- return numSymbols;
- }
- parseIdentifierDefinition(idExpr) {
- let name = null;
- if (this.getToken().is(_.Token.bare_identifier)) {
- name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.bare_identifier);
- } else if (this.getToken().is(_.Token.inttype)) {
- name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.inttype);
- } else if (this.getToken().isKeyword()) {
- name = this.getToken().getSpelling().str();
- this.consumeToken();
- } else {
- throw new mlir.Error(`Expected bare identifier ${this.location()}`);
- }
- for (const entry of this.dimsAndSymbols) {
- if (entry.name === name) {
- throw new mlir.Error(`Redefinition of identifier '${name}' ${this.location()}`);
- }
- }
- this.dimsAndSymbols.push({ name, expr: idExpr });
- }
- parseAffineMapRange(numDims, numSymbols) {
- const exprs = [];
- this.parseToken(_.Token.l_paren, "Expected '('");
- if (this.getToken().isNot(_.Token.r_paren)) {
- do {
- const expr = this.parseAffineExpr();
- if (!expr) {
- throw new mlir.Error(`Failed to parse affine expression ${this.location()}`);
- }
- exprs.push(expr);
- } while (this.consumeIf(_.Token.comma));
- }
- this.parseToken(_.Token.r_paren, "Expected ')'");
- return _.AffineMap.get(numDims, numSymbols, exprs);
- }
- parseIntegerSetConstraints(numDims, numSymbols) {
- const constraints = [];
- const isEqs = [];
- this.parseToken(_.Token.l_paren, "Expected '('");
- if (this.getToken().isNot(_.Token.r_paren)) {
- do {
- const { expr, isEq } = this.parseAffineConstraint();
- constraints.push(expr);
- isEqs.push(isEq);
- } while (this.consumeIf(_.Token.comma));
- }
- this.parseToken(_.Token.r_paren, "Expected ')'");
- // If no constraints, return degenerate true set (0 == 0)
- if (constraints.length === 0) {
- const zero = this.getAffineConstantExpr(0);
- return _.IntegerSet.get(numDims, numSymbols, [zero], [true]);
- }
- return _.IntegerSet.get(numDims, numSymbols, constraints, isEqs);
- }
- parseAffineConstraint() {
- const lhsExpr = this.parseAffineExpr();
- if (!lhsExpr) {
- throw new mlir.Error(`Expected affine expression ${this.location()}`);
- }
- if (this.consumeIf(_.Token.greater) && this.consumeIf(_.Token.equal)) {
- const rhsExpr = this.parseAffineExpr();
- return { expr: this.makeSubExpr(lhsExpr, rhsExpr), isEq: false };
- }
- if (this.consumeIf(_.Token.less) && this.consumeIf(_.Token.equal)) {
- const rhsExpr = this.parseAffineExpr();
- return { expr: this.makeSubExpr(rhsExpr, lhsExpr), isEq: false };
- }
- if (this.consumeIf(_.Token.equal) && this.consumeIf(_.Token.equal)) {
- const rhsExpr = this.parseAffineExpr();
- return { expr: this.makeSubExpr(lhsExpr, rhsExpr), isEq: true };
- }
- throw new mlir.Error(`Expected '>=', '<=', or '==' after affine expression ${this.location()}`);
- }
- makeSubExpr(lhs, rhs) {
- const negOne = this.getAffineConstantExpr(-1);
- const negRhs = this.getAffineBinaryOpExpr(_.AffineExprKind.Mul, negOne, rhs);
- return this.getAffineBinaryOpExpr(_.AffineExprKind.Add, lhs, negRhs);
- }
- parseAffineExpr() {
- return this.parseAffineLowPrecOpExpr(null, null);
- }
- parseAffineLowPrecOpExpr(llhs, llhsOp) {
- const lhs = this.parseAffineOperandExpr(llhs);
- if (!lhs) {
- return null;
- }
- const lOp = this.consumeIfLowPrecOp();
- if (lOp) {
- if (llhs) {
- const sum = this.makeAffineBinaryOp(llhsOp, llhs, lhs);
- return this.parseAffineLowPrecOpExpr(sum, lOp);
- }
- return this.parseAffineLowPrecOpExpr(lhs, lOp);
- }
- const hOp = this.consumeIfHighPrecOp();
- if (hOp) {
- const highRes = this.parseAffineHighPrecOpExpr(lhs, hOp);
- if (!highRes) {
- return null;
- }
- const expr = llhs ? this.makeAffineBinaryOp(llhsOp, llhs, highRes) : highRes;
- const nextOp = this.consumeIfLowPrecOp();
- if (nextOp) {
- return this.parseAffineLowPrecOpExpr(expr, nextOp);
- }
- return expr;
- }
- if (llhs) {
- return this.makeAffineBinaryOp(llhsOp, llhs, lhs);
- }
- return lhs;
- }
- parseAffineHighPrecOpExpr(llhs, llhsOp) {
- const lhs = this.parseAffineOperandExpr(llhs);
- if (!lhs) {
- return null;
- }
- const op = this.consumeIfHighPrecOp();
- if (op) {
- if (llhs) {
- const expr = this.makeAffineBinaryOp(llhsOp, llhs, lhs);
- return this.parseAffineHighPrecOpExpr(expr, op);
- }
- return this.parseAffineHighPrecOpExpr(lhs, op);
- }
- if (llhs) {
- return this.makeAffineBinaryOp(llhsOp, llhs, lhs);
- }
- return lhs;
- }
- consumeIfLowPrecOp() {
- switch (this.getToken().kind) {
- case _.Token.plus:
- this.consumeToken(_.Token.plus);
- return _.AffineLowPrecOp.Add;
- case _.Token.minus:
- this.consumeToken(_.Token.minus);
- return _.AffineLowPrecOp.Sub;
- default:
- return _.AffineLowPrecOp.LNoOp;
- }
- }
- consumeIfHighPrecOp() {
- switch (this.getToken().kind) {
- case _.Token.star:
- this.consumeToken(_.Token.star);
- return _.AffineHighPrecOp.Mul;
- case _.Token.kw_floordiv:
- this.consumeToken(_.Token.kw_floordiv);
- return _.AffineHighPrecOp.FloorDiv;
- case _.Token.kw_ceildiv:
- this.consumeToken(_.Token.kw_ceildiv);
- return _.AffineHighPrecOp.CeilDiv;
- case _.Token.kw_mod:
- this.consumeToken(_.Token.kw_mod);
- return _.AffineHighPrecOp.Mod;
- default:
- return _.AffineHighPrecOp.HNoOp;
- }
- }
- // Create affine binary op expression
- makeAffineBinaryOp(op, lhs, rhs) {
- if (op === 'Sub') {
- // Subtraction: lhs - rhs = lhs + (-1 * rhs)
- const negOne = this.getAffineConstantExpr(-1);
- const negRhs = this.getAffineBinaryOpExpr(_.AffineExprKind.Mul, negOne, rhs);
- return this.getAffineBinaryOpExpr(_.AffineExprKind.Add, lhs, negRhs);
- }
- return this.getAffineBinaryOpExpr(op, lhs, rhs);
- }
- parseAffineOperandExpr(lhs) {
- if (this.getToken().is(_.Token.kw_symbol)) {
- return this.parseSymbolSSAIdExpr();
- }
- if (this.getToken().is(_.Token.percent_identifier)) {
- return this.parseSSAIdExpr(false);
- }
- if (this.getToken().is(_.Token.integer)) {
- return this.parseIntegerExpr();
- }
- if (this.getToken().is(_.Token.l_paren)) {
- return this.parseParentheticalExpr();
- }
- if (this.getToken().is(_.Token.minus)) {
- return this.parseNegateExpression(lhs);
- }
- if (this.getToken().isAny(_.Token.bare_identifier, _.Token.inttype)) {
- return this.parseBareIdExpr();
- }
- if (this.getToken().isKeyword()) {
- return this.parseBareIdExpr();
- }
- if (this.getToken().isAny(_.Token.plus, _.Token.star)) {
- if (lhs) {
- throw new mlir.Error(`Missing right operand of binary operator ${this.location()}`);
- }
- throw new mlir.Error(`Missing left operand of binary operator ${this.location()}`);
- }
- if (lhs) {
- throw new mlir.Error(`Missing right operand of binary operator ${this.location()}`);
- }
- throw new mlir.Error(`Expected affine expression ${this.location()}`);
- }
- parseSymbolSSAIdExpr() {
- this.parseToken(_.Token.kw_symbol);
- this.parseToken(_.Token.l_paren);
- const symbolExpr = this.parseSSAIdExpr(true);
- this.parseToken(_.Token.r_paren);
- return symbolExpr;
- }
- parseSSAIdExpr(isSymbol) {
- if (!this.allowParsingSSAIds) {
- throw new mlir.Error(`Unexpected SSA identifier ${this.location()}`);
- }
- if (this.getToken().isNot(_.Token.percent_identifier)) {
- throw new mlir.Error(`Expected SSA identifier ${this.location()}`);
- }
- const name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.percent_identifier);
- for (const entry of this.dimsAndSymbols) {
- if (entry.name === name) {
- return entry.expr;
- }
- }
- if (this.parseElement) {
- this.parseElement(isSymbol);
- }
- const idExpr = isSymbol
- ? this.getAffineSymbolExpr(this.numSymbolOperands++)
- : this.getAffineDimExpr(this.numDimOperands++);
- this.dimsAndSymbols.push({ name, expr: idExpr });
- return idExpr;
- }
- // Parse integer literal
- parseIntegerExpr() {
- const val = this.getToken().getUInt64IntegerValue();
- this.consumeToken(_.Token.integer);
- return this.getAffineConstantExpr(Number(val));
- }
- // Parse parenthesized expression
- parseParentheticalExpr() {
- this.parseToken(_.Token.l_paren, "Expected '('");
- if (this.getToken().is(_.Token.r_paren)) {
- throw new mlir.Error(`No expression inside parentheses ${this.location()}`);
- }
- const expr = this.parseAffineExpr();
- this.parseToken(_.Token.r_paren, "Expected ')'");
- return expr;
- }
- parseNegateExpression(lhs) {
- this.parseToken(_.Token.minus, "Expected '-'");
- const operand = this.parseAffineOperandExpr(lhs);
- if (!operand) {
- throw new mlir.Error(`Missing operand of negation ${this.location()}`);
- }
- const negOne = this.getAffineConstantExpr(-1);
- return this.getAffineBinaryOpExpr(_.AffineExprKind.Mul, negOne, operand);
- }
- parseBareIdExpr() {
- if (!this.isCurrentTokenAKeyword()) {
- throw new mlir.Error(`Expected bare identifier ${this.location()}`);
- }
- const name = this.getTokenSpelling().str();
- this.consumeToken();
- for (const entry of this.dimsAndSymbols) {
- if (entry.name === name) {
- return entry.expr;
- }
- }
- throw new mlir.Error(`Use of undeclared identifier '${name}' ${this.location()}`);
- }
- parseAffineMapOfSSAIds(delimiter = 'Paren') {
- const exprs = [];
- const open = delimiter === 'Paren' ? _.Token.l_paren : _.Token.l_square;
- const close = delimiter === 'Paren' ? _.Token.r_paren : _.Token.r_square;
- this.parseToken(open, `Expected '${open === _.Token.l_paren ? '(' : '['}'`);
- if (this.getToken().isNot(close)) {
- do {
- const expr = this.parseAffineExpr();
- exprs.push(expr);
- } while (this.consumeIf(_.Token.comma));
- }
- this.parseToken(close, `Expected '${close === _.Token.r_paren ? ')' : ']'}'`);
- return _.AffineMap.get(this.numDimOperands, this.dimsAndSymbols.length - this.numDimOperands, exprs);
- }
- parseAffineExprOfSSAIds() {
- return this.parseAffineExpr();
- }
- getAffineDimExpr(position) {
- return new _.AffineDimExpr(position);
- }
- getAffineSymbolExpr(position) {
- return new _.AffineSymbolExpr(position);
- }
- getAffineConstantExpr(constant) {
- return new _.AffineConstantExpr(constant);
- }
- getAffineBinaryOpExpr(kind, lhs, rhs) {
- return new _.AffineBinaryOpExpr(kind, lhs, rhs);
- }
- };
- _.OperationParser = class extends _.Parser {
- constructor(state) {
- super(state);
- this.isolatedNameScopes = [];
- this.pushSSANameScope(/* isIsolated */ true);
- this._redirect = new Map([
- ['builtin.func', 'func.func'],
- ['builtin.constant', 'arith.constant'],
- ['builtin.return', 'func.return'],
- ['builtin.select', 'arith.select'],
- ['scf.select', 'arith.select'],
- ['scf.call', 'func.call'],
- ['builtin.view', 'memref.view'],
- ['builtin.dealloc', 'memref.dealloc'], ['func.dealloc', 'memref.dealloc'],
- // Arith operations (from both builtin and func default dialects)
- ['builtin.addi', 'arith.addi'], ['func.addi', 'arith.addi'],
- ['builtin.subi', 'arith.subi'], ['func.subi', 'arith.subi'],
- ['builtin.muli', 'arith.muli'], ['func.muli', 'arith.muli'],
- ['builtin.divi_signed', 'arith.divsi'], ['func.divi_signed', 'arith.divsi'],
- ['builtin.divi_unsigned', 'arith.divui'], ['func.divi_unsigned', 'arith.divui'],
- ['builtin.divsi', 'arith.divsi'], ['func.divsi', 'arith.divsi'],
- ['builtin.divui', 'arith.divui'], ['func.divui', 'arith.divui'],
- ['builtin.remi_signed', 'arith.remsi'], ['func.remi_signed', 'arith.remsi'],
- ['builtin.remi_unsigned', 'arith.remui'], ['func.remi_unsigned', 'arith.remui'],
- ['builtin.andi', 'arith.andi'], ['func.andi', 'arith.andi'],
- ['builtin.ori', 'arith.ori'], ['func.ori', 'arith.ori'],
- ['builtin.xori', 'arith.xori'], ['func.xori', 'arith.xori'],
- ['builtin.shli', 'arith.shli'], ['func.shli', 'arith.shli'],
- ['builtin.shrsi', 'arith.shrsi'], ['func.shrsi', 'arith.shrsi'],
- ['builtin.shrui', 'arith.shrui'], ['func.shrui', 'arith.shrui'],
- ['builtin.addf', 'arith.addf'], ['func.addf', 'arith.addf'],
- ['builtin.subf', 'arith.subf'], ['func.subf', 'arith.subf'],
- ['builtin.mulf', 'arith.mulf'], ['func.mulf', 'arith.mulf'],
- ['builtin.divf', 'arith.divf'], ['func.divf', 'arith.divf'],
- ['builtin.cmpi', 'arith.cmpi'], ['func.cmpi', 'arith.cmpi'],
- ['builtin.cmpf', 'arith.cmpf'], ['func.cmpf', 'arith.cmpf'],
- ['builtin.index_cast', 'arith.index_cast'], ['func.index_cast', 'arith.index_cast'],
- ['builtin.sitofp', 'arith.sitofp'], ['func.sitofp', 'arith.sitofp'],
- ['builtin.fptosi', 'arith.fptosi'], ['func.fptosi', 'arith.fptosi'],
- ['builtin.truncf', 'arith.truncf'], ['func.truncf', 'arith.truncf'],
- ['builtin.extf', 'arith.extf'], ['func.extf', 'arith.extf'],
- ['builtin.splat', 'vector.splat'],
- ['func.splat', 'vector.splat'],
- ['scf.splat', 'vector.splat'],
- // Memref operations
- ['builtin.alloc', 'memref.alloc'], ['func.alloc', 'memref.alloc'],
- ['builtin.load', 'memref.load'], ['func.load', 'memref.load'],
- ['builtin.store', 'memref.store'], ['func.store', 'memref.store'],
- ['builtin.subview', 'memref.subview'], ['func.subview', 'memref.subview'],
- ['builtin.dim', 'memref.dim'], ['func.dim', 'memref.dim'],
- ['builtin.view', 'memref.view'], ['func.view', 'memref.view'],
- // Control flow operations
- ['builtin.cond_br', 'cf.cond_br'], ['func.cond_br', 'cf.cond_br'],
- ['builtin.br', 'cf.br'], ['func.br', 'cf.br'],
- ['builtin.switch', 'cf.switch'], ['func.switch', 'cf.switch'],
- ['builtin.assert', 'cf.assert'], ['func.assert', 'cf.assert'],
- // Other redirects
- ['flow.constant', 'flow.tensor.constant'],
- ['util.initializer.return', 'util.return']
- ]);
- }
- finalize(block) {
- const attributeAliases = this.state.attributeAliasDefinitions;
- const deferredRefs = this.state.deferredLocsReferences;
- const resolveLocation = (opOrArgument) => {
- const fwdLoc = opOrArgument.loc;
- if (fwdLoc instanceof _.OpaqueLoc) {
- const locInfo = deferredRefs[fwdLoc.index];
- const identifier = `#${locInfo.identifier}`;
- const attr = attributeAliases.get(identifier);
- if (!attr) {
- throw new mlir.Error(`Operation location alias '${locInfo.identifier}' was never defined.`);
- }
- if (attr instanceof _.LocationAttr === false) {
- throw new mlir.Error(`Expected location but found '${attr}'.`);
- }
- opOrArgument.loc = attr;
- }
- };
- // Walk all operations and resolve locations on ops and block arguments
- const walk = (operations) => {
- for (const op of operations) {
- if (op.loc) {
- resolveLocation(op);
- }
- if (op.body && op.body.blocks) {
- for (const blk of op.body.blocks) {
- if (blk.arguments) {
- for (const arg of blk.arguments) {
- if (arg.loc) {
- resolveLocation(arg);
- }
- }
- }
- if (blk.operations) {
- walk(blk.operations);
- }
- }
- }
- }
- };
- walk(block.operations);
- }
- parseOperation() {
- const loc = this.getToken().loc.copy();
- const resultIDs = [];
- let numExpectedResults = 0;
- if (this.getToken().is(_.Token.percent_identifier)) {
- const parseNextResult = () => {
- const name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.percent_identifier);
- let expectedSubResults = 1;
- if (this.consumeIf(_.Token.colon)) {
- if (this.getToken().isNot(_.Token.integer)) {
- throw new mlir.Error(`Expected integer number of results ${this.location()}`);
- }
- const val = parseInt(this.getToken().getSpelling().str(), 10);
- if (!Number.isInteger(val) || val < 1) {
- throw new mlir.Error(`Expected named operation to have at least 1 result ${this.location()}`);
- }
- this.consumeToken(_.Token.integer);
- expectedSubResults *= val;
- }
- resultIDs.push([name, expectedSubResults, null]);
- numExpectedResults += expectedSubResults;
- return true;
- };
- this.parseCommaSeparatedList('none', parseNextResult);
- this.parseToken(_.Token.equal, "Expected '=' after SSA name");
- }
- let op = null;
- // Reference Parser.cpp:1305: nameTok.is(Token::bare_identifier) || nameTok.isKeyword()
- const nameTok = this.getToken();
- if (nameTok.is(_.Token.bare_identifier) || nameTok.isKeyword()) {
- op = this.parseCustomOperation(resultIDs);
- } else if (nameTok.is(_.Token.string)) {
- op = this.parseGenericOperation();
- } else {
- throw new mlir.Error(`${this.getToken().is(_.Token.eof) ? 'Unexpected end of input' : `Unexpected operation name '${this.getTokenSpelling().str()}'`} ${this.location()}`);
- }
- if (!op) {
- throw new mlir.Error(`Failed to parse operation ${this.location()}`);
- }
- if (resultIDs.length > 0 && op.results.length === 0) {
- throw new mlir.Error(`Cannot name an operation with no results ${loc.toString()}`);
- }
- if (resultIDs.length > 0 && op.results.length !== numExpectedResults) {
- throw new mlir.Error(`Operation '${op.name.getStringRef()}' defines '${op.results.length}' results but was provided '${numExpectedResults}' to bind ${loc.toString()}`);
- }
- let index = 0;
- for (const resIt of resultIDs) {
- for (let subRes = 0; subRes < resIt[1]; subRes++) {
- const result = op.results[index++];
- if (result.name) {
- throw new mlir.Error(`Result '${result.name}' already has name ${this.location()}`);
- }
- // Workaround: Visualization-specific addition is to store name on result for display
- result.name = subRes === 0 ? resIt[0] : `${resIt[0]}.${subRes}`;
- this.addDefinition(new _.UnresolvedOperand(resIt[2], resIt[0], subRes), result);
- }
- }
- return op;
- }
- parseCustomOperation(resultIDs) {
- const opLoc = this.getToken().loc.copy();
- const opNameInfo = this.parseCustomOperationName();
- const opState = new _.OperationState(opLoc, opNameInfo);
- delete opNameInfo.identifier; // Workaround
- const defaultDialect = (opNameInfo && opNameInfo.metadata && opNameInfo.metadata.defaultDialect) || this.state.defaultDialectStack[this.state.defaultDialectStack.length - 1];
- this.state.defaultDialectStack.push(defaultDialect);
- const customParser = new _.CustomOpAsmParser(resultIDs, this);
- if (!opNameInfo.dialect.parseOperation(customParser, opState)) {
- this.state.defaultDialectStack.pop();
- throw new mlir.Error(`Unsupported custom operation '${opState.identifier}' ${this.location()}`);
- }
- if (!opNameInfo.metadata.hasParser && !opNameInfo.metadata.hasCustomAssemblyFormat && opNameInfo.metadata.assemblyFormat && opState.compatibility !== true) {
- throw new mlir.Error(`Operation '${opState.identifier}' has assembly format but was handled by custom dialect code.`);
- }
- opState.loc = this.parseTrailingLocationSpecifier() || {};
- this.state.defaultDialectStack.pop();
- return _.Operation.create(opState);
- }
- parseCustomOperationName() {
- const nameTok = this.getToken();
- if (nameTok.kind !== _.Token.bare_identifier && !nameTok.isKeyword()) {
- throw new mlir.Error(`Expected bare identifier or keyword ${this.location()}`);
- }
- let identifier = nameTok.getSpelling().str(); // Workaround: keep the original source file identifier
- this.consumeToken();
- let index = identifier.indexOf('.');
- if (index === -1) {
- const dialect = this.state.defaultDialectStack[this.state.defaultDialectStack.length - 1];
- identifier = `${dialect}.${identifier}`;
- }
- let opName = identifier;
- if (opName === 'func.constant' && this.getToken().isNot(_.Token.at_identifier)) {
- // Workaround: old std.constant should be arith.constant
- opName = 'arith.constant';
- } else if (opName.startsWith('check.')) {
- // Workaround: Handle conflicting dialects from stablehlo and iree
- const dialect = this.getToken().is(_.Token.l_paren) || this.getToken().is(_.Token.less) ? 'iree' : 'stablehlo';
- opName = opName.replace('check.', `check.<${dialect}>.`);
- } else if (this._redirect.has(opName)) {
- opName = this._redirect.get(opName);
- }
- index = opName.indexOf('.');
- if (index === -1) {
- throw new mlir.Error(`No dialect found '${opName}' ${this.location()}`);
- }
- const dialectName = opName.substring(0, index);
- const dialect = this.context.getOrLoadDialect(dialectName);
- this.context.checkDialect(dialect, dialectName, 'operation');
- // Normalize operation name to canonical dialect name for metadata lookup
- // (e.g., spv.Load -> spirv.Load when dialect.name is spirv)
- opName = dialectName === dialect.name ? opName : opName.replace(`${dialectName}.`, `${dialect.name}.`);
- const name = _.RegisteredOperationName.lookup(opName, this.context);
- if (!name) {
- throw new mlir.Error(`Unsupported operation '${opName}'.`);
- }
- name.identifier = identifier; // Workaround
- return name;
- }
- parseGenericOperation() {
- const srcLocation = this.getToken().loc.copy();
- const name = this.getToken().getStringValue();
- if (name.length === 0) {
- throw new mlir.Error(`Empty operation name is invalid ${this.location()}`);
- }
- if (name.indexOf('\0') !== -1) {
- throw new mlir.Error(`Null character not allowed in operation name ${this.location()}`);
- }
- this.consumeToken(_.Token.string);
- let opName = _.RegisteredOperationName.lookup(name, this.context);
- if (!opName) {
- opName = new _.OperationName(null, name);
- }
- opName.identifier = name; // Workaround
- const state = new _.OperationState(srcLocation, opName);
- this.parseGenericOperationAfterOpName(state);
- return _.Operation.create(state);
- }
- parseGenericOperationAfterOpName(result) {
- this.parseToken(_.Token.l_paren, "Expected '(' to start operand list");
- const unresolvedOperands = [];
- this.parseOptionalSSAUseList(unresolvedOperands);
- this.parseToken(_.Token.r_paren, "Expected ')' to end operand list");
- if (this.getToken().is(_.Token.l_square)) {
- result.successors = [];
- this.parseSuccessors(result.successors);
- }
- if (this.consumeIf(_.Token.less)) {
- result.propertiesAttr = this.parseAttribute();
- if (!result.propertiesAttr) {
- throw new mlir.Error(`Expected attribute as properties ${this.location()}`);
- }
- this.parseToken(_.Token.greater, "Expected '>' to close properties");
- }
- if (this.consumeIf(_.Token.l_paren)) {
- do {
- const region = result.addRegion();
- this.parseRegion(region, undefined, false);
- } while (this.consumeIf(_.Token.comma));
- this.parseToken(_.Token.r_paren, "Expected ')' to end region list");
- }
- if (this.getToken().is(_.Token.l_brace)) {
- this.parseAttributeDict(result.attributes);
- }
- this.parseToken(_.Token.colon, "Expected ':' followed by operation type");
- const fnType = this.parseType();
- if (fnType instanceof _.FunctionType === false) {
- throw new mlir.Error(`Expected function type ${this.location()}`);
- }
- const operandTypes = fnType.inputs;
- if (operandTypes.length !== unresolvedOperands.length) {
- throw new mlir.Error(`Expected ${unresolvedOperands.length} operand type${unresolvedOperands.length === 1 ? '' : 's'} but had ${operandTypes.length} ${this.location()}`);
- }
- for (let i = 0; i < unresolvedOperands.length; i++) {
- const unresolvedOperand = unresolvedOperands[i];
- const type = operandTypes[i] || null;
- const value = this.resolveSSAUse(unresolvedOperand, type);
- result.operands.push(value);
- }
- result.addTypes(fnType.results);
- result.loc = this.parseTrailingLocationSpecifier();
- }
- pushSSANameScope(isIsolated) {
- if (isIsolated) {
- this.isolatedNameScopes.push(new _.IsolatedSSANameScope());
- }
- this.isolatedNameScopes[this.isolatedNameScopes.length - 1].pushSSANameScope();
- }
- popSSANameScope() {
- const currentNameScope = this.isolatedNameScopes[this.isolatedNameScopes.length - 1];
- if (currentNameScope.definitionsPerScope.length === 1) {
- this.isolatedNameScopes.pop();
- } else {
- currentNameScope.popSSANameScope();
- }
- }
- addDefinition(useInfo, value) {
- const entries = this.getSSAValueEntry(useInfo.name);
- if (entries.length <= useInfo.number) {
- entries.length = useInfo.number + 1;
- }
- entries[useInfo.number] = { value, loc: useInfo.location };
- this.recordDefinition(useInfo.name);
- }
- recordDefinition(def) {
- this.isolatedNameScopes[this.isolatedNameScopes.length - 1].recordDefinition(def);
- }
- parseOptionalSSAUseList(results) {
- if (this.getToken().isNot(_.Token.percent_identifier)) {
- return;
- }
- this.parseCommaSeparatedList('none', () => {
- const result = this.parseSSAUse();
- results.push(result);
- });
- }
- parseSSAUse(allowResultNumber = true) {
- const name = this.getTokenSpelling().str();
- this.parseToken(_.Token.percent_identifier, "Expected SSA operand");
- let number = 0;
- if (this.getToken().is(_.Token.hash_identifier)) {
- if (!allowResultNumber) {
- throw new mlir.Error(`Result number not allowed in argument list ${this.location()}`);
- }
- const value = this.getTokenSpelling().str().substring(1);
- this.consumeToken(_.Token.hash_identifier);
- number = parseInt(value, 10);
- if (isNaN(number)) {
- throw new mlir.Error(`Invalid SSA value result number '${value}' ${this.location()}`);
- }
- }
- return new _.UnresolvedOperand(null, name, number);
- }
- resolveSSAUse(unresolvedOperand, type) {
- if (type !== null && type instanceof _.Type === false) {
- throw new mlir.Error(`Type expected instead of '${type}'.`);
- }
- if (unresolvedOperand instanceof _.UnresolvedOperand) {
- const entries = this.getSSAValueEntry(unresolvedOperand.name);
- if (unresolvedOperand.number < entries.length && entries[unresolvedOperand.number]) {
- const entry = entries[unresolvedOperand.number];
- if (type && entry.value) {
- entry.value.type = type;
- }
- return entry.value;
- }
- const value = new _.Value(unresolvedOperand.toString(), type);
- if (entries.length <= unresolvedOperand.number) {
- entries.length = unresolvedOperand.number + 1;
- }
- entries[unresolvedOperand.number] = { value, loc: unresolvedOperand.location };
- return value;
- }
- // Handle literal operands (e.g., integer literals for buildable types)
- if (unresolvedOperand && unresolvedOperand.literal) {
- return new _.Value(unresolvedOperand.name, type);
- }
- throw new mlir.Error(`UnresolvedOperand expected, got '${JSON.stringify(unresolvedOperand)}' ${this.location()}`);
- }
- getSSAValueEntry(name) {
- const scope = this.isolatedNameScopes[this.isolatedNameScopes.length - 1];
- if (!scope.values.has(name)) {
- scope.values.set(name, []);
- }
- return scope.values.get(name);
- }
- parseBlock(block) {
- block.operations = Array.isArray(block.operations) ? block.operations : [];
- block.arguments = Array.isArray(block.arguments) ? block.arguments : [];
- this.parseToken(_.Token.l_brace, "expected '{' to begin a region");
- if (this.getToken().kind === _.Token.caret_identifier || (this.getToken().kind === _.Token.bare_identifier && this.getToken().getSpelling().str() && this.getToken().getSpelling().str().startsWith('^'))) {
- if (this.getToken().kind === _.Token.caret_identifier) {
- block.name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.caret_identifier);
- } else {
- block.name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.bare_identifier);
- }
- if (this.consumeIf(_.Token.l_paren)) {
- while (!this.consumeIf(_.Token.r_paren) && this.getToken().isNot(_.Token.caret_identifier)) {
- const value = this.getTokenSpelling().str();
- this.consumeToken(_.Token.percent_identifier);
- this.consumeToken(_.Token.colon);
- const type = this.parseType();
- const arg = { value, type };
- const loc = this.parseTrailingLocationSpecifier();
- if (loc) {
- arg.loc = loc;
- }
- block.arguments.push(arg);
- this.consumeIf(_.Token.comma);
- }
- }
- if (block.name && block.name.endsWith(':')) {
- block.name = block.name.slice(0, -1);
- } else {
- this.consumeToken(_.Token.colon);
- }
- }
- while (!this.consumeIf(_.Token.r_brace)) {
- if (this.getToken().kind === _.Token.caret_identifier || (this.getToken().kind === _.Token.bare_identifier && this.getToken().getSpelling().str() && this.getToken().getSpelling().str().startsWith('^'))) {
- break;
- }
- const op = this.parseOperation();
- block.operations.push(op);
- }
- block.loc = this.parseTrailingLocationSpecifier();
- return block;
- }
- parseRegion(region, entryArguments, isIsolatedNameScope) {
- // Push a new name scope for this region
- this.pushSSANameScope(isIsolatedNameScope || false);
- region.blocks = Array.isArray(region.blocks) ? region.blocks : [];
- // Register SSA entries for entry arguments BEFORE parsing the block
- // This ensures operations that reference %arg0 find the pre-registered entries
- const resolvedEntryArgs = [];
- if (entryArguments && entryArguments.length > 0) {
- for (let i = 0; i < entryArguments.length; i++) {
- const arg = entryArguments[i];
- // Use explicit name if provided, otherwise generate %arg0, %arg1, etc.
- const name = arg.name || `%arg${i}`;
- const operand = new _.Value(name, arg.type);
- // Register in SSA scope so operations can find it
- this.addDefinition({ name, number: 0, location: arg.loc || null }, operand);
- resolvedEntryArgs.push(operand);
- }
- }
- const block = {};
- this.parseBlock(block);
- if (resolvedEntryArgs.length > 0) {
- if (block.arguments.length === 0) {
- block.arguments = resolvedEntryArgs;
- } else if (block.arguments.length !== resolvedEntryArgs.length) {
- throw new mlir.Error(`Entry block has ${block.arguments.length} arguments, but function signature has ${resolvedEntryArgs.length} ${this.location()}`);
- }
- }
- region.blocks.push(block);
- let hasMultipleBlocks = false;
- while ((this.getToken().kind === _.Token.caret_identifier || (this.getToken().kind === _.Token.bare_identifier && this.getToken().getSpelling().str() && this.getToken().getSpelling().str().startsWith('^'))) && this.getToken().isNot(_.Token.r_brace)) {
- hasMultipleBlocks = true;
- const nextBlock = {};
- nextBlock.operations = [];
- nextBlock.arguments = [];
- if (this.getToken().kind === _.Token.caret_identifier) {
- nextBlock.name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.caret_identifier);
- } else {
- nextBlock.name = this.getTokenSpelling().str();
- this.consumeToken(_.Token.bare_identifier);
- }
- if (this.consumeIf(_.Token.l_paren)) {
- while (!this.consumeIf(_.Token.r_paren)) {
- const value = this.getTokenSpelling().str();
- this.consumeToken(_.Token.percent_identifier);
- this.consumeToken(_.Token.colon);
- const type = this.parseType();
- const arg = { value, type };
- const loc = this.parseTrailingLocationSpecifier();
- if (loc) {
- arg.loc = loc;
- }
- nextBlock.arguments.push(arg);
- this.consumeIf(_.Token.comma);
- }
- }
- if (nextBlock.name && nextBlock.name.endsWith(':')) {
- nextBlock.name = nextBlock.name.slice(0, -1);
- } else {
- this.consumeToken(_.Token.colon);
- }
- while (!(this.getToken().kind === _.Token.caret_identifier || (this.getToken().kind === _.Token.bare_identifier && this.getTokenSpelling().str() && this.getTokenSpelling().str().startsWith('^'))) && this.getToken().isNot(_.Token.r_brace)) {
- const op = this.parseOperation();
- nextBlock.operations.push(op);
- }
- region.blocks.push(nextBlock);
- }
- if (hasMultipleBlocks) {
- this.consumeIf(_.Token.r_brace);
- }
- this.popSSANameScope();
- return region;
- }
- parseTrailingLocationSpecifier() {
- if (!this.consumeIf(_.Token.kw_loc)) {
- return null;
- }
- this.parseToken(_.Token.l_paren, "expected '(' in location");
- const tok = this.getToken();
- let directLoc = null;
- if (tok.is(_.Token.hash_identifier) && !tok.getSpelling().str().includes('.')) {
- directLoc = this.parseLocationAlias();
- } else {
- directLoc = this.parseLocationInstance();
- }
- this.parseToken(_.Token.r_paren, "expected ')' in location");
- return directLoc;
- }
- };
- _.AsmParser = class extends _.Parser {
- getNameLoc() {
- return this.nameLoc;
- }
- parseKeyword(keyword) {
- if (!this.parser.isCurrentTokenAKeyword()) {
- throw new mlir.Error(`expected '${keyword || 'keyword'}' ${this.location()}`);
- }
- if (keyword && this.parser.getTokenSpelling().str() !== keyword) {
- throw new mlir.Error(`expected '${keyword}' ${this.location()}`);
- }
- const spelling = this.parser.getTokenSpelling().str();
- this.consumeToken();
- return spelling;
- }
- parseOptionalKeyword(arg) {
- if (typeof arg === 'string') {
- const keyword = arg;
- if (!this.parser.isCurrentTokenAKeyword() || this.parser.getTokenSpelling().str() !== keyword) {
- return false;
- }
- this.consumeToken();
- return true;
- }
- if (Array.isArray(arg)) {
- const allowedValues = arg;
- if (this.getToken().is(_.Token.bare_identifier) || this.getToken().isKeyword() || this.getToken().is(_.Token.inttype)) {
- const value = this.getTokenSpelling().str();
- if (allowedValues === undefined || allowedValues.some((v) => v === value)) {
- this.consumeToken();
- return value;
- }
- }
- return null;
- }
- if (arg === undefined) {
- return this.parser.parseOptionalKeyword();
- }
- throw new mlir.Error(`Invalid optional keyword ${this.location()}`);
- }
- parseKeywordType(keyword) {
- this.parseKeyword(keyword);
- return this.parseType();
- }
- parseTypeList() {
- return this.parseTypeListNoParens();
- }
- parseColonType() {
- this.parseToken(_.Token.colon, "expected ':'");
- return this.parseType();
- }
- parseColonTypeList() {
- this.parseToken(_.Token.colon, "expected ':'");
- return this.parseTypeList();
- }
- parseOptionalColonTypeList() {
- if (this.consumeIf(_.Token.colon)) {
- return this.parseTypeList();
- }
- return [];
- }
- parseArrowTypeList() {
- this.parseArrow();
- return this.parseFunctionResultTypes();
- }
- parseOptionalArrowTypeList() {
- if (this.consumeIf(_.Token.arrow)) {
- return this.parseFunctionResultTypes();
- }
- return [];
- }
- parseArrow() {
- this.parseToken(_.Token.arrow, "expected '->'");
- }
- parseOptionalArrow() {
- return this.consumeIf(_.Token.arrow);
- }
- parseEqual() {
- this.parseToken(_.Token.equal, "expected '='");
- }
- parseOptionalEqual() {
- return this.consumeIf(_.Token.equal);
- }
- parseLBrace() {
- this.parseToken(_.Token.l_brace, "expected '{'");
- }
- parseRBrace() {
- this.parseToken(_.Token.r_brace, "expected '}'");
- }
- parseOptionalLBrace() {
- return this.consumeIf(_.Token.l_brace);
- }
- parseOptionalRBrace() {
- return this.consumeIf(_.Token.r_brace);
- }
- parseComma() {
- this.parseToken(_.Token.comma, "expected ','");
- }
- parseOptionalComma() {
- return this.consumeIf(_.Token.comma);
- }
- parseLParen() {
- this.parseToken(_.Token.l_paren, "expected '('");
- }
- parseRParen() {
- this.parseToken(_.Token.r_paren, "expected ')'");
- }
- parseOptionalLParen() {
- return this.consumeIf(_.Token.l_paren);
- }
- parseOptionalRParen() {
- return this.consumeIf(_.Token.r_paren);
- }
- parseLSquare() {
- return this.parseToken(_.Token.l_square, "expected '['");
- }
- parseRSquare() {
- return this.parseToken(_.Token.r_square, "expected ']'");
- }
- parseOptionalLSquare() {
- return this.consumeIf(_.Token.l_square);
- }
- parseOptionalRSquare() {
- return this.consumeIf(_.Token.r_square);
- }
- parseColon() {
- this.parseToken(_.Token.colon, "expected ':'");
- }
- parseOptionalColon() {
- return this.consumeIf(_.Token.colon);
- }
- parseLess() {
- this.parseToken(_.Token.less, "expected '<'");
- }
- parseGreater() {
- this.parseToken(_.Token.greater, "expected '>'");
- }
- parseOptionalLess() {
- return this.consumeIf(_.Token.less);
- }
- parseOptionalGreater() {
- return this.consumeIf(_.Token.greater);
- }
- parseOptionalAttrDict(attributes) {
- if (this.parser.getToken().is(_.Token.l_brace)) {
- this.parser.parseAttributeDict(attributes);
- }
- }
- parseOptionalAttrDictWithKeyword(attributes) {
- if (this.parseOptionalKeyword('attributes')) {
- this.parser.parseAttributeDict(attributes);
- }
- }
- parseSymbolName(attrName, attrs) {
- const result = this.parseOptionalSymbolName();
- if (result === null) {
- throw new mlir.Error(`Expected valid '@'-identifier for symbol name ${this.location()}`);
- }
- attrs.set(attrName, result);
- }
- parseOptionalSymbolName() {
- const atToken = this.parser.getToken();
- if (atToken.isNot(_.Token.at_identifier)) {
- return null;
- }
- const result = new _.StringAttr(atToken.getSymbolReference());
- this.parser.consumeToken();
- if (this.parser.state.asmState) {
- // parser.state.asmState.addUses(SymbolRefAttr::get(result), atToken.getLocRange())
- }
- return result;
- }
- };
- _.OpAsmParser = class extends _.AsmParser {
- parseOptionalLocationSpecifier() {
- if (!this.consumeIf(_.Token.kw_loc)) {
- return null;
- }
- this.parseToken(_.Token.l_paren, "expected '(' in location");
- const tok = this.getToken();
- let directLoc = null;
- if (tok.is(_.Token.hash_identifier) && !tok.getSpelling().str().includes('.')) {
- directLoc = this.parseLocationAlias();
- } else {
- directLoc = this.parseLocationInstance();
- }
- this.parseToken(_.Token.r_paren, "expected ')' in location");
- return directLoc;
- }
- parseFunctionOp(op, allowVariadic) {
- this.parseOptionalVisibilityKeyword(op.attributes);
- this.parseSymbolName('sym_name', op.attributes);
- const sig = this.parseFunctionSignatureWithArguments(allowVariadic);
- const argTypes = [];
- for (const arg of sig.arguments) {
- if (arg.name !== '...') {
- argTypes.push(arg.type);
- }
- }
- const type = new _.FunctionType(argTypes, sig.resultTypes);
- op.addAttribute('function_type', new _.TypeAttrOf(type));
- if (sig.resultAttrs.some((a) => a !== null)) {
- op.addAttribute('res_attrs', sig.resultAttrs);
- }
- const argAttrs = sig.arguments.filter((a) => a.name !== '...').map((a) => a.attrs || null);
- if (argAttrs.some((a) => a !== null)) {
- op.addAttribute('arg_attrs', argAttrs);
- }
- this.parseOptionalAttrDictWithKeyword(op.attributes);
- if (this.getToken().is(_.Token.l_brace)) {
- const region = op.addRegion();
- // Functions are IsolatedFromAbove, so pass true for isIsolatedNameScope
- this.parseRegion(region, sig.arguments, /* isIsolatedNameScope */ true);
- }
- }
- parseFunctionSignature(argOperands) {
- const argTypes = [];
- const argAttrs = [];
- const resultTypes = [];
- const resultAttrs = [];
- this.parseToken(_.Token.l_paren, "expected '(' in function signature");
- if (this.getToken().isNot(_.Token.r_paren)) {
- this.parseTypeAndAttrList(argTypes, argAttrs, argOperands);
- }
- this.parseToken(_.Token.r_paren, "expected ')' in function signature");
- if (this.consumeIf(_.Token.arrow)) {
- this.parseFunctionResultList(resultTypes, resultAttrs);
- }
- return { argTypes, argAttrs, resultTypes, resultAttrs };
- }
- parseFunctionSignatureWithArguments(allowVariadic) {
- const argResult = this.parseFunctionArgumentList(allowVariadic);
- const resultTypes = [];
- const resultAttrs = [];
- if (this.consumeIf(_.Token.arrow)) {
- this.parseFunctionResultList(resultTypes, resultAttrs);
- }
- return { arguments: argResult.arguments, isVariadic: argResult.isVariadic, resultTypes, resultAttrs };
- }
- parseFunctionResultList(types, attrs) {
- if (this.consumeIf(_.Token.l_paren)) {
- if (this.consumeIf(_.Token.r_paren)) {
- return;
- }
- this.parseTypeAndAttrList(types, attrs);
- this.parseToken(_.Token.r_paren, "expected ')' in function result list");
- } else {
- const type = this.parseType();
- types.push(type);
- attrs.push(null);
- }
- }
- // Returns { arguments: Array<OpAsmParser.Argument>, isVariadic: boolean }
- parseFunctionArgumentList(allowVariadic) {
- const inputs = [];
- let isVariadic = false;
- if (this.consumeIf(_.Token.l_paren)) {
- while (!this.consumeIf(_.Token.r_paren)) {
- if (this.getToken().is(_.Token.r_paren)) {
- break;
- }
- if (allowVariadic && this.consumeIf(_.Token.ellipsis)) {
- isVariadic = true;
- this.parseToken(_.Token.r_paren, "expected ')' after '...'");
- break;
- }
- if (this.getToken().is(_.Token.percent_identifier)) {
- const ssaName = this.parseOperand();
- this.consumeToken(_.Token.colon);
- const type = this.parseType();
- let attrs = null;
- if (this.getToken().is(_.Token.l_brace)) {
- attrs = new Map();
- this.parseAttributeDict(attrs);
- }
- const loc = this.parseOptionalLocationSpecifier();
- inputs.push(new _.OpAsmParser.Argument(ssaName, type, attrs, loc));
- } else {
- // Type-only argument (no explicit name like %arg0)
- // Don't generate a name - let the region/SSA system handle it
- const type = this.parseType();
- let attrs = null;
- if (this.getToken().is(_.Token.l_brace)) {
- attrs = new Map();
- this.parseAttributeDict(attrs);
- }
- inputs.push(new _.OpAsmParser.Argument(null, type, attrs, null));
- }
- if (this.getToken().isNot(_.Token.r_paren)) {
- if (!this.consumeIf(_.Token.comma)) {
- break;
- }
- if (this.getToken().is(_.Token.r_paren)) {
- break;
- }
- }
- }
- }
- return { arguments: inputs, isVariadic };
- }
- parseTypeAndAttrList(types, attrs, operands) {
- let index = 0;
- this.parseCommaSeparatedList('none', () => {
- const type = this.parseType();
- types.push(type);
- // Parse optional attribute dict after each type
- if (this.getToken().is(_.Token.l_brace)) {
- const attrList = new Map();
- this.parseAttributeDict(attrList);
- attrs.push(attrList);
- // Associate attrs with operand if available
- if (operands && index < operands.length) {
- operands[index].attributes = attrList;
- }
- } else {
- attrs.push(null);
- }
- index++;
- return true;
- });
- }
- parseDenseI64ArrayAttr(attributeName, attributes) {
- this.parseKeyword(attributeName);
- this.parseEqual();
- const value = this.skip('[');
- attributes.set(attributeName, value);
- }
- parseOptionalVisibilityKeyword(attributes) {
- const visibility = this.parseOptionalKeyword(['public', 'private', 'nested']);
- if (visibility) {
- attributes.set('sym_visibility', visibility);
- }
- }
- parseOperands(operands) {
- this.parseCommaSeparatedList('none', () => {
- operands.push(this.parseOperand());
- });
- }
- parseShapedResultList(operands, operandTypes, resultTypes, tiedOperands) {
- const tiedOperandIndices = [];
- do {
- let type = null;
- let tiedOperandIndex = -1;
- const tiedResult = this.parseOptionalOperand();
- if (tiedResult) {
- tiedOperandIndex = this.findTiedOperand(tiedResult, operands);
- if (this.parseOptionalKeyword('as')) {
- type = this.parseType();
- } else if (tiedOperandIndex >= 0 && tiedOperandIndex < operandTypes.length) {
- type = operandTypes[tiedOperandIndex];
- }
- } else {
- type = this.parseType();
- }
- if (this.getToken().is(_.Token.l_brace)) {
- this.skip(_.Token.l_brace);
- }
- if (type) {
- resultTypes.push(type);
- }
- tiedOperandIndices.push(tiedOperandIndex);
- } while (this.parseOptionalComma());
- if (tiedOperands && tiedOperandIndices.length > 0) {
- tiedOperands.push(...tiedOperandIndices);
- }
- }
- findTiedOperand(tiedResult, operands) {
- for (let i = 0; i < operands.length; i++) {
- if (operands[i].name === tiedResult.name && operands[i].number === tiedResult.number) {
- return i;
- }
- }
- return -1;
- }
- };
- _.OpAsmParser.Argument = class {
- constructor(ssaName, type, attrs, loc) {
- this.ssaName = ssaName;
- this.type = type;
- this.attrs = attrs;
- this.loc = loc;
- }
- get name() {
- return this.ssaName ? this.ssaName.name : null;
- }
- get value() {
- return this.ssaName ? this.ssaName.toString() : null;
- }
- };
- _.CustomOpAsmParser = class extends _.OpAsmParser {
- constructor(resultIDs, parser) {
- super(parser.state);
- this.resultIDs = resultIDs || [];
- this.parser = parser;
- this.nameLoc = parser.getToken().loc.copy();
- }
- parseOperation() {
- return this.parser.parseOperation();
- }
- getNumResults() {
- let count = 0;
- for (const entry of this.resultIDs) {
- count += entry[1];
- }
- return count;
- }
- getResultName(index) {
- if (index < this.resultIDs.length) {
- return this.resultIDs[index][0];
- }
- return null;
- }
- parseArgument(allowType, allowAttrs) {
- const ssaName = this.parseOperand();
- let type = null;
- let attrs = null;
- let loc = null;
- if (allowType) {
- type = this.parseColonType();
- }
- if (allowAttrs) {
- attrs = {};
- this.parseOptionalAttrDict(attrs);
- if (Object.keys(attrs).length === 0) {
- attrs = null;
- }
- }
- loc = this.parseOptionalLocationSpecifier();
- return new _.OpAsmParser.Argument(ssaName, type, attrs, loc);
- }
- parseArgumentList(delimiter, allowType = false, allowAttrs = false) {
- delimiter = delimiter || 'none';
- if (delimiter === 'none') {
- if (this.getToken().isNot(_.Token.percent_identifier)) {
- return [];
- }
- }
- const parseOneArgument = () => {
- if (this.getToken().is(_.Token.percent_identifier)) {
- return this.parseArgument(allowType, allowAttrs);
- }
- return null;
- };
- return this.parseCommaSeparatedList(delimiter, parseOneArgument);
- }
- parseRegion(region, entryArguments, enableNameShadowing) {
- return this.parser.parseRegion(region, entryArguments, enableNameShadowing);
- }
- parseOperand(allowResultNumber = true) {
- return this.parser.parseSSAUse(allowResultNumber);
- }
- parseOptionalOperand(allowResultNumber = true) {
- if (this.getToken().is(_.Token.percent_identifier)) {
- return this.parseOperand(allowResultNumber);
- }
- return null;
- }
- parseOperandList(delimiter) {
- delimiter = delimiter || 'none';
- if (delimiter === 'none') {
- if (this.getToken().isNot(_.Token.percent_identifier)) {
- return [];
- }
- }
- const parseOneOperand = () => {
- if (this.getToken().is(_.Token.percent_identifier)) {
- return this.parseOperand();
- }
- return null;
- };
- return this.parseCommaSeparatedList(delimiter, parseOneOperand);
- }
- resolveOperand(operand, type, result) {
- const resolved = this.parser.resolveSSAUse(operand, type);
- if (result) {
- result.push(resolved);
- }
- return resolved;
- }
- resolveOperands(operands, types, result) {
- if (!Array.isArray(operands)) {
- throw new mlir.Error(`Unexpected operands type '${typeof operands}'.`);
- }
- if (!Array.isArray(types)) {
- return;
- }
- const count = Math.min(operands.length, types.length);
- if (result) {
- for (let i = 0; i < count; i++) {
- const operand = operands[i];
- const type = types[i];
- const resolved = this.parser.resolveSSAUse(operand, type);
- result.push(resolved);
- }
- } else {
- for (let i = 0; i < count; i++) {
- const operand = operands[i];
- const type = types[i];
- if (operand && type) {
- if (operand instanceof _.Value) {
- operand.type = type;
- } else if (typeof operand === 'object') {
- operand.type = type;
- }
- }
- }
- }
- }
- parseGenericOperation() {
- return this.parser.parseGenericOperation();
- }
- parseCustomOperationName() {
- return this.parser.parseCustomOperationName();
- }
- parseOptionalLocationSpecifier() {
- // Ref impl: CustomOpAsmParser::parseOptionalLocationSpecifier (Parser.cpp:2002)
- // Separate implementation from parseTrailingLocationSpecifier
- // If there is a 'loc' we parse a trailing location.
- if (!this.parser.consumeIf(_.Token.kw_loc)) {
- return null;
- }
- this.parser.parseToken(_.Token.l_paren, "expected '(' in location");
- const tok = this.parser.getToken();
- let directLoc = null;
- // Check to see if we are parsing a location alias. We are parsing a
- // location alias if the token is a hash identifier *without* a dot in it -
- // the dot signifies a dialect attribute. Otherwise, we parse the location
- // directly.
- if (tok.is(_.Token.hash_identifier) && !tok.getSpelling().str().includes('.')) {
- directLoc = this.parser.parseLocationAlias();
- } else {
- directLoc = this.parser.parseLocationInstance();
- }
- this.parser.parseToken(_.Token.r_paren, "expected ')' in location");
- return directLoc;
- }
- };
- _.DialectAsmParser = class extends _.AsmParser {
- };
- _.CustomDialectAsmParser = class extends _.DialectAsmParser {
- constructor(fullSpec, parser) {
- super(parser.state);
- this.fullSpec = fullSpec;
- this.parser = parser;
- this.nameLoc = parser.getToken().loc.copy();
- }
- };
- _.TensorLiteralParser = class {
- constructor(parser) {
- this._parser = parser;
- // Reference stores pairs of (isNegative, token) for proper type validation
- // We store {isNegative, value, kind} objects
- this._storage = [];
- this._shape = [];
- this._hexStorage = null;
- }
- parse(allowHex) {
- if (allowHex && this._parser.getToken().is(_.Token.string)) {
- const hexStr = this._parser.getToken().getStringValue();
- this._parser.consumeToken(_.Token.string);
- if (hexStr.startsWith('0x')) {
- const cleanHex = hexStr.substring(2);
- const data = new Uint8Array(cleanHex.length >> 1);
- for (let i = 0; i < data.length; i++) {
- const index = i << 1;
- data[i] = parseInt(cleanHex.substring(index, index + 2), 16);
- }
- this._hexStorage = data;
- return { storage: data, shape: null };
- }
- // Non-hex string element
- this._storage.push({ isNegative: false, value: hexStr, kind: 'string' });
- return { storage: this._storage, shape: this._shape };
- }
- if (this._parser.getToken().is(_.Token.l_square)) {
- this.parseList(this._shape);
- } else {
- this.parseElement();
- // Single element parsed without list - shape stays empty (splat)
- }
- return { storage: this._storage, shape: this._shape };
- }
- parseList(dims) {
- this._parser.consumeToken(_.Token.l_square);
- let first = true;
- let newDims = [];
- let size = 0;
- while (!this._parser.consumeIf(_.Token.r_square)) {
- const thisDims = [];
- if (this._parser.getToken().is(_.Token.l_square)) {
- this.parseList(thisDims);
- } else {
- this.parseElement();
- }
- size++;
- if (!first) {
- const compareDims = (a, b) =>{
- if (a.length !== b.length) {
- return false;
- }
- for (let i = 0; i < a.length; i++) {
- if (a[i] !== b[i]) {
- return false;
- }
- }
- return true;
- };
- // Verify consistent dimensions (reference checks prevDims == newDims)
- const dimsMatch = compareDims(thisDims, newDims);
- if (!dimsMatch) {
- throw new mlir.Error(`Tensor literal is invalid; ranks are not consistent between elements ${this._parser.location()}`);
- }
- }
- newDims = thisDims;
- first = false;
- this._parser.consumeIf(_.Token.comma);
- }
- dims.length = 0;
- dims.push(size);
- dims.push(...newDims);
- }
- parseElement() {
- switch (this._parser.getToken().kind) {
- // Parse a boolean element.
- case _.Token.kw_true:
- case _.Token.kw_false: {
- const value = this._parser.getTokenSpelling().str();
- this._parser.consumeToken(this._parser.getToken().kind);
- this._storage.push({ isNegative: false, value, kind: 'boolean' });
- break;
- }
- case _.Token.floatliteral:
- case _.Token.integer: {
- const value = this._parser.getTokenSpelling().str();
- const kind = this._parser.getToken().kind === _.Token.floatliteral ? 'float' : 'int';
- this._parser.consumeToken(this._parser.getToken().kind);
- this._storage.push({ isNegative: false, value, kind });
- break;
- }
- case _.Token.minus: {
- this._parser.consumeToken(_.Token.minus);
- if (!this._parser.getToken().isAny(_.Token.floatliteral, _.Token.integer)) {
- throw new mlir.Error(`Expected integer or floating point literal ${this._parser.location()}`);
- }
- const value = this._parser.getTokenSpelling().str();
- const kind = this._parser.getToken().kind === _.Token.floatliteral ? 'float' : 'int';
- this._parser.consumeToken(this._parser.getToken().kind);
- this._storage.push({ isNegative: true, value, kind });
- break;
- }
- case _.Token.string: {
- const value = this._parser.getToken().getStringValue();
- this._parser.consumeToken(_.Token.string);
- this._storage.push({ isNegative: false, value, kind: 'string' });
- break;
- }
- case _.Token.l_paren:
- this._parser.consumeToken(_.Token.l_paren);
- this.parseElement();
- this._parser.parseToken(_.Token.comma, "expected ',' between complex elements");
- this.parseElement();
- this._parser.parseToken(_.Token.r_paren, "expected ')' after complex elements");
- break;
- default:
- throw new mlir.Error(`Expected element literal of primitive type ${this._parser.location()}`);
- }
- }
- getShape() {
- return this._shape;
- }
- getAttr(type) {
- // Handle hex storage directly
- if (this._hexStorage instanceof Uint8Array) {
- return this._hexStorage;
- }
- const elementType = type && type.elementType ? type.elementType : null;
- const numElements = type && type.getNumElements ? type.getNumElements() : 0;
- const isComplex = elementType instanceof _.ComplexType;
- const baseElemType = isComplex && elementType.elementType ? elementType.elementType : elementType;
- // Determine element type properties for validation
- let isUnsigned = false;
- let isInteger = false;
- let isI1 = false;
- let bitWidth = 0;
- if (baseElemType) {
- const typeStr = baseElemType.toString();
- isUnsigned = typeStr.startsWith('ui');
- isI1 = typeStr === 'i1';
- const intMatch = typeStr.match(/^[su]?i(\d+)$/);
- if (intMatch) {
- isInteger = true;
- bitWidth = parseInt(intMatch[1], 10);
- } else if (typeStr === 'index') {
- isInteger = true;
- bitWidth = 64;
- }
- }
- // Validate and convert storage elements
- const convertElement = (elem) => {
- const { isNegative, value, kind } = elem;
- if (isNegative && isUnsigned) {
- throw new mlir.Error(`Expected unsigned integer elements, but parsed negative value`);
- }
- if (kind === 'boolean') {
- if (isInteger && !isI1) {
- throw new mlir.Error(`Expected i1 type for 'true' or 'false' values`);
- }
- return value === true || value === 'true' ? 1 : 0;
- }
- if (kind === 'float' && isInteger) {
- throw new mlir.Error(`Expected integer elements, but parsed floating-point`);
- }
- let result = null;
- if (kind === 'int') {
- if (bitWidth >= 64) {
- result = BigInt(value);
- if (isNegative) {
- result = -result;
- }
- } else {
- result = parseInt(value, 10);
- if (isNegative) {
- result = -result;
- }
- }
- } else if (kind === 'float') {
- result = parseFloat(value);
- if (isNegative) {
- result = -result;
- }
- } else {
- result = value;
- }
- return result;
- };
- // Handle zero-element tensors (e.g., tensor<2x0x3xi4>)
- if (numElements === 0) {
- return [];
- }
- // Limit splat expansion to avoid memory issues with huge tensors
- const maxSplatExpansion = 10000000;
- // Handle splats - Reference: if shape.empty() and storage has elements, it's a splat
- const isSplat = this._shape.length === 0 && this._storage.length > 0;
- if (isSplat && numElements > 1) {
- if (numElements > maxSplatExpansion) {
- // Too large to expand - return null to indicate we can't provide the data
- return null;
- }
- if (isComplex && this._storage.length === 2) {
- // Complex splat: storage has 2 elements (real, imag)
- const result = [];
- const real = convertElement(this._storage[0]);
- const imag = convertElement(this._storage[1]);
- for (let i = 0; i < numElements; i++) {
- result.push(new base.Complex(real, imag));
- }
- return result;
- }
- // Regular splat: replicate single value
- const converted = convertElement(this._storage[0]);
- return new Array(numElements).fill(converted);
- }
- // Non-splat complex: convert pairs to base.Complex objects
- if (isComplex && Array.isArray(this._storage)) {
- const result = [];
- for (let i = 0; i < this._storage.length; i += 2) {
- result.push(new base.Complex(convertElement(this._storage[i]), convertElement(this._storage[i + 1])));
- }
- return result;
- }
- return this._storage.map(convertElement);
- }
- };
- _.EncodingReader = class {
- constructor(data) {
- if (data instanceof Uint8Array) {
- this._data = data;
- this._reader = base.BinaryReader.open(data);
- } else {
- this._data = null;
- this._reader = data;
- }
- }
- get length() {
- return this._reader.length;
- }
- get position() {
- return this._reader.position;
- }
- empty() {
- return this.position >= this.length;
- }
- size() {
- return this._reader.length - this._reader.position;
- }
- seek(offset) {
- this._reader.seek(offset);
- }
- skipBytes(length) {
- this._reader.skip(length);
- }
- parseBytes(length) {
- return this._reader.read(length);
- }
- parseByte() {
- return this._reader.byte();
- }
- peek() {
- const position = this._reader.position;
- const value = this._reader.byte();
- this._reader.seek(position);
- return value;
- }
- parseVarInt() {
- let result = this._reader.byte();
- if (result & 1) {
- return BigInt(result >> 1);
- }
- if (result === 0) {
- return this._reader.uint64();
- }
- result = BigInt(result);
- let mask = 1n;
- let numBytes = 0n;
- let shift = 8n;
- while (result > 0n && (result & mask) === 0n) {
- if (numBytes >= 7n) {
- throw new mlir.Error('Invalid varint.');
- }
- result |= (BigInt(this._reader.byte()) << shift);
- mask <<= 1n;
- shift += 8n;
- numBytes++;
- }
- if (numBytes === 0n || numBytes > 7n) {
- throw new mlir.Error(`Invalid varint.`);
- }
- result >>= numBytes + 1n;
- return result;
- }
- parseSignedVarInt() {
- const n = this.parseVarInt();
- return (n >> 1n) ^ -(n & 1n);
- }
- parseVarIntWithFlag() {
- const result = this.parseVarInt();
- return [result >> 1n, (result & 1n) === 1n];
- }
- parseNullTerminatedString() {
- const reader = this._reader;
- let result = '';
- let value = -1;
- const maxLength = reader.length - reader.position;
- let bytesRead = 0;
- for (; ;) {
- if (bytesRead >= maxLength) {
- throw new mlir.Error('Malformed null-terminated string, no null character found.');
- }
- value = reader.byte();
- bytesRead++;
- if (value === 0x00) {
- break;
- }
- result += String.fromCharCode(value);
- }
- return result;
- }
- parseEntry(entries, entryStr) {
- const entryIdx = this.parseVarInt().toNumber();
- return this.resolveEntry(entries, entryIdx, entryStr);
- }
- resolveEntry(entries, entryIdx, entryStr) {
- if (entryIdx >= entries.length) {
- throw new mlir.Error(`Invalid '${entryStr}' index.`);
- }
- return entries[entryIdx];
- }
- parseSection(alignmentValidator) {
- const sectionIDAndHasAlignment = this.parseByte();
- const length = this.parseVarInt().toNumber();
- const sectionID = sectionIDAndHasAlignment & 0x7F;
- const hasAlignment = sectionIDAndHasAlignment & 0x80;
- if (sectionID >= 9) { // kNumSections
- throw new mlir.Error(`Invalid section ID: ${sectionID}.`);
- }
- if (hasAlignment) {
- const alignment = this.parseVarInt().toNumber();
- alignmentValidator(alignment);
- this.alignTo(alignment);
- }
- const sectionData = this.parseBytes(length);
- return [sectionID, sectionData];
- }
- parseBlobAndAlignment() {
- const alignment = this.parseVarInt().toNumber();
- const dataSize = this.parseVarInt().toNumber();
- this.alignTo(alignment);
- const data = this.parseBytes(dataSize);
- return [data, alignment];
- }
- alignTo(alignment) {
- if ((alignment & (alignment - 1)) !== 0) {
- throw new mlir.Error('Expected alignment to be a power-of-two.');
- }
- while ((this.position & (alignment - 1)) !== 0) {
- const padding = this.parseByte();
- if (padding !== 0xCB) {
- throw new mlir.Error(`Expected alignment byte (0xCB), but got: 0x${padding.toString(16)}.`);
- }
- }
- }
- };
- _.BytecodeDialect = class {
- load(reader, ctx) {
- this.dialect = ctx.getOrLoadDialect(this.name);
- if (!this.dialect) {
- throw new mlir.Error(`Dialect '${this.name}' is unknown.`);
- }
- this.interface = this.dialect;
- if (this.versionBuffer) {
- const encReader = new _.EncodingReader(this.versionBuffer, reader.loc);
- const versionReader = reader.withEncodingReader(encReader);
- const loadedVersion = this.interface.readVersion(versionReader);
- if (!loadedVersion) {
- return false;
- }
- }
- return true;
- }
- getLoadedDialect() {
- return this.dialect;
- }
- };
- _.DialectBytecodeReader = class {
- };
- _.DialectReader = class extends _.DialectBytecodeReader {
- constructor(attrTypeReader, stringReader, resourceReader, dialectsMap, reader, bytecodeVersion, depth = 0) {
- super();
- this.attrTypeReader = attrTypeReader;
- this.stringReader = stringReader;
- this.resourceReader = resourceReader;
- this.dialectsMap = dialectsMap;
- this.reader = reader;
- this.bytecodeVersion = bytecodeVersion;
- this.depth = depth;
- this._floatBuffer = new ArrayBuffer(8);
- this._floatView = new DataView(this._floatBuffer);
- this._floatBitWidths = { f16: 16, bf16: 16, f32: 32, f64: 64, f80: 80, f128: 128, tf32: 19, f4E2M1FN: 4, f6E2M3FN: 6, f6E3M2FN: 6, f8E5M2: 8, f8E4M3: 8, f8E4M3FN: 8, f8E5M2FNUZ: 8, f8E4M3FNUZ: 8, f8E4M3B11FNUZ: 8, f8E3M4: 8, f8E8M0FNU: 8 };
- }
- readType() {
- const index = this.reader.parseVarInt().toNumber();
- if (this.attrTypeReader.isResolving() && this.depth > this.maxAttrTypeDepth) {
- const existing = this.attrTypeReader.getTypeOrSentinel(index);
- if (!existing) {
- throw new mlir.Error(`Exceeded maximum type depth at '${index}'.`);
- }
- return existing;
- }
- return this.attrTypeReader.readType(index, this.depth + 1);
- }
- readAttribute() {
- const index = this.reader.parseVarInt().toNumber();
- if (this.attrTypeReader.isResolving() && this.depth > this.maxAttrTypeDepth) {
- const existing = this.attrTypeReader.getAttributeOrSentinel(index);
- if (!existing) {
- throw new mlir.Error(`Exceeded maximum attribute depth at '${index}'.`);
- }
- return existing;
- }
- return this.attrTypeReader.readAttribute(index, this.depth + 1);
- }
- readString() {
- return this.stringReader.parseString(this.reader);
- }
- readVarInt() {
- return this.reader.parseVarInt().toNumber();
- }
- readSignedVarInt() {
- return this.reader.parseSignedVarInt();
- }
- readByte() {
- return this.reader.parseByte();
- }
- readBlob() {
- const size = this.reader.parseVarInt().toNumber();
- return this.reader.parseBytes(size);
- }
- readResourceHandle() {
- return this.resourceReader.parseResourceHandle(this.reader);
- }
- readSignedVarInts() {
- const count = this.reader.parseVarInt().toNumber();
- const result = new Array(count);
- for (let i = 0; i < count; i++) {
- result[i] = this.reader.parseSignedVarInt().toNumber();
- }
- return result;
- }
- readAPIntWithKnownWidth(bitWidth) {
- if (bitWidth <= 8) {
- return BigInt(this.reader.parseByte());
- }
- if (bitWidth <= 64) {
- return this.reader.parseSignedVarInt();
- }
- const numWords = this.reader.parseVarInt().toNumber();
- let value = 0n;
- for (let i = 0; i < numWords; i++) {
- const word = this.reader.parseSignedVarInt();
- value |= (word << BigInt(i * 64));
- }
- return value;
- }
- readAPFloatWithKnownSemantics(type) {
- const bitWidth = this._floatBitWidths[type.name];
- if (!bitWidth) {
- throw new mlir.Error(`Unsupported float type '${type.name}'.`);
- }
- const bits = this.readAPIntWithKnownWidth(bitWidth);
- if (bitWidth <= 32) {
- this._floatView.setUint32(0, Number(bits), true);
- } else {
- this._floatView.setBigUint64(0, bits, true);
- }
- switch (type.name) {
- case 'f16': return this._floatView.getFloat16(0, true);
- case 'f32': return this._floatView.getFloat32(0, true);
- case 'f64': return this._floatView.getFloat64(0, true);
- case 'bf16': return this._floatView.getBfloat16(0, true);
- default: throw new mlir.Error(`Unsupported float type '${type.name}'.`);
- }
- }
- withEncodingReader(encodingReader) {
- return new _.DialectReader(this.attrTypeReader, this.stringReader, this.resourceReader, this.dialectsMap, encodingReader, this.bytecodeVersion, this.depth);
- }
- };
- _.AttrTypeReader = class {
- constructor(stringReader, resourceReader, dialectsMap, bytecodeVersion, fileLoc, config) {
- this.stringReader = stringReader;
- this.resourceReader = resourceReader;
- this.dialectsMap = dialectsMap;
- this.fileLoc = fileLoc;
- this.bytecodeVersion = bytecodeVersion;
- this.parserConfig = config;
- this.attributes = [];
- this.types = [];
- this.maxAttrTypeDepth = 5;
- this._resolving = false;
- }
- initialize(dialects, sectionData, offsetSectionData) {
- this._sectionData = sectionData;
- const offsetReader = new _.EncodingReader(offsetSectionData);
- const numAttributes = offsetReader.parseVarInt().toNumber();
- const numTypes = offsetReader.parseVarInt().toNumber();
- this.attributes = new Array(numAttributes);
- this.types = new Array(numTypes);
- let currentOffset = 0;
- const parseEntries = (entries) => {
- let currentIndex = 0;
- const endIndex = entries.length;
- while (currentIndex < endIndex) {
- const dialectIndex = offsetReader.parseVarInt().toNumber();
- const dialect = dialects[dialectIndex];
- const numEntries = offsetReader.parseVarInt().toNumber();
- for (let i = 0; i < numEntries; i++) {
- const entry = {};
- const entrySizeWithFlag = offsetReader.parseVarInt();
- entry.hasCustomEncoding = (entrySizeWithFlag & 1n) === 1n;
- const entrySize = (entrySizeWithFlag >> 1n).toNumber();
- // Store offset like old code (don't bound reading to entry size)
- entry.offset = currentOffset;
- entry.size = entrySize;
- entry.dialect = dialect;
- entry.resolved = null;
- currentOffset += entrySize;
- entries[currentIndex++] = entry;
- }
- }
- };
- // Process attributes, then types
- parseEntries(this.attributes);
- parseEntries(this.types);
- }
- isResolving() {
- return this._resolving;
- }
- getAttributeOrSentinel(index) {
- if (index >= this.attributes.length) {
- return null;
- }
- return this.attributes[index].resolved;
- }
- getTypeOrSentinel(index) {
- if (index >= this.types.length) {
- return null;
- }
- return this.types[index].resolved;
- }
- readAttribute(index, depth = 0) {
- return this.readEntry(this.attributes, index, 'attr', depth);
- }
- readType(index, depth = 0) {
- return this.readEntry(this.types, index, 'type', depth);
- }
- resolveAttribute(index, depth = 0) {
- // Resolve an attribute at the given index (equivalent to resolveEntry in C++)
- return this.resolveEntry(this.attributes, index, 'attr', depth);
- }
- resolveType(index, depth = 0) {
- // Resolve a type at the given index (equivalent to resolveEntry in C++)
- return this.resolveEntry(this.types, index, 'type', depth);
- }
- parseAttribute(reader) {
- // Parse an attribute reference from the reader (varint index)
- const index = reader.parseVarInt().toNumber();
- return this.resolveAttribute(index);
- }
- parseType(reader) {
- // Parse a type reference from the reader (varint index)
- const index = reader.parseVarInt().toNumber();
- return this.resolveType(index);
- }
- resolveEntry(entries, index, entryType, depth = 0) {
- // Simplified version of C++ resolveEntry - doesn't handle deferred worklists
- // but does track recursion depth
- const oldResolving = this._resolving;
- this._resolving = true;
- try {
- const result = this.readEntry(entries, index, entryType, depth);
- return result;
- } finally {
- this._resolving = oldResolving;
- }
- }
- readEntry(entries, index, entryType, depth = 0) {
- if (index >= entries.length) {
- throw new mlir.Error(`Invalid '${entryType}' index '${index}'.`);
- }
- const entry = entries[index];
- if (entry.resolved !== null) {
- return entry.resolved;
- }
- if (depth > this.maxAttrTypeDepth) {
- throw new mlir.Error(`Exceeded maximum '${entryType}' depth.`);
- }
- entry.resolved = { name: 'pending', value: `<${entryType} ${index}>` };
- const entryData = this._sectionData.subarray(entry.offset, entry.offset + entry.size);
- const reader = new _.EncodingReader(entryData);
- const startPosition = reader.position;
- if (entry.hasCustomEncoding) {
- entry.resolved = this.parseCustomEntry(entry, reader, entryType, index, depth);
- } else {
- entry.resolved = this.parseAsmEntry(reader, entryType);
- }
- const bytesRead = reader.position - startPosition;
- if (bytesRead > entry.size) {
- throw new mlir.Error(`Read ${bytesRead} bytes but entry size is ${entry.size} bytes.`);
- }
- return entry.resolved;
- }
- parseCustomEntry(entry, reader, entryType, index, depth) {
- // Lazy load the dialect interface (like BytecodeDialect::load)
- const context = this.fileLoc.context;
- if (entry.dialect.interface === undefined) {
- entry.dialect.interface = context.getOrLoadDialect(entry.dialect.name);
- }
- context.checkDialect(entry.dialect.interface, entry.dialect.name, 'custom entry');
- const dialect = entry.dialect.interface;
- const dialectReader = new _.DialectReader(this, this.stringReader, this.resourceReader, this.dialectsMap, reader, this.version, depth);
- switch (entryType) {
- case 'type':
- return dialect.readType(dialectReader);
- case 'attr':
- return dialect.readAttribute(dialectReader);
- default:
- throw new mlir.Error(`Unknown entry type '${entryType}'.`);
- }
- }
- parseAsmEntry(reader, entryType) {
- const asmStr = reader.parseNullTerminatedString();
- const context = this.fileLoc.context;
- if (entryType === 'type') {
- const { type, numRead } = _.Parser.parseType(asmStr, context);
- if (!type || numRead !== asmStr.length) {
- throw new mlir.Error(`Failed to parse type '${asmStr}'.`);
- }
- return type;
- }
- const { attribute, numRead } = _.Parser.parseAttribute(asmStr, context, null);
- if (!attribute || numRead !== asmStr.length) {
- throw new mlir.Error(`Failed to parse attribute '${asmStr}'.`);
- }
- return attribute;
- }
- };
- _.StringSectionReader = class {
- initialize(sectionData) {
- const decoder = new TextDecoder('utf-8');
- const stringReader = new _.EncodingReader(sectionData);
- const numStrings = stringReader.parseVarInt().toNumber();
- this.strings = new Array(numStrings);
- let stringDataEndOffset = sectionData.length;
- for (let i = numStrings - 1; i >= 0; i--) {
- const stringSize = stringReader.parseVarInt().toNumber();
- if (stringDataEndOffset < stringSize) {
- throw new mlir.Error('String size exceeds the available data size.');
- }
- const stringOffset = stringDataEndOffset - stringSize;
- const buffer = sectionData.subarray(stringOffset, stringOffset + stringSize - 1);
- this.strings[i] = decoder.decode(buffer);
- stringDataEndOffset = stringOffset;
- }
- if ((sectionData.length - stringReader.size()) !== stringDataEndOffset) {
- throw new mlir.Error('Unexpected trailing data between the offsets for strings and their data.');
- }
- }
- parseString(reader) {
- return reader.parseEntry(this.strings, 'string');
- }
- parseStringWithFlag(reader) {
- const [entryIdx, flag] = reader.parseVarIntWithFlag();
- const str = this.parseStringAtIndex(reader, entryIdx.toNumber());
- return [str, flag];
- }
- parseStringAtIndex(reader, index) {
- if (index >= this.strings.length) {
- throw new mlir.Error('Invalid string index.');
- }
- return this.strings[index];
- }
- };
- _.PropertiesSectionReader = class {
- constructor() {
- this._properties = new Map();
- }
- initialize(sectionData) {
- const reader = new _.EncodingReader(sectionData);
- const count = reader.parseVarInt().toNumber();
- this._properties = new Array(count);
- for (let i = 0; i < this._properties.length; i++) {
- const size = reader.parseVarInt().toNumber();
- const data = reader.parseBytes(size);
- this._properties[i] = data;
- }
- }
- read(fileLoc, dialectReader, opName, opState) {
- const propIdx = dialectReader.readVarInt();
- if (propIdx < this._properties.length) {
- const propData = this._properties[propIdx];
- if (propData.length > 0) {
- const propReader = dialectReader.withEncodingReader(new _.EncodingReader(propData));
- // Reference: https://github.com/llvm/llvm-project/blob/main/mlir/lib/Bytecode/Reader/BytecodeReader.cpp
- // Function operations: sym_name (required), function_type (required), arg_attrs (optional), res_attrs (optional)
- if (opName.getStringRef().endsWith('.func') || /\.func_v\d+$/.test(opName.getStringRef())) {
- if (propReader.reader.position < propData.length) {
- const symNameAttr = propReader.readAttribute();
- if (symNameAttr && symNameAttr.value !== undefined) {
- const name = typeof symNameAttr.value === 'string' ? symNameAttr.value : String(symNameAttr.value);
- opState.addAttribute('sym_name', new _.StringAttr(name));
- }
- }
- if (propReader.reader.position < propData.length) {
- const funcTypeAttr = propReader.readAttribute();
- if (funcTypeAttr instanceof _.TypeAttrOf && funcTypeAttr.type instanceof _.FunctionType) {
- opState.addAttribute('function_type', funcTypeAttr);
- }
- }
- if (propReader.reader.position < propData.length) {
- const argAttrs = propReader.readAttribute();
- if (argAttrs) {
- opState.addAttribute('arg_attrs', argAttrs);
- }
- }
- if (propReader.reader.position < propData.length) {
- const resAttrs = propReader.readAttribute();
- if (resAttrs) {
- opState.addAttribute('res_attrs', resAttrs);
- }
- }
- return;
- }
- if (opName.getStringRef().includes('.constant') || opName.getStringRef().endsWith('.const')) {
- if (propReader.reader.position < propData.length) {
- const attr = propReader.readAttribute();
- if (attr !== null && attr !== undefined) {
- opState.addAttribute('value', attr);
- }
- }
- }
- }
- }
- }
- };
- _.AsmResourceEntryKind = {
- Blob: 0,
- Bool: 1,
- String: 2
- };
- _.ParsedResourceEntry.Bytecode = class {
- constructor(key, kind, reader, stringReader) {
- this.key = key;
- this.kind = kind;
- this.reader = reader;
- this.stringReader = stringReader;
- }
- parseAsBlob() {
- if (this.kind === _.AsmResourceEntryKind.Blob) {
- const [data] = this.reader.parseBlobAndAlignment();
- return data;
- }
- return null;
- }
- };
- _.ResourceSectionReader = class {
- constructor() {
- this.dialectResources = [];
- this.dialectResourceHandleRenamingMap = new Map();
- }
- initialize(fileLoc, config, dialects, stringReader, sectionData, offsetSectionData, dialectReader) {
- const resourceReader = new _.EncodingReader(sectionData);
- const offsetReader = new _.EncodingReader(offsetSectionData);
- const numExternalResourceGroups = offsetReader.parseVarInt().toNumber();
- const parseGroup = (handler, allowEmpty = false, processKeyFn = null) => {
- const resolveKey = (key) => {
- const remapped = this.dialectResourceHandleRenamingMap.get(key);
- return remapped === undefined ? key : remapped;
- };
- return this.parseResourceGroup(fileLoc, allowEmpty, offsetReader, resourceReader, stringReader, handler, resolveKey, processKeyFn);
- };
- for (let i = 0; i < numExternalResourceGroups; i++) {
- const key = stringReader.parseString(offsetReader);
- const handler = config && config.getResourceParser ? config.getResourceParser(key) : null;
- if (!parseGroup(handler)) {
- return false;
- }
- }
- const ctx = fileLoc.context;
- while (!offsetReader.empty()) {
- const dialect = offsetReader.parseEntry(dialects, "dialect");
- dialect.load(dialectReader, ctx);
- const handler = dialect.getLoadedDialect();
- if (!handler) {
- throw new mlir.Error(`Unknown dialect '${dialect.name}'.`);
- }
- const processResourceKeyFn = (key) => {
- const handle = handler.declareResource(key);
- if (!handle) {
- throw new mlir.Error(`Unknown 'resource' key '${key}' for dialect '${dialect.name}'.`);
- }
- this.dialectResourceHandleRenamingMap.set(key, handler.getResourceKey(handle));
- this.dialectResources.push(handler);
- return true;
- };
- if (!parseGroup(handler, true, processResourceKeyFn)) {
- return false;
- }
- }
- return true;
- }
- parseResourceGroup(fileLoc, allowEmpty, offsetReader, resourceReader, stringReader, handler, remapKey, processKeyFn) {
- const numResources = offsetReader.parseVarInt().toNumber();
- for (let i = 0; i < numResources; i++) {
- const key = stringReader.parseString(offsetReader);
- const resourceOffset = offsetReader.parseVarInt().toNumber();
- const kind = offsetReader.parseByte();
- const data = resourceReader.parseBytes(resourceOffset);
- if (processKeyFn && !processKeyFn(key)) {
- return false;
- }
- if (allowEmpty && data.length === 0) {
- continue;
- }
- const entryReader = new _.EncodingReader(data);
- const resolvedKey = remapKey ? remapKey(key) : key;
- const entry = new _.ParsedResourceEntry.Bytecode(resolvedKey, kind, entryReader, stringReader);
- if (!handler) {
- continue;
- }
- handler.parseResource(entry);
- if (!entryReader.empty()) {
- throw new mlir.Error(`Unexpected trailing bytes in resource entry '${resolvedKey}'.`);
- }
- }
- return true;
- }
- parseResourceHandle(reader) {
- return reader.parseEntry(this.dialectResources, "resource handle");
- }
- };
- _.Location = class {
- constructor(context) {
- this.context = context;
- }
- };
- _.LocationAttr = class {
- constructor() {
- this.name = 'loc';
- }
- get value() {
- return 'unknown';
- }
- };
- _.UnknownLoc = class extends _.LocationAttr {
- get value() {
- return 'unknown';
- }
- };
- _.FileLineColLoc = class extends _.LocationAttr {
- constructor(filename, line, col) {
- super();
- this.filename = filename;
- this.line = line;
- this.col = col;
- }
- get value() {
- return `${this.filename}:${this.line}:${this.col}`;
- }
- };
- _.FileLineColRange = class extends _.LocationAttr {
- constructor(filename, startLine, startCol, endLine, endCol) {
- super();
- this.filename = filename;
- this.startLine = startLine;
- this.startCol = startCol;
- this.endLine = endLine;
- this.endCol = endCol;
- }
- get value() {
- if (this.endLine !== undefined && this.endCol !== undefined) {
- return `${this.filename}:${this.startLine}:${this.startCol} to ${this.endLine}:${this.endCol}`;
- }
- if (this.endCol !== undefined) {
- return `${this.filename}:${this.startLine}:${this.startCol} to :${this.endCol}`;
- }
- if (this.startCol !== undefined) {
- return `${this.filename}:${this.startLine}:${this.startCol}`;
- }
- return `${this.filename}:${this.startLine}`;
- }
- };
- _.NameLoc = class extends _.LocationAttr {
- constructor(name, childLoc) {
- super();
- this._name = name;
- this.childLoc = childLoc;
- }
- get value() {
- if (this.childLoc) {
- const childStr = this.childLoc.value || this.childLoc;
- return `"${this._name}"(${childStr})`;
- }
- return `"${this._name}"`;
- }
- };
- _.CallSiteLoc = class extends _.LocationAttr {
- constructor(callee, caller) {
- super();
- this.callee = callee;
- this.caller = caller;
- }
- get value() {
- const calleeStr = this.callee && this.callee.value ? this.callee.value : this.callee;
- const callerStr = this.caller && this.caller.value ? this.caller.value : this.caller;
- return `callsite(${calleeStr} at ${callerStr})`;
- }
- };
- _.FusedLoc = class extends _.LocationAttr {
- constructor(locations, metadata) {
- super();
- this.locations = locations;
- this.metadata = metadata;
- }
- get value() {
- const locStrs = this.locations.map((loc) => loc && loc.value ? loc.value : loc);
- if (this.metadata) {
- const metaStr = this.metadata.value === undefined ? this.metadata : this.metadata.value;
- return `fused<${metaStr}>[${locStrs.join(', ')}]`;
- }
- return `fused[${locStrs.join(', ')}]`;
- }
- };
- _.XtenNNDictLoc = class extends _.LocationAttr {
- constructor(dictContent) {
- super();
- this.dictContent = dictContent;
- }
- get value() {
- return `#xten_nn<dict_loc(${this.dictContent}`;
- }
- };
- _.OpaqueLoc = class extends _.LocationAttr {
- constructor(index, identifier, fallbackLoc) {
- super();
- this.index = index;
- this.identifier = identifier;
- this.fallbackLoc = fallbackLoc || new _.UnknownLoc();
- }
- get value() {
- return this.fallbackLoc.value;
- }
- };
- _.BytecodeReader = class {
- constructor(reader, config) {
- this.reader = new _.EncodingReader(reader);
- this.fileLoc = new _.Location(config.context);
- this.valueScopes = [];
- this.opNames = [];
- this.dialects = [];
- this.dialectsMap = new Map();
- this.resourceReader = new _.ResourceSectionReader();
- this.stringReader = new _.StringSectionReader();
- this.propertiesReader = new _.PropertiesSectionReader();
- this.attrTypeReader = new _.AttrTypeReader(this.stringReader, this.resourceReader, this.dialectsMap, this.version, this.fileLoc, this.config);
- // Store buffer start address for alignment validation
- this._bufferStart = reader instanceof Uint8Array ? reader.byteOffset : 0;
- }
- read() {
- const reader = this.reader;
- const signature = reader.parseBytes(4);
- if (String.fromCharCode(...signature) !== 'ML\xEFR') {
- throw new mlir.Error('Invalid MLIR bytecode signature.');
- }
- this.parseVersion(reader);
- this.producer = reader.parseNullTerminatedString();
- const sectionDatas = new Map();
- while (reader.position < reader.length) {
- const sectionIDAndHasAlignment = reader.parseByte();
- const sectionID = sectionIDAndHasAlignment & 0x7F;
- const length = reader.parseVarInt().toNumber();
- const hasAlignment = sectionIDAndHasAlignment & 0x80;
- if (sectionID >= 9) {
- throw new mlir.Error(`Unsupported section identifier '${sectionID}'.`);
- }
- if (hasAlignment) {
- const alignment = reader.parseVarInt().toNumber();
- this.checkSectionAlignment(alignment);
- reader.alignTo(alignment);
- }
- const sectionData = reader.parseBytes(length);
- sectionDatas.set(sectionID, sectionData);
- }
- for (let sectionID = 0; sectionID < 9; sectionID++) {
- if (sectionDatas.has(sectionID)) {
- continue;
- }
- if (sectionID <= 4 || (sectionID === 8 && this.version >= 5)) { // kString, kDialect, kAttrType, kAttrTypeOffset, kIR, kProperties + kNativePropertiesEncoding
- throw new mlir.Error(`Missing section '${sectionID}'.`);
- }
- }
- this.stringReader.initialize(sectionDatas.get(0));
- if (sectionDatas.has(8)) {
- this.propertiesReader.initialize(sectionDatas.get(8));
- }
- this.parseDialectSection(sectionDatas.get(1));
- this.parseResourceSection(sectionDatas.get(5), sectionDatas.get(6));
- this.attrTypeReader.initialize(this.dialects, sectionDatas.get(2), sectionDatas.get(3));
- return this.parseIRSection(sectionDatas.get(4));
- }
- checkSectionAlignment(alignment) {
- // Check that the bytecode buffer meets the requested section alignment.
- // In JavaScript, we validate the byteOffset within the ArrayBuffer.
- // This ensures the buffer offset is aligned to the requested alignment.
- if (!Number.isInteger(alignment) || alignment <= 0 || (alignment & (alignment - 1)) !== 0) {
- throw new mlir.Error(`Invalid alignment value: ${alignment} (must be a power of 2).`);
- }
- const isGloballyAligned = (this._bufferStart & (alignment - 1)) === 0;
- if (!isGloballyAligned) {
- throw new mlir.Error(`Expected section alignment ${alignment} but bytecode buffer offset 0x${this._bufferStart.toString(16)} is not aligned.`);
- }
- }
- parseVersion(reader) {
- this.version = reader.parseVarInt().toNumber();
- const kVersion = 6;
- const kMinSupportedVersion = 0;
- if (this.version < kMinSupportedVersion || this.version > kVersion) {
- throw new mlir.Error(`Unsupported MLIR bytecode version '${this.version}'.`);
- }
- }
- parseDialectSection(sectionData) {
- const sectionReader = new _.EncodingReader(sectionData);
- const numDialects = sectionReader.parseVarInt().toNumber();
- const checkSectionAlignment = (alignment) => {
- this.checkSectionAlignment(alignment);
- };
- for (let i = 0; i < numDialects; i++) {
- this.dialects.push(new _.BytecodeDialect());
- if (this.version < 1) { // kDialectVersioning
- this.dialects[i].name = this.stringReader.parseString();
- continue;
- }
- const [dialectNameIdx, versionAvailable] = sectionReader.parseVarIntWithFlag();
- this.dialects[i].name = this.stringReader.parseStringAtIndex(sectionReader, dialectNameIdx.toNumber());
- if (versionAvailable) {
- const [sectionID, versionBuffer] = sectionReader.parseSection(checkSectionAlignment);
- this.dialects[i].versionBuffer = versionBuffer;
- if (sectionID !== 7) { // kDialectVersion
- throw new mlir.Error(`Expected dialect version section.`);
- }
- }
- }
- let index = 0;
- let numOps = -1;
- const parseOpName = (dialect) => {
- const opName = {};
- opName.dialect = dialect;
- if (this.version < 5) { // kNativePropertiesEncoding
- opName.name = this.stringReader.parseString(sectionReader);
- } else {
- [opName.name, opName.wasRegistered] = this.stringReader.parseStringWithFlag(sectionReader);
- }
- if (numOps < 0) {
- this.opNames.push(opName);
- } else {
- this.opNames[index++] = opName;
- }
- };
- if (this.version >= 4) { // kElideUnknownBlockArgLocation
- numOps = sectionReader.parseVarInt().toNumber();
- this.opNames = new Array(numOps);
- }
- while (sectionReader.position < sectionData.length) {
- _.BytecodeReader.parseDialectGrouping(sectionReader, this.dialects, parseOpName);
- }
- }
- parseResourceSection(resourceData, resourceOffsetData) {
- if (!resourceOffsetData) {
- return true;
- }
- const dialectReader = new _.DialectReader(this.attrTypeReader, this.stringReader, this.resourceReader, this.dialectsMap, new _.EncodingReader(resourceData || new Uint8Array(0)), this.version);
- return this.resourceReader.initialize(this.fileLoc, this.config, this.dialects, this.stringReader, resourceData, resourceOffsetData, dialectReader);
- }
- parseIRSection(sectionData) {
- const reader = new _.EncodingReader(sectionData);
- const block = { operations: [] };
- this.valueScopes = [[]];
- const regionStack = [{
- block,
- curRegion: 0,
- numRegions: 1,
- curBlock: 0,
- numBlocks: 1,
- numOpsRemaining: 0,
- numValues: 0,
- blocks: [block],
- nextValueIdx: 0,
- isTopLevel: true,
- reader
- }];
- const firstBlockHeader = this.parseBlockHeader(reader);
- regionStack[0].numOpsRemaining = firstBlockHeader.numOps;
- if (firstBlockHeader.hasArgs) {
- const [scope] = this.valueScopes;
- this.parseBlockArguments(reader, block, scope, 0);
- regionStack[0].nextValueIdx = block.arguments ? block.arguments.length : 0;
- }
- // Iteratively parse regions until everything has been resolved.
- while (regionStack.length > 0) {
- this.parseRegions(regionStack, regionStack[regionStack.length - 1]);
- }
- return block;
- }
- parseRegion(readState) {
- const reader = readState.reader;
- const numBlocks = reader.parseVarInt().toNumber();
- if (numBlocks === 0) {
- return false;
- }
- const numValues = reader.parseVarInt().toNumber();
- readState.numValues = numValues;
- readState.numBlocks = numBlocks;
- // Create the blocks within this region.
- const blocks = [];
- for (let j = 0; j < numBlocks; j++) {
- blocks.push({ operations: [], arguments: [] });
- }
- readState.blocks = blocks;
- readState.region.blocks = blocks;
- // Prepare the current value scope for this region.
- const scope = this.valueScopes[this.valueScopes.length - 1];
- const valueOffset = scope.length;
- readState.valueOffset = valueOffset;
- for (let j = 0; j < numValues; j++) {
- scope.push(null);
- }
- // Parse the entry block of the region.
- readState.curBlock = 0;
- const blockHeader = this.parseBlockHeader(reader);
- readState.numOpsRemaining = blockHeader.numOps;
- if (blockHeader.hasArgs) {
- this.parseBlockArguments(reader, blocks[0], scope, valueOffset);
- }
- const numBlockArgs = blocks[0].arguments ? blocks[0].arguments.length : 0;
- readState.nextValueIdx = valueOffset + numBlockArgs;
- return true;
- }
- parseRegions(regionStack, readState) {
- const reader = readState.reader;
- while (true) {
- while (readState.numOpsRemaining > 0) {
- readState.numOpsRemaining--;
- // Read in the next operation. We don't read its regions directly,
- // we handle those afterwards as necessary.
- const { state: opState, resultNames, isIsolatedFromAbove, resultIndices } = this.parseOpWithoutRegions(reader, readState);
- const op = _.Operation.create(opState);
- // Assign result names for Netron display (reference: names are in parser symbol table)
- // Also update the value scope to replace placeholders with actual OpResult objects
- const scope = this.valueScopes[this.valueScopes.length - 1];
- for (let i = 0; i < resultNames.length && i < op.results.length; i++) {
- op.results[i].name = resultNames[i];
- // Replace placeholder in scope with actual result
- if (resultIndices && i < resultIndices.length) {
- const valueIdx = resultIndices[i];
- if (valueIdx < scope.length) {
- scope[valueIdx] = op.results[i];
- }
- }
- }
- readState.blocks[readState.curBlock].operations.push(op);
- if (op.regions && op.regions.length > 0) {
- for (let i = op.regions.length - 1; i >= 0; i--) {
- const region = op.regions[i];
- const childState = {
- region,
- curRegion: 0,
- numRegions: 1,
- curBlock: 0,
- numBlocks: 0,
- numOpsRemaining: 0,
- numValues: 0,
- blocks: [],
- valueOffset: 0,
- nextValueIdx: 0,
- isIsolated: isIsolatedFromAbove,
- reader,
- owningReader: null
- };
- // Isolated regions are encoded as a section in version 2 and above.
- if (this.version >= 2 && isIsolatedFromAbove) { // kLazyLoading
- const checkSectionAlignment = (alignment) => {
- this.checkSectionAlignment(alignment);
- };
- const [sectionID, sectionData] = reader.parseSection(checkSectionAlignment);
- if (sectionID !== 4) { // kIR
- throw new mlir.Error(`Expected IR section for region.`);
- }
- childState.owningReader = new _.EncodingReader(sectionData);
- childState.reader = childState.owningReader;
- }
- // If the op is isolated from above, push a new value scope.
- if (isIsolatedFromAbove) {
- this.valueScopes.push([]);
- }
- // Parse the region and push to stack if non-empty
- if (this.parseRegion(childState)) {
- regionStack.push(childState);
- return;
- }
- }
- }
- }
- // Move to the next block of the region.
- readState.curBlock++;
- if (readState.curBlock >= readState.numBlocks) {
- break;
- }
- const blockHeader = this.parseBlockHeader(reader);
- readState.numOpsRemaining = blockHeader.numOps;
- if (blockHeader.hasArgs) {
- const scope = this.valueScopes[this.valueScopes.length - 1];
- const argOffset = readState.nextValueIdx === undefined ? 0 : readState.nextValueIdx;
- this.parseBlockArguments(reader, readState.blocks[readState.curBlock], scope, argOffset);
- const numBlockArgs = readState.blocks[readState.curBlock].arguments ? readState.blocks[readState.curBlock].arguments.length : 0;
- if (readState.nextValueIdx !== undefined) {
- readState.nextValueIdx += numBlockArgs;
- }
- }
- }
- if (readState.isIsolated) {
- this.valueScopes.pop();
- }
- regionStack.pop();
- }
- parseBlockHeader(reader) {
- const numOpsAndHasArgs = reader.parseVarInt();
- const numOps = (numOpsAndHasArgs >> 1n).toNumber();
- const hasArgs = (numOpsAndHasArgs & 1n) === 1n;
- return { numOps, hasArgs };
- }
- parseBlockArguments(reader, block, scope, valueOffset) {
- const numArgs = reader.parseVarInt().toNumber();
- block.arguments = [];
- for (let i = 0; i < numArgs; i++) {
- let type = null;
- let location = null;
- if (this.version >= 4) { // kElideUnknownBlockArgLocation
- const typeAndLocation = reader.parseVarInt().toNumber();
- const typeIdx = typeAndLocation >> 1;
- const hasLocation = (typeAndLocation & 1) === 1;
- type = this.attrTypeReader.readType(typeIdx);
- if (hasLocation) {
- const locIdx = reader.parseVarInt().toNumber();
- location = this.attrTypeReader.readAttribute(locIdx);
- }
- } else {
- const typeIdx = reader.parseVarInt().toNumber();
- type = this.attrTypeReader.readType(typeIdx);
- const locIdx = reader.parseVarInt().toNumber();
- location = this.attrTypeReader.readAttribute(locIdx);
- }
- const argName = `%${valueOffset + i}`;
- const arg = new _.Value(argName, type);
- arg.location = location;
- block.arguments.push(arg);
- if (scope && (valueOffset + i) < scope.length) {
- scope[valueOffset + i] = arg;
- }
- }
- // Use-list ordering (version >= 3) - stored after all arguments
- // hasUseListOrders byte is read whenever hasArgs is true (not dependent on numArgs)
- if (this.version >= 3) { // kUseListOrdering
- const hasUseListOrders = reader.parseByte();
- if (hasUseListOrders !== 0) {
- const orders = this.parseUseListOrdersForRange(reader, numArgs);
- // Store ordering information on block arguments
- for (const order of orders) {
- if (order.argIndex < block.arguments.length) {
- block.arguments[order.argIndex].useListOrder = order.ordering;
- }
- }
- }
- }
- }
- parseUseListOrdersForRange(reader, numValues) {
- const orders = [];
- let numToRead = 1;
- if (numValues > 1) {
- numToRead = reader.parseVarInt().toNumber();
- }
- for (let i = 0; i < numToRead; i++) {
- let argIndex = 0;
- if (numValues > 1) {
- argIndex = reader.parseVarInt().toNumber();
- }
- const numUsesAndFlag = reader.parseVarInt();
- const numUses = (numUsesAndFlag >> 1n).toNumber();
- const indexPairEncoding = (numUsesAndFlag & 1n) === 1n;
- const indices = [];
- for (let j = 0; j < numUses; j++) {
- indices.push(reader.parseVarInt().toNumber());
- }
- orders.push({ argIndex, ordering: { indexPairEncoding, indices } });
- }
- return orders;
- }
- parseOpName(reader) {
- const opName = reader.parseEntry(this.opNames, 'operation name');
- if (!opName.opName) {
- const name = `${opName.dialect.name}.${opName.name}`;
- opName.opName = _.RegisteredOperationName.lookup(name, this.fileLoc.context);
- if (!opName.opName) {
- throw new mlir.Error(`Unregistered bytecode operation '${name}'.`);
- }
- }
- return [opName.opName, opName.wasRegistered];
- }
- parseOpWithoutRegions(reader, state) {
- // Parse operation name index
- const [opName, wasRegistered] = this.parseOpName(reader);
- const opMask = reader.parseByte();
- const kHasAttrs = 0x01;
- const kHasResults = 0x02;
- const kHasOperands = 0x04;
- const kHasSuccessors = 0x08;
- const kHasInlineRegions = 0x10;
- const kHasUseListOrders = 0x20;
- const kHasProperties = 0x40;
- const opState = new _.OperationState(null, opName);
- const locIdx = reader.parseVarInt().toNumber();
- opState.location = this.attrTypeReader.readAttribute(locIdx);
- if (opMask & kHasAttrs) {
- const dictAttrIdx = reader.parseVarInt().toNumber();
- const dictAttr = this.attrTypeReader.readAttribute(dictAttrIdx);
- if (dictAttr && dictAttr.value) {
- if (dictAttr.value instanceof Map) {
- // Already parsed as Map from custom-encoded DictionaryAttr
- opState.attributes = dictAttr.value;
- } else if (typeof dictAttr.value === 'string') {
- // Parse dictionary attribute from ASM string format
- opState.attributes = this.parseAttributeDict(dictAttr.value);
- }
- }
- }
- // Parse properties (version >= 5)
- if (opMask & kHasProperties) {
- if (wasRegistered) {
- const dialectReader = new _.DialectReader(this.attrTypeReader, this.stringReader, this.resourceReader, this.dialectsMap, reader, this.version);
- this.propertiesReader.read(this.fileLoc, dialectReader, opName, opState);
- } else {
- // Unregistered operations store properties as a single dictionary attribute
- const propAttrIdx = reader.parseVarInt().toNumber();
- const propAttr = this.attrTypeReader.readAttribute(propAttrIdx);
- if (propAttr && propAttr.value) {
- const propAttrs = this.parseAttributeDict(propAttr.value);
- for (const [key, value] of propAttrs) {
- opState.addAttribute(key, value);
- }
- }
- }
- }
- // Parse results - add types to OperationState.types, track values in scope
- const resultNames = [];
- const resultIndices = [];
- if (opMask & kHasResults) {
- const numResults = reader.parseVarInt().toNumber();
- const scope = this.valueScopes[this.valueScopes.length - 1];
- for (let i = 0; i < numResults; i++) {
- const typeIdx = reader.parseVarInt().toNumber();
- const type = this.attrTypeReader.readType(typeIdx);
- opState.addTypes([type]);
- const valueIdx = state && state.nextValueIdx !== undefined ? state.nextValueIdx++ : scope.length;
- const valueName = `%${valueIdx}`;
- resultNames.push(valueName);
- resultIndices.push(valueIdx);
- // Create placeholder in scope (will be replaced after Operation creation with actual OpResult)
- const placeholder = new _.Value(valueName, type);
- if (valueIdx < scope.length) {
- scope[valueIdx] = placeholder;
- } else {
- scope.push(placeholder);
- }
- }
- }
- if (opMask & kHasOperands) {
- const numOperands = reader.parseVarInt().toNumber();
- for (let i = 0; i < numOperands; i++) {
- const valueIdx = reader.parseVarInt().toNumber();
- const scope = this.valueScopes[this.valueScopes.length - 1];
- if (valueIdx < scope.length && scope[valueIdx]) {
- opState.operands.push(scope[valueIdx]);
- } else {
- opState.operands.push(new _.Value(`%${valueIdx}`, null));
- }
- }
- }
- if (opMask & kHasSuccessors) {
- const numSuccessors = reader.parseVarInt().toNumber();
- opState.successors = [];
- for (let i = 0; i < numSuccessors; i++) {
- const blockIdx = reader.parseVarInt().toNumber();
- opState.successors.push(blockIdx);
- }
- }
- // Parse use-list orders (version >= 3)
- if (this.version >= 3 && (opMask & kHasUseListOrders)) { // kUseListOrdering
- const numResults = opState.types.length;
- const orders = this.parseUseListOrdersForRange(reader, numResults);
- if (orders.length > 0) {
- opState.useListOrders = orders;
- }
- }
- let isIsolatedFromAbove = false;
- if (opMask & kHasInlineRegions) {
- const numRegionsAndIsIsolated = reader.parseVarInt();
- const numRegions = (numRegionsAndIsIsolated >> 1n).toNumber();
- isIsolatedFromAbove = (numRegionsAndIsIsolated & 1n) === 1n;
- for (let i = 0; i < numRegions; i++) {
- opState.regions.push({ blocks: [] });
- }
- }
- return { state: opState, resultNames, isIsolatedFromAbove, resultIndices };
- }
- parseAttributeDict(str) {
- const attrs = new Map();
- // Parse dictionary attribute format: {key = value, ...}
- if (!str.startsWith('{') || !str.endsWith('}')) {
- return attrs;
- }
- const content = str.slice(1, -1).trim();
- if (content.length === 0) {
- return attrs;
- }
- let i = 0;
- while (i < content.length) {
- while (i < content.length && /\s/.test(content[i])) {
- i++;
- }
- if (i >= content.length) {
- break;
- }
- // Read key (alphanumeric/underscore)
- const keyStart = i;
- while (i < content.length && /[a-zA-Z0-9_]/.test(content[i])) {
- i++;
- }
- const key = content.slice(keyStart, i);
- if (!key) {
- break;
- }
- while (i < content.length && /\s/.test(content[i])) {
- i++;
- }
- if (i >= content.length || content[i] !== '=') {
- break;
- }
- i++; // skip '='
- while (i < content.length && /\s/.test(content[i])) {
- i++;
- }
- // Read value (until balanced comma or end)
- const valueStart = i;
- let depth = 0;
- let inString = false;
- while (i < content.length) {
- const c = content[i];
- if (inString) {
- if (c === '"' && content[i - 1] !== '\\') {
- inString = false;
- }
- i++;
- continue;
- }
- if (c === '"') {
- inString = true;
- i++;
- continue;
- }
- if (c === '{' || c === '[' || c === '(' || c === '<') {
- depth++;
- } else if (c === '}' || c === ']' || c === ')' || c === '>') {
- depth--;
- } else if (c === ',' && depth === 0) {
- break;
- }
- i++;
- }
- const value = content.slice(valueStart, i).trim();
- attrs.set(key, { value, toString: () => value });
- if (i < content.length && content[i] === ',') {
- i++;
- }
- }
- return attrs;
- }
- static parseDialectGrouping(reader, dialects, entryCallback) {
- const dialect = reader.parseEntry(dialects, 'dialect');
- const numEntries = reader.parseVarInt().toNumber();
- for (let i = 0; i < numEntries; i++) {
- entryCallback(dialect);
- }
- }
- };
- _.Constraint = class {
- constructor(name, args, values) {
- this.name = name;
- this.args = args || null;
- this.values = values || null;
- }
- static parse(value) {
- const tokens = _.Constraint.tokenize(value);
- const result = _.Constraint.parseTokens(tokens, 0);
- const { name, args, values } = result.value;
- return new _.Constraint(name, args, values);
- }
- toString() {
- if (this.args && this.args.length > 0) {
- const args = this.args.map((arg) => arg.name).join(', ');
- return `${this.name}<${args}>`;
- }
- if (this.values && this.values.length > 0) {
- return `${this.name}{${this.values.join(' | ')}}`;
- }
- return this.name;
- }
- static tokenize(str) {
- const tokens = [];
- let i = 0;
- while (i < str.length) {
- const ch = str[i];
- if (/\s/.test(ch)) {
- i++;
- continue;
- }
- if ('<>={}[](),|'.indexOf(ch) !== -1) {
- tokens.push({ type: ch, value: ch, pos: i });
- i++;
- continue;
- }
- if (ch === ':' && i + 1 < str.length && str[i + 1] === ':') {
- tokens.push({ type: '::', value: '::', pos: i });
- i += 2;
- continue;
- }
- if (ch === ':') {
- i++;
- continue;
- }
- if (ch === '"' || ch === "'") {
- const quote = ch;
- let j = i + 1;
- while (j < str.length && str[j] !== quote) {
- if (str[j] === '\\' && j + 1 < str.length) {
- j += 2;
- } else {
- j++;
- }
- }
- if (j < str.length) {
- tokens.push({ type: 'string', value: str.substring(i + 1, j), pos: i });
- i = j + 1;
- } else {
- tokens.push({ type: 'ident', value: str.substring(i), pos: i });
- break;
- }
- continue;
- }
- if (/[a-zA-Z_0-9-]/.test(ch)) {
- let j = i;
- while (j < str.length && /[a-zA-Z_0-9-]/.test(str[j])) {
- j++;
- }
- const ident = str.substring(i, j);
- tokens.push({ type: 'ident', value: ident, pos: i });
- i = j;
- continue;
- }
- i++;
- }
- return tokens;
- }
- static parseTokens(tokens, pos) {
- if (pos >= tokens.length) {
- return null;
- }
- const token = tokens[pos];
- if (token.type === '::') {
- let name = '';
- let nextPos = pos;
- while (nextPos < tokens.length) {
- if (tokens[nextPos].type === '::') {
- name += '::';
- nextPos++;
- } else if (tokens[nextPos].type === 'ident') {
- name += tokens[nextPos].value;
- nextPos++;
- } else {
- break;
- }
- }
- if (!name) {
- return null;
- }
- if (nextPos < tokens.length) {
- const nextToken = tokens[nextPos];
- if (nextToken.type === '{') {
- return _.Constraint._parseEnum(tokens, pos, name);
- }
- if (nextToken.type === '<') {
- return _.Constraint._parseGeneric(tokens, pos, name);
- }
- }
- return { value: { name }, nextPos };
- }
- if (token.type !== 'ident') {
- return null;
- }
- let name = token.value;
- let nextPos = pos + 1;
- while (nextPos < tokens.length && tokens[nextPos].type === '::') {
- nextPos++;
- if (nextPos < tokens.length && tokens[nextPos].type === 'ident') {
- const value = tokens[nextPos++].value;
- name += `::${value}`;
- } else {
- break;
- }
- }
- if (nextPos >= tokens.length) {
- return { value: { name }, nextPos };
- }
- const nextToken = tokens[nextPos];
- if (nextToken.type === '{') {
- return _.Constraint._parseEnum(tokens, pos, name);
- }
- if (nextToken.type === '<') {
- return _.Constraint._parseGeneric(tokens, pos, name);
- }
- return { value: { name }, nextPos };
- }
- static _parseEnum(tokens, startPos, name) {
- let pos = startPos;
- while (pos < tokens.length && (tokens[pos].type === 'ident' || tokens[pos].type === '::')) {
- pos++;
- }
- if (pos >= tokens.length || tokens[pos].type !== '{') {
- return null;
- }
- pos++;
- const values = [];
- let currentValue = '';
- while (pos < tokens.length && tokens[pos].type !== '}') {
- const token = tokens[pos];
- if (token.type === '|') {
- if (currentValue.trim()) {
- values.push(currentValue.trim());
- currentValue = '';
- }
- pos++;
- } else if (token.type === 'ident') {
- if (currentValue) {
- currentValue += ' ';
- }
- currentValue += token.value;
- pos++;
- } else if (token.type === '::') {
- currentValue += '::';
- pos++;
- } else {
- pos++;
- }
- }
- if (currentValue.trim()) {
- values.push(currentValue.trim());
- }
- if (pos < tokens.length && tokens[pos].type === '}') {
- pos++;
- }
- return { value: { name, values }, nextPos: pos };
- }
- static _parseGeneric(tokens, startPos, name) {
- let pos = startPos;
- while (pos < tokens.length && (tokens[pos].type === 'ident' || tokens[pos].type === '::')) {
- pos++;
- }
- if (pos >= tokens.length || tokens[pos].type !== '<') {
- return null;
- }
- pos++;
- const args = [];
- let angleDepth = 1;
- let bracketDepth = 0;
- let currentArg = [];
- while (pos < tokens.length && (angleDepth > 0 || bracketDepth > 0)) {
- const token = tokens[pos];
- if (token.type === '<') {
- angleDepth++;
- currentArg.push(token);
- pos++;
- } else if (token.type === '>') {
- angleDepth--;
- if (angleDepth === 0 && bracketDepth === 0) {
- if (currentArg.length > 0) {
- const parsed = _.Constraint.parseArgumentTokens(currentArg);
- if (parsed !== null) {
- args.push(parsed);
- }
- }
- pos++;
- break;
- }
- currentArg.push(token);
- pos++;
- } else if (token.type === '[') {
- bracketDepth++;
- currentArg.push(token);
- pos++;
- } else if (token.type === ']') {
- bracketDepth--;
- currentArg.push(token);
- pos++;
- } else if (token.type === ',' && angleDepth === 1 && bracketDepth === 0) {
- if (currentArg.length > 0) {
- const parsed = _.Constraint.parseArgumentTokens(currentArg);
- if (parsed !== null) {
- args.push(parsed);
- }
- currentArg = [];
- }
- pos++;
- } else {
- currentArg.push(token);
- pos++;
- }
- }
- return { value: { name, args }, nextPos: pos };
- }
- static parseArgumentTokens(tokens) {
- if (!tokens || tokens.length === 0) {
- return null;
- }
- tokens = tokens.filter((t) => t.type !== undefined);
- if (tokens[0].type === '[') {
- return _.Constraint.parseListArgument(tokens);
- }
- if (tokens[0].type === 'string') {
- return tokens[0].value;
- }
- if (tokens[0].type === 'ident' || tokens[0].type === '::') {
- const result = _.Constraint.parseTokens(tokens, 0);
- if (result && result.nextPos === tokens.length) {
- return result.value;
- }
- }
- let literal = '';
- for (const token of tokens) {
- if (token.type === 'ident' || token.type === 'string') {
- if (literal && !/^[,[\]():.]$/.test(literal[literal.length - 1])) {
- literal += ' ';
- }
- literal += token.value;
- } else if (token.type === '::') {
- literal += '::';
- } else if ('{}[](),.'.indexOf(token.type) !== -1) {
- literal += token.value;
- }
- }
- return literal.trim() || null;
- }
- static parseListArgument(tokens) {
- if (!tokens || tokens.length === 0 || tokens[0].type !== '[') {
- return null;
- }
- let pos = 1;
- const items = [];
- let bracketDepth = 1;
- let angleDepth = 0;
- let currentItem = [];
- while (pos < tokens.length && (bracketDepth > 0 || angleDepth > 0)) {
- const token = tokens[pos];
- if (token.type === '[') {
- bracketDepth++;
- currentItem.push(token);
- pos++;
- } else if (token.type === ']') {
- bracketDepth--;
- if (bracketDepth === 0 && angleDepth === 0) {
- if (currentItem.length > 0) {
- const parsed = _.Constraint.parseArgumentTokens(currentItem);
- if (parsed !== null) {
- items.push(parsed);
- }
- }
- break;
- }
- currentItem.push(token);
- pos++;
- } else if (token.type === '<') {
- angleDepth++;
- currentItem.push(token);
- pos++;
- } else if (token.type === '>') {
- angleDepth--;
- currentItem.push(token);
- pos++;
- } else if (token.type === ',' && bracketDepth === 1 && angleDepth === 0) {
- if (currentItem.length > 0) {
- const parsed = _.Constraint.parseArgumentTokens(currentItem);
- if (parsed !== null) {
- items.push(parsed);
- }
- currentItem = [];
- }
- pos++;
- } else {
- currentItem.push(token);
- pos++;
- }
- }
- return items;
- }
- };
- _.AssemblyFormatParser = class {
- constructor(metadata) {
- this._metadata = metadata;
- this._buffer = metadata.assemblyFormat || '';
- this._pos = 0;
- }
- match(char) {
- return this._pos < this._buffer.length && this._buffer[this._pos] === char;
- }
- accept(str) {
- if (str.length === 1) {
- if (this.match(str)) {
- this._pos++;
- return true;
- }
- return false;
- }
- const remaining = this._buffer.substring(this._pos);
- if (remaining.startsWith(str)) {
- // Check that keyword is not followed by alphanumeric (to avoid "type" in "typename")
- const nextChar = this._buffer[this._pos + str.length];
- if (nextChar && /[a-zA-Z0-9_-]/.test(nextChar)) {
- return false;
- }
- this._pos += str.length;
- return true;
- }
- return false;
- }
- expect(char) {
- if (!this.match(char)) {
- throw new mlir.Error(`Expected '${char}'.`);
- }
- this._pos++;
- }
- parse() {
- const directives = [];
- this._skipWhitespace();
- while (this._pos < this._buffer.length) {
- const directive = this.parseDirective();
- directives.push(directive);
- this._skipWhitespace();
- }
- return directives;
- }
- parseDirective() {
- const ch = this._buffer[this._pos];
- if (!ch || this._pos >= this._buffer.length) {
- throw new mlir.Error(`Unexpected end of format string.`);
- }
- // Parenthesized group: can be optional (...)? or conditional (...):(...) or just grouping (...)
- if (this.match('(')) {
- this.accept('(');
- const elements = [];
- let anchorElement = null;
- this._skipWhitespace();
- while (!this.match(')')) {
- const elem = this.parseDirective();
- if (elem.type === 'anchor') {
- // Standalone anchor - applies to the previous element
- if (elements.length > 0) {
- const prev = elements[elements.length - 1];
- anchorElement = prev.name || prev.type;
- }
- } else {
- if (elem.anchor) {
- anchorElement = elem.name || elem.type;
- }
- elements.push(elem);
- }
- this._skipWhitespace();
- }
- this.expect(')');
- this._skipWhitespace();
- // Check what follows to determine the group type
- if (this.accept('?')) {
- // Optional group: (...)?
- return { type: 'optional_group', elements, anchor: anchorElement };
- }
- if (this.accept(':')) {
- // Conditional alternative: (...):(...)?
- this._skipWhitespace();
- const secondAlt = [];
- let isSecondOptional = false;
- if (this.accept('(')) {
- this._skipWhitespace();
- while (!this.match(')')) {
- const elem = this.parseDirective();
- secondAlt.push(elem);
- this._skipWhitespace();
- }
- this.expect(')');
- this._skipWhitespace();
- if (this.accept('?')) {
- isSecondOptional = true;
- }
- }
- return { type: 'conditional_alternative', firstAlt: elements, secondAlt, secondOptional: isSecondOptional };
- }
- return { type: 'group', elements };
- }
- // Literal: `keyword`
- if (this.accept('`')) {
- const value = this.parseUntil('`');
- this.expect('`');
- // MLIR reference: Empty literals (`` or ` `) are whitespace, not literals
- if (value.length === 0 || value === ' ' || value === '\\n') {
- return { type: 'whitespace', value }; // Return whitespace as a directive
- }
- return { type: 'literal', value };
- }
- if (this.accept('$')) {
- const name = this.parseIdentifier();
- const anchor = this.accept('^');
- const metadata = this._metadata;
- // Determine variable type from metadata first - matches reference implementation
- // Check each metadata category in priority order
- if (metadata.successors && metadata.successors.some((a) => a.name === name)) {
- return { type: 'successor_ref', name, anchor };
- }
- if (metadata.attributes && metadata.attributes.some((a) => a.name === name)) {
- return { type: 'attribute_ref', name, anchor };
- }
- if (metadata.operands && metadata.operands.some((a) => a.name === name)) {
- return { type: 'operand_ref', name, anchor };
- }
- if (metadata.regions && metadata.regions.some((a) => a.name === name)) {
- return { type: 'region_ref', name, anchor };
- }
- throw new mlir.Error(`Unknown variable '$${name}' in assembly format.`);
- }
- if (this.accept('type')) {
- const args = this.parseParenList();
- const anchor = this.accept('^');
- return { type: 'type', args, anchor };
- }
- if (this.accept('qualified')) {
- const args = this.parseParenList();
- const anchor = this.accept('^');
- return { type: 'qualified', args, anchor };
- }
- if (this.accept('attr-dict-with-keyword')) {
- return { type: 'attr_dict_with_keyword' };
- }
- if (this.accept('attr-dict')) {
- return { type: 'attr_dict' };
- }
- if (this.accept('prop-dict')) {
- return { type: 'prop_dict' };
- }
- if (this.accept('functional-type')) {
- const args = this.parseParenList();
- const anchor = this.accept('^');
- return { type: 'functional_type', args, anchor };
- }
- if (this.accept('params')) {
- return { type: 'params' };
- }
- if (this.accept('struct')) {
- this.expect('(');
- const args = [];
- while (!this.match(')')) {
- this._skipWhitespace();
- if (this.match(')')) {
- break;
- }
- const arg = this.parseDirective();
- args.push(arg);
- this._skipWhitespace();
- this.accept(',');
- }
- this.expect(')');
- return { type: 'struct', args };
- }
- if (this.accept('ref')) {
- this.expect('(');
- const arg = this.parseDirective();
- this._skipWhitespace();
- this.expect(')');
- return { type: 'ref', arg };
- }
- if (this.accept('custom')) {
- this.expect('<');
- const parser = this.parseUntil('>');
- this.expect('>');
- const args = this.parseParenList();
- const anchor = this.accept('^');
- return { type: 'custom', parser, args, anchor };
- }
- if (this.accept('oilist')) {
- this._skipWhitespace();
- this.expect('(');
- let content = '';
- let depth = 1;
- while (this._pos < this._buffer.length && depth > 0) {
- const ch = this._buffer[this._pos];
- if (ch === '(') {
- depth++;
- content += ch;
- this._pos++;
- } else if (ch === ')') {
- depth--;
- if (depth > 0) {
- content += ch;
- }
- this._pos++;
- } else {
- content += ch;
- this._pos++;
- }
- }
- return { type: 'oilist', content };
- }
- if (this.accept('operands')) {
- return { type: 'operands' };
- }
- if (this.accept('results')) {
- return { type: 'results' };
- }
- if (this.accept('regions')) {
- return { type: 'regions' };
- }
- if (this.accept('successors')) {
- return { type: 'successors' };
- }
- if (ch === '^') {
- this._pos++;
- return { type: 'anchor' };
- }
- if (/^[:()[\]{}<>,=|]/.test(ch)) {
- this._pos++;
- return { type: 'literal', value: ch };
- }
- const context = this._buffer.substring(Math.max(0, this._pos - 10), Math.min(this._buffer.length, this._pos + 10));
- throw new mlir.Error(`Unexpected '${ch}' in assembly format '${context}...'.`);
- }
- parseIdentifier() {
- let name = '';
- while (this._pos < this._buffer.length) {
- const ch = this._buffer[this._pos];
- if (/[a-zA-Z0-9_]/.test(ch)) {
- name += ch;
- this._pos++;
- } else {
- break;
- }
- }
- return name;
- }
- parseUntil(terminator) {
- let value = '';
- while (this._pos < this._buffer.length && this._buffer[this._pos] !== terminator) {
- value += this._buffer[this._pos];
- this._pos++;
- }
- return value;
- }
- parseParenList() {
- this._skipWhitespace();
- if (!this.accept('(')) {
- return [];
- }
- this._skipWhitespace();
- if (this.accept(')')) {
- return [];
- }
- const items = [];
- const parseElement = () => {
- let element = '';
- let depth = 0;
- while (this._pos < this._buffer.length) {
- this._skipWhitespace();
- if (this.accept('"')) {
- // String literal - consume as a unit
- element += '"';
- element += this.parseUntil('"');
- element += '"';
- this.expect('"');
- } else if (this.accept('$')) {
- element += '$';
- const id = this.parseIdentifier();
- element += id;
- } else if (this.accept('(')) {
- // Nested parentheses - include in element (e.g., type($list))
- element += '(';
- depth++;
- } else if (this.match(')')) {
- if (depth === 0) {
- // End of this element
- break;
- }
- element += ')';
- this.accept(')');
- depth--;
- } else if (this.match(',') && depth === 0) {
- // Comma at top level - end of this element
- break;
- } else if (this.match('-')) {
- // Handle hyphenated identifiers like attr-dict, functional-type
- element += '-';
- this.accept('-');
- } else {
- // Plain identifier (e.g., "type" in type($list))
- const id = this.parseIdentifier();
- if (!id) {
- throw new mlir.Error(`Unexpected '${this._buffer[this._pos]}' in assembly format directive list.`);
- }
- element += id;
- }
- }
- return element.trim();
- };
- const first = parseElement();
- if (!first) {
- throw new mlir.Error('Expected element.');
- }
- items.push(first);
- this._skipWhitespace();
- while (this.accept(',')) {
- this._skipWhitespace();
- const elem = parseElement();
- if (!elem) {
- throw new mlir.Error('Expected element after comma');
- }
- items.push(elem);
- this._skipWhitespace();
- }
- this.expect(')');
- return items;
- }
- _skipWhitespace() {
- while (this._pos < this._buffer.length && /\s/.test(this._buffer[this._pos])) {
- this._pos++;
- }
- }
- };
- _.DialectContext = class {
- constructor(metadata) {
- const operations = metadata.operations;
- this._dialects = new Map();
- this._dialects.set('builtin', new _.BuiltinDialect(operations));
- this._dialects.set('std', new _.Dialect(operations, 'std'));
- this._dialects.set('bufferization', new _.BufferizationDialect(operations));
- this._dialects.set('stablehlo', new _.StableHLODialect(operations));
- this._dialects.set('vhlo', new _.VhloDialect(operations));
- this._dialects.set('interpreter', new _.InterpreterDialect(operations));
- this._dialects.set('affine', new _.AffineDialect(operations));
- this._dialects.set('asuka', new _.AsukaDialect(operations));
- this._dialects.set('arith', new _.ArithDialect(operations));
- this._dialects.set('async', new _.async.AsyncDialect(operations));
- this._dialects.set('cf', new _.CFDialect(operations));
- this._dialects.set('emitc', new _.EmitCDialect(operations));
- this._dialects.set('complex', new _.Dialect(operations, 'complex'));
- this._dialects.set('index', new _.Dialect(operations, 'index'));
- this._dialects.set('pdl', new _.pdl.PDLDialect(operations));
- this._dialects.set('ptr', new _.ptr.PtrDialect(operations));
- this._dialects.set('ub', new _.Dialect(operations, 'ub'));
- this._dialects.set('amdgpu', new _.AMDGPUDialect(operations));
- this._dialects.set('nvgpu', new _.NVGPUDialect(operations));
- this._dialects.set('nvvm', new _.NVVMDialect(operations));
- this._dialects.set('rocdl', new _.ROCDLDialect(operations));
- this._dialects.set('nvws', new _.NVWSDialect(operations));
- this._dialects.set('tti', new _.Dialect(operations, 'tti'));
- this._dialects.set('omp', new _.OpenMPDialect(operations));
- this._dialects.set('proton', new _.ProtonDialect(operations));
- this._dialects.set('proton_gpu', new _.Dialect(operations, 'proton_gpu'));
- this._dialects.set('arm_sme', new _.ArmSMEDialect(operations));
- this._dialects.set('arm_neon', new _.ArmNeonDialect(operations));
- this._dialects.set('arm_sve', new _.ArmSVEDialect(operations));
- this._dialects.set('shard', new _.ShardDialect(operations));
- this._dialects.set('amx', new _.Dialect(operations, 'amx'));
- this._dialects.set('smt', new _.smt.SMTDialect(operations));
- this._dialects.set('lagrad', new _.Dialect(operations, 'lagrad'));
- this._dialects.set('iree_codegen', new _.IREECodegenDialect(operations));
- this._dialects.set('iree_encoding', new _.Dialect(operations, 'iree_encoding'));
- this._dialects.set('test', new _.TestDialect(operations));
- this._dialects.set('scf', new _.SCFDialect(operations));
- this._dialects.set('shape', new _.ShapeDialect(operations));
- this._dialects.set('sparse_tensor', new _.SparseTensorDialect(operations));
- this._dialects.set('func', new _.FuncDialect(operations));
- this._dialects.set('gpu', new _.GpuDialect(operations));
- this._dialects.set('llvm', new _.LLVM.LLVMDialect(operations));
- this._dialects.set('xegpu', new _.XeGPUDialect(operations));
- this._dialects.set('memref', new _.MemRefDialect(operations));
- this._dialects.set('vector', new _.VectorDialect(operations));
- this._dialects.set('x86vector', new _.Dialect(operations, 'x86vector'));
- this._dialects.set('onnx', new _.ONNXDialect(operations));
- this._dialects.set('krnl', new _.KrnlDialect(operations));
- this._dialects.set('torch', new _.TorchDialect(operations));
- this._dialects.set('torch_c', new _.Dialect(operations, 'torch_c'));
- this._dialects.set('hal', new _.HALDialect(operations));
- this._dialects.set('hal_loader', new _.HALLoaderDialect(operations));
- this._dialects.set('hal_inline', new _.Dialect(operations, 'hal_inline'));
- this._dialects.set('util', new _.UtilDialect(operations));
- this._dialects.set('mhlo', new _.MhloDialect(operations));
- this._dialects.set('chlo', new _.ChloDialect(operations));
- this._dialects.set('thlo', new _.THLODialect(operations));
- this._dialects.set('flow', new _.FlowDialect(operations));
- this._dialects.set('stream', new _.StreamDialect(operations));
- this._dialects.set('iree_vector_ext', new _.IREEVectorExtDialect(operations));
- this._dialects.set('iree_tensor_ext', new _.IREETensorExtDialect(operations));
- this._dialects.set('linalg', new _.LinalgDialect(operations));
- this._dialects.set('iree_linalg_ext', new _.Dialect(operations, 'iree_linalg_ext'));
- this._dialects.set('linalg_ext', this._dialects.get('iree_linalg_ext'));
- this._dialects.set('quant', new _.QuantDialect(operations));
- this._dialects.set('tensor', new _.TensorDialect(operations));
- this._dialects.set('tosa', new _.TosaDialect(operations));
- this._dialects.set('tf', new _.TFDialect(operations));
- this._dialects.set('tf_saved_model', new _.Dialect(operations, 'tf_saved_model'));
- this._dialects.set('tf_type', new _.TFTypeDialect(operations));
- this._dialects.set('tf_device', new _.TFDeviceDialect(operations));
- this._dialects.set('tf_executor', new _.TFExecutorDialect(operations));
- this._dialects.set('tf_framework', new _.TFFrameworkDialect(operations));
- this._dialects.set('tfr', new _.TFRDialect(operations));
- this._dialects.set('corert', new _.CoreRTDialect(operations));
- this._dialects.set('tfrt', new _.TFRTDialect(operations));
- this._dialects.set('tfrt_fallback', new _.Dialect(operations, 'tfrt_fallback'));
- this._dialects.set('tfrt_fallback_async', new _.TFRTFallbackAsyncDialect(operations));
- this._dialects.set('tfl', new _.TFLDialect(operations));
- this._dialects.set('stdx', new _.StdxDialect(operations));
- this._dialects.set('vm', new _.VMDialect(operations));
- this._dialects.set('math', new _.MathDialect(operations));
- this._dialects.set('tm_tensor', new _.TMTensorDialect(operations));
- this._dialects.set('ml_program', new _.MLProgramDialect(operations));
- this._dialects.set('iree_gpu', new _.IREEGPUDialect(operations));
- this._dialects.set('tile', new _.TileDialect(operations));
- this._dialects.set('pxa', new _.PXADialect(operations));
- this._dialects.set('irdl', new _.IRDLDialect(operations));
- this._dialects.set('transform', new _.TransformDialect(operations));
- this._dialects.set('wasmssa', new _.WasmSSADialect(operations));
- this._dialects.set('spirv', new _.spirv.SPIRVDialect(operations));
- this._dialects.set('spv', this._dialects.get('spirv'));
- this._dialects.set('toy', new _.ToyDialect(operations));
- this._dialects.set('top', new _.Dialect(operations, 'top'));
- this._dialects.set('tpu', new _.Dialect(operations, 'tpu'));
- this._dialects.set('sdfg', new _.SdfgDialect(operations));
- this._dialects.set('sdir', this._dialects.get('sdfg'));
- this._dialects.set('check', new _.Dialect(operations, 'check'));
- this._dialects.set('tt', new _.triton.TritonDialect(operations));
- this._dialects.set('ttg', new _.triton.gpu.TritonGPUDialect(operations));
- this._dialects.set('triton_gpu', this._dialects.get('ttg'));
- this._dialects.set('gluon', new _.GluonDialect(operations));
- this._dialects.set('ttng', new _.TritonNvidiaGPUDialect(operations));
- this._dialects.set('nvidia_gpu', this._dialects.get('ttng'));
- this._dialects.set('amdg', new _.TritonAMDGPUDialect(operations));
- this._dialects.set('amd_gpu', this._dialects.get('amdg'));
- this._dialects.set('michelson', new _.MichelsonDialect(operations));
- this._dialects.set('tensorrt', new _.TensorRTDialect(operations));
- this._dialects.set('executor', new _.executor.ExecutorDialect(operations));
- this._dialects.set('exec', this._dialects.get('executor'));
- this._dialects.set('tfrt_test', new _.TFRTTestDialect(operations));
- this._dialects.set('xevm', new _.XeVMDialect(operations));
- this._dialects.set('vmvx', new _.VMVXDialect(operations));
- this._dialects.set('mlrt', new _.MLRTDialect(operations));
- this._dialects.set('tfrt_tensor', new _.TFRTTensorDialect(operations));
- this._dialects.set('tfrt_dht', new _.TFRTDHTDialect(operations));
- this._dialects.set('coo', new _.Dialect(operations, 'coo'));
- this._dialects.set('tfd', new _.TFDDialect(operations));
- this._dialects.set('acc', new _.ACCDialect(operations));
- this._dialects.set('cuda', new _.Dialect(operations, 'cuda'));
- this._dialects.set('trtrt', new _.Dialect(operations, 'trtrt'));
- this._dialects.set('plan', new _.PlanDialect(operations));
- this._dialects.set('kernel', new _.KernelDialect(operations));
- this._dialects.set('nvg', new _.Dialect(operations, 'nvg'));
- this._dialects.set('mpi', new _.Dialect(operations, 'mpi'));
- this._dialects.set('pdl_interp', new _.PDLInterpDialect(operations));
- this._dialects.set('standalone', new _.Dialect(operations, 'standalone'));
- this._dialects.set('custom', new _.Dialect(operations, 'custom'));
- this._dialects.set('layer', new _.Dialect(operations, 'layer'));
- this._dialects.set('foo', new _.Dialect(operations, 'foo'));
- this._dialects.set('some', new _.Dialect(operations, 'some'));
- this._dialects.set('ts', new _.Dialect(operations, 'ts'));
- this._dialects.set('tf_mlrt', new _.Dialect(operations, 'tf_mlrt'));
- this._dialects.set('io_parameters', new _.IOParametersDialect(operations));
- this._dialects.set('pcf', new _.PCFDialect(operations));
- this._dialects.set('linalgx', new _.Dialect(operations, 'linalgx'));
- this._dialects.set('xsmm', new _.XSMMDialect(operations));
- this._dialects.set('sdy', new _.SdyDialect(operations));
- this._dialects.set('mpmd', new _.MPMDDialect(operations));
- this._dialects.set('tfg', new _.TFGDialect(operations));
- this._dialects.set('vt', new _.Dialect(operations, 'vt'));
- this._dialects.set('testd', new _.Dialect(operations, 'testd'));
- this._dialects.set('cmath', new _.Dialect(operations, 'cmath'));
- this._dialects.set('bytecode', new _.Dialect(operations, 'bytecode'));
- this._dialects.set('test_irdl_to_cpp', new _.Dialect(operations, 'test_irdl_to_cpp'));
- this._dialects.set('iree_unregistered', new _.Dialect(operations, 'iree_unregistered'));
- this._dialects.set('cir', new _.Dialect(operations, 'cir'));
- this._dialects.set('migraphx', new _.Dialect(operations, 'migraphx'));
- this._dialects.set('xla', new _.XlaDialect(operations));
- this._dialects.set('xla_gpu', new _.XlaGpuDialect(operations));
- this._dialects.set('xla_cpu', new _.Dialect(operations, 'xla_cpu'));
- this._dialects.set('xla_framework', new _.Dialect(operations, 'xla_framework'));
- this._dialects.set('ifrt', new _.Dialect(operations, 'ifrt'));
- this._dialects.set('vifrt', new _.Dialect(operations, 'vifrt'));
- this._dialects.set('nir', new _.Dialect(operations, 'nir'));
- this._dialects.set('triton_xla', new _.TritonXlaDialect(operations));
- this._dialects.set('xtile', new _.XTileDialect(operations));
- this._dialects.set('xten_nn', new _.XtenNNDialect(operations));
- this._dialects.set('ensemble', new _.EnsembleDialect(operations));
- this._dialects.set('poly', new _.PolyDialect(operations));
- this._dialects.set('noisy', new _.NoisyDialect(operations));
- this._dialects.set('plugin', new _.Dialect(operations, 'plugin'));
- this._dialects.set('ttcore', new _.Dialect(operations, 'ttcore'));
- this._dialects.set('ttir', new _.Dialect(operations, 'ttir'));
- this._dialects.set('ttnn', new _.Dialect(operations, 'ttnn'));
- this._dialects.set('ttkernel', new _.Dialect(operations, 'ttkernel'));
- this._dialects.set('ttmetal', new _.Dialect(operations, 'ttmetal'));
- }
- getOrLoadDialect(name) {
- return this._dialects.get(name);
- }
- getLoadedDialect(name) {
- return this._dialects.has(name);
- }
- checkDialect(dialect, dialectName, context) {
- if (!dialect) {
- switch (dialectName) {
- case 'bstnnx':
- case 'common':
- case 'dxgml':
- case 'dxgml_pattern':
- case 'gim':
- case 'nir':
- case 'nn':
- case 'torq_hl':
- case 'xir':
- case 'xth':
- throw new mlir.Error(`Undocumented ${context} dialect '${dialectName}'.`);
- default:
- throw new mlir.Error(`Unsupported ${context} dialect '${dialectName}'.`);
- }
- }
- }
- redirect(name) {
- return this._redirect.has(name) ? this._redirect.get(name) : name;
- }
- };
- _.Dialect = class {
- constructor(operations, name) {
- this._name = name;
- this._operations = new Map();
- this._customDirectives = new Map();
- this._customTypes = new Map();
- this._customAttributes = new Map();
- this.registerCustomDirective('DynamicIndexList', this.parseDynamicIndexList.bind(this));
- this.registerCustomDirective('Offsets', this.parseOffsets.bind(this));
- this.registerCustomDirective('SymbolVisibility', this.parseSymbolVisibility.bind(this));
- this.registerCustomDirective('TypeOrAttr', this.parseTypeOrAttr.bind(this));
- this.registerCustomDirective('CopyOpRegion', this.parseCopyOpRegion.bind(this));
- this.registerCustomDirective('SizeAwareType', this.parseSizeAwareType.bind(this));
- this.registerCustomDirective('ResultTypeList', this.parseResultTypeList.bind(this));
- this.registerCustomDirective('CmdParameterGatherOperations', this.parseCmdParameterGatherOperations.bind(this));
- this.registerCustomDirective('AsyncParameterGatherOperations', this.parseAsyncParameterGatherOperations.bind(this));
- this.registerCustomDirective('CmdParameterScatterOperations', this.parseCmdParameterScatterOperations.bind(this));
- this.registerCustomDirective('AsyncParameterScatterOperations', this.parseAsyncParameterScatterOperations.bind(this));
- this.registerCustomAttribute('TypedAttrInterface', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('VM_ConstantIntegerValueAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('Util_AnySerializableAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('DenseElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('I32ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('I64ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('F64ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('IndexElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('AnyI32ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('AnyIntElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('StringElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('RankedF32ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('RankedF64ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('RankedI32ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('RankedI64ElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('SignlessIntElementsAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('AnyAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('UnitAttr', this.parseUnitAttr.bind(this));
- this.registerCustomAttribute('UnitProp', this.parseUnitAttr.bind(this));
- this.registerCustomAttribute('SymbolNameAttr', this.parseSymbolNameAttr.bind(this));
- this.registerCustomAttribute('SymbolRefAttr', this.parseSymbolRefAttr.bind(this));
- this.registerCustomAttribute('FlatSymbolRefAttr', this.parseFlatSymbolRefAttr.bind(this));
- this.registerCustomAttribute('OptionalAttr', this.parseOptionalAttr.bind(this));
- this.registerCustomAttribute('OptionalProp', this.parseOptionalAttr.bind(this));
- this.registerCustomAttribute('DefaultValuedOptionalAttr', this.parseDefaultValuedOptionalAttr.bind(this));
- this.registerCustomAttribute('DefaultValuedAttr', this.parseDefaultValuedAttr.bind(this));
- this.registerCustomAttribute('DefaultValuedEnumAttr', this.parseDefaultValuedAttr.bind(this));
- this.registerCustomAttribute('DefaultValuedProp', this.parseDefaultValuedAttr.bind(this));
- this.registerCustomAttribute('ConfinedAttr', this.parseConfinedAttr.bind(this));
- this.registerCustomAttribute('IntValidAlignment', this.parseConfinedAttr.bind(this));
- this.registerCustomAttribute('TypeAttrOf', this.parseTypeAttrOf.bind(this));
- this.registerCustomAttribute('AnyAttrOf', this.parseAnyAttrOf.bind(this));
- this.registerCustomAttribute('ArrayAttr', this.parseArrayAttr.bind(this));
- this.registerCustomAttribute('TypedArrayAttrBase', this.parseArrayAttr.bind(this));
- this.registerCustomAttribute('DenseI64ArrayAttr', _.DenseI64ArrayAttr.parse);
- this.registerCustomAttribute('I64Attr', (parser) => parser.parseAttribute(new _.IntegerType('i64')));
- this.registerCustomAttribute('I32Attr', (parser) => parser.parseAttribute(new _.IntegerType('i32')));
- this.registerCustomAttribute('I16Attr', (parser) => parser.parseAttribute(new _.IntegerType('i16')));
- this.registerCustomAttribute('I8Attr', (parser) => parser.parseAttribute(new _.IntegerType('i8')));
- this.registerCustomAttribute('I1Attr', (parser) => parser.parseAttribute(new _.IntegerType('i1')));
- this.registerCustomAttribute('SI64Attr', (parser) => parser.parseAttribute(new _.IntegerType('si64')));
- this.registerCustomAttribute('SI32Attr', (parser) => parser.parseAttribute(new _.IntegerType('si32')));
- this.registerCustomAttribute('UI64Attr', (parser) => parser.parseAttribute(new _.IntegerType('ui64')));
- this.registerCustomAttribute('UI32Attr', (parser) => parser.parseAttribute(new _.IntegerType('ui32')));
- this.registerCustomAttribute('IndexAttr', (parser) => parser.parseAttribute(new _.IntegerType('index')));
- this.registerCustomAttribute('I64Prop', (parser) => parser.parseAttribute(new _.IntegerType('i64')));
- this.registerCustomAttribute('I32Prop', (parser) => parser.parseAttribute(new _.IntegerType('i32')));
- this.registerCustomAttribute('AlignmentProp', (parser) => parser.parseAttribute(new _.IntegerType('i64')));
- this.registerCustomAttribute('F64Attr', (parser) => parser.parseAttribute(new _.FloatType('f64')));
- this.registerCustomAttribute('F32Attr', (parser) => parser.parseAttribute(new _.FloatType('f32')));
- this.registerCustomAttribute('F16Attr', (parser) => parser.parseAttribute(new _.FloatType('f16')));
- this.registerCustomAttribute('BF16Attr', (parser) => parser.parseAttribute(new _.FloatType('bf16')));
- this.registerCustomAttribute('StrAttr', (parser) => parser.parseAttribute(new _.PrimitiveType('string')));
- this.registerCustomAttribute('TypedStrAttr', this.parseTypedAttrInterface.bind(this));
- this.registerCustomAttribute('LevelAttr', (parser) => parser.parseAttribute(new _.IntegerType('index')));
- this.registerCustomType('Optional', this.parseOptional.bind(this));
- for (const metadata of operations.get(name) || []) {
- this.registerOperandName(metadata.name, metadata);
- }
- }
- get name() {
- return this._name;
- }
- readType(/* reader */) {
- throw new mlir.Error(`Bytecode dialect '${this._name}' does not support bytecode types.`);
- }
- readAttribute(/* reader */) {
- throw new mlir.Error(`Dialect '${this._name}' does not support bytecode attributes.`);
- }
- parseResource(/* entry */) {
- throw new mlir.Error(`Dialect '${this._name}' does not support resources.`);
- }
- getOperation(opName) {
- const op = this._operations.get(opName);
- if (op && !op.metadata._) {
- if (Array.isArray(op.metadata.operands)) {
- for (const input of op.metadata.operands) {
- if (input && input.type) {
- input.type = _.Constraint.parse(input.type);
- }
- }
- }
- if (Array.isArray(op.metadata.results)) {
- for (const output of op.metadata.results) {
- if (output && output.type) {
- output.type = _.Constraint.parse(output.type);
- }
- }
- }
- if (Array.isArray(op.metadata.attributes)) {
- for (const attribute of op.metadata.attributes) {
- if (attribute && attribute.type) {
- attribute.type = _.Constraint.parse(attribute.type);
- }
- }
- }
- if (Array.isArray(op.metadata.regions)) {
- for (const region of op.metadata.regions) {
- if (region && region.type) {
- region.type = _.Constraint.parse(region.type);
- }
- }
- }
- if (Array.isArray(op.metadata.traits)) {
- for (const trait of op.metadata.traits) {
- if (trait && trait.type) {
- trait.type = _.Constraint.parse(trait.type);
- }
- }
- }
- op.metadata._ = true;
- }
- return op || null;
- }
- registerOperandName(opName, metadata) {
- const opInfo = new _.RegisteredOperationName(this, opName, metadata);
- this._operations.set(opName, opInfo);
- }
- registerCustomDirective(name, parserFn) {
- this._customDirectives.set(name, parserFn);
- }
- registerCustomType(name, parserFn) {
- this._customTypes.set(name, parserFn);
- }
- registerCustomAttribute(name, parserFn) {
- this._customAttributes.set(name, parserFn);
- }
- createBuildableType(constraint) {
- switch (constraint.name) {
- case 'AMDGPU_TDMDescriptorType': return new _.Type('!amdgpu.tdm.descriptor');
- case 'AnyIRModule': return new _.Type('!transform.any_op');
- case 'AnySignlessIntegerOrIndex': return new _.IndexType();
- case 'Async_CoroHandleType': return new _.Type('!async.coro.handle');
- case 'Async_CoroIdType': return new _.Type('!async.coro.id');
- case 'Async_CoroStateType': return new _.Type('!async.coro.state');
- case 'Async_GroupType': return new _.async.GroupType();
- case 'Async_TokenType': return new _.async.TokenType();
- case 'Async_ValueType': return new _.async.ValueType(new _.MemRefType([-1], new _.Type('f32')));
- case 'BF16': return new _.Type('bf16');
- case 'BoolType': return new _.Type('!smt.bool');
- case 'CanonicalLoopInfoType': return new _.Type('!omp.canonical_loop_info');
- case 'ControlType': return new _.Type('!tf_type.control');
- case 'CoreRT_OpHandlerType': return new _.Type('!corert.ophandler');
- case 'CoreRT_StringType': return new _.Type('!corert.string');
- case 'CoreRT_TensorHandleType': return new _.Type('!corert.tensorhandle');
- case 'CUDA_BlasGemmAlgorithm': return new _.Type('!cuda.blas.gemm_algorithm');
- case 'CUDA_BlasHandle': return new _.Type('!cuda.blas.handle');
- case 'CUDA_Device': return new _.Type('!cuda.device');
- case 'CUDA_Event': return new _.Type('!cuda.event');
- case 'CUDA_Function': return new _.Type('!cuda.function');
- case 'CUDA_Module': return new _.Type('!cuda.module');
- case 'CUDA_Stream': return new _.Type('!cuda.stream');
- case 'EmitC_PointerType': return new _.Type('!emitc.ptr<i32>');
- case 'Executor_HostPtr': return new _.Type('!executor.ptr<host>');
- case 'Executor_StrLiteral': return new _.Type('!executor.str_literal');
- case 'Executor_Table': return new _.Type('!executor.table<>');
- case 'F128': return new _.Type('f128');
- case 'F16': return new _.Type('f16');
- case 'F32': return new _.Type('f32');
- case 'F4E2M1FN': return new _.Type('f4E2M1FN');
- case 'F64': return new _.Type('f64');
- case 'F6E2M3FN': return new _.Type('f6E2M3FN');
- case 'F6E3M2FN': return new _.Type('f6E3M2FN');
- case 'F80': return new _.Type('f80');
- case 'F8E3M4': return new _.Type('f8E3M4');
- case 'F8E4M3': return new _.Type('f8E4M3');
- case 'F8E4M3B11FNUZ': return new _.Type('f8E4M3B11FNUZ');
- case 'F8E4M3FN': return new _.Type('f8E4M3FN');
- case 'F8E4M3FNUZ': return new _.Type('f8E4M3FNUZ');
- case 'F8E5M2': return new _.Type('f8E5M2');
- case 'F8E5M2FNUZ': return new _.Type('f8E5M2FNUZ');
- case 'F8E8M0FNU': return new _.Type('f8E8M0FNU');
- case 'FLOW_Channel': return new _.Type('!flow.channel');
- case 'GPU_AsyncToken': return new _.Type('!gpu.async.token');
- case 'GPU_SparseDnTensorHandle': return new _.Type('!gpu.sparse.dntensor_handle');
- case 'GPU_SparseSpGEMMOpHandle': return new _.Type('!gpu.sparse.spgemmop_handle');
- case 'GPU_SparseSpMatHandle': return new _.Type('!gpu.sparse.spmat_handle');
- case 'HostBufferType': return new _.Type('!tfrt_ht.host_buffer');
- case 'I1': return new _.IntegerType('i1');
- case 'I128': return new _.IntegerType('i128');
- case 'I16': return new _.IntegerType('i16');
- case 'I32': return new _.IntegerType('i32');
- case 'I64': return new _.IntegerType('i64');
- case 'I8': return new _.IntegerType('i8');
- case 'Ifrt_ArrayType': return new _.Type('!ifrt.array');
- case 'Ifrt_ControlType': return new _.Type('!ifrt.control');
- case 'Index': return new _.IndexType();
- case 'IntType': return new _.Type('!smt.int');
- case 'IRDL_AttributeType': return new _.Type('!irdl.attribute');
- case 'IRDL_RegionType': return new _.Type('!irdl.region');
- case 'IREE_Input_GlobalRefAttr': return new _.Type('!iree_input.global.ref');
- case 'LLVM_DefaultPointer': return new _.Type('!llvm.ptr');
- case 'LLVM_PointerGeneric': return new _.Type('!llvm.ptr');
- case 'LLVM_PointerGlobal': return new _.Type('!llvm.ptr<1>');
- case 'LLVM_PointerShared': return new _.Type('!llvm.ptr<3>');
- case 'LLVM_PointerSharedCluster': return new _.Type('!llvm.ptr<7>');
- case 'LLVM_PointerTensor': return new _.Type('!llvm.ptr<6>');
- case 'LLVM_PointerTensorMemory': return new _.Type('!llvm.ptr<6>');
- case 'LLVM_TokenType': return new _.Type('!llvm.token');
- case 'MLProgram_TokenType': return new _.Type('!ml_program.token');
- case 'MlrtAsyncHandleType': return new _.Type('!mlrt.async_handle');
- case 'MlrtFutureType': return new _.Type('!mlrt.future');
- case 'MlrtPromiseType': return new _.Type('!mlrt.promise');
- case 'MPI_Comm': return new _.Type('!mpi.comm');
- case 'MPI_Request': return new _.Type('!mpi.request');
- case 'MPI_Retval': return new _.Type('!mpi.retval');
- case 'Noisy_I32': return new _.Type('!noisy.i32');
- case 'NoneType': return new _.Type('none');
- case 'NullPointer': return new _.Type('!iree_codegen.null_pointer');
- case 'NVGPU_DeviceAsyncToken': return new _.Type('!nvgpu.device.async.token');
- case 'NVGPU_MmaSparseSyncMetadataType': return new _.VectorType([2], new _.IntegerType('i16'));
- case 'OMP_MapBoundsType': return new _.Type('!omp.map.bounds');
- case 'OpaqueTensorType': return new _.Type('!tf_type.tensor');
- case 'OpenACC_DataBoundsType': return new _.Type('!acc.data_bounds');
- case 'OpenACC_DeclareTokenType': return new _.Type('!acc.declare_token');
- case 'OpenMP_MapBoundsType': return new _.Type('!omp.map.bounds');
- case 'PDL_Attribute': return new _.pdl.AttributeType();
- case 'PDL_Operation': return new _.pdl.OperationType();
- case 'PDL_OperationType': return new _.pdl.OperationType();
- case 'PDL_RangeOf': return new _.pdl.RangeType(this.createBuildableType(constraint.args[0]));
- case 'PDL_RangeType': return new _.Type('!pdl.range<type>');
- case 'PDL_Type': return new _.pdl.TypeType();
- case 'PDL_TypeType': return new _.pdl.TypeType();
- case 'PDL_Value': return new _.pdl.ValueType();
- case 'PDL_ValueType': return new _.pdl.ValueType();
- case 'Ptr_PtrType': return new _.Type('!ptr.ptr');
- case 'ROCDL_V16BF16Type': return new _.VectorType([16], new _.Type('bf16'));
- case 'ROCDL_V16F16Type': return new _.VectorType([16], new _.Type('f16'));
- case 'ROCDL_V16F32Type': return new _.VectorType([16], new _.Type('f32'));
- case 'ROCDL_V2BF16Type': return new _.VectorType([2], new _.Type('bf16'));
- case 'ROCDL_V2F16Type': return new _.VectorType([2], new _.Type('f16'));
- case 'ROCDL_V2F32Type': return new _.VectorType([2], new _.Type('f32'));
- case 'ROCDL_V2I16Type': return new _.VectorType([2], new _.IntegerType('i16'));
- case 'ROCDL_V2I32Type': return new _.VectorType([2], new _.IntegerType('i32'));
- case 'ROCDL_V32BF16Type': return new _.VectorType([32], new _.Type('bf16'));
- case 'ROCDL_V32F16Type': return new _.VectorType([32], new _.Type('f16'));
- case 'ROCDL_V32F32Type': return new _.VectorType([32], new _.Type('f32'));
- case 'ROCDL_V3I32Type': return new _.VectorType([3], new _.IntegerType('i32'));
- case 'ROCDL_V6I32Type': return new _.VectorType([6], new _.IntegerType('i32'));
- case 'ROCDL_V8BF16Type': return new _.VectorType([8], new _.Type('bf16'));
- case 'ROCDL_V8F16Type': return new _.VectorType([8], new _.Type('f16'));
- case 'ROCDL_V8F32Type': return new _.VectorType([8], new _.Type('f32'));
- case 'Shape_ExtentTensorType': return new _.Type('tensor<?xindex>');
- case 'Shape_ShapeType': return new _.Type('!shape.shape');
- case 'Shape_SizeType': return new _.Type('!shape.size');
- case 'Shape_ValueShapeType': return new _.Type('!shape.value_shape');
- case 'Shape_WitnessType': return new _.Type('!shape.witness');
- case 'StaticShapeTensorOf': return new _.UnrankedTensorType(this.createBuildableType(constraint.args[0][0]));
- case 'SI16': return new _.IntegerType('si16');
- case 'SI32': return new _.IntegerType('si32');
- case 'SI64': return new _.IntegerType('si64');
- case 'SI8': return new _.IntegerType('si8');
- case 'SMTFuncType': return new _.Type('!smt.func<() -> ()>');
- case 'SPIRV_AnyPtr': return new _.Type('!spirv.ptr<i32, StorageBuffer>');
- case 'SPIRV_BFloat16KHR': return new _.FloatType('bf16');
- case 'SPIRV_Bool': return new _.IntegerType('i1');
- case 'SPIRV_Float16': return new _.FloatType('f16');
- case 'SPIRV_Float32': return new _.FloatType('f32');
- case 'SPIRV_Int8': return new _.IntegerType('i8');
- case 'SPIRV_Int16': return new _.IntegerType('i16');
- case 'SPIRV_Int32': return new _.IntegerType('i32');
- case 'SPIRV_Int32Vec4': return new _.VectorType([4], new _.IntegerType('i32'));
- case 'Stream_Channel': return new _.Type('!stream.channel');
- case 'Stream_Dim': return new _.IndexType();
- case 'Stream_Offset': return new _.IndexType();
- case 'Stream_Size': return new _.IndexType();
- case 'TensorRTRuntime_Context': return new _.Type('!trtrt.context');
- case 'TensorRTRuntime_Engine': return new _.Type('!trtrt.engine');
- case 'TensorType': return new _.Type('!tfrt_tensor.tensor');
- case 'TF_MLRT_FutureType': return new _.Type('!tf_mlrt.tensor');
- case 'TF32': return new _.Type('tf32');
- case 'TFAllocatorType': return new _.Type('!tfrt_fallback.tf_allocator');
- case 'TFDeviceType': return new _.Type('!tf_mlrt.device');
- case 'TfeControlType': return new _.Type('!tfe.control');
- case 'TfeTokenType': return new _.Type('!tfe.token');
- case 'TFFramework_JITCallableType': return new _.Type('!tf_framework.jit_callable');
- case 'TFFramework_OpKernelContextType': return new _.Type('!tf_framework.op_kernel_context');
- case 'TFL_Control': return new _.Type('!tfl.control');
- case 'TFL_Quint8': return new _.Type('!quant.uniform<u8:f32, 1.0>');
- case 'TFL_Str': return new _.Type('!tf.string');
- case 'TFR_AttrType': return new _.Type('!tfr.attr');
- case 'TFR_TensorListType': return new _.Type('!tfr.tensor_list');
- case 'TFR_TensorType': return new _.Type('!tfr.tensor');
- case 'TFRT_ChainType': return new _.Type('!tfrt.chain');
- case 'TFRT_DeviceType': return new _.Type('!tfrt.device');
- case 'TFRT_Fallback_TFTensorType': return new _.Type('!tfrt_fallback.tf_tensor');
- case 'TFRT_StringType': return new _.Type('!tfrt.string');
- case 'TFRT_TensorTypeType': return new _.Type('!tfrt.tensor_type');
- case 'TFTensorType': return new _.Type('!tfrt_fallback.tf_tensor');
- case 'Torch_BoolType': return new _.Type('!torch.bool');
- case 'Torch_DeviceType': return new _.Type('!torch.Device');
- case 'Torch_FloatType': return new _.Type('!torch.float');
- case 'Torch_GeneratorType': return new _.Type('!torch.Generator');
- case 'Torch_IntType': return new _.Type('!torch.int');
- case 'Torch_LinearParamsType': return new _.Type('!torch.LinearParams');
- case 'Torch_NoneType': return new _.Type('!torch.none');
- case 'Torch_NumberType': return new _.Type('!torch.number');
- case 'Torch_StringType': return new _.Type('!torch.str');
- case 'Transform_AffineMapParamType': return new _.Type('!transform.affine_map');
- case 'Transform_AnyOpType': return new _.Type('!transform.any_op');
- case 'Transform_AnyParamType': return new _.Type('!transform.any_param');
- case 'Transform_AnyValue': return new _.Type('!transform.any_value');
- case 'Transform_AnyValueType': return new _.Type('!transform.any_value');
- case 'Transform_ParamType': return new _.Type('!transform.param<i64>');
- case 'Transform_TypeParamType': return new _.Type('!transform.type');
- case 'TransformHandleTypeInterface': return new _.Type('!transform.any_op');
- case 'TestTransformTestDialectParamType': return new _.Type('!transform.test_dialect_param');
- case 'TS_FixedRankShape': return new _.OpaqueType('ts', `.fixed_rank_shape.${constraint.args[0].name}`);
- case 'TS_PartialShape': return new _.Type('!ts.partial_shape');
- case 'TS_Shape': return new _.Type('!ts.shape');
- case 'TTG_AsyncToken': return new _.Type('!ttg.async.token');
- case 'UI16': return new _.IntegerType('ui16');
- case 'UI32': return new _.IntegerType('ui32');
- case 'UI64': return new _.IntegerType('ui64');
- case 'UI8': return new _.IntegerType('ui8');
- case 'Util_BufferType': return new _.Type('!util.buffer');
- case 'Util_Offset': return new _.IndexType();
- case 'Util_Size': return new _.IndexType();
- case 'VM_CondValue': return new _.IntegerType('i32');
- case 'VM_RefType': return new _.Type('!vm.ref<?>');
- case 'XeGPU_Nbarrier': return new _.Type('!xegpu.nbarrier');
- case 'XLA_BufferType': return new _.Type('!xla_framework.buffer');
- case 'XLAFramework_BufferType': return new _.Type('!xla_framework.buffer');
- default: return null;
- }
- }
- inferResultTypes(op, vars) {
- const metadata = op.name.getRegisteredInfo()?.metadata;
- if (!metadata?.results) {
- return;
- }
- // If type(results) was used in the format, all result types were explicitly parsed
- if (vars.has('results') && vars.get('results').types.length > 0) {
- if (op.types.length === 0) {
- op.addTypes(vars.get('results').types);
- }
- return;
- }
- const operandNames = metadata.operands?.map((o) => o.name) || [];
- // Build result types in metadata order
- const orderedTypes = [];
- // Track current offset in op.types for custom directive handling
- let customDirectiveOffset = 0;
- // Iterate over each result in metadata order
- for (const resultInfo of metadata.results) {
- // If this result was parsed from assembly format, use those types
- if (vars.has(resultInfo.name) && vars.get(resultInfo.name).types.length > 0) {
- const varTypes = vars.get(resultInfo.name).types;
- orderedTypes.push(...varTypes);
- customDirectiveOffset += varTypes.length;
- continue;
- }
- const typeName = typeof resultInfo.type === 'string' ? resultInfo.type : resultInfo.type?.name;
- const isVariadicOrOptional = typeName === 'Variadic' || typeName === 'Optional';
- // For Variadic/Optional not in vars, check if custom directive added types to op.types
- if (isVariadicOrOptional && op.types.length > customDirectiveOffset) {
- // Count how many types belong to this variadic result
- // (all types from customDirectiveOffset to end, minus types for later results in vars)
- let typesForLaterResults = 0;
- for (let i = metadata.results.indexOf(resultInfo) + 1; i < metadata.results.length; i++) {
- const laterResult = metadata.results[i];
- if (vars.has(laterResult.name) && vars.get(laterResult.name).types.length > 0) {
- typesForLaterResults += vars.get(laterResult.name).types.length;
- }
- }
- const variadicCount = op.types.length - customDirectiveOffset - typesForLaterResults;
- if (variadicCount > 0) {
- for (let i = 0; i < variadicCount; i++) {
- orderedTypes.push(op.types[customDirectiveOffset + i]);
- }
- customDirectiveOffset += variadicCount;
- continue;
- }
- }
- // Try to resolve from traits
- let resolved = false;
- if (metadata.traits) {
- for (const trait of metadata.traits) {
- if (resolved) {
- break;
- }
- // Handle both string and object trait types
- const traitName = trait.type?.name;
- const traitArgs = trait.type?.args;
- // SameOperandsAndResultType: result type = first operand type
- if (traitName === 'SameOperandsAndResultType') {
- if (op.operands.length > 0 && op.operands[0].type) {
- orderedTypes.push(op.operands[0].type);
- resolved = true;
- }
- }
- // AllTypesMatch: result type = matched operand/result/attribute type
- if (traitName === 'AllTypesMatch') {
- const names = traitArgs?.[0];
- if (Array.isArray(names) && (names.includes(resultInfo.name) || (resultInfo.name === 'result' && names.includes('result')))) {
- const sourceTypes = [];
- for (const argName of names) {
- if (argName === resultInfo.name) {
- continue;
- }
- // Check if source is in vars (another parsed result)
- if (vars.get(argName) && vars.get(argName).types.length > 0) {
- sourceTypes.push(...vars.get(argName).types);
- break;
- }
- // Check if source is an operand
- const operandMetaIdx = operandNames.indexOf(argName);
- if (operandMetaIdx >= 0) {
- const operandMeta = metadata.operands[operandMetaIdx];
- const operandTypeObj = operandMeta?.type;
- const operandTypeName = typeof operandTypeObj === 'string' ? operandTypeObj : (operandTypeObj?.name || '');
- const isVariadic = operandTypeName === 'Variadic' || operandTypeName.startsWith('Variadic<');
- if (isVariadic) {
- for (let j = 0; j < op.operands.length; j++) {
- if (op.operands[j].type) {
- sourceTypes.push(op.operands[j].type);
- }
- }
- } else if (operandMetaIdx < op.operands.length && op.operands[operandMetaIdx].type) {
- sourceTypes.push(op.operands[operandMetaIdx].type);
- }
- if (sourceTypes.length > 0) {
- break;
- }
- }
- // Check if source is an attribute with a type
- if (metadata.attributes) {
- const attrMeta = metadata.attributes.find((a) => a.name === argName);
- if (attrMeta) {
- const attr = op.attributes.get(argName);
- if (attr && attr.type) {
- sourceTypes.push(attr.type);
- break;
- }
- }
- }
- }
- if (sourceTypes.length > 0) {
- orderedTypes.push(...sourceTypes);
- resolved = true;
- }
- }
- }
- // TypesMatchWith: result type = transform(source type)
- if (traitName === 'TypesMatchWith') {
- const [lhsArg, rhsArg, transformer] = traitArgs || [];
- let resultArg = null;
- let sourceArg = null;
- if ((rhsArg === resultInfo.name || (resultInfo.name === 'result' && rhsArg === 'result'))) {
- resultArg = rhsArg;
- sourceArg = lhsArg;
- } else if ((lhsArg === resultInfo.name || (resultInfo.name === 'result' && lhsArg === 'result'))) {
- resultArg = lhsArg;
- sourceArg = rhsArg;
- }
- if (resultArg && sourceArg) {
- const sourceTypes = [];
- // Check if source is in vars (parsed result or operand type from type($var))
- if (vars.has(sourceArg) && vars.get(sourceArg).types.length > 0) {
- sourceTypes.push(...vars.get(sourceArg).types);
- }
- // Check if source is an operand - fall back to operand's value type
- if (sourceTypes.length === 0) {
- const operands = metadata.operands || [];
- let actualOperandIdx = 0;
- for (let i = 0; i < operands.length; i++) {
- const operandTypeObj = operands[i].type;
- const operandType = typeof operandTypeObj === 'string' ? operandTypeObj : (operandTypeObj?.name || '');
- const isEnum = /\{[^}]+\}/.test(operandType) || operandType.startsWith('AtomicBinOp') || operandType.startsWith('AtomicOrdering') || operandType.startsWith('LLVM_AtomicOrdering') || operandType.includes('Enum');
- if (operands[i].name === sourceArg) {
- const isVariadicSource = operandType === 'Variadic';
- if (isVariadicSource) {
- for (let j = actualOperandIdx; j < op.operands.length; j++) {
- if (op.operands[j].type) {
- sourceTypes.push(op.operands[j].type);
- }
- }
- } else if (actualOperandIdx < op.operands.length && op.operands[actualOperandIdx].type) {
- sourceTypes.push(op.operands[actualOperandIdx].type);
- }
- break;
- }
- if (!isEnum) {
- actualOperandIdx++;
- }
- }
- }
- if (sourceTypes.length > 0) {
- for (const sourceType of sourceTypes) {
- let resultType = sourceType;
- if (transformer === '::getI1SameShape($_self)') {
- if (sourceType instanceof _.VectorType) {
- resultType = new _.VectorType(sourceType.dimensions, new _.IntegerType('i1'), sourceType.scalableDims);
- } else if (sourceType instanceof mlir.TensorType) {
- resultType = new mlir.TensorType(sourceType.dimensions, new _.IntegerType('i1'));
- } else {
- resultType = new _.IntegerType('i1');
- }
- } else if (transformer && transformer.includes('.getElementType()')) {
- if (sourceType && sourceType.elementType) {
- resultType = sourceType.elementType;
- } else {
- throw new mlir.Error(`Cannot get element type.`);
- }
- } else if (transformer && transformer.includes('getPointeeType')) {
- if (sourceType && sourceType.pointeeType) {
- resultType = sourceType.pointeeType;
- } else {
- throw new mlir.Error(`Cannot get pointee type.`);
- }
- } else if (transformer && transformer.includes('getValAndBoolStructType')) {
- const typeStr = sourceType.toString ? sourceType.toString() : String(sourceType);
- resultType = new _.Type(`!llvm.struct<(${typeStr}, i1)>`);
- } else if (transformer && transformer.includes('.getResults()')) {
- // Extract results from a FunctionType
- if (sourceType instanceof _.FunctionType && sourceType.results) {
- orderedTypes.push(...sourceType.results);
- resolved = true;
- continue; // Already added results, skip single push below
- }
- }
- if (resultType) {
- orderedTypes.push(resultType);
- resolved = true;
- }
- }
- }
- }
- }
- // FirstAttrDerivedResultType: result type = first attribute's type
- if (traitName === 'FirstAttrDerivedResultType') {
- const firstAttr = metadata.attributes?.[0];
- if (firstAttr) {
- const attr = op.attributes.get(firstAttr.name);
- if (attr && attr.type) {
- orderedTypes.push(attr.type);
- resolved = true;
- }
- }
- }
- }
- }
- // For single variadic result matching single variadic operand with same element type, infer from operands
- if (!resolved && resultInfo.type?.name === 'Variadic' &&
- metadata.results?.length === 1 && metadata.operands?.length === 1) {
- const resultElementType = resultInfo.type.args?.[0]?.name;
- const operandInfo = metadata.operands[0];
- if (operandInfo.type?.name === 'Variadic' &&
- operandInfo.type.args?.[0]?.name === resultElementType) {
- for (const operand of op.operands) {
- if (operand.type) {
- orderedTypes.push(operand.type);
- }
- }
- resolved = orderedTypes.length > 0;
- }
- }
- // Fallback: try buildable type (skip for variadic/optional)
- if (!resolved && resultInfo.type && !isVariadicOrOptional) {
- const type = this.createBuildableType(resultInfo.type);
- if (type) {
- orderedTypes.push(type);
- }
- }
- }
- // Replace op.types with correctly ordered types
- if (orderedTypes.length > 0) {
- op.types.length = 0;
- op.addTypes(orderedTypes);
- }
- }
- parseDirective(directive, parser, op, opInfo, directives, i, vars) {
- const isVariadic = (type) => {
- if (type.name === 'Variadic' || type.name === 'VariadicOfVariadic') {
- return true;
- }
- if (Array.isArray(type.args) && type.args.length > 0) {
- return isVariadic(type.args[0]);
- }
- return false;
- };
- const isVariadicOfVariadic = (type) => {
- if (type.name === 'VariadicOfVariadic') {
- return true;
- }
- if (Array.isArray(type.args) && type.args.length > 0) {
- return isVariadicOfVariadic(type.args[0]);
- }
- return false;
- };
- const isOptional = (type) => {
- if (type.name === 'Optional') {
- return true;
- }
- if (Array.isArray(type.args) && type.args.length > 0) {
- return isOptional(type.args[0]);
- }
- return false;
- };
- switch (directive.type) {
- case 'whitespace':
- // Skip whitespace directives - they're just formatting hints
- break;
- case 'literal':
- // Make '[' optional when followed by a variadic operand with buildable type
- if (directive.value === '[' && parser.getToken().isNot(_.Token.l_square)) {
- const nextDir = directives[i + 1];
- if (nextDir && nextDir.type === 'operand_ref') {
- const operandMeta = opInfo.metadata?.operands?.find((o) => o.name === nextDir.name);
- if (operandMeta && isVariadic(operandMeta.type)) {
- // Skip '[' and mark to skip ']' later
- vars.set('_skipClosingBracket', true);
- break;
- }
- }
- }
- if (directive.value === ']' && vars.get('_skipClosingBracket')) {
- vars.delete('_skipClosingBracket');
- break;
- }
- parser.consumeToken();
- break;
- case 'region_ref': {
- const regionMeta = opInfo.metadata && opInfo.metadata.regions && opInfo.metadata.regions.find((r) => r.name === directive.name);
- const isVariadicRegion = regionMeta && regionMeta.type && regionMeta.type.name === 'VariadicRegion';
- const isIsolated = op.name.getRegisteredInfo().hasTrait('IsolatedFromAbove');
- if (isVariadicRegion) {
- if (parser.getToken().is(_.Token.l_brace)) {
- do {
- const region = op.addRegion();
- parser.parseRegion(region, undefined, isIsolated);
- } while (parser.parseOptionalComma() && parser.getToken().is(_.Token.l_brace));
- }
- } else {
- const region = op.addRegion();
- parser.parseRegion(region, undefined, isIsolated);
- }
- break;
- }
- case 'successor_ref': {
- if (!op.successors) {
- op.successors = [];
- }
- // Check if this successor is variadic from metadata or context
- const refName = directive.name;
- let isVariadicSuccessor = false;
- if (opInfo.metadata && opInfo.metadata.successors) {
- const successorMeta = opInfo.metadata.successors.find((s) => s.name === refName);
- if (successorMeta && successorMeta.type) {
- // Check for VariadicSuccessor type
- const typeStr = typeof successorMeta.type === 'string' ? successorMeta.type : successorMeta.type.name;
- isVariadicSuccessor = typeStr && typeStr.startsWith('VariadicSuccessor');
- }
- }
- // Also check context: if next directive is ')' literal, we're inside parentheses
- const nextDir = i + 1 < directives.length ? directives[i + 1] : null;
- const isVariadicContext = isVariadicSuccessor || (nextDir && nextDir.type === 'literal' && nextDir.value === ')');
- // Check if next directive is a custom directive that handles the parentheses (like ResultTypeList)
- const nextDirHandlesParens = nextDir && nextDir.type === 'custom' &&
- (nextDir.parser === 'ResultTypeList' || nextDir.parser === 'TypeList');
- const parseOneSuccessor = () => {
- const successor = {};
- successor.label = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- // Don't consume '(' if next directive handles it (e.g., custom<ResultTypeList>)
- if (!nextDirHandlesParens && parser.parseOptionalLParen()) {
- successor.arguments = [];
- while (parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- successor.arguments.push(parser.parseOperand());
- parser.parseOptionalComma();
- } else {
- break;
- }
- }
- parser.resolveOperands(successor.arguments, parser.parseOptionalColonTypeList());
- parser.parseOptionalRParen();
- }
- op.successors.push(successor);
- };
- if (isVariadicContext) {
- // Variadic successors: parse 0 or more successors
- while (parser.getToken().is(_.Token.caret_identifier)) {
- parseOneSuccessor();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- } else {
- parseOneSuccessor();
- }
- break;
- }
- case 'attribute_ref': {
- const refName = directive.name;
- if (op.attributes.has(refName)) {
- break;
- }
- const attrInfo = opInfo.metadata && opInfo.metadata.attributes && opInfo.metadata.attributes.find((attr) => attr.name === refName);
- const attrType = attrInfo ? attrInfo.type : null;
- let attrValue = null;
- // Pass type to suppress : type suffix parsing (it's a separate directive in assembly format)
- if (attrType && attrType !== 'Attribute') {
- attrValue = this.parseCustomAttributeWithFallback(parser, attrType);
- } else {
- attrValue = parser.parseAttribute(attrType || 'Attribute');
- }
- if (attrValue) {
- op.addAttribute(refName, attrValue);
- }
- break;
- }
- case 'operand_ref': {
- const name = directive.name;
- const input = opInfo.metadata?.operands?.find((inp) => inp.name === name);
- const isVariadicOp = input ? isVariadic(input.type) : false;
- const isVariadicOfVariadicOp = input ? isVariadicOfVariadic(input.type) : false;
- const isOptionalOp = input ? isOptional(input.type) : false;
- // Check for buildable types using createType
- let buildableType = null;
- if (isVariadicOp && input?.type?.args?.[0]) {
- buildableType = this.createBuildableType(input.type.args[0]);
- } else if (input?.type) {
- buildableType = this.createBuildableType(input.type);
- }
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- const entry = vars.get(name);
- if (isVariadicOfVariadicOp) {
- // Parse grouped operands: (op, op), (), (op)
- do {
- if (!parser.parseOptionalLParen()) {
- break;
- }
- while (parser.getToken().is(_.Token.percent_identifier)) {
- entry.operands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- } while (parser.parseOptionalComma());
- } else if (isVariadicOp) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.r_square) && parser.getToken().isNot(_.Token.r_brace) && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace) && parser.getToken().isNot(_.Token.equal)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- entry.operands.push(parser.parseOperand());
- if (buildableType) {
- entry.types.push(buildableType);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else if (buildableType && (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus))) {
- // Handle integer literals for buildable integer types (e.g., I32)
- let value = '';
- if (parser.consumeIf(_.Token.minus)) {
- value = '-';
- }
- value += parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- entry.operands.push({ name: value, literal: true });
- if (buildableType) {
- entry.types.push(buildableType);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else {
- break;
- }
- }
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- entry.operands.push(parser.parseOperand());
- if (buildableType) {
- entry.types.push(buildableType);
- }
- } else if (buildableType && (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus))) {
- // Handle integer literals for buildable integer types (e.g., I32)
- let value = '';
- if (parser.consumeIf(_.Token.minus)) {
- value = '-';
- }
- value += parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- entry.operands.push({ name: value, literal: true });
- entry.types.push(buildableType);
- } else if (parser.getToken().is(_.Token.l_brace)) {
- // Check if this is a region, not an operand
- const isActualOperand = opInfo.metadata?.operands?.some((inp) => inp.name === name);
- if (!isActualOperand) {
- const regionMeta = opInfo.metadata?.regions?.find((r) => r.name === name);
- const isVariadicRegion = regionMeta?.type?.name === 'VariadicRegion';
- const isIsolated = op.name.hasTrait('IsolatedFromAbove');
- if (isVariadicRegion) {
- do {
- parser.parseRegion(op.addRegion(), undefined, isIsolated);
- } while (parser.parseOptionalComma() && parser.getToken().is(_.Token.l_brace));
- } else {
- parser.parseRegion(op.addRegion(), undefined, isIsolated);
- }
- }
- } else if (parser.getToken().is(_.Token.at_identifier)) {
- const attrVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- op.addAttribute(name, attrVal);
- } else if (!isOptionalOp && parser.getToken().is(_.Token.bare_identifier)) {
- // Check if this is an enum type that should be an attribute
- // Enum types have a values array after type parsing
- const inputType = input?.type;
- if (inputType && Array.isArray(inputType.values)) {
- const attrVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- op.addAttribute(name, attrVal);
- } else {
- throw new mlir.Error(`Variable '${name}' has incorrect metadata (expected attribute, got operand).`);
- }
- } else if (!isOptionalOp && parser.getToken().is(_.Token.integer)) {
- const attrVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- op.addAttribute(name, attrVal);
- } else if (!isOptionalOp && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.r_square) && parser.getToken().isNot(_.Token.r_brace) && parser.getToken().isNot(_.Token.eof)) {
- const attr = parser.parseAttribute();
- if (attr) {
- op.addAttribute(name, attr);
- }
- }
- break;
- }
- case 'operands': {
- if (!vars.has('operands')) {
- vars.set('operands', { operands: [], types: [] });
- }
- const operandsList = parser.parseOperandList();
- vars.get('operands').operands.push(...operandsList);
- break;
- }
- case 'results': {
- const args = parser.parseArgumentList('none', true);
- const types = args.map((a) => a.type).filter((t) => t);
- op.addTypes(types);
- break;
- }
- case 'type':
- case 'qualified': {
- if (!directive.args || directive.args.length === 0) {
- // Bare type directive - parse types for operands
- const types = parser.parseTypeListNoParens();
- if (vars.has('operands')) {
- vars.get('operands').types.push(...types);
- }
- break;
- }
- const arg = directive.args[0] === 'type' && directive.args.length > 1 ? directive.args[1] : directive.args[0];
- // Handle qualified($attr) - attribute reference
- if (directive.type === 'qualified' && arg.startsWith('$') && !arg.startsWith('$results') && !arg.startsWith('$operands')) {
- if (!arg.startsWith('type($')) {
- const attrName = arg.substring(1);
- const attr = parser.parseAttribute();
- if (attr) {
- // Store full attribute object to preserve .type for FirstAttrDerivedResultType
- op.addAttribute(attrName, attr);
- }
- break;
- }
- }
- // Extract name from $name or type($name) or type(operands) or type(results)
- let name = null;
- if (arg.startsWith('type($') && arg.endsWith(')')) {
- name = arg.substring(6, arg.length - 1);
- } else if (arg.startsWith('type(') && arg.endsWith(')')) {
- // Handle type(operands) or type(results)
- name = arg.substring(5, arg.length - 1);
- } else if (arg.startsWith('$')) {
- name = arg.substring(1);
- } else if (arg === 'results' || arg === 'operands') {
- name = arg;
- }
- if (!name) {
- break;
- }
- // Check if it's a result or operand
- const resultMeta = opInfo.metadata?.results?.find((r) => r.name === name);
- const operandMeta = opInfo.metadata?.operands?.find((o) => o.name === name);
- const isResult = Boolean(resultMeta) && !operandMeta;
- let isVariadicType = false;
- if (resultMeta) {
- isVariadicType = isVariadic(resultMeta.type);
- } else if (operandMeta) {
- isVariadicType = isVariadic(operandMeta.type);
- }
- const isVariadicOfVariadicType = operandMeta ? isVariadicOfVariadic(operandMeta.type) : false;
- let isOptionalType = false;
- if (operandMeta) {
- isOptionalType = isOptional(operandMeta.type);
- } else if (resultMeta) {
- isOptionalType = isOptional(resultMeta.type);
- }
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- const entry = vars.get(name);
- if (isResult || name === 'results') {
- // Result type - add to op.types and track in entry.types
- if (isVariadicType || name === 'results') {
- const types = parser.parseTypeListNoParens();
- entry.types.push(...types);
- op.addTypes(types);
- } else if (isOptionalType && op.types.length === 0) {
- const type = parser.parseOptionalType();
- if (type) {
- entry.types.push(type);
- op.addTypes([type]);
- }
- } else {
- const type = this.parseCustomTypeWithFallback(parser, resultMeta?.type);
- entry.types.push(type);
- op.addTypes([type]);
- }
- } else if (isVariadicOfVariadicType) {
- // Parse grouped types: (type, type), (), (type)
- do {
- if (!parser.parseOptionalLParen()) {
- break;
- }
- if (parser.getToken().isNot(_.Token.r_paren)) {
- entry.types.push(...parser.parseTypeListNoParens());
- }
- parser.parseRParen();
- } while (parser.parseOptionalComma());
- } else if (isVariadicType || name === 'operands') {
- // Variadic operand - parse type list
- entry.types.push(...parser.parseTypeListNoParens());
- } else if (entry.operands.length > 0) {
- // Single operand - parse one type per operand
- // Clear any previously inferred (buildable) types since explicit type is parsed
- const type = this.parseCustomTypeWithFallback(parser, operandMeta?.type);
- entry.types = [];
- for (let j = 0; j < entry.operands.length; j++) {
- entry.types.push(type);
- }
- } else if (isOptionalType) {
- // Optional operand - parse type if present
- const type = parser.parseOptionalType();
- if (type) {
- entry.types.push(type);
- }
- } else {
- // No operands collected yet, just parse and store type
- // Use custom type parser if available for the operand/result constraint
- const typeConstraint = operandMeta?.type || resultMeta?.type;
- const type = this.parseCustomTypeWithFallback(parser, typeConstraint);
- entry.types.push(type);
- }
- break;
- }
- case 'attr_dict_with_keyword':
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(op.attributes);
- }
- break;
- case 'attr_dict':
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(op.attributes);
- }
- break;
- case 'prop_dict':
- if (parser.parseOptionalLess()) {
- op.propertiesAttr = parser.parseAttribute();
- parser.parseGreater();
- }
- break;
- case 'regions': {
- const isIsolated = op.name.hasTrait('IsolatedFromAbove');
- while (parser.getToken().is(_.Token.l_brace)) {
- const region = op.addRegion();
- parser.parseRegion(region, undefined, isIsolated);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- break;
- }
- case 'successors': {
- op.successors = op.successors || [];
- if (parser.getToken().is(_.Token.caret_identifier)) {
- op.successors.push({ label: parser.getToken().getSpelling().str() });
- parser.consumeToken(_.Token.caret_identifier);
- while (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.caret_identifier)) {
- op.successors.push({ label: parser.getToken().getSpelling().str() });
- parser.consumeToken(_.Token.caret_identifier);
- } else {
- break;
- }
- }
- }
- break;
- }
- case 'functional_type': {
- const type = parser.parseFunctionType();
- if (!(type instanceof _.FunctionType)) {
- throw new mlir.Error('Invalid functional-type function type.');
- }
- // Distribute input types to operands in metadata order
- let typeIndex = 0;
- for (const operandMeta of opInfo.metadata?.operands || []) {
- if (vars.has(operandMeta.name)) {
- const entry = vars.get(operandMeta.name);
- for (let j = 0; j < entry.operands.length && typeIndex < type.inputs.length; j++) {
- entry.types.push(type.inputs[typeIndex]);
- typeIndex++;
- }
- }
- }
- op.addTypes(type.results);
- // Track result types - second arg is typically the result (e.g., $outputs or results)
- if (directive.args && directive.args.length > 1) {
- const resultArg = directive.args[1];
- if (resultArg) {
- // Handle both $name and name (without $)
- const resultName = resultArg.startsWith('$') ? resultArg.substring(1) : resultArg;
- if (resultName === 'results') {
- // 'results' is a keyword meaning all results - create vars entry
- if (!vars.has('results')) {
- vars.set('results', { types: [] });
- }
- vars.get('results').types.push(...type.results);
- } else if (vars.has(resultName)) {
- vars.get(resultName).types.push(...type.results);
- }
- }
- }
- break;
- }
- case 'custom': {
- const fn = this._customDirectives.get(directive.parser);
- if (!fn) {
- throw new mlir.Error(`Custom directive parser '${directive.parser}' not implemented.`);
- }
- // Parse args and pass resolved arrays
- // Always pass parser and op first, then resolved args
- const callArgs = [parser, op];
- for (const arg of directive.args || []) {
- if (arg.startsWith('ref(type($') && arg.endsWith('))')) {
- // ref(type($name)) - reference to the type of an operand
- // 'ref(type($' is 10 characters, so slice from index 10 to skip -2 for '))'
- const name = arg.slice(10, -2);
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- callArgs.push(vars.get(name).types);
- } else if (arg.startsWith('ref($') && arg.endsWith(')')) {
- const name = arg.slice(5, -1);
- // Check if this is an attribute reference
- const isAttribute = opInfo.metadata?.attributes?.some((a) => a.name === name);
- if (isAttribute) {
- // Get attribute value from op.attributes
- const attrValue = op.attributes.get(name);
- callArgs.push(attrValue);
- } else {
- // Operand reference - get from ctx
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- callArgs.push(vars.get(name).operands);
- }
- } else if (arg.startsWith('type($') && arg.endsWith(')')) {
- const name = arg.slice(6, -1);
- // Always use vars entry for the specific name (result or operand)
- // This ensures Optional results like asyncToken are tracked separately
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- callArgs.push(vars.get(name).types);
- } else if (arg.startsWith('$')) {
- const name = arg.slice(1);
- // Could be operand ref or attribute name
- const isOperand = opInfo.metadata?.operands?.some((o) => o.name === name);
- if (isOperand) {
- if (!vars.has(name)) {
- vars.set(name, { operands: [], types: [] });
- }
- callArgs.push(vars.get(name).operands);
- } else {
- // Pass attribute name for custom directive to handle
- callArgs.push(name);
- }
- } else {
- // Warn if a $-prefixed arg wasn't resolved - indicates missing metadata
- if (typeof arg === 'string' && arg.startsWith('$')) {
- throw new mlir.Error(`Custom directive '${directive.parser}' received unresolved arg '${arg}' for op '${opInfo.metadata.name}'. Check metadata for missing operand/attribute definition.`);
- }
- callArgs.push(arg);
- }
- }
- fn(...callArgs);
- break;
- }
- case 'oilist': {
- const clauses = directive.content.split('|').map((c) => c.trim());
- const parsedClauses = [];
- for (const clauseStr of clauses) {
- const clauseParser = new _.AssemblyFormatParser({ ...opInfo.metadata, assemblyFormat: clauseStr });
- const elements = clauseParser.parse();
- parsedClauses.push({ elements, parsed: false, clauseStr });
- }
- // Helper to check if a clause's variables are used by later custom directives
- const isHandledByCustomDirective = (clauseStr) => {
- const varMatches = clauseStr.matchAll(/\$(\w+)/g);
- const clauseVars = [...varMatches].map((m) => m[1]);
- if (clauseVars.length === 0) {
- return false;
- }
- for (let j = i + 1; j < directives.length; j++) {
- const laterDir = directives[j];
- if (laterDir.type === 'custom' && laterDir.args && Array.isArray(laterDir.args)) {
- const customVarNames = [];
- for (const arg of laterDir.args) {
- const argVarMatches = arg.matchAll(/\$(\w+)/g);
- for (const match of argVarMatches) {
- customVarNames.push(match[1]);
- }
- }
- if (clauseVars.some((v) => customVarNames.includes(v))) {
- return true;
- }
- }
- }
- return false;
- };
- let progress = true;
- while (progress) {
- progress = false;
- for (const clause of parsedClauses) {
- if (clause.parsed) {
- continue;
- }
- if (clause.elements.length === 0) {
- continue;
- }
- if (isHandledByCustomDirective(clause.clauseStr)) {
- clause.parsed = true;
- continue;
- }
- const [firstElem] = clause.elements;
- let matches = false;
- if (firstElem.type === 'literal') {
- if (firstElem.value.length === 1 && /[(){}[\],:<>=]/.test(firstElem.value)) {
- matches = parser.getToken().is(firstElem.value);
- } else {
- matches = parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === firstElem.value || parser.getToken().is(`kw_${firstElem.value}`);
- }
- }
- if (matches) {
- for (const elem of clause.elements) {
- this.parseDirective(elem, parser, op, opInfo, directives, i, vars);
- }
- clause.parsed = true;
- progress = true;
- }
- }
- }
- break;
- }
- case 'optional_group': {
- let shouldParse = false;
- const firstElem = directive.elements.find((elem) => elem.type !== 'whitespace');
- if (firstElem) {
- if (firstElem.type === 'literal') {
- if (firstElem.value.length === 1 && /[(){}[\],:<>=?]/.test(firstElem.value)) {
- shouldParse = parser.getToken().is(firstElem.value);
- } else if (firstElem.value === '->') {
- shouldParse = parser.getToken().is(_.Token.arrow);
- } else if (firstElem.value === '...') {
- shouldParse = parser.getToken().is(_.Token.ellipsis);
- } else {
- shouldParse = parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === firstElem.value || parser.getToken().is(`kw_${firstElem.value}`);
- }
- } else if (firstElem.type === 'attribute_ref') {
- const attrInfo = opInfo.metadata && opInfo.metadata.attributes && opInfo.metadata.attributes.find((attr) => attr.name === firstElem.name);
- const attrType = attrInfo ? attrInfo.type : null;
- // Check if attribute type is an array (TypedArrayAttrBase, ArrayAttr, etc.)
- // Handle both string types ("OptionalAttr<StrAttr>") and object types ({name: "OptionalAttr", args: [...]})
- let elementType = attrType;
- const getTypeName = (t) => typeof t === 'string' ? t : t?.name;
- const typeContains = (t, pattern) => {
- const name = typeof t === 'string' ? t : t?.name;
- return name && name.includes(pattern);
- };
- // Unwrap wrapper types
- if (typeContains(elementType, 'OptionalAttr')) {
- elementType = typeof elementType === 'string' ? elementType.replace(/OptionalAttr<(.+)>/, '$1') : elementType.args?.[0];
- }
- if (typeContains(elementType, 'DefaultValuedAttr') || typeContains(elementType, 'DefaultValuedStrAttr')) {
- elementType = typeof elementType === 'string' ? elementType.replace(/DefaultValuedAttr<(.+?)(?:,.*?)?>/, '$1') : elementType.args?.[0];
- }
- if (typeContains(elementType, 'ConfinedAttr')) {
- elementType = typeof elementType === 'string' ? elementType.replace(/ConfinedAttr<(.+?)(?:,.*?)?>/, '$1') : elementType.args?.[0];
- }
- const typeName = getTypeName(elementType);
- let shouldTryParse = false;
- if (typeContains(elementType, 'TypedArrayAttrBase') || typeContains(elementType, 'ArrayAttr')) {
- shouldTryParse = parser.getToken().is(_.Token.l_square);
- } else if (typeName && /^[SU]?I\d+Attr$|^IntegerAttr$|^IndexAttr$/.test(typeName)) {
- shouldTryParse = parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus);
- } else if (typeContains(elementType, 'ElementsAttr')) {
- // ElementsAttr values start with specific keywords: dense, sparse, array, dense_resource
- shouldTryParse = parser.getToken().is(_.Token.kw_dense) || parser.getToken().is(_.Token.kw_sparse) ||
- parser.getToken().is(_.Token.kw_array) || parser.getToken().is(_.Token.kw_dense_resource);
- } else if (typeName === 'StrAttr' || typeName === 'StringAttr') {
- shouldTryParse = parser.getToken().is(_.Token.string);
- } else {
- shouldTryParse = parser.getToken().is(_.Token.bare_identifier) || parser.getToken().isKeyword() || parser.getToken().is(_.Token.hash_identifier) || parser.getToken().is(_.Token.at_identifier) || parser.getToken().is(_.Token.string) || parser.getToken().is(_.Token.l_square) || parser.getToken().is(_.Token.integer);
- }
- if (shouldTryParse) {
- let result = null;
- if (attrType && attrType !== 'Attribute') {
- result = this.parseCustomAttributeWithFallback(parser, attrType);
- } else {
- result = parser.parseOptionalAttribute(attrType || 'Attribute');
- }
- if (result !== null) {
- op.addAttribute(firstElem.name, result);
- shouldParse = true;
- }
- }
- } else if (firstElem.type === 'successor_ref') {
- shouldParse = parser.getToken().is(_.Token.caret_identifier);
- } else if (firstElem.type === 'region_ref') {
- shouldParse = parser.getToken().is(_.Token.l_brace);
- } else if (firstElem.type === 'operand_ref') {
- let isKeywordInput = false;
- if (opInfo.metadata && opInfo.metadata.operands) {
- const inputInfo = opInfo.metadata.operands.find((inp) => inp.name === firstElem.name);
- if (inputInfo) {
- const inputType = inputInfo.type;
- if (typeof inputType === 'string' &&
- (inputType.includes('Prop') || inputType.endsWith('Predicate') ||
- inputType.includes('Flags') || inputType.includes('Enum'))) {
- isKeywordInput = true;
- }
- }
- }
- if (isKeywordInput) {
- shouldParse = parser.getToken().is(_.Token.bare_identifier);
- } else {
- shouldParse = parser.getToken().is(_.Token.percent_identifier);
- }
- } else if (firstElem.type === 'operands') {
- shouldParse = parser.getToken().is(_.Token.l_paren) || parser.getToken().is(_.Token.percent_identifier);
- } else if (firstElem.type === 'custom') {
- const fn = this._customDirectives.get(firstElem.parser);
- if (fn) {
- const resolvedArgs = (firstElem.args || []).map((arg) => typeof arg === 'string' && arg.startsWith('$') ? arg.slice(1) : arg);
- const result = fn(parser, op, ...resolvedArgs);
- if (result === null) {
- shouldParse = false;
- } else {
- shouldParse = 'skip_first';
- }
- }
- } else if (firstElem.type === 'qualified') {
- if (firstElem.args && firstElem.args.length > 0) {
- const [arg] = firstElem.args;
- if (arg.startsWith('$')) {
- shouldParse = parser.getToken().is(_.Token.hash_identifier);
- } else if (arg.startsWith('type($')) {
- shouldParse = parser.getToken().is(_.Token.exclamation_identifier) || parser.getToken().is(_.Token.bare_identifier);
- }
- }
- }
- }
- if (shouldParse) {
- // Recursively parse nested elements using the same parseDirective method
- // If shouldParse === 'skip_first', the custom directive already parsed the first element
- const startIdx = shouldParse === 'skip_first' ? 1 : 0;
- for (let elemIdx = startIdx; elemIdx < directive.elements.length; elemIdx++) {
- this.parseDirective(directive.elements[elemIdx], parser, op, opInfo, directive.elements, elemIdx, vars);
- }
- }
- break;
- }
- case 'conditional_alternative': {
- const checkMatch = (elem) => {
- if (elem.type === 'literal') {
- if (elem.value.length === 1 && /[(){}[\],:<>=?]/.test(elem.value)) {
- return parser.getToken().is(elem.value);
- }
- return parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === elem.value || parser.getToken().is(`kw_${elem.value}`);
- }
- if (elem.type === 'operand_ref') {
- return parser.getToken().is(_.Token.percent_identifier);
- }
- if (elem.type === 'attribute_ref') {
- return parser.getToken().is(_.Token.bare_identifier) || parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.floatliteral) || parser.getToken().is(_.Token.l_square) || parser.getToken().is(_.Token.at_identifier) || parser.getToken().is(_.Token.hash_identifier);
- }
- if (elem.type === 'region_ref') {
- return parser.getToken().is(_.Token.l_brace);
- }
- if (elem.type === 'successor_ref') {
- return parser.getToken().is(_.Token.caret_identifier);
- }
- if (elem.type === 'custom') {
- // Custom directives can start with various tokens including negative integers
- return parser.getToken().is(_.Token.bare_identifier) || parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus) || parser.getToken().is(_.Token.percent_identifier) || parser.getToken().is(_.Token.l_square) || parser.getToken().is(_.Token.l_paren) || parser.getToken().is(_.Token.question);
- }
- return false;
- };
- const firstElem = directive.firstAlt.find((e) => e.type !== 'whitespace');
- let matchedFirst = firstElem && checkMatch(firstElem);
- let customDirectiveHandledFirst = false;
- // For custom directives, try calling them and check if they return null
- if (matchedFirst && firstElem.type === 'custom') {
- const fn = this._customDirectives.get(firstElem.parser);
- if (fn) {
- const resolvedArgs = (firstElem.args || []).map((arg) => {
- if (typeof arg === 'string' && arg.startsWith('$')) {
- return arg.slice(1); // Strip $ prefix to get attribute name
- }
- return arg;
- });
- const result = fn(parser, op, ...resolvedArgs);
- if (result === null) {
- matchedFirst = false;
- } else {
- customDirectiveHandledFirst = true;
- }
- }
- }
- if (matchedFirst) {
- const startIdx = customDirectiveHandledFirst ? 1 : 0;
- for (let elemIdx = startIdx; elemIdx < directive.firstAlt.length; elemIdx++) {
- this.parseDirective(directive.firstAlt[elemIdx], parser, op, opInfo, directive.firstAlt, elemIdx, vars);
- }
- } else if (directive.secondOptional) {
- const secondElem = directive.secondAlt.find((e) => e.type !== 'whitespace');
- const matchedSecond = secondElem && checkMatch(secondElem);
- if (matchedSecond) {
- for (const elem of directive.secondAlt) {
- this.parseDirective(elem, parser, op, opInfo, directive.secondAlt, 0, vars);
- }
- }
- } else if (directive.secondAlt && directive.secondAlt.length > 0) {
- for (const elem of directive.secondAlt) {
- this.parseDirective(elem, parser, op, opInfo, directive.secondAlt, 0, vars);
- }
- }
- break;
- }
- default: {
- throw new mlir.Error(`Unsupported directive type '${directive.type}' ${parser.location()}.`);
- }
- }
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (!opInfo) {
- return false;
- }
- if ((opInfo.metadata.hasParser || opInfo.metadata.hasCustomAssemblyFormat) && !opInfo.metadata.assemblyFormat) {
- throw new mlir.Error(`Operation parser '${result.op}' not implemented.`);
- }
- // Mark as using assembly format parsing (bypasses validation check)
- if (result.compatibility === undefined && opInfo.metadata.assemblyFormat) {
- result.compatibility = true;
- }
- const vars = new Map();
- for (const input of opInfo.metadata?.operands || []) {
- vars.set(input.name, { operands: [], types: [] });
- }
- for (const resultInfo of opInfo.metadata?.results || []) {
- vars.set(resultInfo.name, { types: [] });
- }
- const directives = opInfo.directives || [];
- for (let i = 0; i < directives.length; i++) {
- this.parseDirective(directives[i], parser, result, opInfo, directives, i, vars);
- }
- for (const [, v] of vars) {
- if (v.operands?.length > 0 && v.types?.length > 0) {
- parser.resolveOperands(v.operands, v.types, result.operands);
- } else if (v.operands?.length > 0 && result.types.length > 0) {
- // SameOperandsAndResultType: use result type for operands
- const types = v.operands.map(() => result.types[0]);
- parser.resolveOperands(v.operands, types, result.operands);
- } else if (v.operands?.length > 0) {
- // No explicit type - resolve from scope (type from definition)
- for (const operand of v.operands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- }
- // AttrSizedOperandSegments trait
- if (vars && opInfo.metadata?.operands && result.name.hasTrait('AttrSizedOperandSegments')) {
- const segmentSizes = [];
- for (const operandMeta of opInfo.metadata.operands) {
- const entry = vars.get(operandMeta.name);
- segmentSizes.push(entry?.operands?.length || 0);
- }
- result.addAttribute('operandSegmentSizes', segmentSizes);
- }
- this.inferResultTypes(result, vars);
- return true;
- }
- parseType(parser, dialect) {
- const mnemonic = parser.parseOptionalKeyword();
- if (mnemonic) {
- let type = `!${dialect}.${mnemonic}`;
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- return null;
- }
- parseAttribute(/* parser, type */) {
- return null;
- }
- parseCustomTypeWithFallback(parser, type) {
- const optionalType = parser.parseOptionalType();
- if (optionalType) {
- return optionalType;
- }
- if (type && this._customTypes.has(type.name)) {
- let typeT = this._customTypes.get(type.name);
- if (typeof typeT !== 'function') {
- typeT = { type, name: typeT };
- }
- return parser.parseCustomTypeWithFallback(typeT);
- }
- throw new mlir.Error(`Unsupported type constraint '${type}'.`);
- }
- parseCustomAttributeWithFallback(parser, type) {
- if (type && this._customAttributes.has(type.name)) {
- const attrT = this._customAttributes.get(type.name);
- return parser.parseCustomAttributeWithFallback(attrT, type);
- }
- if (type && Array.isArray(type.values)) {
- const value = parser.parseOptionalKeyword(type.values);
- if (value !== null) {
- return new _.TypedAttr(value, null);
- }
- }
- return parser.parseOptionalAttribute(null);
- }
- parseOptionalAttr(parser, type) {
- if (!Array.isArray(type.args) || type.args.length === 0) {
- throw new mlir.Error(`Invalid 'OptionalAttr' type.`);
- }
- const [elementType] = type.args;
- return this.parseCustomAttributeWithFallback(parser, elementType);
- }
- parseDefaultValuedAttr(parser, type) {
- if (!Array.isArray(type.args) || type.args.length === 0) {
- throw new mlir.Error(`Invalid 'DefaultValuedAttr' type.`);
- }
- const [elementType] = type.args;
- return this.parseCustomAttributeWithFallback(parser, elementType);
- }
- parseDefaultValuedOptionalAttr(parser, type) {
- if (!Array.isArray(type.args) || type.args.length === 0) {
- throw new mlir.Error(`Invalid 'DefaultValuedOptionalAttr' type.`);
- }
- const [elementType] = type.args;
- return this.parseCustomAttributeWithFallback(parser, elementType);
- }
- parseTypeAttrOf(parser, type) {
- if (!Array.isArray(type.args) || type.args.length === 0) {
- throw new mlir.Error(`Invalid 'TypeAttrOf' type.`);
- }
- const parsedType = parser.parseOptionalType();
- if (parsedType) {
- return { value: parsedType, type: 'type' };
- }
- return null;
- }
- parseAnyAttrOf(parser) {
- // This allows parseAttribute to handle the full syntax including `: type` suffix
- return parser.parseOptionalAttribute(null);
- }
- parseArrayAttr(parser) {
- if (parser.getToken().is(_.Token.l_square)) {
- return parser.parseOptionalAttribute();
- }
- // Handle attribute alias references that resolve to arrays
- if (parser.getToken().is(_.Token.hash_identifier)) {
- return parser.parseAttribute();
- }
- return null;
- }
- parseConfinedAttr(parser, type) {
- if (!Array.isArray(type.args) || type.args.length === 0) {
- throw new mlir.Error(`Invalid ConfinedAttr type.`);
- }
- const [baseType] = type.args;
- return this.parseCustomAttributeWithFallback(parser, baseType);
- }
- parseTypedAttrInterface(parser) {
- return parser.parseAttribute();
- }
- parseUnitAttr(parser) {
- parser.consumeIf(_.Token.kw_unit);
- return new _.UnitAttr();
- }
- parseSymbolNameAttr(parser) {
- const value = parser.parseOptionalSymbolName();
- if (value) {
- return new _.StringAttr(value);
- }
- return null;
- }
- parseSymbolRefAttr(parser) {
- return parser.parseAttribute();
- }
- parseFlatSymbolRefAttr(parser) {
- return this.parseSymbolRefAttr(parser);
- }
- // custom<DynamicIndexList>($dynamic_operands, $static_attr, $scalable_attr?, "Delimiter::Paren"?)
- parseDynamicIndexList(parser, op, operandsAttr, staticAttrName, scalableAttrName, delimiterSpec) {
- // Determine delimiter from delimiterSpec
- let openDelim = '[';
- let closeDelim = ']';
- if (typeof delimiterSpec === 'string' && delimiterSpec.includes('Paren')) {
- openDelim = '(';
- closeDelim = ')';
- }
- const unresolvedOperands = [];
- const staticValues = [];
- const scalableFlags = [];
- if (parser.consumeIf(openDelim)) {
- while (parser.getToken().isNot(closeDelim)) {
- const isScalable = parser.parseOptionalLSquare();
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- staticValues.push(_.ShapedType.kDynamic);
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- } else {
- const intVal = parser.parseOptionalInteger();
- if (intVal === null) {
- break;
- }
- staticValues.push(intVal);
- }
- scalableFlags.push(isScalable);
- if (isScalable) {
- if (!parser.parseOptionalRSquare()) {
- throw new mlir.Error(`Expected ']' for scalable index ${parser.location()}`);
- }
- }
- parser.parseOptionalComma();
- }
- parser.consumeToken(closeDelim);
- }
- const indexType = new _.IndexType();
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => indexType), op.operands);
- if (staticAttrName && staticValues.length > 0) {
- op.addAttribute(staticAttrName, staticValues);
- }
- if (scalableFlags.length > 0 && scalableFlags.some((f) => f)) {
- if (!scalableAttrName) {
- throw new mlir.Error(`Scalable indices found but no scalable attribute name provided ${parser.location()}`);
- }
- op.addAttribute(scalableAttrName, scalableFlags);
- }
- }
- parseOffsets(parser, op, attrName) {
- const values = [];
- while (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus)) {
- if (parser.consumeIf(_.Token.minus)) {
- if (parser.getToken().is(_.Token.integer)) {
- values.push(-parser.parseInteger());
- } else {
- throw new mlir.Error(`Expected integer after '-' in offsets ${parser.location()}`);
- }
- } else {
- values.push(parser.parseInteger());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (attrName) {
- op.addAttribute(attrName, values);
- }
- }
- parseSymbolVisibility(parser, op) {
- const visibility = parser.parseOptionalKeyword(['public', 'private', 'nested']);
- if (visibility) {
- op.addAttribute('sym_visibility', visibility);
- }
- }
- parseTypeOrAttr(parser, op, typeArg, attrArg) {
- if (parser.parseOptionalEqual()) {
- const attr = parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- op.addAttribute(typeArg, type);
- attr.type = type;
- } else if (attr && attr.type) {
- op.addAttribute(typeArg, attr.type);
- }
- op.addAttribute(attrArg, attr);
- return;
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- op.addAttribute(typeArg, type);
- if (parser.parseOptionalEqual()) {
- const attr = parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- const attrType = parser.parseType();
- attr.type = attrType;
- }
- op.addAttribute(attrArg, attr);
- }
- return;
- }
- throw new mlir.Error(`Expected ':' or '=' in TypeOrAttr ${parser.location()}`);
- }
- parseSizeAwareType(parser, op, typeArg) {
- const type = parser.parseType();
- parser.parseLBrace();
- const operand = parser.parseOperand();
- parser.parseRBrace();
- if (!Array.isArray(typeArg)) {
- throw new mlir.Error(`Invalid argument 'typeArg'.`);
- }
- if (typeArg.length === 0) {
- typeArg.push(type);
- } else {
- typeArg[0] = type;
- }
- if (operand) {
- parser.resolveOperand(operand, null, op.operands);
- }
- }
- parseCopyOpRegion(parser, result) {
- result.regions.push({ blocks: [] });
- }
- parseResultTypeList(parser, op) {
- parser.parseLParen();
- const types = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- types.push(parser.parseType());
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- if (types.length > 0) {
- op.addAttribute('result_types', types.map((t) => t.toString()));
- }
- }
- parseCmdParameterGatherOperations(parser, op, sourceScopeOps, sourceKeysOps, sourceOffsetsOps, targetOps, targetTypes, targetSizeOps, targetOffsetsOps, targetLengthsOps) {
- do {
- const firstOperand = parser.parseOperand();
- let key = firstOperand;
- if (parser.parseOptionalColon()) {
- sourceScopeOps.push(firstOperand);
- parser.parseColon();
- key = parser.parseOperand();
- }
- sourceKeysOps.push(key);
- parser.parseLSquare();
- sourceOffsetsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseArrow();
- const rowTarget = parser.parseOperand();
- if (targetOps.length === 0) {
- targetOps.push(rowTarget);
- }
- parser.parseLSquare();
- targetOffsetsOps.push(parser.parseOperand());
- parser.parseKeyword('for');
- targetLengthsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseColon();
- const rowType = parser.parseType();
- if (targetTypes.length === 0) {
- targetTypes.push(rowType);
- }
- parser.parseLBrace();
- const rowSize = parser.parseOperand();
- if (targetSizeOps.length === 0) {
- targetSizeOps.push(rowSize);
- }
- parser.parseRBrace();
- } while (parser.parseOptionalComma());
- }
- // custom<AsyncParameterGatherOperations>(...)
- // Parses: %scope::%key[%offset] -> %target[%offset to %end for %length] : type{%size}, ...
- parseAsyncParameterGatherOperations(parser, op, sourceScopeOps, sourceKeysOps, sourceOffsetsOps, targetOps, targetTypes, targetSizeOps, targetOffsetsOps, targetEndsOps, targetLengthsOps) {
- do {
- const firstOperand = parser.parseOperand();
- let key = firstOperand;
- if (parser.parseOptionalColon()) {
- sourceScopeOps.push(firstOperand);
- parser.parseColon();
- key = parser.parseOperand();
- }
- sourceKeysOps.push(key);
- parser.parseLSquare();
- sourceOffsetsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseArrow();
- const rowTarget = parser.parseOperand();
- if (targetOps.length === 0) {
- targetOps.push(rowTarget);
- }
- parser.parseLSquare();
- targetOffsetsOps.push(parser.parseOperand());
- parser.parseKeyword('to');
- targetEndsOps.push(parser.parseOperand());
- parser.parseKeyword('for');
- targetLengthsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseColon();
- const rowType = parser.parseType();
- if (targetTypes.length === 0) {
- targetTypes.push(rowType);
- }
- parser.parseLBrace();
- const rowSize = parser.parseOperand();
- if (targetSizeOps.length === 0) {
- targetSizeOps.push(rowSize);
- }
- parser.parseRBrace();
- } while (parser.parseOptionalComma());
- }
- // custom<CmdParameterScatterOperations>(...)
- // Parses: %source[%offset for %length] : type{%size} -> %scope::%key[%offset], ...
- parseCmdParameterScatterOperations(parser, op, sourceOps, sourceTypes, sourceSizeOps, sourceOffsetsOps, sourceLengthsOps, targetScopeOps, targetKeysOps, targetOffsetsOps) {
- do {
- const rowSource = parser.parseOperand();
- if (sourceOps.length === 0) {
- sourceOps.push(rowSource);
- }
- parser.parseLSquare();
- sourceOffsetsOps.push(parser.parseOperand());
- parser.parseKeyword('for');
- sourceLengthsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseColon();
- const rowType = parser.parseType();
- if (sourceTypes.length === 0) {
- sourceTypes.push(rowType);
- }
- parser.parseLBrace();
- const rowSize = parser.parseOperand();
- if (sourceSizeOps.length === 0) {
- sourceSizeOps.push(rowSize);
- }
- parser.parseRBrace();
- // Parse -> %scope::%key[%offset]
- parser.parseArrow();
- const firstOperand = parser.parseOperand();
- let key = firstOperand;
- if (parser.parseOptionalColon()) {
- targetScopeOps.push(firstOperand);
- parser.parseColon();
- key = parser.parseOperand();
- }
- targetKeysOps.push(key);
- parser.parseLSquare();
- targetOffsetsOps.push(parser.parseOperand());
- parser.parseRSquare();
- } while (parser.parseOptionalComma());
- }
- // custom<AsyncParameterScatterOperations>(...)
- // Parses: %source[%offset to %end for %length] : type{%size} -> %scope::%key[%offset], ...
- parseAsyncParameterScatterOperations(parser, op, sourceOps, sourceTypes, sourceSizeOps, sourceOffsetsOps, sourceEndsOps, sourceLengthsOps, targetScopeOps, targetKeysOps, targetOffsetsOps) {
- do {
- const rowSource = parser.parseOperand();
- if (sourceOps.length === 0) {
- sourceOps.push(rowSource);
- }
- parser.parseLSquare();
- sourceOffsetsOps.push(parser.parseOperand());
- parser.parseKeyword('to');
- sourceEndsOps.push(parser.parseOperand());
- parser.parseKeyword('for');
- sourceLengthsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseColon();
- const rowType = parser.parseType();
- if (sourceTypes.length === 0) {
- sourceTypes.push(rowType);
- }
- parser.parseLBrace();
- const rowSize = parser.parseOperand();
- if (sourceSizeOps.length === 0) {
- sourceSizeOps.push(rowSize);
- }
- parser.parseRBrace();
- // Parse -> %scope::%key[%offset]
- parser.parseArrow();
- const firstOperand = parser.parseOperand();
- let key = firstOperand;
- if (parser.parseOptionalColon()) {
- targetScopeOps.push(firstOperand);
- parser.parseColon();
- key = parser.parseOperand();
- }
- targetKeysOps.push(key);
- parser.parseLSquare();
- targetOffsetsOps.push(parser.parseOperand());
- parser.parseRSquare();
- } while (parser.parseOptionalComma());
- }
- parseEnumFlags(parser, type, separator) {
- const flags = [];
- do {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (!type.values.includes(value)) {
- throw new mlir.Error(`Invalid enum value '${value}' ${parser.location()}`);
- }
- flags.push(value);
- } while (parser.consumeIf(separator));
- return new _.TypedAttr(flags.join(', '));
- }
- parseEnumFlagsAngleBracketComma(parser, type) {
- if (parser.parseOptionalLess()) {
- const value = this.parseEnumFlags(parser, type, ',');
- parser.parseGreater();
- return value;
- }
- return parser.parseOptionalAttribute();
- }
- parseEnumFlagsAngleBracketPipe(parser, type) {
- if (parser.parseOptionalLess()) {
- const value = this.parseEnumFlags(parser, type, '|');
- parser.parseGreater();
- return value;
- }
- return parser.parseOptionalAttribute();
- }
- parseOptional(parser) {
- return parser.parseOptionalType();
- }
- };
- _.HLODialect = class extends _.Dialect {
- constructor(operations, name) {
- super(operations, name);
- this.registerCustomDirective('SameOperandsAndResultType', this.parseSameOperandsAndResultType.bind(this));
- this.registerCustomDirective('VariadicSameOperandsAndResultType', this.parseVariadicSameOperandsAndResultType.bind(this));
- this.registerCustomDirective('ComplexOpType', this.parseComplexOpType.bind(this));
- this.registerCustomDirective('SelectOpType', this.parseSelectOpType.bind(this));
- this.registerCustomDirective('TupleOpType', this.parseTupleOpType.bind(this));
- this.registerCustomDirective('PairwiseOpType', this.parsePairwiseOpType.bind(this));
- this.registerCustomDirective('ConvolutionDimensions', this.parseConvolutionDimensions.bind(this));
- this.registerCustomDirective('DotDimensionNumbers', this.parseDotDimensionNumbers.bind(this));
- this.registerCustomDirective('PrecisionConfig', this.parsePrecisionConfig.bind(this));
- this.registerCustomDirective('PrecisionConfigAndAlgorithm', this.parsePrecisionConfigAndAlgorithm.bind(this));
- this.registerCustomDirective('WindowAttributes', this.parseWindowAttributes.bind(this));
- this.registerCustomDirective('SliceRanges', this.parseSliceRanges.bind(this));
- this.registerCustomDirective('CustomCallTarget', this.parseCustomCallTarget.bind(this));
- this.registerCustomDirective('VariadicOperandWithAttribute', this.parseVariadicOperandWithAttribute.bind(this));
- }
- assignFromFunctionType(operandArrays, resultArray, fnType) {
- for (let i = 0; i < operandArrays.length && i < fnType.inputs.length; i++) {
- operandArrays[i].push(fnType.inputs[i]);
- }
- if (fnType.results.length === 1) {
- resultArray.push(fnType.results[0]);
- }
- }
- parseSameOperandsAndResultType(parser, op, ...typeArrays) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- const operandArrays = typeArrays.slice(0, -1);
- const resultArray = typeArrays[typeArrays.length - 1];
- this.assignFromFunctionType(operandArrays, resultArray, type);
- } else {
- for (const arr of typeArrays) {
- arr.push(type);
- }
- }
- }
- // custom<VariadicSameOperandsAndResultType>(ref($inputs), type($inputs), type($result))
- parseVariadicSameOperandsAndResultType(parser, op, operands, operandTypes, resultTypes) {
- const type = parser.parseType();
- // All operands get the same type
- for (let i = 0; i < operands.length; i++) {
- operandTypes.push(type);
- }
- // Result also gets the same type
- resultTypes.push(type);
- }
- // custom<ComplexOpType>(ref($operands), type($operands), type($result))
- parseComplexOpType(parser, op, operands, operandTypes, resultTypes) {
- const type = parser.parseType();
- for (let i = 0; i < operands.length; i++) {
- operandTypes.push(type);
- }
- resultTypes.push(type);
- }
- // custom<SelectOpType>(type($pred), type($on_true), type($on_false), type($result))
- parseSelectOpType(parser, op, predTypes, onTrueTypes, onFalseTypes, resultTypes) {
- const firstType = parser.parseType();
- if (parser.parseOptionalComma()) {
- const secondType = parser.parseType();
- predTypes.push(firstType);
- onTrueTypes.push(secondType);
- onFalseTypes.push(secondType);
- resultTypes.push(secondType);
- } else {
- predTypes.push(firstType);
- onTrueTypes.push(firstType);
- onFalseTypes.push(firstType);
- resultTypes.push(firstType);
- }
- }
- // custom<TupleOpType>(type($operands), type($result))
- parseTupleOpType(parser, op, operandTypes, resultTypes) {
- const type = parser.parseType();
- operandTypes.push(type);
- resultTypes.push(type);
- }
- // custom<PairwiseOpType>(type($operands), type($results))
- parsePairwiseOpType(parser, op, operandTypes, resultTypes) {
- while (true) {
- const type = parser.parseType();
- if (!type) {
- break;
- }
- operandTypes.push(type);
- resultTypes.push(type);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parseDims(parser) {
- const dims = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- dims.push(parseInt(parser.getToken().getSpelling().str(), 10));
- parser.consumeToken();
- } else if (parser.getToken().is(_.Token.bare_identifier)) {
- dims.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.bare_identifier);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalRSquare();
- }
- return dims;
- }
- parseConvolutionDimensions(parser, op, attrName) {
- const input = this.parseDims(parser);
- parser.parseKeyword('x');
- const kernel = this.parseDims(parser);
- parser.parseArrow();
- const output = this.parseDims(parser);
- op.addAttribute(attrName, new _.ConvDimensionNumbersAttr(input, kernel, output));
- }
- parseWindowAttributes(parser, op, stridesAttr, paddingAttr, lhsDilationAttr, rhsDilationAttr, reversalAttr) {
- const windowAttrs = {
- stride: [],
- pad: [],
- lhs_dilate: [],
- rhs_dilate: [],
- window_reversal: []
- };
- const parseArray = () => {
- return parser.parseCommaSeparatedList('square', () => {
- if (parser.getToken().is(_.Token.l_square)) {
- return parseArray();
- }
- const intValue = parser.parseOptionalInteger();
- if (intValue !== null) {
- return intValue;
- }
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const retVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- return retVal;
- }
- return null;
- });
- };
- while (parser.getToken().isNot(_.Token.r_brace)) {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const key = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalEqual()) {
- windowAttrs[key] = parseArray();
- }
- parser.parseOptionalComma();
- } else {
- break;
- }
- }
- if (stridesAttr && windowAttrs.stride.length > 0) {
- op.addAttribute(stridesAttr, windowAttrs.stride);
- }
- if (paddingAttr && windowAttrs.pad.length > 0) {
- op.addAttribute(paddingAttr, windowAttrs.pad);
- }
- if (lhsDilationAttr && windowAttrs.lhs_dilate.length > 0) {
- op.addAttribute(lhsDilationAttr, windowAttrs.lhs_dilate);
- }
- if (rhsDilationAttr && windowAttrs.rhs_dilate.length > 0) {
- op.addAttribute(rhsDilationAttr, windowAttrs.rhs_dilate);
- }
- if (reversalAttr && windowAttrs.window_reversal.length > 0) {
- op.addAttribute(reversalAttr, windowAttrs.window_reversal);
- }
- }
- parseDotDimensionNumbers(parser, op, attrName = 'dot_dimension_numbers') {
- const dimensions = {
- lhs_batching_dimensions: [],
- rhs_batching_dimensions: [],
- lhs_contracting_dimensions: [],
- rhs_contracting_dimensions: []
- };
- const parseIntArray = () => {
- return parser.parseCommaSeparatedList('optionalSquare', () => {
- const value = parser.parseOptionalInteger();
- if (value !== null) {
- return value;
- }
- parser.consumeToken();
- return null;
- });
- };
- const parsePair = () => {
- const first = parseIntArray();
- let second = [];
- if (parser.parseOptionalKeyword('x')) {
- second = parseIntArray();
- }
- return { first, second };
- };
- if (parser.parseOptionalKeyword('batching_dims') || parser.parseOptionalKeyword('batch_dims')) {
- parser.parseEqual();
- const pair = parsePair();
- dimensions.lhs_batching_dimensions = pair.first;
- dimensions.rhs_batching_dimensions = pair.second;
- parser.parseOptionalComma();
- }
- parser.parseKeyword('contracting_dims');
- parser.parseEqual();
- const pair = parsePair();
- dimensions.lhs_contracting_dimensions = pair.first;
- dimensions.rhs_contracting_dimensions = pair.second;
- op.addAttribute(attrName, dimensions);
- }
- parsePrecisionConfig(parser, op /*, args */) {
- parser.parseOptionalComma();
- if (!(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'precision')) {
- return;
- }
- parser.parseKeyword('precision');
- parser.parseEqual();
- const precision = parser.parseCommaSeparatedList('square', () => {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const retVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- return retVal;
- }
- parser.consumeToken();
- return null;
- });
- if (precision.length > 0) {
- op.addAttribute('precision_config', precision);
- }
- }
- parsePrecisionConfigAndAlgorithm(parser, op /*, args */) {
- if (!parser.parseOptionalComma()) {
- return;
- }
- if (parser.parseOptionalKeyword('algorithm')) {
- parser.parseOptionalEqual();
- const algorithm = parser.parseAttribute();
- op.addAttribute('algorithm', algorithm);
- return;
- }
- if (parser.parseOptionalKeyword('precision')) {
- parser.parseOptionalEqual();
- const precision = parser.parseCommaSeparatedList('optionalSquare', () => {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const retVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- return retVal;
- }
- parser.consumeToken();
- return null;
- });
- if (precision.length > 0) {
- op.addAttribute('precision_config', precision);
- }
- if (parser.parseOptionalComma()) {
- if (parser.parseOptionalKeyword('algorithm')) {
- parser.parseOptionalEqual();
- const algorithm = parser.parseAttribute();
- op.addAttribute('algorithm', algorithm);
- }
- }
- }
- }
- parseSliceRanges(parser, op /*, args */) {
- const ranges = {
- start_indices: [],
- limit_indices: [],
- strides: []
- };
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- ranges.start_indices.push(parser.parseInteger());
- }
- parser.parseOptionalColon();
- if (parser.getToken().is(_.Token.integer)) {
- ranges.limit_indices.push(parser.parseInteger());
- }
- if (parser.parseOptionalColon()) {
- if (parser.getToken().is(_.Token.integer)) {
- ranges.strides.push(parser.parseInteger());
- }
- } else {
- ranges.strides.push(1);
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalRSquare();
- }
- op.addAttribute('start_indices', ranges.start_indices);
- op.addAttribute('limit_indices', ranges.limit_indices);
- op.addAttribute('strides', ranges.strides);
- }
- // custom<CustomCallTarget>($call_target_name)
- parseCustomCallTarget(parser, op, attrName) {
- let target = null;
- if (parser.getToken().is(_.Token.at_identifier)) {
- target = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- } else if (parser.getToken().is(_.Token.string)) {
- target = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else {
- throw new mlir.Error(`Expected '@' or string for CustomCallTarget at ${parser.location()}`);
- }
- op.addAttribute(attrName || 'call_target_name', target);
- }
- // custom<VariadicOperandWithAttribute>($inputs)
- parseVariadicOperandWithAttribute(parser, op, operands) {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- if (parser.getToken().is(_.Token.l_brace)) {
- operand.attributes = new Map();
- parser.parseAttributeDict(operand.attributes);
- }
- operands.push(operand);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parseReduceOp(parser, result, createDimensions, returnOpName) {
- const unresolvedOperands = [];
- const unresolvedInitOperands = [];
- while (true) {
- if (!parser.parseOptionalLParen()) {
- break;
- }
- const operand = parser.parseOperand();
- parser.parseKeyword('init');
- parser.parseColon();
- const initOperand = parser.parseOperand();
- parser.parseRParen();
- unresolvedOperands.push(operand);
- unresolvedInitOperands.push(initOperand);
- parser.parseOptionalComma();
- }
- const allUnresolved = unresolvedOperands.concat(unresolvedInitOperands);
- if (parser.parseOptionalKeyword('applies')) {
- const innerOpName = parser.parseCustomOperationName();
- parser.parseKeyword('across');
- parser.parseKeyword('dimensions');
- parser.parseEqual();
- parser.parseLSquare();
- const dimensions = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- dimensions.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- } else {
- throw new mlir.Error(`Expected integer dimension in reduce operation ${parser.location()}`);
- }
- if (!parser.parseOptionalComma() && parser.getToken().isNot(_.Token.r_square)) {
- throw new mlir.Error(`Expected ',' or ']' in dimensions list ${parser.location()}`);
- }
- }
- parser.parseRSquare();
- result.addAttribute('dimensions', createDimensions(dimensions));
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(allUnresolved, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- const types = Array.isArray(type) ? type : [type];
- parser.resolveOperands(allUnresolved, allUnresolved.map((_, i) => types[i] || types[0]), result.operands);
- if (parser.parseOptionalArrow()) {
- result.addTypes(parser.parseFunctionResultTypes());
- }
- }
- } else {
- parser.resolveOperands(allUnresolved, allUnresolved.map(() => null), result.operands);
- }
- const region = { blocks: [] };
- const block = { operations: [], arguments: [] };
- let tensorType = null;
- if (result.operands.length > 0 && result.operands[0].type) {
- const operandType = result.operands[0].type;
- if (operandType instanceof _.RankedTensorType || operandType instanceof _.UnrankedTensorType) {
- tensorType = new _.RankedTensorType([], operandType.elementType);
- }
- }
- block.arguments.push({ value: '%lhs', type: tensorType });
- block.arguments.push({ value: '%rhs', type: tensorType });
- const innerOp = new _.OperationState(null, innerOpName);
- innerOp.operands.push(new _.Value('%lhs', tensorType));
- innerOp.operands.push(new _.Value('%rhs', tensorType));
- if (tensorType) {
- innerOp.addTypes([tensorType]);
- }
- block.operations.push(_.Operation.create(innerOp));
- const returnOp = new _.OperationState(null, this.getOperation(returnOpName));
- returnOp.operands.push(new _.Value('%0', tensorType));
- block.operations.push(_.Operation.create(returnOp));
- region.blocks.push(block);
- result.regions.push(region);
- return true;
- }
- // Non-compact syntax: parse "across dimensions = [...] : type reducer (args) { ... }"
- parser.parseKeyword('across');
- parser.parseKeyword('dimensions');
- parser.parseEqual();
- parser.parseLSquare();
- const dimensions = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- dimensions.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- result.addAttribute('dimensions', createDimensions(dimensions));
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(allUnresolved, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- const types = Array.isArray(type) ? type : [type];
- parser.resolveOperands(allUnresolved, types, result.operands);
- }
- }
- parser.parseKeyword('reducer');
- // Parse block arguments: (%lhs : type, %rhs : type) (%linit : type, %rinit : type)
- const regionArgs = [];
- while (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- regionArgs.push({ value, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- const region = result.addRegion();
- parser.parseRegion(region, regionArgs);
- return true;
- }
- // Parse scan operation format: scan (%inputs) inits (%inits) dimension=0 { body } : types
- parseScanOp(parser, result /*, returnOpName */) {
- parser.parseLParen();
- const unresolvedInputs = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedInputs.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseKeyword('inits');
- parser.parseLParen();
- const unresolvedInits = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedInits.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseKeyword('dimension');
- parser.parseOptionalEqual();
- const dimension = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('dimension', dimension);
- // Parse optional attributes: is_reverse=true, is_associative=true
- while (parser.parseOptionalComma() || parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'is_reverse' || parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'is_associative') {
- if (parser.parseOptionalKeyword('is_reverse')) {
- parser.parseOptionalEqual();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('is_reverse', value === 'true');
- } else if (parser.parseOptionalKeyword('is_associative')) {
- parser.parseOptionalEqual();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('is_associative', value === 'true');
- }
- }
- const region = result.addRegion();
- parser.parseRegion(region, undefined, true); // IsolatedFromAbove
- // Parse types: : (input types, init types) -> (output types, carry types)
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- const allUnresolved = unresolvedInputs.concat(unresolvedInits);
- parser.resolveOperands(allUnresolved, type.inputs, result.operands);
- result.addTypes(type.results);
- }
- }
- return true;
- }
- };
- _.StableHLODialect = class extends _.HLODialect {
- constructor(operations) {
- super(operations, 'stablehlo');
- this.registerCustomDirective('ExponentMantissa', this.parseExponentMantissa.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'stablehlo.constant') {
- if (parser.parseOptionalLParen() && parser.parseOptionalRParen()) {
- if (parser.parseOptionalLess()) {
- result.propertiesAttr = parser.parseAttribute();
- parser.parseGreater();
- }
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- parser.parseLParen();
- parser.parseRParen();
- parser.parseArrow();
- const type = parser.parseType();
- result.addTypes([type]);
- } else {
- // Custom form: {attrs} value : type
- parser.parseOptionalAttrDict(result.attributes);
- const value = parser.parseAttribute();
- if (value) {
- result.addAttribute('value', value);
- }
- const types = parser.parseOptionalColonTypeList();
- if (types.length > 0) {
- result.addTypes([types[0]]);
- } else if (value && value.type) {
- result.addTypes([value.type]);
- }
- }
- return true;
- }
- if (result.op === 'stablehlo.while' && parser.getToken().is(_.Token.l_paren)) {
- const unresolvedOperands = [];
- parser.parseOptionalLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOperand(); // Skip block argument name
- if (parser.parseOptionalEqual()) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- if (parser.parseOptionalColon()) {
- const types = [];
- while (!(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'cond') && !(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'attributes')) {
- types.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- for (const type of types) {
- result.addTypes([type]);
- }
- }
- if (parser.parseOptionalKeyword('attributes')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (parser.parseOptionalKeyword('cond')) {
- const cond = result.addRegion();
- parser.parseRegion(cond);
- }
- if (parser.parseOptionalKeyword('do')) {
- const body = result.addRegion();
- parser.parseRegion(body);
- }
- return true;
- }
- if (result.op === 'stablehlo.reduce' && parser.getToken().is(_.Token.l_paren)) {
- // stablehlo uses DenseI64ArrayAttr for dimensions (like b.getDenseI64ArrayAttr in ref impl)
- return super.parseReduceOp(parser, result, (dims) => dims, 'stablehlo.return');
- }
- if (result.op === 'stablehlo.scan' && parser.getToken().is(_.Token.l_paren)) {
- return super.parseScanOp(parser, result, 'stablehlo.return');
- }
- return super.parseOperation(parser, result);
- }
- parseType(parser, dialect) {
- const mnemonic = parser.parseOptionalKeyword();
- if (mnemonic === 'token') {
- return new _.Type(`!${dialect}.token`);
- }
- throw new mlir.Error(`Unknown '${dialect}' type '${mnemonic}' ${parser.getNameLoc()}`);
- }
- parseExponentMantissa(parser, op, exponentAttr, mantissaAttr) {
- const keyword = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- const match = /^e(\d+)m(\d+)$/.exec(keyword);
- if (!match) {
- throw new mlir.Error(`Expected exponent mantissa in format e#m#, got '${keyword}'`);
- }
- const exponent = parseInt(match[1], 10);
- const mantissa = parseInt(match[2], 10);
- op.addAttribute(exponentAttr, exponent);
- op.addAttribute(mantissaAttr, mantissa);
- }
- };
- _.VhloDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'vhlo');
- this.registerCustomDirective('FunctionBody', this.parseFunctionBody.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'vhlo.constant_v1') {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const value = parser.parseAttribute();
- if (value) {
- result.addAttribute('value', value);
- }
- result.addTypes(parser.parseOptionalColonTypeList());
- return true;
- }
- if (result.op === 'vhlo.return_v1') {
- const unresolvedOperands = parser.parseOperandList();
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseFunctionBody(parser, result) {
- parser.parseFunctionOp(result, false);
- }
- };
- _.InterpreterDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'interpreter');
- }
- };
- _.AffineDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'affine');
- }
- parseOperation(parser, result) {
- if (result.op === 'affine.parallel') {
- return this.parseParallelOp(parser, result);
- }
- // Special handling for affine.for - similar to scf.for but with affine expressions
- if (result.op === 'affine.for') {
- return this.parseForOp(parser, result);
- }
- // Special handling for affine.if - has condition before region
- if (result.op === 'affine.if') {
- // affine.if #set(dims)[symbols] [-> (type)] { region }
- // Or: affine.if affine_set<(d0) : (constraint)>(dims)[symbols]
- if (parser.getToken().is(_.Token.hash_identifier)) {
- const condition = parser.parseAttribute();
- result.addAttribute('condition', condition);
- } else if (parser.consumeIf(_.Token.kw_affine_set)) {
- const content = parser.skip('<');
- result.addAttribute('condition', `affine_set${content}`);
- }
- const indexType = new _.IndexType();
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, indexType, result.operands);
- }
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, indexType, result.operands);
- }
- parser.parseOptionalComma();
- }
- }
- result.addTypes(parser.parseOptionalArrowTypeList());
- const region = result.addRegion();
- parser.parseRegion(region);
- if (parser.parseOptionalKeyword('else')) {
- const elseRegion = {};
- parser.parseRegion(elseRegion);
- result.regions.push(elseRegion);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- // Special handling for affine.apply, affine.min, and affine.max
- if (result.op === 'affine.apply' || result.op === 'affine.min' || result.op === 'affine.max') {
- if (parser.getToken().is(_.Token.hash_identifier) || parser.getToken().is(_.Token.kw_affine_map) || parser.getToken().is(_.Token.kw_affine_set)) {
- const value = parser.parseAttribute();
- result.addAttribute('map', value);
- }
- const indexType = new _.IndexType();
- if (parser.getToken().is(_.Token.l_paren)) {
- const unresolvedDims = parser.parseOperandList('paren');
- parser.resolveOperands(unresolvedDims, unresolvedDims.map(() => indexType), result.operands);
- }
- const unresolvedSyms = parser.parseOperandList('optionalSquare');
- parser.resolveOperands(unresolvedSyms, unresolvedSyms.map(() => indexType), result.operands);
- parser.parseOptionalAttrDict(result.attributes);
- result.addTypes([indexType]);
- return true;
- }
- if (result.op === 'affine.store') {
- return this.parseStoreOp(parser, result);
- }
- if (result.op === 'affine.load') {
- return this.parseLoadOp(parser, result);
- }
- if (result.op === 'affine.vector_load') {
- return this.parseVectorLoadOp(parser, result);
- }
- if (result.op === 'affine.vector_store') {
- return this.parseVectorStoreOp(parser, result);
- }
- if (result.op === 'affine.prefetch') {
- const memref = parser.parseOperand();
- parser.skip('[');
- parser.parseComma();
- const rwSpecifier = parser.parseOptionalKeyword();
- result.addAttribute('isWrite', rwSpecifier === 'write');
- parser.parseComma();
- parser.parseKeyword('locality');
- parser.parseLess();
- const locality = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('localityHint', locality);
- parser.parseGreater();
- parser.parseComma();
- const cacheType = parser.parseOptionalKeyword();
- result.addAttribute('isDataCache', cacheType === 'data');
- parser.parseOptionalAttrDict(result.attributes);
- const type = parser.parseColonType();
- parser.resolveOperand(memref, type, result.operands);
- return true;
- }
- // C++-only operation: affine.dma_start
- // Defined in mlir/lib/Dialect/Affine/IR/AffineOps.cpp
- if (result.op === 'affine.dma_start') {
- const indexType = new _.IndexType();
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalAttrDict(result.attributes);
- const types = [];
- if (parser.parseOptionalColon()) {
- do {
- types.push(parser.parseType());
- } while (parser.parseOptionalComma());
- }
- // Resolve operands with types, use index for any operands beyond type count
- const resolveTypes = unresolvedOperands.map((_, i) => i < types.length ? types[i] : indexType);
- parser.resolveOperands(unresolvedOperands, resolveTypes, result.operands);
- return true;
- }
- if (result.op === 'affine.dma_wait') {
- const indexType = new _.IndexType();
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalAttrDict(result.attributes);
- let memrefType = null;
- if (parser.parseOptionalColon()) {
- memrefType = parser.parseType();
- }
- // First operand is tag (memref type), rest are indices (index type)
- const resolveTypes = unresolvedOperands.map((_, i) => i === 0 ? memrefType : indexType);
- parser.resolveOperands(unresolvedOperands, resolveTypes, result.operands);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseForOp(parser, result) {
- const inductionVar = parser.parseOperand();
- parser.parseOptionalLocationSpecifier();
- parser.parseEqual();
- this.parseAffineBound(parser, result, 'lowerBound');
- parser.parseKeyword('to');
- this.parseAffineBound(parser, result, 'upperBound');
- if (parser.parseOptionalKeyword('step')) {
- if (parser.getToken().is(_.Token.integer)) {
- const step = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('step', step);
- }
- }
- if (parser.parseOptionalKeyword('iter_args')) {
- const unresolvedIterOperands = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // iter arg (block arg)
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedIterOperands.push(parser.parseOperand());
- } else {
- // Non-SSA values like constants - skip as they're not operands
- parser.parseAttribute();
- }
- }
- parser.parseOptionalComma();
- }
- }
- const iterTypes = parser.parseOptionalArrowTypeList();
- result.addTypes(iterTypes);
- parser.resolveOperands(unresolvedIterOperands, iterTypes, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- if (region.blocks && region.blocks.length > 0) {
- if (!region.blocks[0].arguments) {
- region.blocks[0].arguments = [];
- }
- if (region.blocks[0].arguments.length > 0) {
- region.blocks[0].arguments[0] = { value: inductionVar };
- } else {
- region.blocks[0].arguments.push({ value: inductionVar });
- }
- }
- result.regions.push(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseAffineBound(parser, op, boundName) {
- // Parse affine bound following reference implementation in AffineOps.cpp parseBound()
- // All affine operands have type index
- const indexType = new _.IndexType();
- // Try parsing SSA value first (shorthand for identity map)
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolved = parser.parseOperand();
- parser.resolveOperands([unresolved], [indexType], op.operands);
- const mapAttrName = boundName === 'lowerBound' ? 'lowerBoundMap' : 'upperBoundMap';
- op.addAttribute(mapAttrName, 'symbol_identity');
- return;
- }
- // Try parsing integer literal (shorthand for constant map)
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus)) {
- const negate = parser.consumeIf(_.Token.minus);
- let value = parser.parseInteger();
- if (negate) {
- value = -value;
- }
- const mapAttrName = boundName === 'lowerBound' ? 'lowerBoundMap' : 'upperBoundMap';
- op.addAttribute(mapAttrName, value);
- return;
- }
- if (!parser.parseOptionalKeyword('min')) {
- parser.parseOptionalKeyword('max');
- }
- if (parser.getToken().is(_.Token.hash_identifier) || parser.getToken().is(_.Token.kw_affine_map)) {
- const mapValue = parser.parseAttribute();
- if (mapValue) {
- const mapAttrName = boundName === 'lowerBound' ? 'lowerBoundMap' : 'upperBoundMap';
- op.addAttribute(mapAttrName, mapValue);
- // Parse dim and symbol operands in ()[...] or (...)
- const unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => indexType), op.operands);
- return;
- }
- }
- throw new mlir.Error(`Expected loop bound (SSA value, integer, or affine map) in affine.for ${parser.location()}`);
- }
- parseStoreOp(parser, result) {
- let unresolvedValue = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedValue = parser.parseOperand();
- }
- // Note: attribute values are not operands
- if (!parser.parseOptionalKeyword('to')) {
- parser.parseOptionalComma();
- }
- const unresolvedAddress = parser.parseOperand();
- parser.skip('[');
- if (parser.parseOptionalColon()) {
- const memrefType = parser.parseType();
- // Value type is element type of memref
- const valueType = memrefType.elementType || memrefType;
- if (unresolvedValue) {
- parser.resolveOperands([unresolvedValue], [valueType], result.operands);
- }
- parser.resolveOperands([unresolvedAddress], [memrefType], result.operands);
- }
- return true;
- }
- parseLoadOp(parser, result) {
- const memref = parser.parseOperand();
- parser.skip('[');
- parser.parseOptionalAttrDict(result.attributes);
- const type = parser.parseColonType();
- parser.resolveOperand(memref, type, result.operands);
- // Result type is element type of memref
- result.addTypes([type.elementType || type]);
- return true;
- }
- parseVectorLoadOp(parser, result) {
- const memref = parser.parseOperand();
- parser.skip('[');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const memrefType = parser.parseType();
- parser.resolveOperand(memref, memrefType, result.operands);
- parser.parseComma();
- const vectorType = parser.parseType();
- result.addTypes([vectorType]);
- } else {
- parser.resolveOperand(memref, null, result.operands);
- }
- return true;
- }
- parseVectorStoreOp(parser, result) {
- const value = parser.parseOperand();
- parser.parseComma();
- const memref = parser.parseOperand();
- parser.skip('[');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const memrefType = parser.parseType();
- parser.parseComma();
- const vectorType = parser.parseType();
- // Resolve operands: value first, then memref
- parser.resolveOperand(value, vectorType, result.operands);
- parser.resolveOperand(memref, memrefType, result.operands);
- } else {
- parser.resolveOperand(value, null, result.operands);
- parser.resolveOperand(memref, null, result.operands);
- }
- return true;
- }
- parseParallelOp(parser, result) {
- const ivArgs = parser.parseArgumentList('paren', false);
- const indexType = new _.IndexType();
- for (const iv of ivArgs) {
- iv.type = indexType;
- }
- if (!parser.parseOptionalEqual()) {
- return false;
- }
- parser.skip('(');
- if (!parser.parseOptionalKeyword('to')) {
- return false;
- }
- parser.skip('(');
- if (parser.parseOptionalKeyword('step')) {
- parser.skip('(');
- }
- if (parser.parseOptionalKeyword('reduce')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.string)) {
- parser.consumeToken(_.Token.string);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- const resultTypes = [];
- const resultAttrs = [];
- parser.parseFunctionResultList(resultTypes, resultAttrs);
- result.addTypes(resultTypes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- // Pass iv arguments to parseRegion - they become block arguments
- const region = result.addRegion();
- parser.parseRegion(region, ivArgs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- };
- _.MemRefDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'memref');
- this.registerCustomDirective('GlobalMemrefOpTypeAndInitialValue', this.parseGlobalMemrefOpTypeAndInitialValue.bind(this));
- // AtomicRMWKindAttr can appear as bare id (addi) or string ("addi") in different test files
- this.registerCustomAttribute('AtomicRMWKindAttr', this.parseAtomicRMWKindAttr.bind(this));
- }
- parseAtomicRMWKindAttr(parser, type) {
- // Accept both bare identifier (addi) and string literal ("addi")
- if (parser.getToken().is(_.Token.string)) {
- const retVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- return retVal;
- }
- if (parser.getToken().is(_.Token.bare_identifier) && type.values && type.values.includes(parser.getTokenSpelling().str())) {
- const retVal = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- return retVal;
- }
- return null;
- }
- parseGlobalMemrefOpTypeAndInitialValue(parser, op, typeAttr = 'type', initialValueAttr = 'initial_value') {
- const type = parser.parseType();
- op.addAttribute(typeAttr, { value: type, type: 'type' });
- if (parser.parseOptionalEqual()) {
- if (parser.parseOptionalKeyword('uninitialized')) {
- op.addAttribute(initialValueAttr, 'uninitialized');
- } else {
- // Pass the type to parseAttribute to suppress : type suffix parsing
- const initialValue = parser.parseAttribute(type);
- op.addAttribute(initialValueAttr, initialValue);
- }
- }
- }
- parseOperation(parser, result) {
- if (result.op === 'memref.tensor_load') {
- const unresolvedOperands = parser.parseOperandList();
- const types = parser.parseOptionalColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- // Infer result tensor type from memref operand type
- if (result.operands.length > 0 && result.operands[0].type) {
- const memrefType = result.operands[0].type;
- // Convert memref type to tensor type
- if (memrefType instanceof _.MemRefType) {
- result.addTypes([new _.RankedTensorType(memrefType.shape, memrefType.elementType, null)]);
- } else if (memrefType instanceof _.UnrankedMemRefType) {
- result.addTypes([new _.UnrankedTensorType(memrefType.elementType)]);
- }
- }
- return true;
- }
- if (result.op === 'memref.store') {
- result.compatibility = true;
- return this.parseStoreOp(parser, result);
- }
- if (result.op === 'memref.alloca_scope') {
- return this.parseAllocaScopeOp(parser, result);
- }
- if (result.op === 'memref.transpose') {
- return this.parseTransposeOp(parser, result);
- }
- if (result.op === 'memref.generic_atomic_rmw') {
- return this.parseGenericAtomicRMWOp(parser, result);
- }
- if (result.op === 'memref.prefetch') {
- const memref = parser.parseOperand();
- const indices = parser.parseOperandList('square');
- parser.parseComma();
- const readOrWrite = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('isWrite', readOrWrite === 'write');
- parser.parseComma();
- parser.parseKeyword('locality');
- parser.parseLess();
- const localityHint = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- result.addAttribute('localityHint', localityHint);
- parser.parseGreater();
- parser.parseComma();
- const cacheType = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('isDataCache', cacheType === 'data');
- const type = parser.parseColonType();
- parser.resolveOperand(memref, type, result.operands);
- const indexType = new _.IndexType();
- const indexTypes = indices.map(() => indexType);
- parser.resolveOperands(indices, indexTypes, result.operands);
- return true;
- }
- if (result.op === 'memref.dma_start') {
- const srcMemRef = parser.parseOperand();
- const srcIndices = parser.parseOperandList('square');
- parser.parseComma();
- const dstMemRef = parser.parseOperand();
- const dstIndices = parser.parseOperandList('square');
- parser.parseComma();
- const numElements = parser.parseOperand();
- parser.parseComma();
- const tagMemRef = parser.parseOperand();
- const tagIndices = parser.parseOperandList('square');
- const strideInfo = [];
- while (parser.parseOptionalComma() && parser.getToken().is(_.Token.percent_identifier)) {
- strideInfo.push(parser.parseOperand());
- }
- const types = parser.parseColonTypeList();
- const indexType = new _.IndexType();
- parser.resolveOperand(srcMemRef, types[0], result.operands);
- const srcIndexTypes = srcIndices.map(() => indexType);
- parser.resolveOperands(srcIndices, srcIndexTypes, result.operands);
- parser.resolveOperand(dstMemRef, types[1], result.operands);
- const dstIndexTypes = dstIndices.map(() => indexType);
- parser.resolveOperands(dstIndices, dstIndexTypes, result.operands);
- parser.resolveOperand(numElements, indexType, result.operands);
- parser.resolveOperand(tagMemRef, types[2], result.operands);
- const tagIndexTypes = tagIndices.map(() => indexType);
- parser.resolveOperands(tagIndices, tagIndexTypes, result.operands);
- if (strideInfo.length > 0) {
- const strideTypes = strideInfo.map(() => indexType);
- parser.resolveOperands(strideInfo, strideTypes, result.operands);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseGenericAtomicRMWOp(parser, result) {
- const memref = parser.parseOperand();
- const indices = parser.parseOperandList('square');
- parser.parseColon();
- const memrefType = parser.parseType();
- parser.resolveOperand(memref, memrefType, result.operands);
- const indexType = new _.IndexType();
- const indexTypes = indices.map(() => indexType);
- parser.resolveOperands(indices, indexTypes, result.operands);
- const region = result.addRegion();
- parser.parseRegion(region);
- parser.parseOptionalAttrDict(result.attributes);
- if (memrefType && memrefType.elementType) {
- result.addTypes([memrefType.elementType]);
- }
- return true;
- }
- parseTransposeOp(parser, result) {
- const operand = parser.parseOperand();
- const dims = parser.skip('(');
- parser.parseArrow();
- const results = parser.skip('(');
- const permutation = `affine_map<${dims} -> ${results}>`;
- result.addAttribute('permutation', permutation);
- parser.parseOptionalAttrDict(result.attributes);
- const srcType = parser.parseColonType();
- parser.resolveOperand(operand, srcType, result.operands);
- const dstType = parser.parseKeywordType('to');
- result.addTypes([dstType]);
- return true;
- }
- parseAllocaScopeOp(parser, result) {
- const resultTypes = parser.parseOptionalArrowTypeList();
- result.addTypes(resultTypes);
- const region = result.addRegion();
- parser.parseRegion(region);
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseStoreOp(parser, result) {
- // or old: value to memref[indices] : type
- let valueOperand = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- valueOperand = parser.parseOperand();
- } else {
- // Non-standard: constant value - store as attribute
- const value = parser.parseAttribute();
- result.addAttribute('value', value);
- }
- // Accept either ',' (new) or 'to' (old)
- if (!parser.parseOptionalKeyword('to')) {
- parser.parseOptionalComma();
- }
- const memrefOperand = parser.parseOperand();
- parser.skip('[');
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const memrefType = parser.parseType();
- // Value type is element type of memref
- const valueType = memrefType.elementType || memrefType;
- if (valueOperand) {
- parser.resolveOperand(valueOperand, valueType, result.operands);
- }
- parser.resolveOperand(memrefOperand, memrefType, result.operands);
- }
- return true;
- }
- };
- _.VectorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'vector');
- this.registerCustomAttribute('Vector_CombiningKindAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- this.registerCustomAttribute('Arith_FastMathAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'vector.splat') {
- const unresolvedOperands = parser.parseOperandList();
- const types = parser.parseOptionalColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- result.addTypes(types);
- return true;
- }
- if (result.op === 'vector.contract') {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- } else if (parser.getToken().is(_.Token.hash_identifier)) {
- parser.consumeToken(_.Token.hash_identifier);
- }
- const unresolvedOperands = parser.parseOperandList();
- parser.parseOptionalAttrDict(result.attributes);
- const types = parser.parseColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- const resultType = parser.parseKeywordType('into');
- result.addTypes([resultType]);
- return true;
- }
- if (result.op === 'vector.mask') {
- let mask = null;
- let passthru = null;
- let hasPassthru = false;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- mask = parser.parseOperand();
- }
- if (parser.parseOptionalComma()) {
- hasPassthru = true;
- passthru = parser.parseOperand();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- const [maskType] = parser.parseOptionalColonTypeList();
- const resultTypes = parser.parseOptionalArrowTypeList();
- if (mask) {
- parser.resolveOperand(mask, maskType, result.operands);
- }
- if (hasPassthru && passthru) {
- parser.resolveOperand(passthru, resultTypes[0], result.operands);
- }
- result.addTypes(resultTypes);
- return true;
- }
- if (result.op === 'vector.outerproduct') {
- return this.parseOuterProductOp(parser, result);
- }
- if (result.op === 'vector.transfer_read' || result.op === 'vector.transfer_write') {
- return this.parseTransferOp(parser, result);
- }
- if (result.op === 'vector.extract' && !result.isGeneric) {
- result.compatibility = true; // compatibility
- return this.parseExtractOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseOuterProductOp(parser, result) {
- const operandsInfo = parser.parseOperandList();
- parser.parseOptionalAttrDict(result.attributes);
- const tLHS = parser.parseColonType();
- parser.parseComma();
- const tRHS = parser.parseType();
- const vLHS = tLHS instanceof _.VectorType ? tLHS : null;
- const vRHS = tRHS instanceof _.VectorType ? tRHS : null;
- let resType = null;
- if (vRHS) {
- const scalableDimsRes = [
- vLHS.scalableDims ? vLHS.scalableDims[0] : false,
- vRHS.scalableDims ? vRHS.scalableDims[0] : false
- ];
- resType = new _.VectorType([vLHS.shape[0], vRHS.shape[0]], vLHS.elementType, scalableDimsRes);
- } else {
- // Scalar RHS operand
- const scalableDimsRes = [vLHS.scalableDims ? vLHS.scalableDims[0] : false];
- resType = new _.VectorType([vLHS.shape[0]], vLHS.elementType, scalableDimsRes);
- }
- parser.resolveOperand(operandsInfo[0], tLHS, result.operands);
- parser.resolveOperand(operandsInfo[1], tRHS, result.operands);
- if (operandsInfo.length > 2) {
- parser.resolveOperand(operandsInfo[2], resType, result.operands);
- }
- result.addTypes([resType]);
- return true;
- }
- parseExtractOp(parser, result) {
- // Old syntax (pre-2023): %r = vector.extract %v[0] : vector<4xf32>
- // New syntax: %r = vector.extract %v[0] : f32 from vector<4xf32>
- const unresolvedSource = parser.parseOperand();
- const unresolvedDynIndices = [];
- const indexType = new _.IndexType();
- let numStaticIndices = 0;
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- const staticIndex = parser.parseOptionalInteger();
- if (staticIndex !== null) {
- numStaticIndices++;
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const dynIndex = parser.parseOperand();
- unresolvedDynIndices.push(dynIndex);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalRSquare();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const resultType = parser.parseType();
- if (parser.parseOptionalKeyword('from')) {
- const sourceType = parser.parseType();
- parser.resolveOperand(unresolvedSource, sourceType, result.operands);
- result.addTypes([resultType]);
- } else {
- // Old syntax: the type after ':' is the source type
- // Result type is inferred by removing dimensions based on indices
- parser.resolveOperand(unresolvedSource, resultType, result.operands);
- // Infer result type from source type and number of indices
- const numIndices = numStaticIndices + unresolvedDynIndices.length;
- if (resultType instanceof _.VectorType && numIndices > 0) {
- const shape = resultType.shape.slice(numIndices);
- if (shape.length === 0) {
- // Scalar result
- result.addTypes([resultType.elementType]);
- } else {
- const scalableDims = resultType.scalableDims ? resultType.scalableDims.slice(numIndices) : [];
- result.addTypes([new _.VectorType(shape, resultType.elementType, scalableDims)]);
- }
- } else if (resultType instanceof _.VectorType) {
- // No indices - result is the same as source
- result.addTypes([resultType]);
- }
- }
- parser.resolveOperands(unresolvedDynIndices, unresolvedDynIndices.map(() => indexType), result.operands);
- }
- return true;
- }
- parseTransferOp(parser, result) {
- // or: vector.transfer_read %source[%i, %j, ...], %padding, %mask {attrs} : memref_type, vector_type
- // or: vector.transfer_write %value, %dest[%i, %j, ...] {attrs} : vector_type, memref_type
- // or: vector.transfer_write %value, %dest[%i, %j, ...], %mask {attrs} : vector_type, memref_type
- const unresolvedFirst = parser.parseOperand();
- const hasIndicesAfterFirst = parser.getToken().is(_.Token.l_square);
- if (hasIndicesAfterFirst) {
- parser.skip('[');
- }
- parser.parseOptionalComma();
- const unresolvedSecond = parser.parseOperand();
- if (!hasIndicesAfterFirst && parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- // Optional mask parameter (third operand)
- let unresolvedMask = null;
- if (parser.parseOptionalComma()) {
- unresolvedMask = parser.parseOperand();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type1 = parser.parseType();
- parser.parseOptionalComma();
- const type2 = parser.parseType();
- parser.resolveOperand(unresolvedFirst, type1, result.operands);
- parser.resolveOperand(unresolvedSecond, type2, result.operands);
- if (unresolvedMask) {
- // Mask type would be a vector of i1, use type2 as approximation
- parser.resolveOperand(unresolvedMask, type2, result.operands);
- }
- // For transfer_read, type2 is the result type
- // For transfer_write to tensor (not memref), the result is the updated tensor
- if (result.op === 'vector.transfer_read') {
- result.addTypes([type2]);
- } else if (result.op === 'vector.transfer_write' && type2 instanceof _.RankedTensorType) {
- result.addTypes([type2]);
- }
- }
- return true;
- }
- inferResultTypes(op, vars) {
- if (op.op === 'vector.shuffle') {
- const maskAttr = op.attributes.get('mask');
- if (maskAttr instanceof _.DenseI64ArrayAttr && op.operands.length > 0) {
- const v1Type = op.operands[0].type;
- if (v1Type instanceof _.VectorType) {
- const maskLength = maskAttr.value.length;
- const trailingDims = v1Type.shape.slice(1);
- const resultShape = [maskLength, ...trailingDims];
- const resultType = new _.VectorType(resultShape, v1Type.elementType, v1Type.scalableDims ? [false, ...v1Type.scalableDims.slice(1)] : []);
- op.addTypes([resultType]);
- return;
- }
- }
- }
- if (op.op === 'vector.to_elements' && op.operands.length > 0) {
- const vecType = op.operands[0].type;
- if (vecType instanceof _.VectorType) {
- const elType = vecType.elementType;
- for (let i = 0; i < vecType.getNumElements(); i++) {
- op.addTypes([elType]);
- }
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.TensorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tensor');
- }
- parseOperation(parser, result) {
- if (result.op === 'tensor.expand_shape') {
- // The new tensor.expand_shape format includes 'output_shape':
- // $src $reassociation `output_shape` custom<DynamicIndexList>(...) attr-dict `:` type($src) `into` type($result)
- // Old format (deprecated):
- // $src $reassociation attr-dict `:` type($src) `into` type($result)
- result.compatibility = true;
- const unresolvedOperand = parser.parseOperand();
- const reassociation = parser.parseAttribute();
- result.addAttribute('reassociation', reassociation);
- if (parser.parseOptionalKeyword('output_shape')) {
- // New format: parse output_shape dynamic index list
- this.parseDynamicIndexList(parser, result, ['$output_shape', '$static_output_shape']);
- }
- // Both formats: attr-dict `:` type($src) `into` type($result)
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const srcType = parser.parseType();
- parser.resolveOperands([unresolvedOperand], [srcType], result.operands);
- parser.parseKeyword('into');
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TorchDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'torch');
- this.simpleTypes = new Set([
- 'int', 'float', 'bool', 'str', 'none', 'Device', 'Generator',
- 'qint8', 'quint8', 'qint16', 'qint32', 'quint4x2', 'quint2x4',
- 'LinearParams', 'number', 'any'
- ]);
- }
- parseType(parser, dialect) {
- const mnemonic = parser.parseOptionalKeyword();
- if (!mnemonic) {
- return null;
- }
- let type = `!${dialect}.${mnemonic}`;
- if (this.simpleTypes.has(mnemonic)) {
- return new _.Type(type);
- }
- if (mnemonic === 'vtensor' || mnemonic === 'tensor' || mnemonic === 'list' || mnemonic === 'tuple' || mnemonic === 'union' || mnemonic === 'optional' || mnemonic === 'dict' || mnemonic.startsWith('nn.')) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op.startsWith('torch.constant.')) {
- result.label = 'torch.constant';
- } else if (result.op.startsWith('torch.aten.') || result.op.startsWith('torch.prim.') || result.op.startsWith('torch.prims.') || result.op.startsWith('torch.torchvision.')) {
- result.label = `${result.op.split('.')[1]}.${result.op.split('.')[2]}`;
- }
- if (result.op === 'torch.constant.int') {
- const value = parser.parseOptionalInteger();
- if (value !== null) {
- result.addAttribute('value', value);
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addTypes([new _.Type('!torch.int')]);
- return true;
- }
- if (result.op === 'torch.onnx.rotary_embedding') {
- const unresolvedOperands = parser.parseOperandList();
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- return true;
- }
- if (result.op === 'torch.bind_symbolic_shape') {
- const unresolved = parser.parseOperand();
- parser.parseOptionalComma();
- const shapeSymbols = parser.parseOperandList('square');
- parser.parseOptionalComma();
- const shapeExpr = parser.parseAttribute();
- result.addAttribute('shape_expressions', shapeExpr.value || shapeExpr);
- parser.parseOptionalAttrDict(result.attributes);
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(unresolved, type, result.operands);
- for (const sym of shapeSymbols) {
- parser.resolveOperand(sym, null, result.operands);
- }
- return true;
- }
- if (result.op === 'torch.initialize.global_slots') {
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseLSquare();
- const slotSymNames = [];
- while (!parser.parseOptionalRSquare()) {
- const slotSymName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- slotSymNames.push(slotSymName);
- parser.parseLParen();
- const unresolved = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(unresolved, type, result.operands);
- parser.parseRParen();
- }
- result.addAttribute('slotSymNames', slotSymNames);
- return true;
- }
- const opInfo = result.name.getRegisteredInfo();
- if (opInfo.metadata.hasCustomAssemblyFormat && !opInfo.metadata.assemblyFormats) {
- return this.parseDefaultTorchOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseDefaultTorchOp(parser, result) {
- const unresolvedOperands = parser.parseOperandList();
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- parser.resolveOperands(unresolvedOperands, parser.parseTypeList(), result.operands);
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- if (parser.parseOptionalArrow()) {
- // Handle both -> (type, type) and -> type, type syntaxes
- const types = parser.getToken().is(_.Token.l_paren) ? parser.parseTypeListParens() : parser.parseTypeListNoParens();
- result.addTypes(types);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- if (parser.parseOptionalKeyword('else') && parser.getToken().is(_.Token.l_brace)) {
- const elseRegion = {};
- parser.parseRegion(elseRegion);
- result.regions.push(elseRegion);
- }
- }
- return true;
- }
- };
- _.IREEDialect = class extends _.Dialect {
- constructor(operations, name) {
- super(operations, name);
- this.registerCustomDirective('DispatchEntryPoints', this.parseDispatchEntryPoints.bind(this));
- this.registerCustomDirective('ShapedTiedResult', this.parseShapedTiedResult.bind(this));
- this.registerCustomDirective('SymbolAlias', this.parseSymbolAlias.bind(this));
- this.registerCustomDirective('TypeAlias', this.parseTypeAlias.bind(this));
- this.registerCustomDirective('WorkgroupCountRegion', this.parseWorkgroupCountRegion.bind(this));
- this.registerCustomDirective('ShapedFunctionType', this.parseShapedFunctionType.bind(this));
- }
- parseShapedFunctionType(parser, op, unresolvedArguments /*, otherArgs */) {
- const operandTypes = [];
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const type = parser.parseType();
- if (type) {
- operandTypes.push(type);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- parser.parseArrow();
- const operands = unresolvedArguments || op.operands;
- const resultTypes = [];
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseShapedResultList(operands, operandTypes, resultTypes, null);
- }
- parser.parseRParen();
- } else {
- parser.parseShapedResultList(operands, operandTypes, resultTypes, null);
- }
- op.addTypes(resultTypes);
- }
- parseDispatchEntryPoints(parser, op, attrName = 'entry_points') {
- const entryPoints = [];
- if (parser.parseOptionalLBrace()) {
- do {
- if (parser.getToken().is(_.Token.at_identifier)) {
- let symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- if (parser.getToken().is(_.Token.colon)) {
- const curPointer = parser.getToken().loc.position;
- parser.consumeToken(_.Token.colon);
- if (parser.consumeIf(_.Token.colon)) {
- if (parser.getToken().is(_.Token.at_identifier)) {
- const nested = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- symbol += `::${nested}`;
- }
- } else {
- parser.resetToken(curPointer);
- }
- }
- entryPoints.push(symbol);
- }
- } while (parser.parseOptionalComma());
- parser.parseRBrace();
- } else if (parser.getToken().is(_.Token.at_identifier)) {
- let symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- if (parser.getToken().is(_.Token.colon)) {
- const curPointer = parser.getToken().loc.position;
- parser.consumeToken(_.Token.colon);
- if (parser.consumeIf(_.Token.colon)) {
- if (parser.getToken().is(_.Token.at_identifier)) {
- const nested = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- symbol += `::${nested}`;
- }
- } else {
- parser.resetToken(curPointer);
- }
- }
- entryPoints.push(symbol);
- }
- const value = entryPoints.length === 1 ? entryPoints[0] : entryPoints;
- op.addAttribute(attrName, value);
- }
- parseShapedTiedResult(parser, op /*, args */) {
- // or just: type{dims}
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // tiedOperand - parsed but not stored in OperationState
- parser.parseKeyword('as');
- }
- const resultType = parser.parseType();
- op.types.push(resultType); // Only add the type
- if (parser.parseOptionalLBrace()) {
- const indexType = new _.IndexType();
- while (parser.getToken().isNot(_.Token.r_brace)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const dim = parser.parseOperand();
- parser.resolveOperand(dim, indexType, op.operands);
- parser.parseOptionalComma();
- } else {
- break;
- }
- }
- parser.parseRBrace();
- }
- }
- parseSymbolAlias(parser, op, symNameAttr, aliasAttr) {
- const alias = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- let symName = alias;
- if (parser.parseOptionalKeyword('as')) {
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().is(_.Token.string)) {
- symName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else if (parser.getToken().is(_.Token.at_identifier)) {
- symName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- }
- parser.parseOptionalRParen();
- }
- }
- if (symNameAttr && aliasAttr) {
- op.addAttribute(symNameAttr, symName);
- op.addAttribute(aliasAttr, alias);
- }
- }
- parseTypeAlias(parser, op, encodingAttrName, typeArg) {
- const encodingType = parser.parseType();
- let storageType = encodingType;
- if (parser.parseOptionalKeyword('as')) {
- storageType = parser.parseType();
- }
- if (encodingAttrName) {
- op.addAttribute(encodingAttrName, encodingType);
- }
- if (!Array.isArray(typeArg)) {
- throw new mlir.Error(`Invalid argument 'typeArg'.`);
- }
- if (typeArg.length > 0) {
- typeArg[0] = storageType;
- } else {
- typeArg.push(storageType);
- }
- }
- parseWorkgroupCountRegion(parser, result) {
- if (!(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'workgroups')) {
- return;
- }
- parser.parseKeyword('workgroups');
- const region = { blocks: [] };
- const block = { arguments: [], operations: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- if (parser.parseOptionalColon()) {
- arg.type = parser.parseType();
- }
- block.arguments.push(arg);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalLBrace()) {
- while (parser.getToken().isNot(_.Token.r_brace)) {
- const innerOp = parser.parseOperation();
- if (innerOp) {
- block.operations.push(innerOp);
- }
- if (parser.getToken().is(_.Token.r_brace)) {
- break;
- }
- }
- parser.parseRBrace();
- }
- region.blocks.push(block);
- result.regions.push(region);
- }
- };
- _.HALDialect = class extends _.IREEDialect {
- constructor(operations) {
- super(operations, 'hal');
- this.simpleTypes = new Set(['allocator', 'buffer', 'buffer_view', 'channel', 'command_buffer', 'descriptor_set', 'descriptor_set_layout', 'device', 'event', 'executable', 'executable_layout', 'fence', 'file', 'semaphore']);
- this.registerCustomAttribute('HAL_PipelineLayoutAttr', this.parsePipelineLayoutAttr.bind(this));
- this.registerCustomDirective('ExportConditionRegion', this.parseExportConditionRegion.bind(this));
- this.registerCustomDirective('TargetConditionObjects', this.parseTargetConditionObjects.bind(this));
- this.registerCustomDirective('WorkgroupCountRegion', this.parseWorkgroupCountRegion.bind(this));
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- if (this.simpleTypes.has(typeName)) {
- // Note: !hal.buffer{%size} syntax is handled by custom<SizeAwareType> directive,
- // not by the type parser. The type parser just returns the base type.
- return new _.Type(`!${dialect}.${typeName}`);
- }
- return null;
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (result.op === 'hal.tensor.cast') {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- parser.resolveOperands(unresolvedOperands, [type], result.operands);
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- return true;
- }
- if (result.op === 'hal.constant') {
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- const value = parser.parseAttribute();
- result.addAttribute('value', value.value === undefined ? value : value.value);
- result.addTypes(parser.parseOptionalColonTypeList());
- return true;
- }
- if (result.op === 'hal.device.switch') {
- if (parser.parseOptionalLess()) {
- while (!parser.parseOptionalGreater()) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow() || parser.parseOptionalColon()) {
- const resultType = parser.parseType();
- result.types = [resultType];
- }
- while (parser.getToken().is(_.Token.hash_identifier)) {
- const region = {};
- const caseAttr = parser.parseAttribute();
- region.caseAttribute = caseAttr;
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region);
- }
- result.regions.push(region);
- parser.parseOptionalComma();
- }
- return true;
- }
- if (result.op === 'hal.executable.constant.block') {
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const arg = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(arg, type, result.operands);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType([], resultTypes)));
- }
- if (parser.parseOptionalKeyword('as')) {
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- const keys = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.string)) {
- keys.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.string);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- result.addAttribute('keys', keys);
- } else if (parser.getToken().is(_.Token.string)) {
- const key = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('keys', [key]);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- // Handle hal.executable.create with both old (layouts) and new (affinity) syntax
- if (result.op === 'hal.executable.create') {
- result.compatibility = true;
- const inputNames = new Set(opInfo.metadata.operands.map((input) => input.name));
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.kw_loc)) {
- const paramName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalLParen()) {
- if (inputNames.has(paramName) && parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let operandType = null;
- if (parser.parseOptionalColon()) {
- operandType = parser.parseType();
- }
- parser.parseRParen();
- parser.resolveOperand(operand, operandType, result.operands);
- } else if (inputNames.has(paramName) && parser.getToken().is(_.Token.l_square)) {
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- parser.parseRParen();
- } else {
- let parenDepth = 1;
- let paramValue = '';
- while (parenDepth > 0 && parser.getToken().isNot(_.Token.eof)) {
- if (parser.getToken().is(_.Token.l_paren)) {
- parenDepth++;
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else if (parser.getToken().is(_.Token.r_paren)) {
- parenDepth--;
- if (parenDepth > 0) {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else {
- parser.parseRParen();
- }
- } else {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- }
- }
- // Normalize old 'layouts' parameter to 'affinity' for consistency
- const normalizedName = paramName === 'layouts' ? 'affinity' : paramName;
- result.addAttribute(normalizedName, paramValue);
- }
- } else {
- break;
- }
- }
- result.addTypes(parser.parseOptionalColonTypeList());
- return true;
- }
- // Handle operations with <%operand : type> syntax and/or named parameters
- // e.g., hal.allocator.compute_size<%allocator : !hal.allocator> shape([...]) type(...) encoding(...) : index
- // or hal.executable_layout.lookup device(%device : !hal.device) layouts([[...]]) : !hal.executable_layout
- // Exclude hal.executable, hal.interface, and hal.device.switch which have special handling
- if ((result.op.startsWith('hal.allocator.') || result.op.startsWith('hal.buffer.') || result.op.startsWith('hal.buffer_view.') ||
- result.op.startsWith('hal.command_buffer.') || result.op.startsWith('hal.executable_layout') ||
- result.op.startsWith('hal.executable.') || result.op.startsWith('hal.descriptor_set_layout') ||
- result.op.startsWith('hal.device.')) &&
- result.op !== 'hal.device.allocator' &&
- result.op !== 'hal.buffer_view.buffer' &&
- result.op !== 'hal.executable' &&
- result.op !== 'hal.interface' &&
- result.op !== 'hal.device.switch' &&
- result.op !== 'hal.device.memoize' &&
- result.op !== 'hal.command_buffer.execution_barrier' &&
- result.op !== 'hal.executable.entry_point' &&
- result.op !== 'hal.executable.variant' &&
- result.op !== 'hal.executable.lookup' &&
- result.op !== 'hal.interface.binding' &&
- result.op !== 'hal.executable.create' &&
- result.op !== 'hal.executable.export' &&
- result.op !== 'hal.executable.binary' &&
- result.op !== 'hal.executable.source' &&
- result.op !== 'hal.executable.condition' &&
- result.op !== 'hal.executable.constant.block' &&
- result.op !== 'hal.executable.constant.load') {
- if (result.op === 'hal.allocator.allocate' || result.op === 'hal.command_buffer.create' || result.op === 'hal.buffer_view.create' || result.op === 'hal.command_buffer.device' || result.op === 'hal.command_buffer.dispatch' || result.op === 'hal.device.query') {
- result.compatibility = true;
- }
- if (parser.parseOptionalLess()) {
- while (!parser.parseOptionalGreater()) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- parser.parseOptionalComma();
- }
- }
- // Also handle bracket expressions between parameters like layout(...)[%c0]
- // Stop when we hit a colon (result type) or something that doesn't look like a parameter
- // Named parameters don't have dots, so if we see an id with a dot, it's likely the next operation
- // Also exclude common operation keywords that shouldn't be treated as parameters
- const notParameterNames = new Set(['br', 'cond_br', 'return', 'yield', 'call', 'unreachable', 'assert']);
- while (parser.getToken().is(_.Token.l_square) || (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() !== 'attributes' && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.kw_loc) && parser.getTokenSpelling().str() && parser.getTokenSpelling().str().indexOf('.') === -1 && !notParameterNames.has(parser.getTokenSpelling().str()))) {
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- continue;
- }
- const paramName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalLParen()) {
- // Check if this named parameter is actually an input from the operation metadata
- const inputNames = new Set((opInfo.metadata && opInfo.metadata.operands || []).map((i) => i.name));
- if (inputNames.has(paramName) && parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let operandType = null;
- if (parser.parseOptionalColon()) {
- operandType = parser.parseType();
- }
- parser.parseRParen();
- parser.resolveOperand(operand, operandType, result.operands);
- } else if (inputNames.has(paramName) && parser.getToken().is(_.Token.l_square)) {
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- parser.parseRParen();
- } else {
- let parenDepth = 1;
- let paramValue = '';
- while (parenDepth > 0 && parser.getToken().isNot(_.Token.eof)) {
- if (parser.getToken().is(_.Token.l_paren)) {
- parenDepth++;
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else if (parser.getToken().is(_.Token.r_paren)) {
- parenDepth--;
- if (parenDepth > 0) {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else {
- parser.parseRParen();
- }
- } else {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- }
- }
- result.addAttribute(paramName, paramValue);
- }
- } else {
- // Not a named parameter - we've consumed an id token that doesn't belong to us
- // This shouldn't happen with proper MLIR, but break gracefully
- break;
- }
- }
- result.addTypes(parser.parseOptionalColonTypeList());
- // Handle old IREE format: !hal.buffer{%size} where {%size} follows the type
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- if (parser.parseOptionalEqual()) {
- const value = parser.parseAttribute();
- result.addAttribute('default', value.value);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- return true;
- }
- if (result.op === 'hal.executable.condition' || result.op === 'hal.executable.constant.block') {
- const sig = parser.parseFunctionSignatureWithArguments(false);
- const argTypes = sig.arguments.map((a) => a.type);
- const type = new _.FunctionType(argTypes, sig.resultTypes);
- result.addAttribute('function_type', new _.TypeAttrOf(type));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, sig.arguments);
- }
- return true;
- }
- // Handle operations with visibility + symbol (similar to flow dialect)
- if (result.op === 'hal.executable' || result.op === 'hal.executable.source' || result.op === 'hal.interface' || result.op === 'hal.executable.binary') {
- result.compatibility = true;
- this.parseSymbolVisibility(parser, result);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- // Handle hal.interface.binding.subspan with old syntax (symbol reference)
- // Old syntax: hal.interface.binding.subspan @io::@binding[operand] : type
- // New syntax: hal.interface.binding.subspan layout(...) binding(...) : type
- if (result.op === 'hal.interface.binding.subspan' && parser.getToken().is(_.Token.at_identifier)) {
- result.compatibility = true;
- // Old syntax - parse symbol reference and bracket expression
- const symbolRef = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('layout', symbolRef);
- const unresolvedOperands = [];
- const indexType = new _.IndexType();
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- } else {
- parser.consumeToken();
- }
- parser.parseOptionalComma();
- }
- }
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => indexType), result.operands);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type) {
- result.addTypes([type]);
- }
- if (parser.parseOptionalLBrace()) {
- const dynamicDimOperands = [];
- if (parser.getToken().isNot(_.Token.r_brace)) {
- do {
- const dimOperand = parser.parseOperand();
- dynamicDimOperands.push(dimOperand);
- } while (parser.parseOptionalComma());
- }
- parser.parseRBrace();
- parser.resolveOperands(dynamicDimOperands, dynamicDimOperands.map(() => indexType), result.operands);
- }
- }
- return true;
- }
- // Handle operations with named parameters: hal.interface.binding, hal.executable.variant, etc.
- if (result.op === 'hal.interface.binding' || result.op === 'hal.executable.variant' || result.op === 'hal.executable.entry_point' || result.op === 'hal.executable.export') {
- result.compatibility = true;
- this.parseSymbolVisibility(parser, result);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- parser.parseOptionalComma();
- }
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() !== 'attributes' && parser.getToken().isNot(_.Token.l_brace) && parser.getToken().isNot(_.Token.kw_loc)) {
- const tokenValue = parser.getTokenSpelling().str();
- if (tokenValue && tokenValue.includes('.')) {
- break;
- }
- const paramName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (paramName === 'condition') {
- parser.parseLParen();
- const regionArgs = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- regionArgs.push({ value: arg, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseArrow();
- parser.parseType();
- const conditionRegion = { arguments: regionArgs };
- parser.parseRegion(conditionRegion);
- result.regions.push(conditionRegion);
- continue;
- }
- if (parser.parseOptionalLParen()) {
- let parenDepth = 1;
- let paramValue = '';
- while (parenDepth > 0 && parser.getToken().isNot(_.Token.eof)) {
- if (parser.getToken().is(_.Token.l_paren)) {
- parenDepth++;
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else if (parser.getToken().is(_.Token.r_paren)) {
- parenDepth--;
- if (parenDepth > 0) {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- } else {
- parser.parseRParen();
- }
- } else {
- paramValue += parser.getToken().getSpelling().str();
- parser.consumeToken();
- }
- }
- result.addAttribute(paramName, paramValue);
- parser.parseOptionalComma();
- } else if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.hash_identifier)) {
- const value = parser.parseAttribute();
- result.addAttribute(paramName, value.value);
- } else if (parser.getToken().is(_.Token.string)) {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute(paramName, value);
- } else {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken();
- result.addAttribute(paramName, value);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else {
- break;
- }
- }
- if (parser.parseOptionalArrow()) {
- const resultTypes = [];
- const resultAttrs = [];
- parser.parseFunctionResultList(resultTypes, resultAttrs);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalKeyword('count')) {
- this.parseWorkgroupCountRegion(parser, result);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalKeyword('count')) {
- this.parseWorkgroupCountRegion(parser, result);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parsePipelineLayoutAttr(parser) {
- // HAL_PipelineLayoutAttr format: <constants = N, bindings = [...], flags = ...>
- if (parser.getToken().is(_.Token.less)) {
- return parser.parseAttribute();
- }
- return parser.parseOptionalAttribute();
- }
- parseExportConditionRegion(parser, result) {
- parser.parseLParen();
- const regionArgs = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- regionArgs.push({ value: arg, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseArrow();
- parser.parseType();
- const region = { arguments: regionArgs };
- parser.parseRegion(region);
- result.regions.push(region);
- }
- parseTargetConditionObjects(parser, result) {
- // #target if(...) { region } ordinal(N) = [objects], ...
- do {
- if (parser.getToken().is(_.Token.hash_identifier)) {
- parser.parseAttribute();
- }
- if (parser.parseOptionalKeyword('if')) {
- this.parseTargetConditionRegion(parser, result);
- }
- if (parser.parseOptionalKeyword('ordinal')) {
- parser.parseLParen();
- parser.consumeToken(_.Token.integer);
- parser.parseRParen();
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- }
- } while (parser.parseOptionalComma());
- }
- parseTargetConditionRegion(parser, result) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.eof)) {
- parser.parseOperand();
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- if (parser.parseOptionalArrow()) {
- parser.parseType();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- }
- parseWorkgroupCountRegion(parser, result) {
- // (args) -> (index, index, index) { region }
- const region = { blocks: [] };
- const block = { arguments: [], operations: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.eof)) {
- const arg = parser.parseOperand();
- if (parser.parseOptionalColon()) {
- arg.type = parser.parseType();
- }
- block.arguments.push(arg);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.eof)) {
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- region.blocks.push(block);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region);
- }
- result.regions.push(region);
- }
- };
- _.IREECodegenDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'iree_codegen');
- }
- inferResultTypes(op, vars) {
- if (op.op === 'iree_codegen.inner_tiled') {
- const outputsEntry = vars.get('outputs');
- if (outputsEntry.types.length > 0) {
- op.addTypes(outputsEntry.types);
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- parseOperation(parser, result) {
- if (result.op === 'iree_codegen.workgroup_count_hint') {
- const staticSizes = [];
- const unresolvedSizes = [];
- parser.parseOptionalKeyword('sizes');
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedSizes.push(parser.parseOperand());
- staticSizes.push(-9223372036854775808);
- } else if (parser.getToken().is(_.Token.integer)) {
- const constValue = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- staticSizes.push(parseInt(constValue, 10));
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- // Resolve operands with index type (sizes are typically index)
- const indexType = new _.IndexType();
- for (const unresolved of unresolvedSizes) {
- parser.resolveOperand(unresolved, indexType, result.operands);
- }
- if (staticSizes.length > 0) {
- result.addAttribute('static_sizes', staticSizes);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.HALLoaderDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'hal_loader');
- this.registerCustomDirective('DispatchBindings', this.parseDispatchBindings.bind(this));
- }
- parseDispatchBindings(parser, result) {
- const unresolvedBuffers = [];
- const bufferTypes = [];
- const unresolvedOffsets = [];
- const unresolvedLengths = [];
- do {
- parser.parseLParen();
- unresolvedBuffers.push(parser.parseOperand());
- parser.parseColon();
- bufferTypes.push(parser.parseType());
- parser.parseRParen();
- parser.parseLSquare();
- unresolvedOffsets.push(parser.parseOperand());
- parser.parseComma();
- unresolvedLengths.push(parser.parseOperand());
- parser.parseRSquare();
- } while (parser.parseOptionalComma());
- const indexType = new _.IndexType();
- for (let i = 0; i < unresolvedBuffers.length; i++) {
- parser.resolveOperand(unresolvedBuffers[i], bufferTypes[i], result.operands);
- }
- for (const unresolved of unresolvedOffsets) {
- parser.resolveOperand(unresolved, indexType, result.operands);
- }
- for (const unresolved of unresolvedLengths) {
- parser.resolveOperand(unresolved, indexType, result.operands);
- }
- }
- };
- _.UtilDialect = class extends _.IREEDialect {
- constructor(operations) {
- super(operations, 'util');
- this.registerCustomDirective('OperandTypeList', this.parseOperandTypeList.bind(this));
- this.registerCustomDirective('TiedFunctionResultList', this.parseTiedFunctionResultList.bind(this));
- this.registerCustomDirective('TypeAlias', this.parseTypeAlias.bind(this));
- this.registerCustomDirective('TypedValueList', this.parseTypedValueList.bind(this));
- this.registerCustomDirective('RangeList', this.parseRangeList.bind(this));
- this.registerCustomDirective('ListTypeGet', this.parseListTypeGet.bind(this));
- this.registerCustomDirective('ListTypeSet', this.parseListTypeSet.bind(this));
- this.registerCustomDirective('ValueTypeList', this.parseValueTypeList.bind(this));
- this.simpleTypes = new Set(['buffer', 'list', 'object', 'ptr']);
- }
- parseTypeAlias(parser /*, op, args */) {
- parser.parseType();
- if (parser.parseOptionalKeyword('as')) {
- parser.parseType();
- }
- }
- parseTypedValueList(parser, op /*, args */) {
- parser.parseLSquare();
- if (parser.getToken().isNot(_.Token.r_square)) {
- const unresolvedValues = [];
- do {
- unresolvedValues.push(parser.parseOperand());
- } while (parser.parseOptionalComma());
- for (const unresolved of unresolvedValues) {
- parser.resolveOperand(unresolved, null, op.operands);
- }
- }
- parser.parseRSquare();
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- if (this.simpleTypes.has(typeName)) {
- if (typeName === 'list' && parser.parseOptionalLess()) {
- let elementType = null;
- if (parser.consumeIf(_.Token.question)) {
- elementType = new _.util.VariantType();
- } else {
- elementType = parser.parseType();
- }
- parser.parseGreater();
- return new _.util.ListType(elementType);
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- return null;
- }
- parseOperandTypeList(parser, op /*, args */) {
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- let index = 0;
- do {
- const type = parser.parseType();
- if (index < op.operands.length) {
- op.operands[index].type = type;
- }
- index++;
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- }
- parseTiedFunctionResultList(parser, op /*, args */) {
- const parseTiedResultOrType = () => {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const tiedRef = parser.parseOperand();
- let tiedType = null;
- for (let i = 0; i < op.operands.length; i++) {
- if (op.operands[i].value === tiedRef) {
- tiedType = op.operands[i].type;
- break;
- }
- }
- if (parser.parseOptionalKeyword('as')) {
- return parser.parseType();
- }
- if (tiedType) {
- return tiedType;
- }
- return new _.Type('!util.unknown');
- }
- return parser.parseType();
- };
- if (parser.parseOptionalLParen()) {
- let index = 0;
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const type = parseTiedResultOrType();
- if (index < op.types.length) {
- op.types[index] = type;
- } else {
- op.addTypes([type]);
- }
- index++;
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- } else {
- let index = 0;
- do {
- const type = parseTiedResultOrType();
- if (index < op.types.length) {
- op.types[index] = type;
- } else {
- op.addTypes([type]);
- }
- index++;
- } while (parser.parseOptionalComma());
- }
- }
- parseOperation(parser, result) {
- if (result.op === 'util.assume.int') {
- return this.parseAssumeIntOp(parser, result);
- }
- if (result.op === 'util.initializer') {
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'util.unreachable') {
- result.compatibility = true;
- if (parser.getToken().is(_.Token.string)) {
- const message = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('message', message);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- return true;
- }
- if (result.op === 'util.func') {
- this.parseUtilFuncOp(parser, result);
- return true;
- }
- if (result.op === 'util.unfoldable_constant') {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const value = parser.parseAttribute();
- result.addAttribute('value', value);
- // Use type from attribute if present, otherwise parse explicit type
- if (value && value.type) {
- result.addTypes([value.type]);
- } else if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseUtilFuncOp(parser, result) {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- parser.parseSymbolName('sym_name', result.attributes);
- const argResult = parser.parseFunctionArgumentList(false);
- const resultTypes = [];
- const resultAttrs = [];
- const tiedOperandIndices = [];
- // Parse result list which may contain:
- // - Regular type: tensor<...>
- // - Tied reference: %arg1 (inherits type from argument)
- // - Tied with type override: %arg2 as tensor<...>
- const parseTiedResultOrType = () => {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const tiedRef = parser.parseOperand();
- let tiedIndex = -1;
- for (let i = 0; i < argResult.arguments.length; i++) {
- if (argResult.arguments[i].value === tiedRef) {
- tiedIndex = i;
- break;
- }
- }
- tiedOperandIndices.push(tiedIndex);
- // Check for 'as type' override
- if (parser.parseOptionalKeyword('as')) {
- return parser.parseType();
- }
- if (tiedIndex >= 0 && argResult.arguments[tiedIndex].type) {
- return argResult.arguments[tiedIndex].type;
- }
- return new _.Type('!util.unknown');
- }
- tiedOperandIndices.push(-1);
- return parser.parseType();
- };
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- resultTypes.push(parseTiedResultOrType());
- if (parser.getToken().is(_.Token.l_brace)) {
- const attrList = new Map();
- parser.parseAttributeDict(attrList);
- resultAttrs.push(attrList);
- } else {
- resultAttrs.push(null);
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- } else {
- do {
- resultTypes.push(parseTiedResultOrType());
- resultAttrs.push(null);
- } while (parser.parseOptionalComma());
- }
- }
- const argTypes = argResult.arguments.filter((a) => a.value !== '...').map((a) => a.type);
- const type = new _.FunctionType(argTypes, resultTypes);
- result.addAttribute('function_type', new _.TypeAttrOf(type));
- if (tiedOperandIndices.some((i) => i >= 0)) {
- result.addAttribute('tied_operands', tiedOperandIndices);
- }
- if (resultAttrs.some((a) => a !== null)) {
- result.addAttribute('res_attrs', resultAttrs);
- }
- const argAttrs = argResult.arguments.filter((a) => a.value !== '...').map((a) => a.attrs || null);
- if (argAttrs.some((a) => a !== null)) {
- result.addAttribute('arg_attrs', argAttrs);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, argResult.arguments);
- }
- }
- parseAssumeIntOp(parser, result) {
- const allOperandAssumptions = [];
- const unresolvedOperands = [];
- do {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- const operandAssumptions = [];
- if (parser.parseOptionalLSquare()) {
- if (parser.getToken().isNot(_.Token.r_square)) {
- do {
- const assumption = this.parseIntAssumptionAttr(parser);
- operandAssumptions.push(assumption);
- } while (parser.parseOptionalComma());
- }
- parser.parseRSquare();
- } else if (parser.getToken().is(_.Token.less)) {
- const assumption = this.parseIntAssumptionAttr(parser);
- operandAssumptions.push(assumption);
- }
- allOperandAssumptions.push(operandAssumptions);
- } while (parser.parseOptionalComma());
- parser.parseColon();
- const parsedOperandTypes = [];
- do {
- const type = parser.parseType();
- parsedOperandTypes.push(type);
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedOperands, parsedOperandTypes, result.operands);
- for (const type of parsedOperandTypes) {
- result.addTypes([type || null]);
- }
- result.addAttribute('assumptions', allOperandAssumptions);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- return true;
- }
- parseIntAssumptionAttr(parser) {
- parser.parseLess();
- const assumption = {};
- if (parser.getToken().isNot(_.Token.greater)) {
- do {
- const key = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (!parser.parseOptionalEqual()) {
- throw new mlir.Error(`Expected '=' after ${key} ${parser.location()}`);
- }
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- assumption[key] = value;
- } while (parser.parseOptionalComma());
- }
- parser.parseGreater();
- return assumption;
- }
- parseRangeList(parser, op, offsetsAttr) {
- const unresolvedOffsets = [];
- const unresolvedLengths = [];
- do {
- parser.parseLSquare();
- unresolvedOffsets.push(parser.parseOperand());
- parser.parseKeyword('for');
- unresolvedLengths.push(parser.parseOperand());
- parser.parseRSquare();
- } while (parser.parseOptionalComma());
- const indexType = new _.IndexType();
- for (const unresolved of unresolvedOffsets) {
- parser.resolveOperand(unresolved, indexType, op.operands);
- }
- for (const unresolved of unresolvedLengths) {
- parser.resolveOperand(unresolved, indexType, op.operands);
- }
- if (offsetsAttr) {
- op.addAttribute(`${offsetsAttr}_count`, unresolvedOffsets.length);
- }
- }
- parseListTypeGet(parser, op, listTypeArr, resultTypeArr) {
- const listType = parser.parseType();
- let elementType = null;
- if (parser.parseOptionalArrow()) {
- elementType = parser.parseType();
- } else if (listType instanceof _.util.ListType) {
- elementType = listType.elementType;
- }
- if (Array.isArray(listTypeArr) && listType) {
- listTypeArr.push(listType);
- }
- if (Array.isArray(resultTypeArr) && elementType) {
- resultTypeArr.push(elementType);
- }
- }
- parseListTypeSet(parser, op, listTypeArr, valueTypeArr) {
- const leadingType = parser.parseType();
- let listType = null;
- let elementType = null;
- if (parser.parseOptionalArrow()) {
- elementType = leadingType;
- listType = parser.parseType();
- } else if (leadingType instanceof _.util.ListType) {
- listType = leadingType;
- elementType = leadingType.elementType;
- }
- if (Array.isArray(listTypeArr) && listType) {
- listTypeArr.push(listType);
- }
- if (Array.isArray(valueTypeArr) && elementType) {
- valueTypeArr.push(elementType);
- }
- }
- parseValueTypeList(parser, result) {
- parser.parseLSquare();
- if (parser.getToken().isNot(_.Token.r_square)) {
- const unresolvedOperands = [];
- const types = [];
- do {
- unresolvedOperands.push(parser.parseOperand());
- parser.parseColon();
- types.push(parser.parseType());
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- }
- parser.parseRSquare();
- }
- };
- _.FlowDialect = class extends _.IREEDialect {
- constructor(operations) {
- super(operations, 'flow');
- this.registerCustomDirective('DispatchWorkgroupBody', this.parseDispatchWorkgroupBody.bind(this));
- this.registerCustomDirective('DispatchWorkgroupsCountRegion', this.parseDispatchWorkgroupsCountRegion.bind(this));
- this.registerCustomDirective('ShapedFunctionType', this.parseShapedFunctionType.bind(this));
- this.registerCustomDirective('ShapedOperandList', this.parseShapedOperandList.bind(this));
- this.registerCustomDirective('ParameterReference', this.parseParameterReference.bind(this));
- }
- parseParameterReference(parser, op, scopeOperands, keyOperands) {
- const firstOperand = parser.parseOperand();
- if (parser.parseOptionalColon()) {
- parser.parseColon();
- const keyOperand = parser.parseOperand();
- scopeOperands.push(firstOperand);
- keyOperands.push(keyOperand);
- } else {
- keyOperands.push(firstOperand);
- }
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (typeName === 'channel') {
- return new _.Type(type);
- }
- if (typeName === 'dispatch.tensor') {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op === 'flow.ex.stream.fragment') {
- return this.parseDispatchWorkgroupsOp(parser, result);
- }
- if (result.op === 'flow.dispatch.region') {
- return this.parseDispatchRegionOp(parser, result);
- }
- if (result.op === 'flow.dispatch.tensor.load' || result.op === 'flow.dispatch.tensor.store') {
- return this.parseTensorLoadStoreOp(parser, result);
- }
- // Handle operations with visibility + symbol that aren't in schema or need manual parsing
- if (result.op === 'flow.dispatch.entry') {
- this.parseSymbolVisibility(parser, result);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'flow.func') {
- return this.parseFlowFuncOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseFlowFuncOp(parser, result) {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- parser.parseSymbolName('sym_name', result.attributes);
- const argResult = parser.parseFunctionArgumentList();
- const inputs = argResult.arguments.map((a) => a.type);
- const results = [];
- if (parser.parseOptionalArrow()) {
- const hasParens = parser.parseOptionalLParen();
- if (!hasParens || parser.getToken().isNot(_.Token.r_paren)) {
- do {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- if (parser.parseOptionalKeyword('as')) {
- const resultType = parser.parseType();
- results.push(resultType);
- } else {
- results.push(new _.Type('tied'));
- }
- } else {
- const resultType = parser.parseType();
- results.push(resultType);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- if (!hasParens) {
- break;
- }
- } while (parser.parseOptionalComma());
- }
- if (hasParens) {
- parser.parseRParen();
- }
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- parseDispatchRegionOp(parser, result) {
- const workloadOperands = parser.parseOperandList('optionalSquare');
- for (const workload of workloadOperands) {
- parser.resolveOperand(workload, null, result.operands);
- }
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const type = parser.parseType();
- if (parser.parseOptionalLBrace()) {
- while (!parser.parseOptionalRBrace()) {
- const tied = parser.parseOperand();
- parser.resolveOperand(tied, null, result.operands);
- parser.parseOptionalComma();
- }
- }
- result.types.push(type);
- parser.parseOptionalComma();
- }
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- this.parseDispatchWorkgroupsCountRegion(parser, result);
- return true;
- }
- parseDispatchWorkgroupsOp(parser, result) {
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.consumeToken(); // read subscript value
- parser.parseOptionalComma();
- }
- }
- const unresolvedOperands = parser.parseOperandList('paren');
- if (parser.parseOptionalColon()) {
- if (parser.parseOptionalLParen()) {
- const inputTypes = parser.parseTypeListNoParens();
- parser.parseRParen();
- parser.resolveOperands(unresolvedOperands, inputTypes, result.operands);
- parser.parseArrow();
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- } else {
- const types = parser.parseTypeListNoParens();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- }
- } else if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalEqual()) {
- const args = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const arg = parser.parseArgument(true, false);
- args.push(arg);
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow() || parser.parseOptionalKeyword('to')) {
- parser.parseType();
- }
- const region = result.addRegion();
- parser.parseRegion(region, args, /* enableNameShadowing */ true);
- }
- return true;
- }
- parseShapedFunctionType(parser, op, unresolvedArguments /*, otherArgs */) {
- // unresolvedArguments is ref($arguments) - array of unresolved operands with .name and .number
- const operandTypes = [];
- if (parser.parseOptionalLParen()) {
- let index = 0;
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const type = parser.parseType();
- if (type) {
- operandTypes.push(type);
- const startIdx = Math.max(0, op.operands.length - (index + 1));
- if (startIdx + index < op.operands.length && !op.operands[startIdx + index].type) {
- op.operands[startIdx + index].type = type;
- }
- index++;
- }
- if (parser.parseOptionalLBrace()) {
- while (!parser.parseOptionalRBrace()) {
- parser.parseOperand();
- parser.parseOptionalComma();
- }
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- let index = 0;
- const hasParens = parser.parseOptionalLParen();
- if (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.l_brace) && parser.getToken().isNot(_.Token.kw_loc) && parser.getToken().isNot(_.Token.equal)) {
- do {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const tiedResult = parser.parseOperand();
- // Handle optional "as type" for tied results
- if (parser.parseOptionalKeyword('as')) {
- const type = parser.parseType();
- if (type) {
- if (index < op.types.length) {
- op.types[index] = type;
- } else {
- op.addTypes([type]);
- }
- }
- } else {
- // Look up type from tied operand using unresolvedArguments
- // unresolvedArguments contains UnresolvedOperand objects with .name and .number
- const operands = unresolvedArguments || [];
- const tiedOperandIndex = parser.findTiedOperand(tiedResult, operands);
- if (tiedOperandIndex >= 0 && tiedOperandIndex < operandTypes.length) {
- const type = operandTypes[tiedOperandIndex];
- if (type) {
- if (index < op.types.length) {
- op.types[index] = type;
- } else {
- op.addTypes([type]);
- }
- }
- }
- }
- index++;
- } else {
- const type = parser.parseType();
- if (type) {
- if (index < op.types.length) {
- op.types[index] = type;
- } else {
- op.addTypes([type]);
- }
- index++;
- }
- }
- if (parser.parseOptionalLBrace()) {
- while (!parser.parseOptionalRBrace()) {
- parser.parseOperand();
- parser.parseOptionalComma();
- }
- }
- if (!hasParens) {
- break;
- }
- } while (parser.parseOptionalComma());
- }
- if (hasParens) {
- parser.parseRParen();
- }
- }
- }
- parseTensorLoadStoreOp(parser, result) {
- // or: store %26, %arg4, offsets = [...] : type -> type
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- break;
- }
- }
- // Note: first parameter might not need comma-eating if we just broke from operand loop
- let needComma = parser.getToken().isNot(_.Token.bare_identifier); // If we're not at _.Token.bare_identifier, we need to eat commas
- while (needComma ? parser.parseOptionalComma() : true) {
- needComma = true; // After first iteration, always need comma
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const paramName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- } else {
- parser.consumeToken();
- }
- result.addAttribute(paramName, paramName);
- }
- } else {
- break;
- }
- }
- const types = parser.parseOptionalColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- // For tensor.load, there's a -> result type
- // For tensor.store, the -> is followed by the output tensor type (not a result)
- if (parser.parseOptionalArrow() || parser.parseOptionalKeyword('to')) {
- const resultType = parser.parseType();
- if (result.op === 'flow.dispatch.tensor.load' && resultType) {
- result.addTypes([resultType]);
- }
- }
- return true;
- }
- parseDispatchWorkgroupBody(parser, op /*, args */) {
- parser.parseLParen();
- const regionArgs = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const arg = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- regionArgs.push({ name: arg, type: argType });
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- const region = { blocks: [{ arguments: regionArgs, operations: [] }] };
- parser.parseRegion(region);
- op.regions.push(region);
- }
- parseDispatchWorkgroupsCountRegion(parser, op /*, args */) {
- if (!parser.parseOptionalKeyword('count')) {
- return;
- }
- parser.parseLParen();
- const regionArgs = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const arg = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- regionArgs.push({ name: arg, type: argType });
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- parser.parseArrow();
- if (parser.parseOptionalLParen()) {
- parser.parseType();
- parser.parseOptionalComma();
- parser.parseType();
- parser.parseOptionalComma();
- parser.parseType();
- parser.parseRParen();
- } else {
- parser.parseType();
- parser.parseOptionalComma();
- parser.parseType();
- parser.parseOptionalComma();
- parser.parseType();
- }
- const region = { blocks: [{ arguments: regionArgs, operations: [] }] };
- parser.parseRegion(region);
- op.regions.push(region);
- }
- parseShapedOperandList(parser, result) {
- const unresolvedValues = [];
- const valueTypes = [];
- const unresolvedDims = [];
- do {
- unresolvedValues.push(parser.parseOperand());
- parser.parseColon();
- const valueType = parser.parseType();
- valueTypes.push(valueType);
- if (valueType) {
- const typeStr = valueType.toString();
- const dynamicDimCount = (typeStr.match(/\?/g) || []).length;
- if (dynamicDimCount > 0 && parser.parseOptionalLBrace()) {
- for (let i = 0; i < dynamicDimCount; i++) {
- if (i > 0) {
- parser.parseOptionalComma();
- }
- unresolvedDims.push(parser.parseOperand());
- }
- parser.parseRBrace();
- }
- }
- } while (parser.parseOptionalComma());
- for (let i = 0; i < unresolvedValues.length; i++) {
- parser.resolveOperand(unresolvedValues[i], valueTypes[i], result.operands);
- }
- const indexType = new _.IndexType();
- for (const unresolved of unresolvedDims) {
- parser.resolveOperand(unresolved, indexType, result.operands);
- }
- }
- };
- _.StreamDialect = class extends _.IREEDialect {
- constructor(operations, name = 'stream') {
- super(operations, name);
- this.registerCustomDirective('DispatchOperands', this.parseDispatchOperands.bind(this));
- this.registerCustomDirective('DispatchResources', this.parseDispatchResources.bind(this));
- this.registerCustomDirective('ExplicitResourceRegion', this.parseExplicitResourceRegion.bind(this));
- this.registerCustomDirective('ShapedTypeList', this.parseShapedTypeList.bind(this));
- this.registerCustomDirective('ResourceRegion', this.parseResourceRegion.bind(this));
- this.registerCustomDirective('ParameterLoadOperations', this.parseParameterLoadOperations.bind(this));
- this.registerCustomDirective('EncodedResourceOperands', this.parseEncodedResourceOperands.bind(this));
- this.registerCustomDirective('DispatchEntryPoints', this.parseDispatchEntryPoints.bind(this));
- this.registerCustomDirective('ShapedTiedResult', this.parseShapedTiedResult.bind(this));
- this.registerCustomDirective('EncodedShapedFunctionType', this.parseEncodedShapedFunctionType.bind(this));
- this.registerCustomDirective('CollectiveParam', this.parseCollectiveParam.bind(this));
- this.registerCustomDirective('PackSliceRanges', this.parsePackSliceRanges.bind(this));
- this.registerCustomDirective('WorkgroupCountRegion', this.parseWorkgroupCountRegion.bind(this));
- this.registerCustomDirective('DispatchFunctionSignature', this.parseDispatchFunctionSignature.bind(this));
- this.registerCustomDirective('ShapedFunctionSignature', this.parseShapedFunctionSignature.bind(this));
- this.registerCustomDirective('ConstantValueList', this.parseConstantValueList.bind(this));
- this.registerCustomDirective('CmdCallOperands', this.parseCmdCallOperands.bind(this));
- this.registerCustomDirective('ParameterReference', this.parseParameterReference.bind(this));
- this.registerCustomDirective('ParameterGatherOperations', this.parseParameterGatherOperations.bind(this));
- this.registerCustomDirective('ParameterScatterOperations', this.parseParameterScatterOperations.bind(this));
- this.registerCustomDirective('SymbolAlias', this.parseSymbolAlias.bind(this));
- }
- parseDispatchResources(parser, op /*, args */) {
- do {
- const accessMode = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- const unresolvedResource = parser.parseOperand();
- parser.parseLSquare();
- const unresolvedOffset = parser.parseOperand();
- parser.parseKeyword('for');
- const unresolvedLength = parser.parseOperand();
- parser.parseRSquare();
- parser.parseColon();
- const resourceType = parser.parseType();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- op.addAttribute('resource_access', accessMode);
- parser.resolveOperand(unresolvedResource, resourceType, op.operands);
- const indexType = new _.IndexType();
- parser.resolveOperand(unresolvedOffset, indexType, op.operands);
- parser.resolveOperand(unresolvedLength, indexType, op.operands);
- } while (parser.parseOptionalComma());
- }
- parseShapedTypeList(parser, op, ...args) {
- // 2-arg: custom<ShapedTypeList>(type($operands), $sizes)
- // 3-arg: custom<ShapedTypeList>(type($operands), type($results), $sizes)
- let operandTypes = null;
- let resultTypes = null;
- let sizeOperands = null;
- if (args.length === 2) {
- [operandTypes, sizeOperands] = args;
- } else if (args.length >= 3) {
- [operandTypes, resultTypes, sizeOperands] = args;
- } else {
- operandTypes = args[0] || null;
- }
- const indexType = new _.IndexType();
- do {
- const type = parser.parseType();
- if (operandTypes) {
- operandTypes.push(type);
- }
- if (resultTypes) {
- resultTypes.push(type);
- }
- if (parser.parseOptionalLBrace()) {
- do {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const sizeOperand = parser.parseOperand();
- if (sizeOperands) {
- sizeOperands.push(sizeOperand);
- }
- parser.resolveOperand(sizeOperand, indexType, op.operands);
- }
- } while (parser.parseOptionalComma());
- parser.parseRBrace();
- }
- } while (parser.parseOptionalComma());
- }
- parseExplicitResourceRegion(parser, op /*, args */) {
- parser.parseLParen();
- const regionArgs = [];
- const unresolvedOperands = [];
- const operandTypes = [];
- const unresolvedSizes = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseKeyword('as');
- const arg = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- operandTypes.push(argType);
- regionArgs.push({ name: arg, type: argType });
- if (parser.parseOptionalLBrace()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedSizes.push(parser.parseOperand());
- }
- parser.parseRBrace();
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- for (let i = 0; i < unresolvedOperands.length; i++) {
- parser.resolveOperand(unresolvedOperands[i], operandTypes[i] || null, op.operands);
- }
- const indexType = new _.IndexType();
- for (const unresolved of unresolvedSizes) {
- parser.resolveOperand(unresolved, indexType, op.operands);
- }
- const region = { blocks: [{ arguments: regionArgs, operations: [] }] };
- parser.parseRegion(region);
- op.regions.push(region);
- }
- parseResourceRegion(parser, op /*, args */) {
- const regionArgs = [];
- const unresolvedOperands = [];
- const operandTypes = [];
- const unresolvedSizes = [];
- const indexType = new _.IndexType();
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- parser.parseKeyword('as');
- const arg = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- operandTypes.push(argType);
- regionArgs.push({ name: arg, type: argType });
- if (parser.parseOptionalLBrace()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedSizes.push(parser.parseOperand());
- }
- parser.parseRBrace();
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- for (let i = 0; i < unresolvedOperands.length; i++) {
- parser.resolveOperand(unresolvedOperands[i], operandTypes[i], op.operands);
- }
- for (const unresolved of unresolvedSizes) {
- parser.resolveOperand(unresolved, indexType, op.operands);
- }
- const resultSizes = [];
- const parseResultTypeOrTied = () => {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- if (parser.parseOptionalKeyword('as')) {
- const resultType = parser.parseType();
- op.addTypes([resultType]);
- } else {
- op.addTypes([new _.Type('tied')]);
- }
- } else {
- const resultType = parser.parseType();
- op.addTypes([resultType]);
- }
- if (parser.parseOptionalLBrace()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- resultSizes.push(parser.parseOperand());
- }
- parser.parseRBrace();
- }
- };
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- parseResultTypeOrTied();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- } else {
- parseResultTypeOrTied();
- }
- }
- for (const unresolved of resultSizes) {
- parser.resolveOperand(unresolved, indexType, op.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = { blocks: [{ arguments: regionArgs, operations: [] }] };
- parser.parseRegion(region);
- op.regions.push(region);
- }
- }
- parseParameterLoadOperations(parser, op, sourceScopeOps, sourceKeysOps, sourceOffsetsOps, resultTypes, resultSizesOps) {
- do {
- const firstOperand = parser.parseOperand();
- let key = firstOperand;
- if (parser.parseOptionalColon()) {
- sourceScopeOps.push(firstOperand);
- parser.parseColon();
- key = parser.parseOperand();
- }
- sourceKeysOps.push(key);
- parser.parseLSquare();
- sourceOffsetsOps.push(parser.parseOperand());
- parser.parseRSquare();
- parser.parseColon();
- const resultType = parser.parseType();
- op.addTypes([resultType]);
- if (parser.parseOptionalLBrace()) {
- resultSizesOps.push(parser.parseOperand());
- parser.parseRBrace();
- }
- } while (parser.parseOptionalComma());
- }
- parseEncodedResourceOperands(parser /*, op, args */) {
- do {
- parser.parseOperand();
- parser.parseColon();
- parser.parseType();
- parser.skip('{');
- parser.parseKeyword('in');
- parser.parseType();
- parser.skip('{');
- } while (parser.parseOptionalComma());
- }
- parseDispatchEntryPoints(parser, op /*, args */) {
- if (parser.parseOptionalLBrace()) {
- do {
- const symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- op.addAttribute('entry_point', symbol);
- } while (parser.parseOptionalComma());
- parser.parseRBrace();
- } else {
- const symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- op.addAttribute('entry_point', symbol);
- }
- }
- parseShapedTiedResult(parser, op /*, args */) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // tiedOperand - parsed but not stored in OperationState
- parser.parseKeyword('as');
- }
- const type = parser.parseType();
- op.types.push(type);
- if (parser.parseOptionalLBrace()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedSize = parser.parseOperand();
- const indexType = new _.IndexType();
- parser.resolveOperand(unresolvedSize, indexType, op.operands);
- }
- parser.parseRBrace();
- }
- }
- parseEncodedShapedTypeList(parser, types) {
- do {
- const type0 = parser.parseType();
- parser.skip('{');
- if (parser.parseOptionalKeyword('in')) {
- const type1 = parser.parseType();
- parser.skip('{');
- types.push(type1);
- } else {
- types.push(type0);
- }
- } while (parser.parseOptionalComma());
- }
- parseEncodedShapedResultList(parser, operands, operandTypes, resultTypes) {
- do {
- let type0 = null;
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- type0 = parser.parseType();
- parser.skip('{');
- }
- if (!parser.parseOptionalKeyword('in')) {
- if (type0) {
- resultTypes.push(type0);
- }
- continue;
- }
- let resultType = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const tiedResult = parser.parseOperand();
- const tiedOperandIndex = parser.findTiedOperand(tiedResult, operands);
- if (parser.parseOptionalKeyword('as')) {
- resultType = parser.parseType();
- } else if (tiedOperandIndex >= 0 && operandTypes[tiedOperandIndex]) {
- resultType = operandTypes[tiedOperandIndex];
- }
- } else {
- resultType = parser.parseType();
- }
- parser.skip('{');
- if (resultType) {
- resultTypes.push(resultType);
- }
- } while (parser.parseOptionalComma());
- }
- parseEncodedShapedFunctionType(parser, op, operandsRef, operandTypes /* , ... */) {
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- this.parseEncodedShapedTypeList(parser, operandTypes);
- }
- parser.parseRParen();
- parser.parseArrow();
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- this.parseEncodedShapedResultList(parser, operandsRef, operandTypes, op.types);
- }
- parser.parseRParen();
- } else {
- this.parseEncodedShapedResultList(parser, operandsRef, operandTypes, op.types);
- }
- }
- parseCollectiveParam(parser, op /*, args */) {
- for (const keyword of ['source', 'target', 'source_target_pair']) {
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === keyword) {
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseLParen();
- const unresolvedParam = parser.parseOperand();
- parser.resolveOperand(unresolvedParam, null, op.operands);
- parser.parseRParen();
- return;
- }
- }
- }
- parsePackSliceRanges(parser, op, lifetimeIntervalsName, dynamicSliceSizesOperands, packedOffsetsTypes) {
- const indexType = new _.IndexType();
- while (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseComma();
- parser.parseAttribute();
- parser.parseRSquare();
- parser.parseEqual();
- const unresolvedOperand = parser.parseOperand();
- parser.resolveOperand(unresolvedOperand, indexType, op.operands);
- packedOffsetsTypes.push(indexType);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parseWorkgroupCountRegion(parser, op /*, args */) {
- if (!parser.parseOptionalKeyword('workgroups')) {
- return;
- }
- const region = { blocks: [] };
- const block = { arguments: [], operations: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- if (parser.parseOptionalColon()) {
- arg.type = parser.parseType();
- }
- block.arguments.push(arg);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- region.blocks.push(block);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region);
- }
- op.regions.push(region);
- }
- parseDispatchFunctionSignature(parser, op /*, args */) {
- const inputs = [];
- const results = [];
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- parser.parseOperand();
- // skip('[', ']') already handles checking for '[' presence
- parser.skip('[');
- parser.parseColon();
- const type = parser.parseType();
- inputs.push(type);
- parser.skip('{');
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- const parseResultTypeOrTied = () => {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- if (parser.parseOptionalKeyword('as')) {
- return parser.parseType();
- }
- return new _.Type('tied');
- }
- return parser.parseType();
- };
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- results.push(parseResultTypeOrTied());
- parser.skip('{');
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- } else {
- results.push(parseResultTypeOrTied());
- parser.skip('{');
- }
- }
- op.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- }
- parseShapedFunctionSignature(parser, op /*, args */) {
- this.parseDispatchFunctionSignature(parser, op);
- }
- parseConstantValueList(parser, op /*, args */) {
- do {
- const resultType = parser.parseType();
- op.addTypes([resultType]);
- if (parser.parseOptionalLBrace()) {
- // Size is an SSA value like %c4, not an attribute
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolved = parser.parseOperand();
- parser.resolveOperand(unresolved, null, op.operands);
- } else {
- const size = parser.parseAttribute();
- if (size) {
- // If it's an integer literal, create operand directly
- op.operands.push(new _.Value(null, size));
- }
- }
- parser.parseRBrace();
- }
- parser.parseEqual();
- parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- } while (parser.parseOptionalComma());
- }
- parseCmdCallOperands(parser, op /*, args */) {
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- const indexType = new _.IndexType();
- do {
- // Check for access mode keyword (ro, rw, wo)
- const accessMode = parser.parseOptionalKeyword('ro') || parser.parseOptionalKeyword('rw') || parser.parseOptionalKeyword('wo');
- if (accessMode) {
- // Resource operand with offset/length: access operand[offset for length]
- const unresolvedResource = parser.parseOperand();
- parser.parseLSquare();
- const unresolvedOffset = parser.parseOperand();
- parser.parseKeyword('for');
- const unresolvedLength = parser.parseOperand();
- parser.parseRSquare();
- op.addAttribute('resource_access', accessMode);
- parser.resolveOperand(unresolvedResource, null, op.operands);
- parser.resolveOperand(unresolvedOffset, indexType, op.operands);
- parser.resolveOperand(unresolvedLength, indexType, op.operands);
- } else {
- // Primitive/custom operand
- const unresolvedOperand = parser.parseOperand();
- parser.resolveOperand(unresolvedOperand, null, op.operands);
- }
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- }
- parseParameterReference(parser, op, scopeOperands, keyOperands) {
- const firstOperand = parser.parseOperand();
- if (parser.parseOptionalColon()) {
- parser.parseColon();
- const keyOperand = parser.parseOperand();
- if (scopeOperands) {
- scopeOperands.push(firstOperand);
- }
- if (keyOperands) {
- keyOperands.push(keyOperand);
- }
- } else if (keyOperands) {
- keyOperands.push(firstOperand);
- }
- }
- parseParameterGatherOperations(parser /*, op, args */) {
- do {
- // %scope::%key[offset] -> %target[offset for length] : type{size}
- this.parseParameterReference(parser);
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseArrow();
- parser.parseOperand();
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseKeyword('for');
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseColon();
- parser.parseType();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- } while (parser.parseOptionalComma());
- }
- parseParameterScatterOperations(parser /*, op, args */) {
- do {
- // %source[offset for length] : type{size} -> %scope::%key[offset]
- parser.parseOperand();
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseKeyword('for');
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseColon();
- parser.parseType();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- parser.parseArrow();
- this.parseParameterReference(parser);
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- } while (parser.parseOptionalComma());
- }
- parseSymbolAlias(parser, op /*, args */) {
- parser.parseSymbolName('sym_name', op.attributes);
- if (parser.parseOptionalEqual()) {
- const ref = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- op.addAttribute('function_ref', ref);
- }
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- const simpleTypes = ['binding', 'channel', 'timepoint', 'file'];
- if (simpleTypes.includes(typeName)) {
- return new _.Type(type);
- }
- if (typeName === 'resource') {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- // Handle test.fence type (Stream_TestFence in StreamTypes.td)
- if (typeName === 'test') {
- if (parser.consumeIf('.')) {
- const subtype = parser.parseOptionalKeyword();
- if (subtype === 'fence') {
- return new _.Type(`!${dialect}.test.fence`);
- }
- // Handle unknown test.X subtypes generically
- return new _.Type(`!${dialect}.test.${subtype}`);
- }
- // Just "test" without subtype - return as is
- return new _.Type(type);
- }
- // Fallback for unknown stream types - parse generically like base Dialect
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- parseDispatchOperands(parser, op, resourceOperands /* offsets, ends, lengths */) {
- // args are: [$resource_operands, $resource_operand_offsets, $resource_operand_ends, $resource_operand_lengths]
- // resourceOperands is passed by ref so ShapedFunctionType can use it for tied operand lookup
- parser.parseLParen();
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseRParen();
- return;
- }
- const unresolvedOperands = [];
- do {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- if (Array.isArray(resourceOperands)) {
- resourceOperands.push(operand);
- }
- if (parser.parseOptionalLSquare()) {
- unresolvedOperands.push(parser.parseOperand()); // offset
- parser.parseKeyword('to');
- unresolvedOperands.push(parser.parseOperand()); // end
- parser.parseKeyword('for');
- unresolvedOperands.push(parser.parseOperand()); // length
- parser.parseRSquare();
- }
- } while (parser.parseOptionalComma());
- parser.parseRParen();
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, op.operands);
- }
- }
- };
- _.IOParametersDialect = class extends _.StreamDialect {
- constructor(operations) {
- super(operations, 'io_parameters');
- }
- };
- _.PCFDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'pcf');
- this.registerCustomDirective('ParallelExecutionBody', this.parseParallelExecutionBody.bind(this));
- this.registerCustomDirective('InferNumIndexArgs', this.parseInferNumIndexArgs.bind(this));
- }
- parseInferNumIndexArgs() {
- }
- parseParallelExecutionBody(parser, result) {
- const inits = [];
- const dynamicSizes = [];
- const resultTypes = [];
- const isTied = [];
- const regionRefArgs = [];
- const indexArgs = [];
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- result.addAttribute('num_leading_args', (result.attributes.get('num_leading_args') || 0) + 1);
- regionRefArgs.push({ value: arg, type: argType });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseKeyword('execute');
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const refArg = parser.parseOperand();
- regionRefArgs.push({ value: refArg });
- if (parser.parseOptionalEqual()) {
- const initOperand = parser.parseOperand();
- inits.push({ value: initOperand });
- isTied.push(true);
- } else {
- isTied.push(false);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- const indexArg = parser.parseOperand();
- parser.parseColon();
- const indexType = parser.parseType();
- indexArgs.push({ value: indexArg, type: indexType });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- if (regionRefArgs.length > 0 && parser.parseOptionalColon()) {
- parser.parseLParen();
- let refIdx = result.attributes.get('num_leading_args') || 0;
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const refType = parser.parseType();
- if (refIdx < regionRefArgs.length) {
- regionRefArgs[refIdx].type = refType;
- }
- refIdx++;
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseArrow();
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const resType = parser.parseType();
- resultTypes.push(resType);
- result.addTypes([resType]);
- if (parser.parseOptionalLBrace()) {
- while (parser.getToken().isNot(_.Token.r_brace)) {
- const dim = parser.parseOperand();
- dynamicSizes.push({ value: dim });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRBrace();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- for (const init of inits) {
- parser.resolveOperand(init.value, null, result.operands);
- }
- for (const dim of dynamicSizes) {
- parser.resolveOperand(dim.value, null, result.operands);
- }
- if (isTied.length > 0) {
- result.addAttribute('is_tied', isTied);
- }
- const region = { blocks: [{ arguments: [...regionRefArgs, ...indexArgs], operations: [] }] };
- parser.parseRegion(region);
- result.regions.push(region);
- }
- };
- _.IREEVectorExtDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'iree_vector_ext');
- }
- parseOperation(parser, result) {
- if (result.op === 'iree_vector_ext.transfer_gather') {
- const unresolvedSource = parser.parseOperand();
- const unresolvedIndices = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedIndices.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- parser.parseLSquare();
- const indexed = [];
- const unresolvedIndexVecs = [];
- const indexVecTypes = [];
- while (!parser.parseOptionalRSquare()) {
- if (parser.parseOptionalKeyword('None')) {
- indexed.push(false);
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const indexVec = parser.parseOperand();
- parser.parseColon();
- const indexVecType = parser.parseType();
- unresolvedIndexVecs.push(indexVec);
- indexVecTypes.push(indexVecType);
- indexed.push(true);
- }
- parser.parseOptionalComma();
- }
- result.addAttribute('indexed', indexed);
- parser.parseComma();
- const padding = parser.parseAttribute();
- result.addAttribute('padding', padding);
- let unresolvedMask = null;
- if (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedMask = parser.parseOperand();
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- let sourceType = null;
- if (parser.parseOptionalColon()) {
- sourceType = parser.parseType();
- parser.parseComma();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- parser.resolveOperand(unresolvedSource, sourceType, result.operands);
- const indexType = new _.IndexType();
- for (const idx of unresolvedIndices) {
- parser.resolveOperand(idx, indexType, result.operands);
- }
- for (let i = 0; i < unresolvedIndexVecs.length; i++) {
- parser.resolveOperand(unresolvedIndexVecs[i], indexVecTypes[i], result.operands);
- }
- if (unresolvedMask) {
- parser.resolveOperand(unresolvedMask, null, result.operands);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.IREETensorExtDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'iree_tensor_ext');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (typeName === 'dispatch.tensor') {
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- };
- _.LinalgDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'linalg');
- this._namedStructuredOps = new Set([
- 'linalg.matmul', 'linalg.batch_matmul', 'linalg.batch_reduce_matmul',
- 'linalg.matvec', 'linalg.vecmat', 'linalg.dot', 'linalg.batch_matvec',
- 'linalg.conv_1d', 'linalg.conv_1d_ncw_fcw', 'linalg.conv_1d_nwc_wcf',
- 'linalg.conv_2d', 'linalg.conv_2d_nchw_fchw', 'linalg.conv_2d_nchw_fchw_q',
- 'linalg.conv_2d_ngchw_fgchw', 'linalg.conv_2d_ngchw_gfchw', 'linalg.conv_2d_ngchw_gfchw_q',
- 'linalg.conv_2d_nhwc_fhwc', 'linalg.conv_2d_nhwc_fhwc_q',
- 'linalg.conv_2d_nhwc_hwcf', 'linalg.conv_2d_nhwc_hwcf_q',
- 'linalg.conv_2d_nhwgc_gfhwc', 'linalg.conv_2d_nhwgc_gfhwc_q',
- 'linalg.conv_3d', 'linalg.conv_3d_ncdhw_fcdhw', 'linalg.conv_3d_ndhwc_dhwcf', 'linalg.conv_3d_ndhwc_dhwcf_q',
- 'linalg.depthwise_conv_1d_ncw_cw', 'linalg.depthwise_conv_1d_nwc_wc', 'linalg.depthwise_conv_1d_nwc_wcm',
- 'linalg.depthwise_conv_2d_nchw_chw', 'linalg.depthwise_conv_2d_nhwc_hwc', 'linalg.depthwise_conv_2d_nhwc_hwc_q',
- 'linalg.depthwise_conv_2d_nhwc_hwcm', 'linalg.depthwise_conv_2d_nhwc_hwcm_q',
- 'linalg.depthwise_conv_3d_ncdhw_cdhw', 'linalg.depthwise_conv_3d_ndhwc_dhwc', 'linalg.depthwise_conv_3d_ndhwc_dhwcm',
- 'linalg.pooling_nchw_max', 'linalg.pooling_nchw_sum',
- 'linalg.pooling_nhwc_max', 'linalg.pooling_nhwc_max_unsigned', 'linalg.pooling_nhwc_min', 'linalg.pooling_nhwc_min_unsigned', 'linalg.pooling_nhwc_sum',
- 'linalg.pooling_ncw_max', 'linalg.pooling_ncw_sum',
- 'linalg.pooling_nwc_max', 'linalg.pooling_nwc_max_unsigned', 'linalg.pooling_nwc_min', 'linalg.pooling_nwc_min_unsigned', 'linalg.pooling_nwc_sum',
- 'linalg.pooling_ndhwc_max', 'linalg.pooling_ndhwc_min', 'linalg.pooling_ndhwc_sum'
- ]);
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (result.op === 'linalg.generic') {
- return this.parseGenericOp(parser, result);
- }
- if (result.op === 'linalg.init_tensor') {
- if (parser.parseOptionalLSquare()) {
- const dims = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- dims.push(parser.parseOperand().name);
- } else if (parser.getToken().is(_.Token.integer)) {
- dims.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- result.addAttribute('static_sizes', dims);
- }
- result.addTypes(parser.parseOptionalColonTypeList());
- return true;
- }
- if (result.op === 'linalg.fill') {
- // Form 1: ins/outs format - use parseNamedStructuredOp
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'ins' || parser.getToken().is(_.Token.l_brace) || parser.getToken().is(_.Token.less)) {
- return this.parseNamedStructuredOp(parser, result);
- }
- let unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- unresolvedOperands = parser.parseOperandList();
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- return true;
- }
- if (result.op === 'linalg.conv') {
- let unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- unresolvedOperands = parser.parseOperandList();
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- return true;
- }
- if (result.op === 'linalg.yield') {
- const unresolvedOperands = parser.parseOperandList();
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- return true;
- }
- if (result.op === 'linalg.pack') {
- return this.parsePackOp(parser, result);
- }
- if (result.op === 'linalg.unpack') {
- return this.parseUnpackOp(parser, result);
- }
- if (result.op === 'linalg.transpose') {
- return this.parseDstStyleOp(parser, result, (parser, attributes) => {
- parser.parseDenseI64ArrayAttr('permutation', attributes);
- });
- }
- if (result.op === 'linalg.reduce') {
- // Optional short form: { payload_op attr-dict }
- let payloadOpName = null;
- const payloadOpAttrs = new Map();
- if (parser.parseOptionalLBrace()) {
- payloadOpName = parser.parseCustomOperationName();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(payloadOpAttrs);
- }
- parser.parseRBrace();
- }
- if (!this.parseDstStyleOp(parser, result, (parser, attributes) => {
- parser.parseDenseI64ArrayAttr('dimensions', attributes);
- })) {
- return false;
- }
- // Parse block arguments and region (or add body with payload op)
- if (payloadOpName) {
- this.addBodyWithPayloadOp(result, payloadOpName, payloadOpAttrs, true, true);
- } else {
- const regionArgs = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- regionArgs.push({ value, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- const region = result.addRegion();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region, regionArgs);
- }
- }
- return true;
- }
- if (result.op === 'linalg.broadcast') {
- return this.parseDstStyleOp(parser, result, (parser, attributes) => {
- parser.parseDenseI64ArrayAttr('dimensions', attributes);
- });
- }
- if (result.op === 'linalg.elementwise') {
- parser.parseKeyword('kind');
- parser.parseEqual();
- const kind = parser.parseAttribute();
- result.addAttribute('kind', kind.value);
- const indexingMapsAttr = this.parseIndexingMapsAttr(parser);
- if (indexingMapsAttr !== null) {
- result.addAttribute('indexing_maps', indexingMapsAttr);
- }
- return this.parseNamedStructuredOp(parser, result);
- }
- if (result.op === 'linalg.map') {
- return this.parseMapOp(parser, result);
- }
- if (result.op === 'linalg.contract') {
- const indexingMapsAttr = this.parseIndexingMapsAttr(parser);
- if (!indexingMapsAttr) {
- throw new mlir.Error(`Expected 'indexing_maps' attribute ${parser.location()}`);
- }
- result.addAttribute('indexing_maps', indexingMapsAttr);
- return this.parseNamedStructuredOp(parser, result);
- }
- if (this._namedStructuredOps.has(result.op)) {
- const indexingMapsAttr = this.parseIndexingMapsAttr(parser);
- if (indexingMapsAttr) {
- result.addAttribute('indexing_maps', indexingMapsAttr);
- }
- return this.parseNamedStructuredOp(parser, result);
- }
- if (opInfo.metadata && opInfo.metadata.assemblyFormat) {
- return super.parseOperation(parser, result);
- }
- if (parser.getToken().is(_.Token.l_brace) || parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'ins' || parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'outs') {
- const parsed = this.parseCommonStructuredOpParts(parser, result);
- if (!parsed) {
- return false;
- }
- if (parser.parseOptionalKeyword('attrs')) {
- parser.parseEqual();
- parser.parseAttributeDict(result.attributes);
- } else if (parser.getToken().is(_.Token.l_brace) && parser.getTokenSpelling().str() !== '^') {
- const saved = parser.save();
- parser.parseLBrace();
- if (parser.getToken().isNot(_.Token.percent_identifier) && parser.getToken().isNot(_.Token.bare_identifier)) {
- parser.restore(saved);
- } else {
- parser.restore(saved);
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- // Parse region (for generic ops)
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, []);
- }
- return true;
- }
- return false;
- }
- // "Common parsing used for both named structured ops created by ods-gen and by
- // manually defined C++ ops. Does not handle regions."
- // Returns { inputTypes, outputTypes } on success, null on failure.
- parseCommonStructuredOpParts(parser, result, addOperandSegmentSizes = true) {
- if (parser.parseOptionalLess()) {
- result.propertiesAttr = parser.parseAttribute();
- parser.parseGreater();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const inputTypes = [];
- const outputTypes = [];
- if (parser.parseOptionalKeyword('ins')) {
- if (!parser.parseOptionalLParen()) {
- return null;
- }
- const unresolvedIns = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedIns.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- inputTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.resolveOperands(unresolvedIns, inputTypes, result.operands);
- }
- if (!parser.parseOptionalRParen()) {
- return null;
- }
- }
- if (parser.parseOptionalKeyword('outs')) {
- if (!parser.parseOptionalLParen()) {
- return null;
- }
- const unresolvedOuts = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOuts.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- outputTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.resolveOperands(unresolvedOuts, outputTypes, result.operands);
- }
- if (!parser.parseOptionalRParen()) {
- return null;
- }
- }
- if (addOperandSegmentSizes) {
- result.addAttribute('operandSegmentSizes', [inputTypes.length, outputTypes.length]);
- }
- return { inputTypes, outputTypes };
- }
- parseNamedStructuredOp(parser, result) {
- const parsed = this.parseCommonStructuredOpParts(parser, result);
- if (!parsed) {
- return false;
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- return true;
- }
- parseIndexingMapsAttr(parser) {
- if (!parser.parseOptionalKeyword('indexing_maps')) {
- return null;
- }
- parser.parseEqual();
- return parser.parseAttribute();
- }
- parseDstStyleOp(parser, op, parseAttrsFn) {
- const parsed = this.parseCommonStructuredOpParts(parser, op, false);
- if (!parsed) {
- return false;
- }
- for (const outputType of parsed.outputTypes) {
- if (outputType instanceof _.RankedTensorType) {
- op.addTypes([outputType]);
- }
- }
- if (parseAttrsFn) {
- parseAttrsFn(parser, op.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(op.attributes);
- }
- return true;
- }
- addBodyWithPayloadOp(op, payloadOpName, payloadOpAttrs /*, initFirst, mapInit */) {
- const region = op.addRegion();
- const block = { operations: [], arguments: [] };
- for (const operand of op.operands) {
- if (operand && operand.type) {
- const elemType = operand.type.elementType || operand.type;
- block.arguments.push({ value: null, type: elemType });
- }
- }
- const payloadState = new _.OperationState(null,payloadOpName);
- if (op.operands.length > 0) {
- const lastOperand = op.operands[op.operands.length - 1];
- if (lastOperand && lastOperand.type) {
- const elemType = lastOperand.type.elementType || lastOperand.type;
- payloadState.types = [elemType];
- }
- }
- for (const [name, value] of payloadOpAttrs) {
- payloadState.attributes.set(name, value);
- }
- block.operations.push(_.Operation.create(payloadState));
- const yieldState = new _.OperationState(null, this.getOperation('linalg.yield'));
- block.operations.push(_.Operation.create(yieldState));
- region.blocks = [block];
- }
- parseMapOp(parser, result) {
- let payloadOpName = null;
- const payloadOpAttrs = new Map();
- if (parser.parseOptionalLBrace()) {
- payloadOpName = parser.parseCustomOperationName();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(payloadOpAttrs);
- }
- parser.parseRBrace();
- }
- if (!this.parseDstStyleOp(parser, result)) {
- return false;
- }
- // Parse block arguments and region (or add body with payload op)
- if (payloadOpName) {
- if (result.operands.length > 0) {
- this.addBodyWithPayloadOp(result, payloadOpName, payloadOpAttrs, false, false);
- } else {
- result.addRegion();
- }
- } else {
- const regionArgs = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- regionArgs.push({ value, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- const region = result.addRegion();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region, regionArgs);
- }
- }
- return true;
- }
- parseGenericOp(parser, result) {
- if (parser.getToken().is(_.Token.l_brace) || parser.getToken().is(_.Token.hash_identifier)) {
- if (parser.getToken().is(_.Token.hash_identifier)) {
- const attrRef = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.hash_identifier);
- result.addAttribute('trait', attrRef);
- } else {
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalKeyword('ins')) {
- parser.parseLParen();
- const unresolvedIns = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedIns.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- const insTypes = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- insTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.resolveOperands(unresolvedIns, insTypes, result.operands);
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('outs')) {
- parser.parseLParen();
- const unresolvedOuts = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOuts.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- const outsTypes = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- outsTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.resolveOperands(unresolvedOuts, outsTypes, result.operands);
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('attrs')) {
- parser.parseEqual();
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- if (parser.parseOptionalArrow()) {
- const hasParens = parser.getToken().is(_.Token.l_paren);
- const types = hasParens ? parser.parseTypeListParens() : parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- return true;
- }
- parsePackOp(parser, result) {
- const unresolvedSource = parser.parseOperand();
- let unresolvedPadding = null;
- let paddingType = null;
- if (parser.parseOptionalKeyword('padding_value')) {
- parser.parseLParen();
- unresolvedPadding = parser.parseOperand();
- parser.parseColon();
- paddingType = parser.parseType();
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('outer_dims_perm')) {
- parser.parseEqual();
- const outerDimsPerm = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- outerDimsPerm.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('outer_dims_perm', outerDimsPerm.map((v) => BigInt(v)));
- }
- parser.parseKeyword('inner_dims_pos');
- parser.parseEqual();
- const innerDimsPos = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- innerDimsPos.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('inner_dims_pos', innerDimsPos.map((v) => BigInt(v)));
- parser.parseKeyword('inner_tiles');
- parser.parseEqual();
- const staticInnerTiles = [];
- const dynamicTileOperands = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- dynamicTileOperands.push(operand);
- staticInnerTiles.push(_.ShapedType.kDynamic);
- } else if (parser.getToken().is(_.Token.integer)) {
- staticInnerTiles.push(BigInt(parser.getToken().getSpelling().str()));
- parser.consumeToken(_.Token.integer);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('static_inner_tiles', staticInnerTiles);
- parser.parseKeyword('into');
- const unresolvedDest = parser.parseOperand();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- const sourceType = parser.parseType();
- parser.parseArrow();
- const destType = parser.parseType();
- parser.resolveOperand(unresolvedSource, sourceType, result.operands);
- parser.resolveOperand(unresolvedDest, destType, result.operands);
- if (unresolvedPadding) {
- parser.resolveOperand(unresolvedPadding, paddingType, result.operands);
- }
- for (const dynOp of dynamicTileOperands) {
- parser.resolveOperand(dynOp, null, result.operands);
- }
- if (!sourceType.toString().startsWith('memref')) {
- result.addTypes([destType]);
- }
- return true;
- }
- parseUnpackOp(parser, result) {
- const unresolvedSource = parser.parseOperand();
- if (parser.parseOptionalKeyword('outer_dims_perm')) {
- parser.parseEqual();
- const outerDimsPerm = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- outerDimsPerm.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('outer_dims_perm', outerDimsPerm.map((v) => BigInt(v)));
- }
- parser.parseKeyword('inner_dims_pos');
- parser.parseEqual();
- const innerDimsPos = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- innerDimsPos.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('inner_dims_pos', innerDimsPos.map((v) => BigInt(v)));
- parser.parseKeyword('inner_tiles');
- parser.parseEqual();
- const staticInnerTiles = [];
- const dynamicTileOperands = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- dynamicTileOperands.push(operand);
- staticInnerTiles.push(-9223372036854775808n); // ShapedType::kDynamic
- } else if (parser.getToken().is(_.Token.integer)) {
- staticInnerTiles.push(BigInt(parser.getToken().getSpelling().str()));
- parser.consumeToken(_.Token.integer);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('static_inner_tiles', staticInnerTiles);
- parser.parseKeyword('into');
- const unresolvedDest = parser.parseOperand();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- const sourceType = parser.parseType();
- parser.parseArrow();
- const destType = parser.parseType();
- parser.resolveOperand(unresolvedSource, sourceType, result.operands);
- parser.resolveOperand(unresolvedDest, destType, result.operands);
- for (const dynOp of dynamicTileOperands) {
- parser.resolveOperand(dynOp, null, result.operands);
- }
- if (!sourceType.toString().startsWith('memref')) {
- result.addTypes([destType]);
- }
- return true;
- }
- };
- _.ONNXDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'onnx');
- }
- parseOperation(parser, result) {
- if (result.op === 'onnx.Constant') {
- parser.parseOptionalAttrDict(result.attributes);
- if (result.attributes.size === 0) {
- const value = parser.parseOptionalAttribute();
- if (value) {
- const name = value instanceof _.SparseElementsAttr ? 'sparse_value' : 'value';
- result.addAttribute(name, value);
- result.addTypes([value.type]);
- return true;
- }
- }
- result.addTypes([parser.parseColonType()]);
- return true;
- }
- if (result.op === 'onnx.ConstantOfShape') {
- parser.parseLParen();
- const unresolved = parser.parseOperand();
- parser.parseRParen();
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- parser.parseLParen();
- const inputType = parser.parseType();
- parser.resolveOperand(unresolved, inputType, result.operands);
- parser.parseRParen();
- parser.parseArrow();
- const outputType = parser.parseType();
- result.addTypes([outputType]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.KrnlDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'krnl');
- }
- parseOperation(parser, result) {
- if (result.op === 'krnl.define_loops') {
- if (parser.getToken().is(_.Token.integer)) {
- const count = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- result.addAttribute('num_loops', count);
- const loopType = new _.Type('!krnl.loop');
- const types = Array(count).fill(loopType);
- result.addTypes(types);
- }
- return true;
- }
- if (result.op === 'krnl.get_linear_offset_index') {
- const unresolvedOperands = [];
- const staticIndices = [];
- const memref = parser.parseOperand();
- unresolvedOperands.push(memref);
- if (parser.parseOptionalKeyword('at')) {
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- // Indices can be either SSA values (%arg) or integer constants (0, 10, etc.)
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const index = parser.parseOperand();
- unresolvedOperands.push(index);
- staticIndices.push(-9223372036854775808n); // ShapedType::kDynamic marker
- } else if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus)) {
- const value = parser.parseInteger();
- staticIndices.push(BigInt(value));
- }
- if (parser.getToken().isNot(_.Token.r_square)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseRSquare();
- if (staticIndices.length > 0) {
- result.addAttribute('static_indices', staticIndices);
- }
- }
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, type, result.operands);
- }
- result.addTypes([new _.IndexType()]);
- return true;
- }
- if (result.op === 'krnl.prefetch') {
- const unresolvedOperands = [];
- const memref = parser.parseOperand();
- unresolvedOperands.push(memref);
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- const index = parser.parseOperand();
- unresolvedOperands.push(index);
- if (parser.getToken().isNot(_.Token.r_square)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseRSquare();
- }
- parser.parseComma();
- const readOrWrite = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('isWrite', readOrWrite === 'write');
- parser.parseComma();
- parser.parseKeyword('locality');
- parser.parseLess();
- const localityHint = parser.parseInteger();
- result.addAttribute('localityHint', localityHint);
- parser.parseGreater();
- parser.parseComma();
- const cacheType = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('isDataCache', cacheType === 'data');
- parser.parseOptionalAttrDict(result.attributes);
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, type, result.operands);
- }
- return true;
- }
- if (result.op === 'krnl.iterate') {
- const unresolvedOperands = parser.parseOperandList('paren');
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- if (parser.parseOptionalKeyword('with')) {
- parser.parseLParen();
- const numOptimizedLoops = result.operands.length;
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOperand();
- parser.parseArrow();
- parser.parseOperand();
- parser.parseEqual();
- parser.parseOptionalKeyword('max');
- if (parser.getToken().isAny(_.Token.kw_affine_map, _.Token.kw_affine_set)) {
- parser.parseAttribute();
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.skip('(');
- }
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- } else {
- parser.parseAttribute();
- }
- parser.parseKeyword('to');
- parser.parseOptionalKeyword('min');
- if (parser.getToken().isAny(_.Token.kw_affine_map, _.Token.kw_affine_set)) {
- parser.parseAttribute();
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.skip('(');
- }
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- } else {
- parser.parseAttribute();
- }
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- result.addAttribute('num_optimized_loops', numOptimizedLoops);
- }
- if (parser.parseOptionalKeyword('iter_args')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOperand();
- parser.parseEqual();
- parser.parseAttribute();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- result.regions = [region];
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.MhloDialect = class extends _.HLODialect {
- constructor(operations) {
- super(operations, 'mhlo');
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (opInfo.metadata && opInfo.metadata.parser && opInfo.metadata.parser.includes('parseOneResultSameOperandTypeOp')) {
- return this.parseOneResultSameOperandTypeOp(parser, result);
- }
- if (result.op === 'mhlo.constant') {
- if (parser.parseOptionalLParen() && parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- parser.parseLParen();
- parser.parseRParen();
- parser.parseArrow();
- const type = parser.parseType();
- result.addTypes([type]);
- }
- } else {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const value = parser.parseAttribute();
- if (value) {
- result.addAttribute('value', value);
- if (value.type && result.types.length === 0) {
- result.addTypes([value.type]);
- }
- }
- if (result.types.length === 0) {
- result.addTypes(parser.parseOptionalColonTypeList());
- }
- }
- return true;
- }
- if (result.op === 'mhlo.compare') {
- // Use modern assemblyFormat if available, but only if this is modern syntax
- // Old format has 'attributes' keyword, new format has attr-dict directly
- if (opInfo.metadata && opInfo.metadata.assemblyFormat && !(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'attributes')) {
- return super.parseOperation(parser, result);
- }
- // Legacy parser for old mhlo.compare without assembly format
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const comparisonDirection = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('comparison_direction', comparisonDirection);
- parser.parseComma();
- const unresolvedOperands = parser.parseOperandList();
- // Check for optional compare_type
- if (parser.parseOptionalComma() && parser.getToken().is(_.Token.bare_identifier)) {
- const compareType = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('compare_type', compareType);
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedOperands, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- // Single type applied to all operands
- const types = unresolvedOperands.map(() => type);
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- }
- if (result.op === 'mhlo.reduce') {
- // mhlo uses raw array for dimensions (like b.getI64TensorAttr in ref impl)
- return super.parseReduceOp(parser, result, (dims) => dims, 'mhlo.return');
- }
- if (result.op === 'mhlo.scan' && parser.getToken().is(_.Token.l_paren)) {
- return super.parseScanOp(parser, result, 'mhlo.return');
- }
- if (result.op === 'mhlo.while') {
- // mhlo.while always uses parenthesized form with named arguments
- parser.parseLParen();
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const firstOperand = parser.parseOperand();
- let operandToResolve = firstOperand;
- if (parser.parseOptionalEqual()) {
- operandToResolve = parser.parseOperand();
- }
- unresolvedOperands.push(operandToResolve);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- const types = [];
- if (parser.parseOptionalColon()) {
- while (!(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'cond') && !(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'attributes')) {
- types.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- result.addTypes(types);
- if (parser.parseOptionalKeyword('attributes')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (parser.parseOptionalKeyword('cond')) {
- const condRegion = {};
- parser.parseRegion(condRegion);
- result.regions.push(condRegion);
- }
- if (parser.parseOptionalKeyword('do')) {
- const bodyRegion = {};
- parser.parseRegion(bodyRegion);
- result.regions.push(bodyRegion);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseOneResultSameOperandTypeOp(parser, result) {
- const unresolvedOperands = parser.parseOperandList();
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- const types = unresolvedOperands.map(() => type);
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (result.types.length > 0) {
- result.types[0] = type;
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- };
- _.ChloDialect = class extends _.HLODialect {
- constructor(operations) {
- super(operations, 'chlo');
- }
- parseOperation(parser, result) {
- if (result.op === 'chlo.scan' && parser.getToken().is(_.Token.l_paren)) {
- return super.parseScanOp(parser, result, 'stablehlo.return');
- }
- return super.parseOperation(parser, result);
- }
- };
- _.THLODialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'thlo');
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (opInfo.metadata.assemblyFormat) {
- return super.parseOperation(parser, result);
- }
- if (parser.parseOptionalKeyword('ins')) {
- parser.parseLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('outs')) {
- parser.parseLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const blockArguments = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- blockArguments.push({ value, type });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = { blocks: [] };
- const block = { operations: [], arguments: blockArguments };
- parser.parseLBrace();
- while (!parser.parseOptionalRBrace()) {
- const operation = parser.parseOperation();
- block.operations.push(_.Operation.create(operation));
- }
- region.blocks.push(block);
- result.regions.push(region);
- }
- return true;
- }
- };
- _.QuantDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'quant');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (typeName === 'uniform' || typeName === 'calibrated' || typeName === 'any') {
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- };
- _.TosaDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tosa');
- this._customOps = new Set([
- 'tosa.apply_scale', 'tosa.argmax', 'tosa.cast_from_block_scaled',
- 'tosa.cast_to_block_scaled', 'tosa.clamp', 'tosa.conv2d_block_scaled',
- 'tosa.matmul_t_block_scaled', 'tosa.max_pool2d', 'tosa.maximum',
- 'tosa.minimum', 'tosa.reduce_max', 'tosa.reduce_min', 'tosa.rescale',
- 'tosa.resize'
- ]);
- this._regionOps = new Set(['tosa.cond_if', 'tosa.while_loop']);
- this.registerCustomDirective('VariableOpTypeOrInitialValue', this.parseVariableOpTypeOrInitialValue.bind(this));
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (typeName === 'shape') {
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- if (typeName === 'mxint8') {
- return new _.Type(`!${dialect}.mxint8`);
- }
- return null;
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (this._regionOps.has(result.op)) {
- let hasBlockArgs = false;
- const unresolvedCond = [];
- const unresolvedInputs = [];
- const blockArgs = [];
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedCond.push(parser.parseOperand());
- }
- if (parser.parseOptionalLParen()) {
- hasBlockArgs = true;
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- blockArgs.push(parser.parseOperand());
- parser.parseEqual();
- unresolvedInputs.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalColon()) {
- // For tosa.while_loop: no condition operand, has block args, function type after colon
- if (unresolvedCond.length === 0 && hasBlockArgs) {
- const functionType = parser.parseFunctionType();
- if (functionType) {
- parser.resolveOperands(unresolvedInputs, functionType.inputs, result.operands);
- result.addTypes(functionType.results);
- }
- } else {
- const condType = parser.parseType();
- if (unresolvedCond.length > 0) {
- parser.resolveOperands(unresolvedCond, [condType], result.operands);
- }
- // If block args present, parse function type for inputs/outputs
- if (hasBlockArgs && parser.getToken().is(_.Token.l_paren)) {
- const functionType = parser.parseFunctionType();
- if (functionType) {
- parser.resolveOperands(unresolvedInputs, functionType.inputs, result.operands);
- result.addTypes(functionType.results);
- }
- } else if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- }
- } else {
- for (const cond of unresolvedCond) {
- parser.resolveOperand(cond, null, result.operands);
- }
- for (const input of unresolvedInputs) {
- parser.resolveOperand(input, null, result.operands);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- if (parser.parseOptionalKeyword('else') || parser.parseOptionalKeyword('do')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const secondRegion = {};
- parser.parseRegion(secondRegion);
- result.regions.push(secondRegion);
- }
- }
- return true;
- }
- if (this._customOps.has(result.op)) {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.getToken().is(_.Token.l_brace)) {
- // Parse attribute dict but check if any are actually inputs
- const inputNames = new Set((opInfo.metadata && opInfo.metadata.operands || []).map((i) => i.name));
- const tempAttrs = new Map();
- parser.parseAttributeDict(tempAttrs);
- for (const [name, value] of tempAttrs) {
- // If this is an input (like input_zp, output_zp), add as operand
- if (inputNames.has(name) && value && typeof value === 'string' && value.startsWith('%')) {
- const unresolvedOperand = new _.UnresolvedOperand(null,value, 0);
- unresolvedOperands.push(unresolvedOperand);
- } else if (inputNames.has(name) && value && value.value && typeof value.value === 'string' && value.value.startsWith('%')) {
- const unresolvedOperand = new _.UnresolvedOperand(null, value.value, 0);
- unresolvedOperands.push(unresolvedOperand);
- } else {
- result.attributes.set(name, value);
- }
- }
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedOperands, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- const types = unresolvedOperands.map(() => type);
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseVariableOpTypeOrInitialValue(parser, op /*, args */) {
- if (parser.parseOptionalEqual()) {
- const initialValue = parser.parseAttribute();
- op.addAttribute('initial_value', initialValue);
- } else if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- op.addAttribute('type', type);
- }
- }
- };
- _.IRDLDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'irdl');
- this.registerCustomDirective('SingleBlockRegion', this.parseSingleBlockRegion.bind(this));
- this.registerCustomDirective('NamedValueList', this.parseNamedValueList.bind(this));
- this.registerCustomDirective('NamedValueListWithVariadicity', this.parseNamedValueListWithVariadicity.bind(this));
- this.registerCustomDirective('AttributesOp', this.parseAttributesOp.bind(this));
- }
- parseSingleBlockRegion(parser, op, /* args */) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = op.addRegion();
- parser.parseRegion(region);
- }
- }
- parseNamedValueList(parser, op, argsAttrName, namesAttrName) {
- const argValues = [];
- const nameValues = [];
- const parseOne = () => {
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken();
- nameValues.push(name);
- parser.parseColon();
- const value = parser.parseOperand();
- argValues.push(value);
- return value;
- };
- parser.parseCommaSeparatedList('paren', parseOne);
- if (argsAttrName && namesAttrName) {
- // args contains SSA values (%0, %1, ...) - resolve and add as operands
- for (const value of argValues) {
- parser.resolveOperand(value, null, op.operands);
- }
- op.addAttribute(namesAttrName, nameValues);
- }
- }
- parseNamedValueListWithVariadicity(parser, op, argsAttrName, namesAttrName, variadicityAttrName) {
- const argValues = [];
- const nameValues = [];
- const variadicityValues = [];
- const parseOne = () => {
- let variadicity = null;
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const peekValue = parser.getTokenSpelling().str();
- if (peekValue === 'single' || peekValue === 'optional' || peekValue === 'variadic') {
- variadicity = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- }
- }
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken();
- nameValues.push(name);
- parser.parseColon();
- const value = parser.parseOperand();
- argValues.push(value);
- variadicityValues.push(variadicity || 'single');
- return value;
- };
- parser.parseCommaSeparatedList('paren', parseOne);
- if (argsAttrName && namesAttrName) {
- // args contains SSA values (%0, %1, ...) - resolve and add as operands
- for (const value of argValues) {
- parser.resolveOperand(value, null, op.operands);
- }
- op.addAttribute(namesAttrName, nameValues);
- if (variadicityAttrName) {
- op.addAttribute(variadicityAttrName, variadicityValues);
- }
- }
- }
- parseAttributesOp(parser, op, argsAttrName, namesAttrName) {
- const argValues = [];
- const nameValues = [];
- if (parser.parseOptionalLBrace()) {
- while (parser.getToken().isNot(_.Token.r_brace)) {
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- nameValues.push(name);
- parser.parseEqual();
- const value = parser.parseOperand();
- argValues.push(value);
- parser.parseOptionalComma();
- }
- parser.parseRBrace();
- }
- if (argsAttrName && namesAttrName) {
- for (const value of argValues) {
- parser.resolveOperand(value, null, op.operands);
- }
- op.addAttribute(namesAttrName, nameValues);
- }
- }
- };
- _.XeGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xegpu');
- this.registerCustomDirective('OptionalDynamicIndexList', this.parseOptionalDynamicIndexList.bind(this));
- }
- parseOptionalDynamicIndexList(parser, op, dynamicAttrName, staticAttrName) {
- const indices = [];
- const dynamicValues = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.floatliteral)) {
- indices.push(parseInt(parser.getToken().getSpelling().str(), 10));
- parser.consumeToken();
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const value = parser.parseOperand();
- dynamicValues.push(value);
- indices.push(-9223372036854775808);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseOptionalRSquare();
- if (dynamicAttrName && staticAttrName) {
- // Dynamic values are SSA operands (%0, %1, ...) - resolve and add as operands
- for (const value of dynamicValues) {
- parser.resolveOperand(value, null, op.operands);
- }
- // Static indices are compile-time constants - add as attribute
- if (indices.length > 0) {
- op.addAttribute(staticAttrName, indices);
- }
- }
- }
- }
- };
- _.ShardDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'shard');
- this.registerCustomDirective('DimensionList', this.parseDimensionList.bind(this));
- }
- parseDimensionList(parser, op, attrName) {
- const dimensions = [];
- while (true) {
- if (parser.getToken().is(_.Token.question)) {
- parser.consumeToken(_.Token.question);
- dimensions.push(-1);
- } else if (parser.getToken().is(_.Token.integer)) {
- dimensions.push(parser.parseInteger());
- } else {
- break;
- }
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const token = parser.getTokenSpelling().str();
- if (token === 'x') {
- parser.consumeToken(_.Token.bare_identifier);
- continue;
- } else if (token.startsWith('x')) {
- parser.consumeToken(_.Token.bare_identifier);
- const remaining = token.substring(1);
- const parts = remaining.split('x');
- for (const part of parts) {
- if (part === '?') {
- dimensions.push(-1);
- } else if (part !== '') {
- const num = parseInt(part, 10);
- if (!isNaN(num)) {
- dimensions.push(num);
- }
- }
- }
- break;
- }
- break;
- }
- if (parser.getToken().isNot(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.question)) {
- break;
- }
- }
- if (attrName) {
- op.addAttribute(attrName, dimensions);
- }
- }
- };
- _.spirv = {};
- _.spirv.PointerType = class extends _.Type {
- constructor(pointeeType, storageClass) {
- super(null);
- this.pointeeType = pointeeType;
- this.storageClass = storageClass;
- }
- static parse(parser) {
- parser.parseLess();
- const pointeeType = parser.parseType();
- parser.parseComma();
- let storageClass = parser.parseOptionalKeyword();
- if (!storageClass) {
- storageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- }
- parser.parseGreater();
- return new _.spirv.PointerType(pointeeType, storageClass);
- }
- toString() {
- const pointeeStr = this.pointeeType?.toString ? this.pointeeType.toString() : this.pointeeType;
- return `!spirv.ptr<${pointeeStr}, ${this.storageClass}>`;
- }
- };
- _.spirv.SPIRVDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'spirv');
- this.typesWithOptionalParams = new Set(['sampler', 'sampled_image', 'matrix', 'image', 'rtarray', 'ptr', 'array', 'struct', 'coopmatrix']);
- this.registerCustomDirective('ImageOperands', this.parseImageOperands.bind(this));
- this.registerCustomDirective('SwitchOpCases', this.parseSwitchOpCases.bind(this));
- this.registerCustomDirective('SPIRV_I32_1DArmTensor', this.parseSPIRV_I32_1DArmTensor.bind(this));
- this.registerCustomAttribute('SPIRV_ScopeAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_MemorySemanticsAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_MemoryAccessAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_GroupOperationAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_KHR_CooperativeMatrixLayoutAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_KHR_CooperativeMatrixOperandsAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('SPIRV_TosaExtNaNPropagationModeAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- }
- parseSwitchOpCases(parser, result) {
- if (!parser.parseOptionalKeyword('default')) {
- return;
- }
- if (!parser.parseOptionalColon()) {
- return;
- }
- if (parser.getToken().isNot(_.Token.caret_identifier)) {
- return;
- }
- const defaultDestination = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const defaultDest = { label: defaultDestination, arguments: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const value = parser.parseOperand();
- defaultDest.arguments.push({ value });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (idx < defaultDest.arguments.length && parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- if (defaultDest.arguments[idx]) {
- defaultDest.arguments[idx].type = type;
- }
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRParen();
- }
- result.successors = result.successors || [];
- result.successors.push(defaultDest);
- const caseValues = [];
- while (parser.parseOptionalComma()) {
- if (parser.getToken().isNot(_.Token.integer) && parser.getToken().isNot(_.Token.minus)) {
- break;
- }
- const value = parser.parseInteger();
- caseValues.push(value);
- if (!parser.parseOptionalColon()) {
- break;
- }
- if (parser.getToken().isNot(_.Token.caret_identifier)) {
- break;
- }
- const caseDestination = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const caseDest = { label: caseDestination, arguments: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const argValue = parser.parseOperand();
- caseDest.arguments.push({ value: argValue });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (idx < caseDest.arguments.length && parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- if (caseDest.arguments[idx]) {
- caseDest.arguments[idx].type = type;
- }
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRParen();
- }
- result.successors.push(caseDest);
- }
- if (caseValues.length > 0) {
- result.addAttribute('literals', caseValues);
- }
- }
- parseImageOperands(parser /*, op, args */) {
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- }
- parseType(parser, dialect) {
- let mnemonic = parser.parseOptionalKeyword();
- if (mnemonic) {
- if (mnemonic === 'ptr' && parser.getToken().is(_.Token.less)) {
- return _.spirv.PointerType.parse(parser);
- }
- // Handle sub-dialect types like arm.tensor, KHR.CooperativeMatrix, etc.
- while (parser.consumeIf('.')) {
- const subType = parser.parseOptionalKeyword();
- if (subType) {
- mnemonic += `.${subType}`;
- } else {
- break;
- }
- }
- let type = `!${dialect}.${mnemonic}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- // spv.Constant value : type
- if (result.op === 'spirv.Constant' || result.op === 'spv.Constant') {
- const value = parser.parseAttribute();
- result.addAttribute('value', value);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- } else if (value && value.type) {
- result.addTypes([value.type]);
- }
- return true;
- }
- // SPIR-V comparison operations: result is i1 (or vector<Nxi1> for vector operands)
- const spirvCompareOps = new Set([
- 'spirv.SLessThan', 'spv.SLessThan',
- 'spirv.SLessThanEqual', 'spv.SLessThanEqual',
- 'spirv.SGreaterThan', 'spv.SGreaterThan',
- 'spirv.SGreaterThanEqual', 'spv.SGreaterThanEqual',
- 'spirv.ULessThan', 'spv.ULessThan',
- 'spirv.ULessThanEqual', 'spv.ULessThanEqual',
- 'spirv.UGreaterThan', 'spv.UGreaterThan',
- 'spirv.UGreaterThanEqual', 'spv.UGreaterThanEqual',
- 'spirv.IEqual', 'spv.IEqual',
- 'spirv.INotEqual', 'spv.INotEqual',
- 'spirv.FOrdEqual', 'spv.FOrdEqual',
- 'spirv.FOrdNotEqual', 'spv.FOrdNotEqual',
- 'spirv.FOrdLessThan', 'spv.FOrdLessThan',
- 'spirv.FOrdLessThanEqual', 'spv.FOrdLessThanEqual',
- 'spirv.FOrdGreaterThan', 'spv.FOrdGreaterThan',
- 'spirv.FOrdGreaterThanEqual', 'spv.FOrdGreaterThanEqual',
- 'spirv.FUnordEqual', 'spv.FUnordEqual',
- 'spirv.FUnordNotEqual', 'spv.FUnordNotEqual',
- 'spirv.FUnordLessThan', 'spv.FUnordLessThan',
- 'spirv.FUnordLessThanEqual', 'spv.FUnordLessThanEqual',
- 'spirv.FUnordGreaterThan', 'spv.FUnordGreaterThan',
- 'spirv.FUnordGreaterThanEqual', 'spv.FUnordGreaterThanEqual'
- ]);
- if (spirvCompareOps.has(result.op)) {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalColon()) {
- const operandType = parser.parseType();
- parser.resolveOperands(unresolvedOperands, [operandType, operandType], result.operands);
- // Result type is i1 for scalar, vector<Nxi1> for vector
- if (operandType instanceof _.VectorType) {
- result.addTypes([new _.VectorType(operandType.shape, new _.IntegerType('i1'), operandType.scalableDims)]);
- } else {
- result.addTypes([new _.IntegerType('i1')]);
- }
- } else {
- parser.resolveOperands(unresolvedOperands, [null, null], result.operands);
- result.addTypes([new _.IntegerType('i1')]);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- // Operations with '->' in their assembly format should use assembly format parsing
- const arrowFormatOps = new Set([
- 'spirv.GL.Distance', 'spirv.GL.FMix', 'spirv.GL.FrexpStruct', 'spirv.GL.Ldexp',
- 'spirv.GL.Length', 'spirv.GL.PackHalf2x16', 'spirv.GL.UnpackHalf2x16',
- 'spirv.GL.PackSnorm4x8', 'spirv.GL.UnpackSnorm4x8',
- 'spirv.GLSL.Distance', 'spirv.GLSL.FMix', 'spirv.GLSL.FrexpStruct', 'spirv.GLSL.Ldexp',
- 'spirv.GLSL.Length', 'spirv.GLSL.PackHalf2x16', 'spirv.GLSL.UnpackHalf2x16',
- 'spirv.GLSL.PackSnorm4x8', 'spirv.GLSL.UnpackSnorm4x8',
- 'spv.GL.Distance', 'spv.GL.FMix', 'spv.GL.FrexpStruct', 'spv.GL.Ldexp',
- 'spv.GL.Length', 'spv.GL.PackHalf2x16', 'spv.GL.UnpackHalf2x16',
- 'spv.GL.PackSnorm4x8', 'spv.GL.UnpackSnorm4x8',
- 'spv.GLSL.Distance', 'spv.GLSL.FMix', 'spv.GLSL.FrexpStruct', 'spv.GLSL.Ldexp',
- 'spv.GLSL.Length', 'spv.GLSL.PackHalf2x16', 'spv.GLSL.UnpackHalf2x16',
- 'spv.GLSL.PackSnorm4x8', 'spv.GLSL.UnpackSnorm4x8'
- ]);
- if ((result.op.startsWith('spirv.GLSL.') || result.op.startsWith('spv.GLSL.') || result.op.startsWith('spirv.GL.') || result.op.startsWith('spv.GL.')) && !arrowFormatOps.has(result.op)) {
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.colon)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, type, result.operands);
- }
- result.types.push(type);
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.SpecConstantComposite' || result.op === 'spv.SpecConstantComposite') {
- parser.parseSymbolName('sym_name', result.attributes);
- parser.parseLParen();
- const constituents = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.at_identifier)) {
- constituents.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.at_identifier);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- result.addAttribute('constituents', constituents);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('type', type.toString());
- }
- return true;
- }
- if (result.op.endsWith('.SpecConstantCompositeReplicate')) {
- parser.parseSymbolName('sym_name', result.attributes);
- parser.parseLParen();
- if (parser.getToken().is(_.Token.at_identifier)) {
- const constituent = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('constituent', constituent);
- }
- parser.parseRParen();
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('type', type.toString());
- }
- return true;
- }
- if (result.op === 'spirv.SpecConstantOperation' || result.op === 'spv.SpecConstantOperation') {
- parser.parseKeyword('wraps');
- const wrappedOp = parser.parseGenericOperation();
- if (wrappedOp) {
- const region = { blocks: [{ operations: [wrappedOp] }] };
- result.regions.push(region);
- if (wrappedOp.results && wrappedOp.results.length > 0) {
- result.addTypes([wrappedOp.results[0].type]);
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (result.op === 'spirv.Constant' || result.op === 'spv.Constant') {
- const value = parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- const valueType = parser.parseType();
- result.addAttribute('value', { ...value, valueType });
- } else {
- result.addAttribute('value', value);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- return true;
- }
- if (result.op === 'spirv.Load' || result.op === 'spv.Load') {
- const storageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('storage_class', storageClass);
- const ptrOperand = parser.parseOperand();
- if (parser.parseOptionalLSquare()) {
- const memoryAccess = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.string)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.string);
- } else if (parser.getToken().is(_.Token.integer)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- if (memoryAccess.length > 0) {
- result.addAttribute('memory_access', memoryAccess.join(', '));
- }
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- parser.resolveOperand(ptrOperand, null, result.operands);
- result.types.push(type);
- } else {
- parser.resolveOperand(ptrOperand, null, result.operands);
- }
- return true;
- }
- if (result.op === 'spirv.CompositeExtract' || result.op === 'spv.CompositeExtract') {
- const compositeOperand = parser.parseOperand();
- if (parser.parseOptionalLSquare()) {
- const indices = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- indices.push(parser.parseInteger());
- }
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('indices', indices);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- parser.resolveOperand(compositeOperand, null, result.operands);
- result.types.push(type);
- } else {
- parser.resolveOperand(compositeOperand, null, result.operands);
- }
- return true;
- }
- // Handle AccessChain with old syntax (no -> for result type)
- // Uses compatibility flag because assemblyFormat includes `->` which is not present in old files
- if (result.op === 'spirv.AccessChain' || result.op === 'spv.AccessChain') {
- result.compatibility = true;
- const unresolvedOperands = [];
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const basePtrType = parser.parseType();
- const types = [basePtrType];
- while (parser.parseOptionalComma()) {
- types.push(parser.parseType());
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- // Check for optional -> result_type (newer syntax)
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- } else {
- // Old syntax without explicit result type - use base pointer type as fallback
- // (In SPIR-V, AccessChain result is a pointer to a sub-element, but
- // for visualization purposes the base ptr type is sufficient)
- result.addTypes([basePtrType]);
- }
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.Variable' || result.op === 'spv.Variable') {
- let unresolvedInit = null;
- if (parser.parseOptionalKeyword('init')) {
- parser.parseLParen();
- unresolvedInit = parser.parseOperand();
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- if (unresolvedInit) {
- // The init value type should match the element type of the pointer result
- parser.resolveOperand(unresolvedInit, null, result.operands);
- }
- } else if (unresolvedInit) {
- parser.resolveOperand(unresolvedInit, null, result.operands);
- }
- return true;
- }
- if (result.op === 'spirv.Store' || result.op === 'spv.Store') {
- const storageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('storage_class', storageClass);
- const unresolvedOperands = [];
- unresolvedOperands.push(parser.parseOperand());
- parser.parseComma();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- const memoryAccess = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.string)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.string);
- } else if (parser.getToken().is(_.Token.integer)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- if (memoryAccess.length > 0) {
- result.addAttribute('memory_access', memoryAccess.join(', '));
- }
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- parser.resolveOperands(unresolvedOperands, [type, type], result.operands);
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.CompositeInsert' || result.op === 'spv.CompositeInsert') {
- const unresolvedOperands = [];
- unresolvedOperands.push(parser.parseOperand());
- parser.parseComma();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- const indices = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- indices.push(parser.parseInteger());
- }
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- result.addAttribute('indices', indices);
- }
- if (parser.parseOptionalColon()) {
- const objType = parser.parseType();
- parser.resolveOperand(unresolvedOperands[0], objType, result.operands);
- if (parser.parseOptionalKeyword('into')) {
- const compositeType = parser.parseType();
- parser.resolveOperand(unresolvedOperands[1], compositeType, result.operands);
- result.types.push(compositeType);
- } else {
- parser.resolveOperand(unresolvedOperands[1], null, result.operands);
- }
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.BranchConditional' || result.op === 'spv.BranchConditional') {
- const conditionOperand = parser.parseOperand();
- parser.resolveOperand(conditionOperand, null, result.operands);
- // Parse optional branch weights [trueWeight, falseWeight]
- if (parser.parseOptionalLSquare()) {
- const weights = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer)) {
- weights.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- if (weights.length > 0) {
- result.addAttribute('branch_weights', weights);
- }
- }
- parser.parseComma();
- if (!result.successors) {
- result.successors = [];
- }
- const trueLabel = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const trueSucc = { label: trueLabel };
- if (parser.parseOptionalLParen()) {
- trueSucc.arguments = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- trueSucc.arguments.push(parser.parseOperand());
- parser.parseOptionalComma();
- } else {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (parser.getToken().isNot(_.Token.r_paren) && idx < trueSucc.arguments.length) {
- trueSucc.arguments[idx].type = parser.parseType();
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- }
- result.successors.push(trueSucc);
- parser.parseComma();
- const falseLabel = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const falseSucc = { label: falseLabel };
- if (parser.parseOptionalLParen()) {
- falseSucc.arguments = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- falseSucc.arguments.push(parser.parseOperand());
- parser.parseOptionalComma();
- } else {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (parser.getToken().isNot(_.Token.r_paren) && idx < falseSucc.arguments.length) {
- falseSucc.arguments[idx].type = parser.parseType();
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- }
- result.successors.push(falseSucc);
- return true;
- }
- if (result.op === 'spirv.CompositeConstruct' || result.op === 'spv.CompositeConstruct') {
- result.compatibility = true;
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.colon)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- if (parser.parseOptionalLParen()) {
- const types = parser.parseTypeList();
- parser.parseRParen();
- parser.parseArrow();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- const type = parser.parseType();
- result.types.push(type);
- } else {
- for (const unresolvedOp of unresolvedOperands) {
- parser.resolveOperand(unresolvedOp, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.SpecConstant' || result.op === 'spv.SpecConstant') {
- parser.parseSymbolName('sym_name', result.attributes);
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'spec_id') {
- parser.parseKeyword('spec_id');
- parser.parseLParen();
- const specId = parser.parseAttribute();
- result.addAttribute('spec_id', specId);
- parser.parseRParen();
- }
- if (parser.parseOptionalEqual()) {
- const defaultValue = parser.parseAttribute();
- result.addAttribute('default_value', defaultValue);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('type', type);
- }
- return true;
- }
- if (result.op === 'spirv.module' || result.op === 'spv.module') {
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const addressingModel = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('addressing_model', addressingModel);
- }
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const memoryModel = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('memory_model', memoryModel);
- }
- if (parser.parseOptionalKeyword('requires')) {
- const vce = parser.parseAttribute();
- result.addAttribute('vce_triple', vce);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'spirv.ARM.Graph') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'spirv.ARM.GraphEntryPoint') {
- const fn = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('fn', fn);
- const interfaceVars = [];
- while (parser.parseOptionalComma()) {
- const varSymbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- interfaceVars.push(varSymbol);
- }
- result.addAttribute('interface', interfaceVars);
- return true;
- }
- if (result.op === 'spirv.func' || result.op === 'spv.func') {
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- let inputs = [];
- const results = [];
- const resultAttrs = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- const argResult = parser.parseFunctionArgumentList();
- inputs = argResult.arguments.map((a) => a.type);
- }
- if (parser.parseOptionalArrow()) {
- parser.parseFunctionResultList(results, resultAttrs);
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- if (parser.getToken().is(_.Token.string)) {
- const control = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('function_control', control);
- }
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- // spirv.func is IsolatedFromAbove
- parser.parseRegion(region, undefined, /* isIsolatedNameScope */ true);
- }
- return true;
- }
- if (result.op === 'spirv.GlobalVariable' || result.op === 'spv.GlobalVariable') {
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- if (parser.parseOptionalKeyword('initializer')) {
- parser.parseLParen();
- const initSymbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- parser.parseRParen();
- result.addAttribute('initializer', initSymbol);
- }
- if (parser.parseOptionalKeyword('built_in')) {
- parser.parseLParen();
- const builtIn = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseRParen();
- result.addAttribute('built_in', builtIn);
- }
- if (parser.parseOptionalKeyword('bind')) {
- parser.parseLParen();
- const binding = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseOptionalComma();
- const set = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseRParen();
- result.addAttribute('descriptor_set', set);
- result.addAttribute('binding', binding);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.types = [type];
- }
- return true;
- }
- if (result.op === 'spirv.EntryPoint' || result.op === 'spv.EntryPoint') {
- if (parser.getToken().is(_.Token.string)) {
- const executionModel = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('execution_model', executionModel);
- }
- result.operands = [];
- while (parser.getToken().is(_.Token.at_identifier)) {
- const symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('fn', new _.SymbolRefAttr(symbol));
- parser.parseOptionalComma();
- }
- return true;
- }
- if (result.op === 'spirv.ExecutionMode' || result.op === 'spv.ExecutionMode') {
- if (parser.getToken().is(_.Token.at_identifier)) {
- const symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('fn', new _.SymbolRefAttr(symbol));
- }
- if (parser.getToken().is(_.Token.string)) {
- const mode = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('execution_mode', mode);
- }
- const params = [];
- while (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.floatliteral) || parser.getToken().is(_.Token.bare_identifier)) {
- const param = parser.getToken().getSpelling().str();
- parser.consumeToken();
- params.push(param);
- } else {
- break;
- }
- }
- if (params.length > 0) {
- result.addAttribute('values', params);
- }
- return true;
- }
- if (result.op === 'spirv.mlir.loop' || result.op === 'spv.mlir.loop' || result.op === 'spirv.mlir.selection' || result.op === 'spv.mlir.selection') {
- if (parser.parseOptionalKeyword('control')) {
- parser.parseLParen();
- const controlValue = parser.parseOptionalKeyword();
- result.addAttribute('selection_control', controlValue);
- parser.parseRParen();
- }
- result.addTypes(parser.parseOptionalArrowTypeList());
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- // spirv.CompositeInsert with 'into' keyword
- if (result.op === 'spirv.CompositeInsert' || result.op === 'spv.CompositeInsert') {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalLSquare()) {
- const indices = [];
- while (!parser.parseOptionalRSquare()) {
- const index = parser.getToken().getSpelling().str();
- parser.consumeToken();
- if (parser.parseOptionalColon()) {
- parser.consumeToken(); // Skip type (e.g., i32)
- }
- indices.push(index);
- parser.parseOptionalComma();
- }
- result.addAttribute('indices', indices);
- }
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- if (parser.parseOptionalKeyword('into')) {
- const resultType = parser.parseType();
- result.types = [resultType];
- }
- return true;
- }
- const arithmeticExtendedOps = new Set([
- 'spirv.IAddCarry', 'spv.IAddCarry',
- 'spirv.ISubBorrow', 'spv.ISubBorrow',
- 'spirv.SMulExtended', 'spv.SMulExtended',
- 'spirv.UMulExtended', 'spv.UMulExtended'
- ]);
- if (arithmeticExtendedOps.has(result.op)) {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalColon()) {
- const resultType = parser.parseType();
- parser.resolveOperands(unresolvedOperands, [resultType, resultType], result.operands);
- result.addTypes([resultType]);
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'spirv.INTEL.SubgroupBlockWrite' || result.op === 'spv.INTEL.SubgroupBlockWrite') {
- const storageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('storage_class', storageClass);
- const ptrUnresolved = parser.parseOperand();
- parser.parseComma();
- const valueUnresolved = parser.parseOperand();
- let ptrType = null;
- let valueType = null;
- if (parser.parseOptionalColon()) {
- valueType = parser.parseType();
- ptrType = new _.Type(`!spirv.ptr<${valueType}, ${storageClass}>`);
- }
- parser.resolveOperand(ptrUnresolved, ptrType, result.operands);
- parser.resolveOperand(valueUnresolved, valueType, result.operands);
- return true;
- }
- if ((result.op === 'spirv.CopyMemory' || result.op === 'spv.CopyMemory')) {
- const targetStorageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('target_storage_class', targetStorageClass);
- const targetUnresolved = parser.parseOperand();
- parser.parseComma();
- const sourceStorageClass = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('source_storage_class', sourceStorageClass);
- const sourceUnresolved = parser.parseOperand();
- if (parser.parseOptionalLSquare()) {
- const memoryAccess = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.string)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.string);
- } else if (parser.getToken().is(_.Token.integer)) {
- memoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- if (memoryAccess.length > 0) {
- result.addAttribute('memory_access', memoryAccess.join(', '));
- }
- }
- if (parser.parseOptionalComma()) {
- if (parser.parseOptionalLSquare()) {
- const sourceMemoryAccess = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.string)) {
- sourceMemoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.string);
- } else if (parser.getToken().is(_.Token.integer)) {
- sourceMemoryAccess.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.integer);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- if (sourceMemoryAccess.length > 0) {
- result.addAttribute('source_memory_access', sourceMemoryAccess.join(', '));
- }
- }
- }
- let targetType = null;
- let sourceType = null;
- if (parser.parseOptionalColon()) {
- const elementType = parser.parseType();
- targetType = new _.spirv.PointerType(elementType, targetStorageClass);
- sourceType = new _.spirv.PointerType(elementType, sourceStorageClass);
- }
- parser.resolveOperand(targetUnresolved, targetType, result.operands);
- parser.resolveOperand(sourceUnresolved, sourceType, result.operands);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseSPIRV_I32_1DArmTensor(parser, op, attrName) {
- const values = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus) || parser.getToken().is(_.Token.floatliteral)) {
- const value = parser.parseInteger();
- values.push(value);
- } else {
- break;
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- op.addAttribute(attrName, values);
- }
- };
- _.WasmSSADialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'wasmssa');
- this.registerCustomType('WasmSSA_LocalRef', this.parseLocalRefType.bind(this));
- this.registerCustomDirective('ElseRegion', this.parseElseRegion.bind(this));
- this.registerCustomAttribute('WasmSSA_ValTypeAttr', this.parseValTypeAttr.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'wasmssa.import_global') {
- const importName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('importName', importName);
- parser.parseKeyword('from');
- const moduleName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('moduleName', moduleName);
- parser.parseKeyword('as');
- parser.parseSymbolName('sym_name', result.attributes);
- if (parser.parseOptionalKeyword('mutable')) {
- result.addAttribute('isMutable', new _.UnitAttr());
- }
- parser.parseColon();
- const type = parser.parseType();
- result.addAttribute('type', type);
- return true;
- }
- if (result.op === 'wasmssa.global') {
- if (parser.parseOptionalKeyword('exported')) {
- result.addAttribute('exported', new _.UnitAttr());
- }
- parser.parseSymbolName('sym_name', result.attributes);
- const type = parser.parseType();
- result.addAttribute('type', type);
- if (parser.parseOptionalKeyword('mutable')) {
- result.addAttribute('isMutable', new _.UnitAttr());
- }
- parser.parseColon();
- const region = result.addRegion();
- parser.parseRegion(region);
- return true;
- }
- if (result.op === 'wasmssa.func') {
- if (parser.parseOptionalKeyword('exported')) {
- result.addAttribute('exported', new _.UnitAttr());
- }
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseType(parser) {
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'local') {
- parser.parseKeyword('local');
- parser.parseKeyword('ref');
- parser.parseKeyword('to');
- const elementType = parser.parseType();
- return new _.Type(`!wasmssa<local ref to ${elementType}>`);
- }
- return null;
- }
- parseLocalRefType(parser) {
- parser.parseKeyword('ref');
- parser.parseKeyword('to');
- const elementType = parser.parseType();
- return new _.Type(`ref to ${elementType}`);
- }
- parseElseRegion(parser, result) {
- if (parser.parseOptionalKeyword('else')) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- }
- inferResultTypes(op, vars) {
- if (op.op === 'wasmssa.local') {
- const typeAttr = op.attributes.get('type');
- if (typeAttr) {
- const elementType = typeAttr.type ? typeAttr.type : typeAttr;
- op.addTypes([new _.Type(`!wasmssa<local ref to ${elementType}>`)]);
- return;
- }
- }
- if (op.op === 'wasmssa.local_get') {
- const localVarEntry = vars.get('localVar');
- if (localVarEntry && localVarEntry.types && localVarEntry.types.length > 0) {
- const localRefType = localVarEntry.types[0];
- if (localRefType) {
- const typeStr = localRefType.toString();
- const match = typeStr.match(/ref\s+to\s+(.+)/);
- if (match) {
- op.addTypes([new _.Type(match[1])]);
- return;
- }
- }
- }
- }
- super.inferResultTypes(op, vars);
- }
- parseValTypeAttr(parser) {
- const type = parser.parseType();
- return new _.TypeAttrOf(type);
- }
- };
- _.CFDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'cf');
- this.registerCustomDirective('SwitchOpCases', this.parseSwitchOpCases.bind(this));
- }
- parseSwitchOpCases(parser, result) {
- if (!parser.parseOptionalKeyword('default')) {
- return false;
- }
- if (!parser.parseOptionalColon()) {
- return false;
- }
- if (parser.getToken().isNot(_.Token.caret_identifier)) {
- return false;
- }
- const defaultDestination = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const defaultDest = { label: defaultDestination, arguments: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const value = parser.parseOperand();
- defaultDest.arguments.push({ value });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (idx < defaultDest.arguments.length && parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- if (defaultDest.arguments[idx]) {
- defaultDest.arguments[idx].type = type;
- }
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRParen();
- }
- result.successors = result.successors || [];
- result.successors.push(defaultDest);
- const caseValues = [];
- const caseOperandSegments = [defaultDest.arguments.length];
- while (parser.parseOptionalComma()) {
- const value = parser.parseOptionalInteger();
- if (value === null) {
- break;
- }
- caseValues.push(value);
- if (!parser.parseOptionalColon()) {
- break;
- }
- if (parser.getToken().isNot(_.Token.caret_identifier)) {
- break;
- }
- const caseDestination = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const caseDest = { label: caseDestination, arguments: [] };
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const operandValue = parser.parseOperand();
- caseDest.arguments.push({ value: operandValue });
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (idx < caseDest.arguments.length && parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- if (caseDest.arguments[idx]) {
- caseDest.arguments[idx].type = type;
- }
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRParen();
- }
- result.successors.push(caseDest);
- caseOperandSegments.push(caseDest.arguments.length);
- }
- if (caseValues.length > 0) {
- result.addAttribute('case_values', caseValues);
- result.addAttribute('case_operand_segments', caseOperandSegments);
- }
- return true;
- }
- };
- _.pdl = {};
- _.pdl.ValueType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!pdl.value';
- }
- };
- _.pdl.TypeType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!pdl.type';
- }
- };
- _.pdl.AttributeType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!pdl.attribute';
- }
- };
- _.pdl.OperationType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!pdl.operation';
- }
- };
- _.pdl.RangeType = class extends _.Type {
- constructor(elementType) {
- super(null);
- this.elementType = elementType;
- }
- static getElementTypeOrSelf(type) {
- if (type instanceof _.pdl.RangeType) {
- return type.elementType;
- }
- return type;
- }
- toString() {
- const elementStr = this.elementType?.toString() || '';
- const match = elementStr.match(/^!pdl\.(.+)$/);
- return `!pdl.range<${match ? match[1] : elementStr}>`;
- }
- };
- _.pdl.PDLDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'pdl');
- this.registerCustomDirective('OperationOpAttributes', this.parseOperationOpAttributes.bind(this));
- this.registerCustomDirective('RangeType', this.parseRangeType.bind(this));
- this.registerCustomDirective('ResultsValueType', this.parseResultsValueType.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'pdl.operation') {
- result.compatibility = true;
- return this.parseOperationOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseOperationOp(parser, result) {
- if (parser.getToken().is(_.Token.string)) {
- const opNameValue = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('opName', opNameValue);
- }
- if (parser.parseOptionalLParen()) {
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- const types = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- types.push(parser.parseType());
- parser.parseOptionalComma();
- }
- }
- for (let i = 0; i < unresolvedOperands.length; i++) {
- parser.resolveOperand(unresolvedOperands[i], types[i] || null, result.operands);
- }
- parser.parseOptionalRParen();
- }
- this.parseOperationOpAttributes(parser, result);
- if (parser.parseOptionalArrow()) {
- parser.parseOptionalLParen();
- const unresolvedTypeValues = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace) && !(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'loc')) {
- unresolvedTypeValues.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- const types = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.l_brace) && !(parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'loc')) {
- types.push(parser.parseType());
- parser.parseOptionalComma();
- }
- }
- for (let i = 0; i < unresolvedTypeValues.length; i++) {
- parser.resolveOperand(unresolvedTypeValues[i], types[i] || null, result.operands);
- }
- parser.parseOptionalRParen();
- }
- result.addTypes([new _.Type('!pdl.operation')]);
- return true;
- }
- parseOperationOpAttributes(parser, result) {
- if (!parser.parseOptionalLBrace()) {
- return true;
- }
- const attributeNames = [];
- while (parser.getToken().isNot(_.Token.r_brace)) {
- const name = parser.parseAttribute();
- if (!parser.parseOptionalEqual()) {
- break;
- }
- const unresolvedValue = parser.parseOperand();
- parser.resolveOperand(unresolvedValue, null, result.operands);
- attributeNames.push(name);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRBrace();
- if (attributeNames.length > 0) {
- result.addAttribute('attributeValueNames', attributeNames);
- }
- return true;
- }
- parseRangeType(parser, op, argumentTypes, resultTypes) {
- if (argumentTypes && argumentTypes.length > 0) {
- const elementType = _.pdl.RangeType.getElementTypeOrSelf(argumentTypes[0]);
- resultTypes.push(new _.pdl.RangeType(elementType));
- return true;
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- resultTypes.push(type);
- }
- return true;
- }
- parseResultsValueType(parser, result) {
- // Parses `-> type` for pdl.results operation
- // If index is present, type can be !pdl.value or !pdl.range<value>
- // If index is absent, type is always !pdl.range<value> (full result range)
- if (parser.parseOptionalArrow()) {
- const type = parser.parseType();
- result.addTypes([type]);
- } else {
- // Default to !pdl.range<value> when no explicit type is given
- result.addTypes([new _.Type('!pdl.range<!pdl.value>')]);
- }
- return true;
- }
- };
- _.PDLInterpDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'pdl_interp');
- this.registerCustomDirective('CreateOperationOpAttributes', this.parseCreateOperationOpAttributes.bind(this));
- this.registerCustomDirective('CreateOperationOpResults', this.parseCreateOperationOpResults.bind(this));
- this.registerCustomDirective('RangeType', this.parseRangeType.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'pdl_interp.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'pdl_interp.foreach') {
- return this.parseForeachOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseForeachOp(parser, result) {
- const loopVar = parser.parseOperand();
- parser.parseColon();
- const loopVarType = parser.parseType();
- parser.parseKeyword('in');
- const range = parser.parseOperand();
- parser.resolveOperand(range, null, result.operands);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- if (region.blocks && region.blocks.length > 0) {
- if (!region.blocks[0].arguments) {
- region.blocks[0].arguments = [];
- }
- region.blocks[0].arguments.push({ value: loopVar, type: loopVarType });
- }
- result.regions.push(region);
- }
- if (parser.parseOptionalArrow()) {
- parser.consumeToken(_.Token.caret_identifier);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseCreateOperationOpAttributes(parser, result) {
- const attrNames = [];
- if (parser.parseOptionalLBrace()) {
- while (parser.getToken().isNot(_.Token.r_brace)) {
- const nameAttr = parser.parseAttribute();
- parser.parseEqual();
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- attrNames.push(nameAttr);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRBrace();
- }
- if (attrNames.length > 0) {
- result.addAttribute('inputAttributeNames', attrNames);
- }
- }
- parseCreateOperationOpResults(parser, result) {
- if (!parser.parseOptionalArrow()) {
- return;
- }
- if (parser.parseOptionalLess()) {
- parser.parseKeyword('inferred');
- parser.parseGreater();
- result.addAttribute('inferredResultTypes', true);
- return;
- }
- parser.parseLParen();
- const unresolvedOperands = [];
- const types = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- do {
- const type = parser.parseType();
- types.push(type);
- } while (parser.parseOptionalComma());
- }
- for (let i = 0; i < unresolvedOperands.length; i++) {
- const type = i < types.length ? types[i] : null;
- parser.resolveOperand(unresolvedOperands[i], type, result.operands);
- }
- parser.parseRParen();
- }
- parseRangeType(parser, op, argumentTypes, resultTypes) {
- if (argumentTypes && argumentTypes.length > 0) {
- const elementType = _.pdl.RangeType.getElementTypeOrSelf(argumentTypes[0]);
- resultTypes.push(new _.pdl.RangeType(elementType));
- return true;
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- resultTypes.push(type);
- }
- return true;
- }
- };
- _.ptr = {};
- _.ptr.PtrType = class extends _.Type {
- constructor(memorySpace) {
- super(null);
- this.memorySpace = memorySpace;
- }
- toString() {
- return `!ptr.ptr<${this.memorySpace}>`;
- }
- };
- _.ptr.PtrDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'ptr');
- this.registerCustomAttribute('EnumProp', this.parseEnumProp.bind(this));
- this.registerCustomAttribute('Ptr_PtrDiffFlags', this.parsePtrDiffFlags.bind(this));
- this.registerCustomType('Ptr_PtrType', this.parsePtrType.bind(this));
- }
- parseEnumProp(parser, type) {
- const [innerType] = type.args;
- return this.parseCustomAttributeWithFallback(parser, innerType);
- }
- parsePtrDiffFlags(parser, type) {
- if (type.values.includes(parser.getTokenSpelling().str())) {
- return this.parseEnumFlags(parser, type, '|');
- }
- return null;
- }
- parsePtrType(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- const memorySpace = content.slice(1, -1);
- return new _.ptr.PtrType(memorySpace);
- }
- return parser.parseType();
- }
- inferResultTypes(op, vars) {
- if (op.op === 'ptr.ptr_add' && op.operands.length >= 2) {
- const baseType = op.operands[0].type;
- const offsetType = op.operands[1].type;
- const offsetIsShaped = offsetType instanceof _.VectorType || offsetType instanceof _.RankedTensorType;
- if (!offsetIsShaped) {
- if (baseType) {
- op.addTypes([baseType]);
- }
- return;
- }
- const baseIsShaped = baseType instanceof _.VectorType || baseType instanceof _.RankedTensorType;
- if (!baseIsShaped) {
- if (offsetType instanceof _.VectorType) {
- op.addTypes([new _.VectorType(offsetType.dimensions, baseType, offsetType.scalableDims)]);
- } else if (offsetType instanceof _.RankedTensorType) {
- op.addTypes([new _.RankedTensorType(offsetType.shape, baseType, offsetType.encoding)]);
- }
- return;
- }
- if (baseType) {
- op.addTypes([baseType]);
- }
- return;
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.EmitCDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'emitc');
- this.registerCustomType('EmitC_LValueType', this.parseLValueType.bind(this));
- this.registerCustomDirective('SwitchCases', this.parseSwitchCases.bind(this));
- this.registerCustomDirective('EmitCGlobalOpTypeAndInitialValue', this.parseTypeAndInitialValue.bind(this));
- this.registerCustomDirective('EmitCFieldOpTypeAndInitialValue', this.parseTypeAndInitialValue.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'emitc.include') {
- if (parser.parseOptionalLess()) {
- const include = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseGreater();
- result.addAttribute('is_standard_include', true);
- result.addAttribute('include', include);
- } else {
- const include = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('include', include);
- }
- return true;
- }
- if (result.op === 'emitc.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'emitc.expression') {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalKeyword('noinline')) {
- result.addAttribute('do_not_inline', true);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- result.addAttribute('type', type);
- result.addTypes(type.results);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'emitc.if') {
- const cond = parser.parseOperand();
- parser.resolveOperand(cond, null, result.operands);
- const thenRegion = {};
- parser.parseRegion(thenRegion);
- result.regions.push(thenRegion);
- if (parser.parseOptionalKeyword('else')) {
- const elseRegion = {};
- parser.parseRegion(elseRegion);
- result.regions.push(elseRegion);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (result.op === 'emitc.do') {
- const bodyRegion = {};
- parser.parseRegion(bodyRegion);
- result.regions.push(bodyRegion);
- parser.parseKeyword('while');
- const condRegion = {};
- parser.parseRegion(condRegion);
- result.regions.push(condRegion);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- return true;
- }
- if (result.op === 'emitc.for') {
- const iterVar = parser.parseOperand();
- parser.parseEqual();
- const lb = parser.parseOperand();
- parser.resolveOperand(lb, null, result.operands);
- parser.parseKeyword('to');
- const ub = parser.parseOperand();
- parser.resolveOperand(ub, null, result.operands);
- parser.parseKeyword('step');
- const step = parser.parseOperand();
- parser.resolveOperand(step, null, result.operands);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('type', type.toString());
- }
- result.addAttribute('iterVar', { value: iterVar, hidden: true });
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseLValueType(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!emitc.lvalue${content}`);
- }
- return null;
- }
- parseSwitchCases(parser, op /*, args */) {
- const caseValues = [];
- while (parser.parseOptionalKeyword('case')) {
- const value = parser.parseInteger();
- caseValues.push(value);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = op.addRegion();
- parser.parseRegion(region);
- }
- }
- if (caseValues.length > 0) {
- op.addAttribute('cases', caseValues);
- }
- }
- parseTypeAndInitialValue(parser, op, typeAttr = 'type', valueAttr = 'initial_value') {
- const type = parser.parseType();
- op.addAttribute(typeAttr, type);
- if (parser.parseOptionalEqual()) {
- const initialValue = parser.parseAttribute(type);
- op.addAttribute(valueAttr, initialValue);
- }
- }
- };
- _.AsukaDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'asuka');
- }
- parseOperation(parser, result) {
- // https://github.com/monellz/FlashTensor/blob/main/bench/ea.mlir
- // uses batch_dims and reduce_dims not valid given the assemblyFormat spec.
- // Custom parsing preserves compatibility with this file.
- if (result.op === 'asuka.dot' || result.op === 'asuka.add' || result.op === 'asuka.split' || result.op === 'asuka.softmax' || result.op === 'asuka.reduce') {
- result.compatibility = true;
- result.operands = parser.parseOperandList();
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace)) {
- const attrName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalEqual()) {
- let attrValue = null;
- if (parser.getToken().is(_.Token.l_square)) {
- attrValue = parser.parseAttribute();
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'x') {
- parser.consumeToken(_.Token.bare_identifier); // consume 'x'
- const secondValue = parser.parseAttribute();
- attrValue = { kind: 'pair', first: attrValue, second: secondValue };
- }
- } else if (parser.getToken().is(_.Token.integer)) {
- attrValue = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- } else {
- attrValue = parser.parseAttribute();
- }
- result.addAttribute(attrName, attrValue);
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- parser.resolveOperands(result.operands, funcType.inputs);
- for (const resultType of funcType.results) {
- result.addTypes([resultType]);
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.async = {};
- _.async.TokenType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!async.token';
- }
- };
- _.async.GroupType = class extends _.Type {
- constructor() {
- super(null);
- }
- toString() {
- return '!async.group';
- }
- };
- _.async.ValueType = class extends _.Type {
- constructor(valueType) {
- super(null);
- this.valueType = valueType;
- }
- toString() {
- const inner = this.valueType?.toString ? this.valueType.toString() : this.valueType;
- return `!async.value<${inner}>`;
- }
- static parse(parser) {
- if (parser.consumeIf(_.Token.less)) {
- const innerType = parser.parseType();
- parser.parseGreater();
- return new _.async.ValueType(innerType);
- }
- return parser.parseType();
- }
- };
- _.async.AsyncDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'async');
- this.registerCustomDirective('AwaitResultType', this.parseAwaitResultType.bind(this));
- this.registerCustomType('Async_ValueType', (parser) => _.async.ValueType.parse(parser));
- }
- parseType(parser, dialect) {
- const mnemonic = parser.parseOptionalKeyword();
- if (mnemonic === 'coro.handle' || mnemonic === 'coro.id' || mnemonic === 'coro.state') {
- return new _.Type(`!${dialect}.${mnemonic}`);
- }
- if (mnemonic === 'token') {
- return new _.async.TokenType();
- }
- if (mnemonic === 'group') {
- return new _.async.GroupType();
- }
- if (mnemonic === 'value') {
- if (parser.consumeIf(_.Token.less)) {
- const innerType = parser.parseType();
- parser.parseGreater();
- return new _.async.ValueType(innerType);
- }
- return new _.async.ValueType(null);
- }
- throw new mlir.Error(`Unknown '${dialect}' type '${mnemonic}' ${parser.getNameLoc()}`);
- }
- parseOperation(parser, result) {
- if (result.op === 'async.execute') {
- return this.parseExecuteOp(parser, result);
- }
- if (result.op === 'async.func') {
- return this.parseFuncOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseExecuteOp(parser, result) {
- const tokenArgs = parser.parseOperandList('optionalSquare');
- const tokenTypes = tokenArgs.map(() => null);
- parser.resolveOperands(tokenArgs, tokenTypes, result.operands);
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- if (parser.parseOptionalKeyword('as')) {
- parser.parseOperand();
- }
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- const valueTypes = [];
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- valueTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- } else {
- valueTypes.push(parser.parseType());
- }
- }
- result.addTypes([new _.async.TokenType()]);
- result.addTypes(valueTypes);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- parseFuncOp(parser, result) {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- parser.parseSymbolName('sym_name', result.attributes);
- const argResult = parser.parseFunctionArgumentList();
- const inputs = argResult.arguments.map((a) => a.type);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- const results = [];
- const resultAttrs = [];
- if (parser.parseOptionalArrow()) {
- parser.parseFunctionResultList(results, resultAttrs);
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- parseAwaitResultType(parser, op, operandTypeArg) {
- const operandType = parser.parseType();
- if (Array.isArray(operandTypeArg)) {
- operandTypeArg.push(operandType);
- }
- if (operandType instanceof _.async.ValueType && operandType.valueType) {
- op.addTypes([operandType.valueType]);
- }
- }
- };
- _.ArithDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'arith');
- this.registerCustomAttribute('Arith_FastMathAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- this.registerCustomAttribute('Arith_IntegerOverflowAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'arith.select') {
- return this.parseSelectOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseSelectOp(parser, result) {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const condType = parser.parseType();
- if (parser.parseOptionalComma()) {
- const resultType = parser.parseType();
- const types = [condType, resultType, resultType];
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (result.types.length > 0) {
- result.types[0] = resultType;
- } else {
- result.addTypes([resultType]);
- }
- } else {
- const types = unresolvedOperands.map(() => condType);
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (result.types.length > 0) {
- result.types[0] = condType;
- } else {
- result.addTypes([condType]);
- }
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- };
- _.BuiltinDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'builtin');
- this.blobManager = new Map();
- }
- parseOperation(parser, result) {
- if (result.op === 'builtin.call' || result.op === 'builtin.call_indirect') {
- parser.parseSymbolName('callee', result.attributes);
- const unresolvedOperands = parser.parseOperandList();
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- readType(reader) {
- const typeCode = reader.readVarInt();
- switch (typeCode) {
- case 0: { // IntegerType
- const widthAndSign = reader.readVarInt();
- const width = widthAndSign >> 2;
- const signedness = widthAndSign & 0x3;
- if (signedness === 0) {
- return new _.IntegerType(`i${width}`);
- }
- if (signedness === 1) {
- return new _.IntegerType(`si${width}`);
- }
- return new _.IntegerType(`ui${width}`);
- }
- case 1: // IndexType
- return new _.IndexType();
- case 2: { // FunctionType
- const numInputs = reader.readVarInt();
- const inputs = [];
- for (let i = 0; i < numInputs; i++) {
- inputs.push(reader.readType());
- }
- const numResults = reader.readVarInt();
- const results = [];
- for (let i = 0; i < numResults; i++) {
- results.push(reader.readType());
- }
- return new _.FunctionType(inputs, results);
- }
- case 3: return new _.FloatType('bf16'); // BFloat16Type
- case 4: return new _.FloatType('f16'); // Float16Type
- case 5: return new _.FloatType('f32'); // Float32Type
- case 6: return new _.FloatType('f64'); // Float64Type
- case 7: return new _.FloatType('f80'); // Float80Type
- case 8: return new _.FloatType('f128'); // Float128Type
- case 9: { // ComplexType
- const elementType = reader.readType();
- return new _.Type(`complex<${elementType.toString()}>`);
- }
- case 10: { // MemRefType
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- reader.readAttribute(); // layout
- return new _.Type(`memref<${shape.join('x')}x${elementType.toString()}>`);
- }
- case 11: { // MemRefTypeWithMemSpace
- reader.readAttribute(); // memorySpace
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- reader.readAttribute(); // layout
- return new _.Type(`memref<${shape.join('x')}x${elementType.toString()}>`);
- }
- case 12: // NoneType
- return new _.NoneType();
- case 13: { // RankedTensorType
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- return new _.RankedTensorType(shape, elementType, null);
- }
- case 14: { // RankedTensorTypeWithEncoding
- const encoding = reader.readAttribute();
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- return new _.RankedTensorType(shape, elementType, encoding);
- }
- case 15: { // TupleType
- const numTypes = reader.readVarInt();
- const types = [];
- for (let i = 0; i < numTypes; i++) {
- types.push(reader.readType());
- }
- return new _.Type(`tuple<${types.map((t) => t.toString()).join(', ')}>`);
- }
- case 16: { // UnrankedMemRefType
- const elementType = reader.readType();
- return new _.Type(`memref<*x${elementType.toString()}>`);
- }
- case 17: { // UnrankedMemRefTypeWithMemSpace
- reader.readAttribute(); // memorySpace
- const elementType = reader.readType();
- return new _.Type(`memref<*x${elementType.toString()}>`);
- }
- case 18: { // UnrankedTensorType
- const elementType = reader.readType();
- return new _.UnrankedTensorType(elementType);
- }
- case 19: { // VectorType
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- return new _.VectorType(shape, elementType);
- }
- case 20: { // VectorTypeWithScalableDims
- const numScalable = reader.readVarInt();
- for (let i = 0; i < numScalable; i++) {
- reader.readByte(); // scalableDims flags
- }
- const shape = reader.readSignedVarInts();
- const elementType = reader.readType();
- return new _.VectorType(shape, elementType);
- }
- default:
- throw new mlir.Error(`Unsupported built-in type code '${typeCode}'.`);
- }
- }
- readAttribute(reader) {
- const typeCode = reader.readVarInt();
- switch (typeCode) {
- case 0: { // ArrayAttr
- const count = reader.readVarInt();
- const elements = [];
- for (let i = 0; i < count; i++) {
- elements.push(reader.readAttribute());
- }
- return new _.ArrayAttr(elements);
- }
- case 1: { // DictionaryAttr
- const count = reader.readVarInt();
- const attrs = new Map();
- for (let i = 0; i < count; i++) {
- const nameAttr = reader.readAttribute();
- const valueAttr = reader.readAttribute();
- const name = nameAttr && nameAttr.value ? nameAttr.value : `attr_${i}`;
- attrs.set(name, valueAttr);
- }
- return { name: 'dictionary', value: attrs };
- }
- case 2: { // StringAttr
- const value = reader.readString();
- return new _.StringAttr(value);
- }
- case 3: { // StringAttrWithType
- const value = reader.readString();
- const type = reader.readType();
- return new _.StringAttr(value, type);
- }
- case 4: { // FlatSymbolRefAttr
- const value = reader.readString();
- return new _.SymbolRefAttr(`@${value}`);
- }
- case 5: { // SymbolRefAttr
- const root = reader.readString();
- const numNested = reader.readVarInt();
- let value = `@${root}`;
- for (let i = 0; i < numNested; i++) {
- value += `::@${reader.readString()}`;
- }
- return new _.SymbolRefAttr(value);
- }
- case 6: { // TypeAttr
- const type = reader.readType();
- return new _.TypeAttrOf(type);
- }
- case 7: // UnitAttr
- return new _.UnitAttr();
- case 8: { // IntegerAttr
- const getIntegerBitWidth = (type) => {
- const str = type ? type.toString() : '';
- const match = str.match(/^[su]?i(\d+)$/);
- if (match) {
- return parseInt(match[1], 10);
- }
- if (str === 'index') {
- return 64;
- }
- throw new mlir.Error(`Unsupported integer type '${str}'.`);
- };
- const type = reader.readType();
- const bitWidth = getIntegerBitWidth(type);
- let value = null;
- if (bitWidth <= 8) {
- value = BigInt(reader.readByte());
- } else if (bitWidth <= 64) {
- value = reader.readSignedVarInt();
- } else {
- const numWords = reader.readVarInt();
- value = 0n;
- for (let i = 0; i < numWords; i++) {
- const word = reader.readSignedVarInt();
- value |= (word << BigInt(i * 64));
- }
- }
- return new _.IntegerAttr(type, value);
- }
- case 9: { // FloatAttr
- const type = reader.readType();
- const value = reader.readAPFloatWithKnownSemantics(type);
- return new _.FloatAttr(type, value);
- }
- case 10: { // CallSiteLoc
- const caller = reader.readAttribute();
- const callee = reader.readAttribute();
- const callerStr = caller && caller.value ? caller.value : '<caller>';
- const calleeStr = callee && callee.value ? callee.value : '<callee>';
- return { name: 'loc', value: `callsite(${callerStr} at ${calleeStr})` };
- }
- case 11: { // FileLineColLoc
- const filename = reader.readString();
- const line = reader.readVarInt();
- const col = reader.readVarInt();
- return { name: 'loc', value: `${filename}:${line}:${col}` };
- }
- case 12: { // FusedLoc
- const count = reader.readVarInt();
- const locations = [];
- for (let i = 0; i < count; i++) {
- const loc = reader.readAttribute();
- locations.push(loc && loc.value ? loc.value : '<loc>');
- }
- return { name: 'loc', value: `fused[${locations.join(', ')}]` };
- }
- case 13: { // FusedLocWithMetadata
- const metadata = reader.readAttribute();
- const count = reader.readVarInt();
- const locations = [];
- for (let i = 0; i < count; i++) {
- const loc = reader.readAttribute();
- locations.push(loc && loc.value ? loc.value : '<loc>');
- }
- const metaStr = metadata && metadata.value !== undefined ? metadata.value : '<meta>';
- return { name: 'loc', value: `fused<${metaStr}>[${locations.join(', ')}]` };
- }
- case 14: { // NameLoc
- const nameAttr = reader.readAttribute();
- const childLoc = reader.readAttribute();
- const nameStr = nameAttr && nameAttr.value !== undefined ? nameAttr.value : '<name>';
- const childStr = childLoc && childLoc.value ? childLoc.value : '<loc>';
- return { name: 'loc', value: `#loc(${nameStr}(${childStr}))` };
- }
- case 15: // UnknownLoc
- return { name: 'loc', value: 'unknown' };
- case 16: { // DenseResourceElementsAttr
- const type = reader.readType();
- const resource = reader.readResourceHandle();
- const blobData = resource && resource.value && resource.value.kind === 'blob' ? resource.value.data : null;
- const handle = new _.DenseResourceElementsHandle(resource ? resource.key : 'unknown', blobData);
- return new _.DenseResourceElementsAttr(type, handle);
- }
- case 17: { // DenseArrayAttr
- const type = reader.readType();
- const size = reader.readVarInt();
- const blob = reader.readBlob();
- return new _.DenseArrayAttr(type, size, blob);
- }
- case 18: { // DenseIntOrFPElementsAttr
- const type = reader.readType();
- const blob = reader.readBlob();
- return new _.DenseElementsAttr(blob, type);
- }
- case 19: { // DenseStringElementsAttr
- const type = reader.readType();
- const isSplat = reader.readVarInt() !== 0;
- const count = reader.readVarInt();
- const strings = [];
- for (let i = 0; i < count; i++) {
- strings.push(reader.readString());
- }
- return { name: 'dense_string', value: strings, type, isSplat };
- }
- case 20: { // SparseElementsAttr
- const type = reader.readType();
- const indices = reader.readAttribute();
- const values = reader.readAttribute();
- return new _.SparseElementsAttr(type, indices, values);
- }
- case 21: { // DistinctAttr
- const referencedAttr = reader.readAttribute();
- return { name: 'distinct', value: referencedAttr };
- }
- case 22: { // FileLineColRange
- const filename = reader.readString();
- const numLocs = reader.readVarInt();
- const locs = [];
- for (let i = 0; i < numLocs; i++) {
- locs.push(reader.readVarInt());
- }
- return { name: 'loc', value: `${filename}:${locs.join(':')}` };
- }
- default:
- return { name: 'builtin', value: `<builtin code ${typeCode}>` };
- }
- }
- declareResource(key) {
- if (!this.blobManager.has(key)) {
- this.blobManager.set(key, new _.DenseResourceElementsHandle(key));
- }
- return this.blobManager.get(key);
- }
- getResourceKey(handle) {
- return handle.key;
- }
- parseResource(entry) {
- const blob = entry.parseAsBlob();
- this.blobManager.get(entry.key).blob = blob;
- }
- };
- _.BufferizationDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'bufferization');
- }
- parseOperation(parser, result) {
- if (result.op === 'bufferization.alloc_tensor') {
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- const unresolvedDynamicDims = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedDynamicDims.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else {
- break;
- }
- }
- parser.parseRParen();
- let unresolvedCopy = null;
- if (parser.parseOptionalKeyword('copy')) {
- parser.parseLParen();
- unresolvedCopy = parser.parseOperand();
- parser.parseRParen();
- }
- let unresolvedSizeHint = null;
- if (parser.parseOptionalKeyword('size_hint')) {
- parser.parseEqual();
- unresolvedSizeHint = parser.parseOperand();
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const resultType = parser.parseType();
- const indexType = new _.IndexType();
- parser.resolveOperands(unresolvedDynamicDims, unresolvedDynamicDims.map(() => indexType), result.operands);
- if (unresolvedCopy) {
- parser.resolveOperand(unresolvedCopy, resultType, result.operands);
- }
- if (unresolvedSizeHint) {
- parser.resolveOperand(unresolvedSizeHint, indexType, result.operands);
- }
- if (result.types.length === 0) {
- result.types.push(resultType);
- } else {
- result.types[0] = resultType;
- }
- }
- return true;
- }
- // bufferization.to_memref %tensor read_only : tensor_type to memref_type
- if (result.op === 'bufferization.to_memref') {
- let unresolvedOperand = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperand = parser.parseOperand();
- }
- if (parser.parseOptionalKeyword('read_only')) {
- result.addAttribute('read_only', true);
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const sourceType = parser.parseType();
- result.addAttribute('source_type', sourceType);
- if (unresolvedOperand) {
- parser.resolveOperand(unresolvedOperand, sourceType, result.operands);
- }
- parser.parseKeyword('to');
- const destType = parser.parseType();
- result.addTypes([destType]);
- } else if (unresolvedOperand) {
- parser.resolveOperand(unresolvedOperand, null, result.operands);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- inferResultTypes(op, vars) {
- if (op.op === 'bufferization.dealloc') {
- // DeallocOp::inferReturnTypes - one i1 per retained memref
- const retainedEntry = vars.get('retained');
- const numRetained = retainedEntry?.operands?.length || 0;
- const i1Type = new _.IntegerType('i1');
- for (let i = 0; i < numRetained; i++) {
- op.addTypes([i1Type]);
- }
- return;
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.SCFDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'scf');
- this.registerCustomDirective('SwitchCases', this.parseSwitchCases.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'scf.for') {
- return this.parseForOp(parser, result);
- }
- if (result.op === 'scf.if') {
- return this.parseIfOp(parser, result);
- }
- if (result.op === 'scf.while') {
- return this.parseWhileOp(parser, result);
- }
- if (result.op === 'scf.forall') {
- return this.parseForallOp(parser, result);
- }
- if (result.op === 'scf.forall.in_parallel') {
- return this.parseInParallelOp(parser, result);
- }
- if (result.op === 'scf.parallel') {
- return this.parseParallelOp(parser, result);
- }
- if (result.op === 'scf.execute_region') {
- return this.parseExecuteRegionOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseForOp(parser, result) {
- if (parser.parseOptionalKeyword('unsigned')) {
- result.addAttribute('unsignedCmp', true);
- }
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- return false;
- }
- const inductionVar = parser.parseOperand();
- if (!parser.parseOptionalEqual()) {
- return false;
- }
- const indexType = new _.IndexType();
- let unresolvedLb = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedLb = parser.parseOperand();
- } else {
- return false;
- }
- if (!parser.parseOptionalKeyword('to')) {
- return false;
- }
- let unresolvedUb = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedUb = parser.parseOperand();
- } else {
- return false;
- }
- if (!parser.parseOptionalKeyword('step')) {
- return false;
- }
- let unresolvedStep = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedStep = parser.parseOperand();
- } else {
- return false;
- }
- parser.resolveOperands([unresolvedLb, unresolvedUb, unresolvedStep], [indexType, indexType, indexType], result.operands);
- let initArgsCount = 0;
- if (parser.parseOptionalKeyword('iter_args')) {
- const unresolvedIterArgs = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // Skip the loop-carried variable name
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedIterArgs.push(parser.parseOperand());
- } else {
- const value = parser.parseAttribute();
- if (value) {
- // Attribute values aren't operands - skip for now
- }
- }
- }
- parser.parseOptionalComma();
- }
- }
- result.addTypes(parser.parseArrowTypeList());
- const iterArgTypes = result.types.map((t) => t || indexType);
- parser.resolveOperands(unresolvedIterArgs, iterArgTypes, result.operands);
- initArgsCount = unresolvedIterArgs.length;
- }
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- if (region.blocks && region.blocks.length > 0) {
- if (!region.blocks[0].arguments) {
- region.blocks[0].arguments = [];
- }
- if (region.blocks[0].arguments.length > 0) {
- region.blocks[0].arguments[0] = { value: inductionVar };
- } else {
- region.blocks[0].arguments.push({ value: inductionVar });
- }
- }
- result.regions.push(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addAttribute('operandSegmentSizes', [1, 1, 1, initArgsCount]);
- return true;
- }
- parseIfOp(parser, result) {
- // Reference impl: condition operand is of type i1
- let unresolvedCond = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedCond = parser.parseOperand();
- } else {
- return false;
- }
- const i1Type = new _.IntegerType('i1');
- parser.resolveOperands([unresolvedCond], [i1Type], result.operands);
- result.addTypes(parser.parseOptionalArrowTypeList());
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- } else {
- return false;
- }
- if (parser.parseOptionalKeyword('else')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseWhileOp(parser, result) {
- const unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // Skip variable name
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- // Note: attribute values are not operands, skip them
- }
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalColon()) {
- const types = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- types.push(parser.parseType());
- parser.parseOptionalComma();
- }
- } else {
- types.push(parser.parseType());
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- result.addTypes(parser.parseOptionalArrowTypeList());
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- if (parser.parseOptionalKeyword('do')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- return true;
- }
- parseForallOp(parser, result) {
- const indexType = new _.IndexType();
- const inductionVars = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- inductionVars.push(parser.parseOperand().name);
- } else {
- return false;
- }
- if (!parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseOptionalRParen();
- break;
- }
- return false;
- }
- }
- const isNormalized = parser.parseOptionalKeyword('in');
- if (!isNormalized && !parser.parseOptionalEqual()) {
- return false;
- }
- // Helper to parse bounds list - only SSA values become operands, integers are static
- const parseBoundsList = () => {
- const bounds = [];
- if (!parser.parseOptionalLParen()) {
- return bounds;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- bounds.push(parser.parseOperand());
- } else if (parser.getToken().is(_.Token.integer)) {
- parser.consumeToken(_.Token.integer); // Skip static bound
- }
- parser.parseOptionalComma();
- }
- return bounds;
- };
- if (isNormalized) {
- // Normalized form: in (bounds)
- const bounds = parseBoundsList();
- parser.resolveOperands(bounds, bounds.map(() => indexType), result.operands);
- } else {
- // Range form: = (lb) to (ub) step (step)
- const lowerBounds = parseBoundsList();
- parser.resolveOperands(lowerBounds, lowerBounds.map(() => indexType), result.operands);
- if (!parser.parseOptionalKeyword('to')) {
- return false;
- }
- const upperBounds = parseBoundsList();
- parser.resolveOperands(upperBounds, upperBounds.map(() => indexType), result.operands);
- if (!parser.parseOptionalKeyword('step')) {
- return false;
- }
- const steps = parseBoundsList();
- parser.resolveOperands(steps, steps.map(() => indexType), result.operands);
- }
- if (parser.parseOptionalKeyword('shared_outs')) {
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // Skip arg name
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- } else {
- parser.parseAttribute(); // Skip attribute value
- }
- }
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const type = parser.parseType();
- result.addTypes([type]);
- parser.parseOptionalComma();
- }
- } else {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- } else {
- return false;
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseParallelOp(parser, result) {
- const indexType = new _.IndexType();
- const inductionVars = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- inductionVars.push(parser.parseOperand().name);
- } else {
- return false;
- }
- parser.parseOptionalComma();
- }
- if (!parser.parseOptionalEqual()) {
- return false;
- }
- const lowerBounds = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- lowerBounds.push(parser.parseOperand());
- } else {
- return false;
- }
- parser.parseOptionalComma();
- }
- parser.resolveOperands(lowerBounds, lowerBounds.map(() => indexType), result.operands);
- if (!parser.parseOptionalKeyword('to')) {
- return false;
- }
- const upperBounds = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- upperBounds.push(parser.parseOperand());
- } else {
- return false;
- }
- parser.parseOptionalComma();
- }
- parser.resolveOperands(upperBounds, upperBounds.map(() => indexType), result.operands);
- if (!parser.parseOptionalKeyword('step')) {
- return false;
- }
- const steps = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- steps.push(parser.parseOperand());
- } else {
- return false;
- }
- parser.parseOptionalComma();
- }
- parser.resolveOperands(steps, steps.map(() => indexType), result.operands);
- if (parser.parseOptionalKeyword('init')) {
- const initVals = [];
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- initVals.push(parser.parseOperand());
- } else {
- const value = parser.parseAttribute();
- if (value) {
- initVals.push(value);
- }
- }
- parser.parseOptionalComma();
- }
- // Init values type inferred from definition
- parser.resolveOperands(initVals, initVals.map(() => null), result.operands);
- }
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const type = parser.parseType();
- result.addTypes([type]);
- parser.parseOptionalComma();
- }
- } else {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- if (region.blocks && region.blocks.length > 0 && inductionVars.length > 0) {
- if (!region.blocks[0].arguments) {
- region.blocks[0].arguments = [];
- }
- for (const iv of inductionVars) {
- region.blocks[0].arguments.push({ value: iv });
- }
- }
- result.regions.push(region);
- } else {
- return false;
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseInParallelOp(parser, result) {
- // scf.forall.in_parallel { region }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- } else {
- return false;
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseSwitchCases(parser, op, casesAttrName) {
- const caseValues = [];
- while (parser.parseOptionalKeyword('case')) {
- if (parser.getToken().isNot(_.Token.integer)) {
- break;
- }
- const value = parser.parseInteger();
- caseValues.push(value);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = op.addRegion();
- parser.parseRegion(region);
- } else {
- break;
- }
- }
- if (casesAttrName) {
- op.addAttribute(casesAttrName, caseValues);
- }
- }
- parseExecuteRegionOp(parser, result) {
- result.addTypes(parser.parseOptionalArrowTypeList());
- if (parser.parseOptionalKeyword('no_inline')) {
- result.addAttribute('no_inline', true);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- };
- _.ShapeDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'shape');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (typeName === 'value' && parser.getToken().is('_')) {
- parser.consumeToken('_');
- const subType = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- type += `_${subType}`;
- }
- const simpleTypes = ['shape', 'witness', 'size', 'value_shape'];
- if (simpleTypes.includes(type.substring(7))) { // Remove "!shape." prefix
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op === 'shape.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'shape.assuming') {
- return this.parseAssumingOp(parser, result);
- }
- if (result.op === 'shape.const_shape') {
- return this.parseConstShapeOp(parser, result);
- }
- if (result.op === 'shape.reduce') {
- return this.parseReduceOp(parser, result);
- }
- if (result.op === 'shape.function_library') {
- return this.parseFunctionLibraryOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseAssumingOp(parser, result) {
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- return false;
- }
- const unresolvedWitness = parser.parseOperand();
- const witnessType = new _.Type('!shape.witness');
- parser.resolveOperand(unresolvedWitness, witnessType, result.operands);
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseConstShapeOp(parser, result) {
- parser.parseOptionalAttrDict(result.attributes);
- const extents = parser.parseAttribute();
- result.addAttribute('shape', extents);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- return true;
- }
- parseReduceOp(parser, result) {
- if (parser.getToken().isNot(_.Token.l_paren)) {
- return false;
- }
- parser.parseOptionalLParen();
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- let shapeType = new _.Type('!shape.shape');
- if (parser.parseOptionalColon()) {
- shapeType = parser.parseType();
- }
- const resultTypes = [];
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- resultTypes.push(...types);
- }
- // First operand is the shape, rest are init values with result types
- if (unresolvedOperands.length > 0) {
- parser.resolveOperand(unresolvedOperands[0], shapeType, result.operands);
- for (let i = 1; i < unresolvedOperands.length; i++) {
- const initType = resultTypes[i - 1] || null;
- parser.resolveOperand(unresolvedOperands[i], initType, result.operands);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseFunctionLibraryOp(parser, result) {
- parser.parseSymbolName('sym_name', result.attributes);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- if (parser.parseOptionalKeyword('mapping')) {
- const mapping = parser.parseAttribute();
- result.addAttribute('mapping', mapping);
- }
- return true;
- }
- };
- _.SparseTensorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'sparse_tensor');
- this.registerCustomDirective('LevelRange', this.parseLevelRange.bind(this));
- }
- parseLevelRange(parser, op, startAttr, endAttr) {
- const loLvl = parser.parseInteger();
- const hiLvl = parser.parseOptionalKeyword('to') ? parser.parseInteger() : loLvl + 1;
- if (startAttr && endAttr) {
- op.addAttribute(startAttr, loLvl);
- op.addAttribute(endAttr, hiLvl);
- }
- }
- parseOperation(parser, result) {
- if (result.op === 'sparse_tensor.iterate') {
- return this.parseIterateOp(parser, result);
- }
- if (result.op === 'sparse_tensor.coiterate') {
- return this.parseCoIterateOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseIterateOp(parser, result) {
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- return false;
- }
- const regionArgs = [];
- const iteratorArg = parser.parseOperand(); // iterator name (block arg)
- regionArgs.push({ name: iteratorArg.name, type: null }); // type determined by tensor
- if (!parser.parseOptionalKeyword('in')) {
- return false;
- }
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- return false;
- }
- const unresolvedTensor = parser.parseOperand();
- const iterArgNames = [];
- const initValues = [];
- if (parser.parseOptionalKeyword('at')) {
- parser.parseOptionalLParen();
- while (parser.getToken().is(_.Token.percent_identifier) || parser.getToken().is(_.Token.bare_identifier)) {
- parser.consumeToken();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- }
- if (parser.parseOptionalKeyword('iter_args')) {
- parser.parseOptionalLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const iterArg = parser.parseOperand();
- iterArgNames.push(iterArg.name);
- if (parser.parseOptionalEqual()) {
- initValues.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- }
- let tensorType = null;
- if (parser.parseOptionalColon()) {
- tensorType = parser.parseType();
- }
- const resultTypes = [];
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- resultTypes.push(...types);
- }
- // Add iter_args to region args with their result types
- for (let i = 0; i < iterArgNames.length; i++) {
- const argType = resultTypes[i] || null;
- regionArgs.push({ name: iterArgNames[i], type: argType });
- }
- parser.resolveOperand(unresolvedTensor, tensorType, result.operands);
- // iter_args block args don't go to operands, but init values do
- for (let i = 0; i < initValues.length; i++) {
- const initType = resultTypes[i] || null;
- parser.resolveOperand(initValues[i], initType, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, regionArgs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseCoIterateOp(parser, result) {
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- const unresolvedTensors = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedTensors.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- if (parser.parseOptionalKeyword('at')) {
- parser.parseOptionalLParen();
- while (parser.getToken().is(_.Token.percent_identifier) || parser.getToken().is(_.Token.bare_identifier)) {
- parser.consumeToken();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- }
- const iterArgNames = [];
- const initValues = [];
- if (parser.parseOptionalKeyword('iter_args')) {
- parser.parseOptionalLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const iterArg = parser.parseOperand(); // block arg name
- iterArgNames.push(iterArg.name);
- if (parser.parseOptionalEqual()) {
- initValues.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- }
- const tensorTypes = [];
- if (parser.parseOptionalColon()) {
- parser.parseOptionalLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- tensorTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseOptionalRParen();
- }
- const resultTypes = [];
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- result.addTypes(types);
- resultTypes.push(...types);
- }
- const regionArgs = [];
- for (let i = 0; i < iterArgNames.length; i++) {
- const argType = resultTypes[i] || null;
- regionArgs.push({ name: iterArgNames[i], type: argType });
- }
- for (let i = 0; i < unresolvedTensors.length; i++) {
- const tensorType = tensorTypes[i] || null;
- parser.resolveOperand(unresolvedTensors[i], tensorType, result.operands);
- }
- for (let i = 0; i < initValues.length; i++) {
- const initType = resultTypes[i] || null;
- parser.resolveOperand(initValues[i], initType, result.operands);
- }
- while (parser.parseOptionalKeyword('case')) {
- const caseArgs = [...regionArgs]; // Start with iter_args
- while (parser.getToken().is(_.Token.percent_identifier) || parser.getToken().is(_.Token.bare_identifier)) {
- const caseArg = parser.getToken().getSpelling().str();
- parser.consumeToken();
- if (caseArg.startsWith('%')) {
- caseArgs.push({ name: caseArg, type: null });
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, caseArgs);
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- };
- _.FuncDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'func');
- }
- parseOperation(parser, result) {
- if (result.op === 'func.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.GpuDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'gpu');
- this.registerCustomDirective('AllReduceOperation', this.parseAllReduceOperation.bind(this));
- this.registerCustomDirective('LaunchFuncOperands', this.parseLaunchFuncOperands.bind(this));
- this.registerCustomDirective('AsyncDependencies', this.parseAsyncDependencies.bind(this));
- this.registerCustomDirective('LaunchDimType', this.parseLaunchDimType.bind(this));
- this.registerCustomDirective('OffloadingHandler', this.parseOffloadingHandler.bind(this));
- }
- parseAllReduceOperation(parser, op, attrName = 'op') {
- const validOps = ['add', 'mul', 'minui', 'minsi', 'minnumf', 'maxui', 'maxsi', 'maxnumf', 'and', 'or', 'xor', 'minimumf', 'maximumf'];
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const opName = parser.getTokenSpelling().str();
- if (validOps.includes(opName)) {
- parser.consumeToken(_.Token.bare_identifier);
- op.addAttribute(attrName, opName);
- }
- }
- }
- parseLaunchDimType(parser, op, typeArg1, typeArg2, clusterTypeArg1, clusterTypeArg2, clusterTypeArg3) {
- // typeArg1 = type($gridSizeX), typeArg2 = ref($clusterSizeX)
- // clusterTypeArg1/2/3 = type($clusterSizeX/Y/Z)
- let dimType = new _.IndexType();
- if (parser.parseOptionalColon()) {
- dimType = parser.parseType();
- }
- // Push type to gridSizeX types array
- if (Array.isArray(typeArg1)) {
- typeArg1.push(dimType);
- }
- // If clusters are present (ref($clusterSizeX) has values), push to cluster type arrays
- const hasCluster = Array.isArray(typeArg2) && typeArg2.length > 0;
- if (hasCluster) {
- if (Array.isArray(clusterTypeArg1)) {
- clusterTypeArg1.push(dimType);
- }
- if (Array.isArray(clusterTypeArg2)) {
- clusterTypeArg2.push(dimType);
- }
- if (Array.isArray(clusterTypeArg3)) {
- clusterTypeArg3.push(dimType);
- }
- }
- }
- parseOperation(parser, result) {
- if (result.op === 'gpu.func') {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- parser.parseSymbolName('sym_name', result.attributes);
- const sig = parser.parseFunctionSignatureWithArguments(false);
- const argTypes = sig.arguments.map((a) => a.type);
- const type = new _.FunctionType(argTypes, sig.resultTypes);
- result.addAttribute('function_type', new _.TypeAttrOf(type));
- const allArgs = [...sig.arguments];
- if (parser.parseOptionalKeyword('workgroup')) {
- const workgroupResult = parser.parseFunctionArgumentList(false);
- allArgs.push(...workgroupResult.arguments);
- }
- if (parser.parseOptionalKeyword('private')) {
- const privateResult = parser.parseFunctionArgumentList(false);
- allArgs.push(...privateResult.arguments);
- }
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'kernel') {
- parser.consumeToken();
- result.addAttribute('gpu.kernel', true);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- // gpu.func is IsolatedFromAbove
- parser.parseRegion(region, allArgs, /* isIsolatedNameScope */ true);
- }
- return true;
- }
- if (result.op === 'gpu.launch') {
- const indexType = new _.IndexType();
- if (parser.parseOptionalKeyword('async')) {
- if (parser.getNumResults() === 0) {
- throw new mlir.Error(`Operation '${result.op}' needs to be named when marked 'async' ${parser.location()}`);
- }
- result.addTypes([new _.Type('!gpu.async.token')]);
- }
- const asyncDeps = parser.parseOperandList('optionalSquare');
- const asyncTypes = asyncDeps.map(() => null);
- parser.resolveOperands(asyncDeps, asyncTypes, result.operands);
- if (parser.parseOptionalKeyword('clusters')) {
- this.parseSizeAssignment(parser, result, indexType);
- parser.parseKeyword('in');
- this.parseSizeAssignment(parser, result, indexType);
- }
- parser.parseKeyword('blocks');
- this.parseSizeAssignment(parser, result, indexType);
- parser.parseKeyword('in');
- this.parseSizeAssignment(parser, result, indexType);
- parser.parseKeyword('threads');
- this.parseSizeAssignment(parser, result, indexType);
- parser.parseKeyword('in');
- this.parseSizeAssignment(parser, result, indexType);
- if (parser.parseOptionalKeyword('dynamic_shared_memory_size')) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, indexType, result.operands);
- }
- if (parser.parseOptionalKeyword('module')) {
- parser.parseLParen();
- const moduleSymbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('module', moduleSymbol);
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('function')) {
- parser.parseLParen();
- const funcSymbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('function', funcSymbol);
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('workgroup')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOperand();
- parser.parseColon();
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('private')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOperand();
- parser.parseColon();
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- // gpu.launch is IsolatedFromAbove
- parser.parseRegion(region, undefined, /* isIsolatedNameScope */ true);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (result.op === 'gpu.warp_execute_on_lane_0') {
- return this.parseWarpExecuteOnLane0Op(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseWarpExecuteOnLane0Op(parser, result) {
- parser.parseLParen();
- const unresolvedLaneId = parser.parseOperand();
- const indexType = new _.IndexType();
- parser.resolveOperand(unresolvedLaneId, indexType, result.operands);
- parser.parseRParen();
- parser.parseLSquare();
- const warpSize = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('warp_size', parseInt(warpSize, 10));
- parser.parseRSquare();
- if (parser.parseOptionalKeyword('args')) {
- parser.parseLParen();
- const unresolvedArgs = parser.parseOperandList('none');
- if (parser.parseOptionalColon()) {
- const types = parser.parseTypeListNoParens();
- parser.resolveOperands(unresolvedArgs, types, result.operands);
- } else {
- for (const arg of unresolvedArgs) {
- parser.resolveOperand(arg, null, result.operands);
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- const types = parser.parseFunctionResultTypes();
- if (result.types.length > 0) {
- result.addTypes(types);
- } else {
- for (const type of types) {
- result.addTypes([type]);
- }
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseSizeAssignment(parser, op, indexType) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand(); // Skip the LHS block arg
- if (parser.parseOptionalEqual()) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, indexType, op.operands);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else {
- break;
- }
- }
- parser.parseRParen();
- }
- parseLaunchFuncOperands(parser, op /*, args */) {
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'args') {
- parser.consumeToken();
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, op.operands);
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseComma();
- }
- }
- parser.parseRParen();
- }
- }
- parseAsyncDependencies(parser, op, asyncTokenTypes, asyncDependencies) {
- // custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
- // If 'async' keyword is present, operation has async result token
- const hasAsync = parser.parseOptionalKeyword('async');
- if (hasAsync && Array.isArray(asyncTokenTypes)) {
- asyncTokenTypes.push(new _.Type('!gpu.async.token'));
- }
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (Array.isArray(asyncDependencies)) {
- asyncDependencies.push(parser.parseOperand());
- } else {
- parser.parseOperand();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- }
- }
- parseOffloadingHandler(parser /*, op, args */) {
- if (parser.parseOptionalLess()) {
- parser.parseAttribute();
- parser.parseGreater();
- }
- }
- };
- _.ArmSMEDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'arm_sme');
- this.registerCustomAttribute('ArmSME_TypeSizeAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- this.registerCustomAttribute('ArmSME_TileSliceLayoutAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- this.registerCustomAttribute('ArmSME_CombiningKindAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- }
- inferResultTypes(op, vars) {
- if (op.op === 'arm_sme.outerproduct' && op.operands.length >= 2) {
- const lhsType = op.operands[0].type;
- if (lhsType instanceof _.VectorType && lhsType.shape.length === 1 && lhsType.scalableDims?.[0]) {
- const size = lhsType.shape[0];
- const tileType = new _.VectorType([size, size], lhsType.elementType, [true, true]);
- op.addTypes([tileType]);
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.ArmNeonDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'arm_neon');
- }
- };
- _.ArmSVEDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'arm_sve');
- }
- };
- _.AMDGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'amdgpu');
- this.registerCustomDirective('MNKDimensionList', this.parseMNKDimensionList.bind(this));
- }
- parseMNKDimensionList(parser, op, mAttr, nAttr, kAttr) {
- const dimInfo = parser.parseDimensionListRanked(false, false);
- const dims = dimInfo.dimensions;
- if (dims.length >= 3 && mAttr && nAttr && kAttr) {
- op.addAttribute(mAttr, dims[0]);
- op.addAttribute(nAttr, dims[1]);
- op.addAttribute(kAttr, dims[2]);
- }
- }
- };
- _.NVGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'nvgpu');
- this.registerCustomType('NVGPU_TensorMapDescriptor', this.parseTensorMapDescriptor.bind(this));
- this.registerCustomType('NVGPU_WarpgroupAccumulator', this.parseWarpgroupAccumulator.bind(this));
- this.registerCustomType('NVGPU_WarpgroupMatrixDescriptor', this.parseWarpgroupMatrixDescriptor.bind(this));
- this.registerCustomType('NVGPU_MBarrierGroup', this.parseMBarrierGroup.bind(this));
- }
- parseTensorMapDescriptor(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!nvgpu.tensormap.descriptor${content}`);
- }
- return null;
- }
- parseWarpgroupAccumulator(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!nvgpu.warpgroup.accumulator${content}`);
- }
- return null;
- }
- parseWarpgroupMatrixDescriptor(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!nvgpu.warpgroup.descriptor${content}`);
- }
- return null;
- }
- parseMBarrierGroup(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!nvgpu.mbarrier.barrier${content}`);
- }
- return null;
- }
- };
- _.NVVMDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'nvvm');
- }
- parseOperation(parser, result) {
- // Reference: NVVMDialect.cpp - parseMmaOperand
- // Helper to parse operand list in the format: name[operands]
- const parseMmaOperand = (name) => {
- parser.parseKeyword(name);
- return parser.parseOperandList('optionalSquare');
- };
- // Reference: NVVMDialect.cpp - MmaOp::parse
- if (result.op === 'nvvm.mma.sync') {
- const fragsA = parseMmaOperand('A');
- const fragsB = parseMmaOperand('B');
- const fragsC = parseMmaOperand('C');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType instanceof _.FunctionType) {
- parser.resolveOperands(fragsA, funcType.inputs, result.operands);
- parser.resolveOperands(fragsB, funcType.inputs, result.operands);
- parser.resolveOperands(fragsC, funcType.inputs, result.operands);
- result.addTypes(funcType.results);
- }
- }
- return true;
- }
- // Reference: NVVMDialect.cpp - MmaSpOp::parse
- if (result.op === 'nvvm.mma.sp.sync') {
- const fragsA = parseMmaOperand('A');
- const fragsB = parseMmaOperand('B');
- const fragsC = parseMmaOperand('C');
- const fragsSparseMetadata = parseMmaOperand('sparseMetadata');
- const fragsSelector = parseMmaOperand('selector');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType instanceof _.FunctionType) {
- parser.resolveOperands(fragsA, funcType.inputs, result.operands);
- parser.resolveOperands(fragsB, funcType.inputs, result.operands);
- parser.resolveOperands(fragsC, funcType.inputs, result.operands);
- parser.resolveOperands(fragsSparseMetadata, funcType.inputs, result.operands);
- parser.resolveOperands(fragsSelector, funcType.inputs, result.operands);
- result.addTypes(funcType.results);
- }
- }
- return true;
- }
- // Reference: NVVMDialect.cpp - MmaBlockScaleOp::parse
- if (result.op === 'nvvm.mma.block_scale') {
- const fragsA = parseMmaOperand('A');
- const fragsB = parseMmaOperand('B');
- const fragsC = parseMmaOperand('C');
- const scaleAOperands = parseMmaOperand('scaleA');
- const scaleBOperands = parseMmaOperand('scaleB');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType instanceof _.FunctionType) {
- parser.resolveOperands(fragsA, funcType.inputs, result.operands);
- parser.resolveOperands(fragsB, funcType.inputs, result.operands);
- parser.resolveOperands(fragsC, funcType.inputs, result.operands);
- parser.resolveOperands(scaleAOperands, funcType.inputs, result.operands);
- parser.resolveOperands(scaleBOperands, funcType.inputs, result.operands);
- result.addTypes(funcType.results);
- }
- }
- return true;
- }
- // Reference: NVVMDialect.cpp - MmaSpBlockScaleOp::parse
- if (result.op === 'nvvm.mma.sp.block_scale') {
- const fragsA = parseMmaOperand('A');
- const fragsB = parseMmaOperand('B');
- const fragsC = parseMmaOperand('C');
- const metadataOperands = parseMmaOperand('sparseMetadata');
- const selectorOperands = parseMmaOperand('selector');
- const scaleAOperands = parseMmaOperand('scaleA');
- const scaleBOperands = parseMmaOperand('scaleB');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType instanceof _.FunctionType) {
- parser.resolveOperands(fragsA, funcType.inputs, result.operands);
- parser.resolveOperands(fragsB, funcType.inputs, result.operands);
- parser.resolveOperands(fragsC, funcType.inputs, result.operands);
- parser.resolveOperands(metadataOperands, funcType.inputs, result.operands);
- parser.resolveOperands(selectorOperands, funcType.inputs, result.operands);
- parser.resolveOperands(scaleAOperands, funcType.inputs, result.operands);
- parser.resolveOperands(scaleBOperands, funcType.inputs, result.operands);
- result.addTypes(funcType.results);
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.NVWSDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'nvws');
- this.registerCustomType('NVWS_ArefType', this.parseArefTypeShorthand.bind(this));
- }
- parseArefTypeShorthand(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!nvws.aref${content}`);
- }
- return parser.parseType();
- }
- parseOperation(parser, result) {
- if (result.op === 'nvws.warp_group') {
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- const numWarps = [];
- let partitionIndex = 0;
- while (parser.parseOptionalKeyword(`partition${partitionIndex}`)) {
- parser.parseKeyword('num_warps');
- parser.parseLParen();
- const n = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- numWarps.push(n);
- parser.parseRParen();
- const region = result.addRegion();
- parser.parseRegion(region);
- partitionIndex++;
- }
- result.addAttribute('numWarps', numWarps);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.OpenMPDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'omp');
- this.registerCustomDirective('MapClause', this.parseMapClause.bind(this));
- this.registerCustomDirective('CaptureType', this.parseCaptureType.bind(this));
- this.registerCustomDirective('MembersIndex', this.parseMembersIndex.bind(this));
- this.registerCustomDirective('PrivateReductionRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('PrivateRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('InReductionPrivateRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('InReductionPrivateReductionRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('TaskReductionRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('UseDeviceAddrUseDevicePtrRegion', this.parsePrivateReductionRegion.bind(this));
- this.registerCustomDirective('TargetOpRegion', this.parseTargetOpRegion.bind(this));
- this.registerCustomDirective('ClauseAttr', this.parseClauseAttr.bind(this));
- this.registerCustomDirective('DependVarList', this.parseDependVarList.bind(this));
- this.registerCustomDirective('LoopTransformClis', this.parseLoopTransformClis.bind(this));
- this.registerCustomDirective('SynchronizationHint', this.parseSynchronizationHint.bind(this));
- this.registerCustomDirective('AlignedClause', this.parseAlignedClause.bind(this));
- this.registerCustomDirective('ScheduleClause', this.parseScheduleClause.bind(this));
- this.registerCustomDirective('AllocateAndAllocator', this.parseAllocateAndAllocator.bind(this));
- this.registerCustomDirective('LinearClause', this.parseLinearClause.bind(this));
- this.registerCustomDirective('UniformClause', this.parseUniformClause.bind(this));
- this.registerCustomDirective('OrderClause', this.parseOrderClause.bind(this));
- this.registerCustomDirective('Copyprivate', this.parseCopyprivate.bind(this));
- this.registerCustomDirective('GrainsizeClause', this.parseGranularityClause.bind(this));
- this.registerCustomDirective('NumTasksClause', this.parseGranularityClause.bind(this));
- this.registerCustomDirective('AffinityClause', this.parseAffinityClause.bind(this));
- this.registerCustomAttribute('DataSharingClauseTypeAttr', this.parseDataSharingClauseTypeAttr.bind(this));
- this.registerCustomAttribute('ClauseCancelConstructTypeAttr', this.parseParenthesizedEnumAttr.bind(this));
- this.registerCustomAttribute('ClauseDependAttr', this.parseParenthesizedEnumAttr.bind(this));
- this.registerCustomAttribute('ClauseOrderingIncludeTypeAttr', this.parseParenthesizedEnumAttr.bind(this));
- this.registerCustomAttribute('ClauseTypeAttr', this.parseParenthesizedEnumAttr.bind(this));
- this.registerCustomAttribute('ClauseDistScheduleTypeAttr', this.parseParenthesizedEnumAttr.bind(this));
- this.registerCustomAttribute('OrderModifierAttr', this.parseParenthesizedEnumAttr.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'omp.loop_nest') {
- return this.parseLoopNestOp(parser, result);
- }
- if (result.op === 'omp.canonical_loop') {
- return this.parseCanonicalLoopOp(parser, result);
- }
- if (result.op === 'omp.unroll_heuristic') {
- return this.parseUnrollHeuristicOp(parser, result);
- }
- if (result.op === 'omp.map.bounds') {
- return this.parseMapBoundsOp(parser, result);
- }
- if (result.op === 'omp.target_allocmem') {
- const unresolvedDevice = parser.parseOperand();
- parser.parseColon();
- const deviceType = parser.parseType();
- parser.resolveOperand(unresolvedDevice, deviceType, result.operands);
- parser.parseComma();
- const inType = parser.parseType();
- result.addAttribute('in_type', { value: inType, type: 'type' });
- const unresolvedTypeparams = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedTypeparams.push(parser.parseOperand());
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseColon();
- const types = parser.parseTypeList();
- parser.resolveOperands(unresolvedTypeparams, types, result.operands);
- parser.parseRParen();
- }
- const unresolvedShape = [];
- while (parser.parseOptionalComma()) {
- unresolvedShape.push(parser.parseOperand());
- }
- const indexType = new _.IndexType();
- for (const s of unresolvedShape) {
- parser.resolveOperand(s, indexType, result.operands);
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addAttribute('operandSegmentSizes', [1, unresolvedTypeparams.length, unresolvedShape.length]);
- result.addTypes([new _.IntegerType('i64')]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseCanonicalLoopOp(parser, result) {
- if (parser.parseOptionalLParen()) {
- const cliOperand = parser.parseOperand();
- // CLI operand is a loop handle, resolve with null type
- parser.resolveOperand(cliOperand, null, result.operands);
- parser.parseRParen();
- }
- const inductionVar = parser.parseOperand();
- parser.parseColon();
- const ivType = parser.parseType();
- parser.parseKeyword('in');
- parser.parseKeyword('range');
- parser.parseLParen();
- const rangeOperand = parser.parseOperand();
- parser.resolveOperand(rangeOperand, null, result.operands);
- parser.parseRParen();
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- // Pass induction variable as region argument
- const regionArgs = [{ name: inductionVar.name, type: ivType }];
- parser.parseRegion(region, regionArgs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseUnrollHeuristicOp(parser, result) {
- parser.parseLParen();
- const applyee = parser.parseOperand();
- // Applyee is a loop handle, resolve with null type
- parser.resolveOperand(applyee, null, result.operands);
- parser.parseRParen();
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- parser.parseRParen();
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseLoopNestOp(parser, result) {
- // Parse CLI operands (loop handles)
- const unresolvedCli = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedCli.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- // Parse types for CLI operands
- const cliTypes = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.equal) && parser.getToken().isNot(_.Token.l_brace)) {
- cliTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.resolveOperands(unresolvedCli, cliTypes, result.operands);
- if (parser.parseOptionalEqual()) {
- if (parser.parseOptionalLParen()) {
- const unresolvedLb = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedLb.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- for (const lb of unresolvedLb) {
- parser.resolveOperand(lb, null, result.operands);
- }
- }
- if (parser.parseOptionalKeyword('to')) {
- if (parser.parseOptionalLParen()) {
- const unresolvedUb = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedUb.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- for (const ub of unresolvedUb) {
- parser.resolveOperand(ub, null, result.operands);
- }
- }
- }
- parser.parseOptionalKeyword('inclusive');
- if (parser.parseOptionalKeyword('step')) {
- if (parser.parseOptionalLParen()) {
- const unresolvedStep = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedStep.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- for (const step of unresolvedStep) {
- parser.resolveOperand(step, null, result.operands);
- }
- }
- }
- }
- if (parser.parseOptionalKeyword('collapse')) {
- parser.parseLParen();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('collapse_num_loops', parseInt(value, 10));
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('tiles')) {
- parser.parseLParen();
- const tiles = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- tiles.push(parseInt(parser.getToken().getSpelling().str(), 10));
- parser.consumeToken(_.Token.integer);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- result.addAttribute('tile_sizes', tiles);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseMapBoundsOp(parser, result) {
- // oilist(lower_bound(...) | upper_bound(...) | extent(...) | stride(...) | start_idx(...)) attr-dict
- result.compatibility = true;
- const operandSegmentSizes = [0, 0, 0, 0, 0]; // lower_bound, upper_bound, extent, stride, start_idx
- const keywords = ['lower_bound', 'upper_bound', 'extent', 'stride', 'start_idx'];
- while (true) {
- let matched = false;
- for (let i = 0; i < keywords.length; i++) {
- if (parser.parseOptionalKeyword(keywords[i])) {
- parser.parseLParen();
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- parser.parseRParen();
- operandSegmentSizes[i] = 1;
- matched = true;
- break;
- }
- }
- if (!matched) {
- break;
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addAttribute('operandSegmentSizes', operandSegmentSizes);
- result.addTypes([new _.Type('!omp.map_bounds_ty')]);
- return true;
- }
- parseParenthesizedEnumAttr(parser) {
- if (parser.parseOptionalLParen()) {
- const value = parser.parseOptionalKeyword();
- parser.parseRParen();
- return new _.TypedAttr(value, null);
- }
- return null;
- }
- parseOrderClause(parser, result) {
- const orderModifiers = ['reproducible', 'unconstrained'];
- const orderKinds = ['concurrent'];
- let orderMod = null;
- let orderKind = null;
- const keyword = parser.parseOptionalKeyword();
- if (orderModifiers.includes(keyword)) {
- orderMod = keyword;
- parser.parseColon();
- orderKind = parser.parseOptionalKeyword();
- } else if (orderKinds.includes(keyword)) {
- orderKind = keyword;
- }
- if (orderKind) {
- result.addAttribute('order_kind', orderKind);
- }
- if (orderMod) {
- result.addAttribute('order_mod', orderMod);
- }
- }
- parseLinearClause(parser, result) {
- const unresolvedLinearVars = [];
- const linearVarTypes = [];
- const unresolvedStepVars = [];
- do {
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- break;
- }
- unresolvedLinearVars.push(parser.parseOperand());
- parser.parseEqual();
- unresolvedStepVars.push(parser.parseOperand());
- parser.parseColon();
- const type = parser.parseType();
- linearVarTypes.push(type);
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedLinearVars, linearVarTypes, result.operands);
- // Step vars typically have same type as linear vars
- parser.resolveOperands(unresolvedStepVars, linearVarTypes, result.operands);
- }
- parseUniformClause(parser, result, uniformVars, uniformTypes) {
- parser.parseCommaSeparatedList('none', () => {
- uniformVars.push(parser.parseOperand());
- parser.parseColon();
- uniformTypes.push(parser.parseType());
- });
- }
- parseCopyprivate(parser, op, varsAttr, typesAttr, symsAttr) {
- const unresolvedVars = [];
- const varTypes = [];
- const copyprivateSyms = [];
- do {
- unresolvedVars.push(parser.parseOperand());
- parser.parseArrow();
- const sym = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- parser.parseColon();
- const type = parser.parseType();
- varTypes.push(type);
- copyprivateSyms.push(sym);
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedVars, varTypes, op.operands);
- if (symsAttr) {
- op.addAttribute(symsAttr, copyprivateSyms);
- }
- }
- parseGranularityClause(parser, op, modAttr) {
- let modifier = null;
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.percent_identifier)) {
- modifier = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseComma();
- }
- const unresolvedOperand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(unresolvedOperand, type, op.operands);
- if (modAttr && modifier) {
- op.addAttribute(modAttr, modifier);
- }
- }
- parseAlignedClause(parser, result) {
- const unresolvedVars = [];
- const varTypes = [];
- const alignments = [];
- do {
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- break;
- }
- unresolvedVars.push(parser.parseOperand());
- parser.parseColon();
- const type = parser.parseType();
- varTypes.push(type);
- parser.parseArrow();
- const alignment = parser.parseAttribute();
- alignments.push(alignment);
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedVars, varTypes, result.operands);
- if (alignments.length > 0) {
- result.addAttribute('alignments', alignments);
- }
- }
- parseAffinityClause(parser, result, affinityVars, affinityTypes) {
- parser.parseCommaSeparatedList('none', () => {
- affinityVars.push(parser.parseOperand());
- parser.parseColon();
- affinityTypes.push(parser.parseType());
- });
- }
- parseScheduleClause(parser, result) {
- const scheduleKinds = ['static', 'dynamic', 'guided', 'auto', 'runtime', 'distribute'];
- let scheduleKind = null;
- for (const kind of scheduleKinds) {
- if (parser.parseOptionalKeyword(kind)) {
- scheduleKind = kind;
- break;
- }
- }
- if (scheduleKind) {
- result.addAttribute('schedule_kind', scheduleKind);
- }
- if (parser.parseOptionalEqual()) {
- let unresolvedChunk = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedChunk = parser.parseOperand();
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (unresolvedChunk) {
- parser.resolveOperand(unresolvedChunk, type, result.operands);
- }
- } else if (unresolvedChunk) {
- parser.resolveOperand(unresolvedChunk, null, result.operands);
- }
- }
- const modifiers = [];
- while (parser.parseOptionalComma()) {
- const mod = parser.parseOptionalKeyword();
- if (mod) {
- modifiers.push(mod);
- }
- }
- if (modifiers.length > 0) {
- result.addAttribute('schedule_modifiers', modifiers);
- }
- }
- parseAllocateAndAllocator(parser, result) {
- const unresolvedAllocators = [];
- const allocatorTypes = [];
- const unresolvedAllocates = [];
- const allocateTypes = [];
- do {
- if (parser.getToken().isNot(_.Token.percent_identifier)) {
- break;
- }
- unresolvedAllocators.push(parser.parseOperand());
- parser.parseColon();
- allocatorTypes.push(parser.parseType());
- parser.parseArrow();
- unresolvedAllocates.push(parser.parseOperand());
- parser.parseColon();
- allocateTypes.push(parser.parseType());
- } while (parser.parseOptionalComma());
- parser.resolveOperands(unresolvedAllocators, allocatorTypes, result.operands);
- parser.resolveOperands(unresolvedAllocates, allocateTypes, result.operands);
- }
- parseSynchronizationHint(parser, op, hintAttr = 'hint') {
- if (parser.parseOptionalKeyword('none')) {
- op.addAttribute(hintAttr, 0);
- return;
- }
- let hint = 0;
- const hints = [];
- while (parser.getToken().is(_.Token.bare_identifier)) {
- const keyword = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- hints.push(keyword);
- if (keyword === 'uncontended') {
- hint |= 1;
- } else if (keyword === 'contended') {
- hint |= 2;
- } else if (keyword === 'nonspeculative') {
- hint |= 4;
- } else if (keyword === 'speculative') {
- hint |= 8;
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- op.addAttribute(hintAttr, hint);
- }
- parseClauseAttr(parser, op, attrName) {
- // Parses a keyword (enum value) and converts to attribute
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const enumValue = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (attrName) {
- op.addAttribute(attrName, enumValue);
- }
- } else if (parser.getToken().is(_.Token.l_brace)) {
- parser.skip('{');
- }
- }
- parseTargetOpRegion(parser, result) {
- const unitAttrKeywords = ['nowait', 'bare'];
- for (const kw of unitAttrKeywords) {
- if (parser.parseOptionalKeyword(kw)) {
- result.addAttribute(kw, true);
- }
- }
- if (parser.parseOptionalKeyword('depend')) {
- parser.skip('(');
- }
- const singleValueKeywords = ['device', 'if', 'thread_limit'];
- for (const kw of singleValueKeywords) {
- if (parser.parseOptionalKeyword(kw)) {
- parser.parseLParen();
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedOperand = parser.parseOperand();
- let opType = null;
- if (parser.parseOptionalColon()) {
- opType = parser.parseType();
- }
- parser.resolveOperand(unresolvedOperand, opType, result.operands);
- } else if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- parser.parseRParen();
- }
- }
- if (parser.parseOptionalKeyword('is_device_ptr')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.parseRParen();
- }
- const keywords = ['has_device_addr', 'host_eval', 'in_reduction', 'map_entries', 'private', 'reduction', 'task_reduction', 'use_device_addr', 'use_device_ptr'];
- let progress = true;
- while (progress) {
- progress = false;
- // Handle private_barrier unit attribute
- if (parser.parseOptionalKeyword('private_barrier')) {
- result.addAttribute('private_needs_barrier', true);
- progress = true;
- continue;
- }
- // Handle list clauses
- if (keywords.some((kw) => parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === kw)) {
- parser.consumeToken(_.Token.bare_identifier);
- progress = true;
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- parser.parseOptionalKeyword('byref');
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.consumeToken(_.Token.at_identifier);
- }
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- }
- if (parser.parseOptionalArrow()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- }
- }
- if (parser.parseOptionalLSquare()) {
- parser.parseKeyword('map_idx');
- parser.parseEqual();
- parser.consumeToken(_.Token.integer);
- parser.parseRSquare();
- }
- if (!parser.parseOptionalComma() || parser.getToken().is(_.Token.colon)) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseType();
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.parseRParen();
- }
- }
- }
- if (parser.getToken().isNot(_.Token.l_brace)) {
- return;
- }
- const region = {};
- parser.parseRegion(region);
- result.regions.push(region);
- }
- parseMapClause(parser, op, attrName = 'map_type') {
- const mapFlags = [];
- do {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const flag = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- mapFlags.push(flag);
- }
- } while (parser.parseOptionalComma());
- if (attrName && mapFlags.length > 0) {
- op.addAttribute(attrName, mapFlags.join(', '));
- }
- }
- parseCaptureType(parser, op, attrName) {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const captureType = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (attrName) {
- op.addAttribute(attrName, captureType);
- }
- }
- }
- parseMembersIndex(parser, op, attrName) {
- const memberIndices = [];
- do {
- if (parser.parseOptionalLSquare()) {
- const indices = [];
- do {
- if (parser.getToken().is(_.Token.integer)) {
- const idx = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- indices.push(idx);
- }
- } while (parser.parseOptionalComma());
- parser.parseRSquare();
- memberIndices.push(indices);
- }
- } while (parser.parseOptionalComma());
- if (attrName && memberIndices.length > 0) {
- op.addAttribute(attrName, memberIndices);
- }
- }
- parsePrivateReductionRegion(parser, result) {
- const singleValueClauses = ['if', 'num_threads', 'thread_limit', 'device', 'safelen', 'simdlen', 'priority', 'grainsize', 'num_tasks', 'final', 'filter'];
- const enumClauses = ['proc_bind', 'order', 'schedule', 'dist_schedule', 'memory_order', 'hint'];
- const listClauses = ['private', 'reduction', 'in_reduction', 'task_reduction', 'copyin', 'copyprivate', 'firstprivate', 'lastprivate', 'shared', 'linear', 'aligned', 'nontemporal', 'inclusive', 'exclusive', 'allocate', 'depend'];
- const unitClauses = ['nowait', 'untied', 'mergeable', 'nogroup', 'simd', 'threads', 'seq_cst', 'acq_rel', 'acquire', 'release', 'relaxed', 'private_barrier'];
- let progress = true;
- while (progress) {
- progress = false;
- // Handle single-value clauses: keyword(value : type)
- for (const kw of singleValueClauses) {
- if (parser.parseOptionalKeyword(kw)) {
- progress = true;
- parser.parseLParen();
- let unresolvedOp = null;
- let opType = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOp = parser.parseOperand();
- } else if (parser.getToken().is(_.Token.integer)) {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute(kw, value);
- }
- if (parser.parseOptionalColon()) {
- opType = parser.parseType();
- }
- if (unresolvedOp) {
- parser.resolveOperand(unresolvedOp, opType, result.operands);
- }
- parser.parseRParen();
- }
- }
- // Handle enum clauses: keyword(enum_value)
- for (const kw of enumClauses) {
- if (parser.parseOptionalKeyword(kw)) {
- progress = true;
- parser.parseLParen();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute(kw, value);
- // Handle modifier syntax like schedule(static, value)
- while (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedOp = parser.parseOperand();
- let opType = null;
- if (parser.parseOptionalColon()) {
- opType = parser.parseType();
- }
- parser.resolveOperand(unresolvedOp, opType, result.operands);
- } else if (parser.getToken().is(_.Token.bare_identifier)) {
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- }
- }
- parser.parseRParen();
- }
- }
- // Handle list clauses: keyword(syms %vals -> %new_vals : types) or keyword(@sym %val : type, ...)
- for (const kw of listClauses) {
- if (parser.parseOptionalKeyword(kw)) {
- progress = true;
- if (parser.parseOptionalLParen()) {
- if (parser.consumeIf(_.Token.kw_mod)) {
- parser.parseColon();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseComma();
- }
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- parser.parseOptionalKeyword('byref');
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.consumeToken(_.Token.at_identifier);
- }
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- if (parser.parseOptionalArrow()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- }
- }
- if (parser.parseOptionalLSquare()) {
- parser.parseKeyword('map_idx');
- parser.parseEqual();
- parser.consumeToken(_.Token.integer);
- parser.parseRSquare();
- }
- if (!parser.parseOptionalComma() || parser.getToken().is(_.Token.colon)) {
- break;
- }
- }
- const types = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- types.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- parser.parseRParen();
- }
- }
- }
- // Handle unit clauses (boolean flags)
- for (const kw of unitClauses) {
- if (parser.parseOptionalKeyword(kw)) {
- progress = true;
- // private_barrier maps to private_needs_barrier attribute
- const attrName = kw === 'private_barrier' ? 'private_needs_barrier' : kw;
- result.addAttribute(attrName, true);
- }
- }
- // Handle map_entries clause: map_entries(%vars : types)
- if (parser.parseOptionalKeyword('map_entries')) {
- progress = true;
- parser.parseLParen();
- const unresolvedMapVars = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- unresolvedMapVars.push(operand);
- }
- if (!parser.parseOptionalComma() || parser.getToken().is(_.Token.colon)) {
- break;
- }
- }
- const mapTypes = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- mapTypes.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- for (let i = 0; i < unresolvedMapVars.length; i++) {
- const type = i < mapTypes.length ? mapTypes[i] : null;
- parser.resolveOperand(unresolvedMapVars[i], type, result.operands);
- }
- parser.parseRParen();
- }
- // Handle num_teams clause: num_teams(lower : type to upper : type) or num_teams(to upper : type)
- if (parser.parseOptionalKeyword('num_teams')) {
- progress = true;
- parser.parseLParen();
- if (parser.parseOptionalKeyword('to')) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const upper = parser.parseOperand();
- let upperType = null;
- if (parser.parseOptionalColon()) {
- upperType = parser.parseType();
- }
- parser.resolveOperand(upper, upperType, result.operands);
- }
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const lower = parser.parseOperand();
- let lowerType = null;
- if (parser.parseOptionalColon()) {
- lowerType = parser.parseType();
- }
- parser.resolveOperand(lower, lowerType, result.operands);
- parser.parseKeyword('to');
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const upper = parser.parseOperand();
- let upperType = null;
- if (parser.parseOptionalColon()) {
- upperType = parser.parseType();
- }
- parser.resolveOperand(upper, upperType, result.operands);
- }
- }
- parser.parseRParen();
- }
- // Handle use_device_addr/use_device_ptr clauses: keyword(%var -> %arg : type, ...)
- for (const kw of ['use_device_addr', 'use_device_ptr', 'has_device_addr', 'host_eval']) {
- if (parser.parseOptionalKeyword(kw)) {
- progress = true;
- parser.parseLParen();
- const unresolvedOperands = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- }
- if (parser.parseOptionalArrow()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- parser.parseOperand();
- }
- }
- if (!parser.parseOptionalComma() || parser.getToken().is(_.Token.colon)) {
- break;
- }
- }
- const types = [];
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- types.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- for (let i = 0; i < unresolvedOperands.length; i++) {
- const type = i < types.length ? types[i] : null;
- parser.resolveOperand(unresolvedOperands[i], type, result.operands);
- }
- parser.parseRParen();
- }
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = {};
- parser.parseRegion(region);
- result.regions.push(region);
- }
- }
- parseDataSharingClauseTypeAttr(parser) {
- if (parser.parseOptionalLBrace()) {
- parser.parseKeyword('type');
- parser.parseEqual();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseRBrace();
- return { value };
- }
- return null;
- }
- parseDependVarList(parser, op, operandAttr, typesAttr, kindAttr) {
- const dependVars = [];
- const dependTypes = [];
- const dependKinds = [];
- do {
- const keyword = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- dependKinds.push(keyword);
- parser.parseArrow();
- const operand = parser.parseOperand();
- dependVars.push(operand);
- parser.parseColon();
- const type = parser.parseType();
- dependTypes.push(type);
- } while (parser.parseOptionalComma());
- if (operandAttr) {
- // depend_vars are SSA operands - resolve and add as operands
- for (let i = 0; i < dependVars.length; i++) {
- const type = i < dependTypes.length ? dependTypes[i] : null;
- parser.resolveOperand(dependVars[i], type, op.operands);
- }
- }
- if (kindAttr) {
- op.addAttribute(kindAttr, dependKinds);
- }
- }
- // Syntax 1: (generatees) <- (applyees) - generatees present (no leading <)
- // Syntax 2: <- (applyees) - generatees omitted (starts with <-)
- parseLoopTransformClis(parser, result) {
- const generatees = [];
- const applyees = [];
- if (!parser.parseOptionalLess()) {
- // Syntax 1: generatees present, parse (generatees) first
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- generatees.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseLess();
- }
- parser.consumeToken(_.Token.minus);
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- applyees.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- for (const g of generatees) {
- parser.resolveOperand(g, null, result.operands);
- }
- for (const a of applyees) {
- parser.resolveOperand(a, null, result.operands);
- }
- }
- };
- _.LLVM = {};
- _.LLVM.LLVMFunctionType = class extends _.Type {
- constructor(returnType, params, varArg = false) {
- super(null);
- this.returnType = returnType;
- this.params = params || [];
- this.varArg = varArg;
- }
- get inputs() {
- return this.params;
- }
- get results() {
- return this.returnType ? [this.returnType] : [];
- }
- toString() {
- const params = this.params.map((t) => t.toString());
- if (this.varArg) {
- params.push('...');
- }
- const returnType = this.returnType ? this.returnType.toString() : 'void';
- return `!llvm.func<${returnType} (${params.join(', ')})>`;
- }
- };
- _.LLVM.LLVMDialect = class extends _.Dialect {
- constructor(operations, name = 'llvm') {
- super(operations, name);
- this.registerCustomDirective('GEPIndices', this.parseGEPIndices.bind(this));
- this.registerCustomDirective('IndirectBrOpSucessors', this.parseIndirectBrOpSucessors.bind(this));
- this.registerCustomDirective('InsertExtractValueElementType', this.parseInsertExtractValueElementType.bind(this));
- this.registerCustomDirective('LLVMLinkage', this.parseLLVMLinkage.bind(this));
- this.registerCustomDirective('OpBundles', this.parseOpBundles.bind(this));
- this.registerCustomDirective('ShuffleType', this.parseShuffleType.bind(this));
- this.registerCustomDirective('SwitchOpCases', this.parseSwitchOpCases.bind(this));
- this.registerCustomAttribute('LLVM_IntegerOverflowFlagsProp', this.parseLLVMIntegerOverflowFlagsProp.bind(this));
- this.registerCustomAttribute('GEPNoWrapFlagsProp', this.parseGEPNoWrapFlagsProp.bind(this));
- this.registerCustomAttribute('LLVM_BlockAddressAttr', this.parseLLVMBlockAddressAttr.bind(this));
- this.registerCustomAttribute('LLVM_BlockTagAttr', this.parseLLVMBlockTagAttr.bind(this));
- this.registerCustomType('LLVM_AnyPointer', this.parseLLVMPointerType.bind(this));
- this.registerCustomType('LLVM_PointerInAddressSpace', this.parseLLVMPointerType.bind(this));
- this.registerCustomType('LLVM_Type', (parser) => parser.parseType());
- }
- parseLLVMIntegerOverflowFlagsProp(parser) {
- if (parser.parseOptionalKeyword('overflow')) {
- return this.parseEnumFlagsAngleBracketComma(parser, { values: ['wrap', 'nuw', 'nsw'] });
- }
- return null;
- }
- parseGEPNoWrapFlagsProp(parser, type) {
- if (type.values.includes(parser.getTokenSpelling().str())) {
- return this.parseEnumFlags(parser, type, '|');
- }
- return null;
- }
- parseLLVMBlockAddressAttr(parser) {
- if (parser.getToken().isNot(_.Token.less)) {
- return null;
- }
- const content = parser.skip('<');
- return { blockaddress: content };
- }
- parseLLVMBlockTagAttr(parser) {
- if (parser.getToken().isNot(_.Token.less)) {
- return null;
- }
- const content = parser.skip('<');
- return { blocktag: content };
- }
- parseLLVMPointerType(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- const inner = content.startsWith('<') && content.endsWith('>') ? content.slice(1, -1) : content;
- if (/^\d+$/.test(inner)) {
- return new _.Type(`!llvm.ptr<${inner}>`);
- }
- }
- return parser.parseType();
- }
- getInsertExtractValueElementType(containerType, position) {
- let llvmType = containerType.toString ? containerType.toString() : String(containerType);
- for (const idx of position) {
- const arrayMatch = llvmType.match(/^!llvm\.array<(\d+)\s*x\s*(.+)>$/);
- if (arrayMatch) {
- llvmType = arrayMatch[2].trim();
- continue;
- }
- const structMatch = llvmType.match(/^!llvm\.struct<(?:"[^"]*",\s*)?\((.+)\)>$/);
- if (structMatch) {
- const elementsStr = structMatch[1];
- const body = [];
- let depth = 0;
- let current = '';
- for (const char of elementsStr) {
- if (char === '<' || char === '(' || char === '[' || char === '{') {
- depth++;
- current += char;
- } else if (char === '>' || char === ')' || char === ']' || char === '}') {
- depth--;
- current += char;
- } else if (char === ',' && depth === 0) {
- body.push(current.trim());
- current = '';
- } else {
- current += char;
- }
- }
- if (current.trim()) {
- body.push(current.trim());
- }
- if (idx >= 0 && idx < body.length) {
- llvmType = body[idx];
- continue;
- }
- }
- break;
- }
- return llvmType;
- }
- parseInsertExtractValueElementType(parser, op, resultTypes, containerTypes, positionArg) {
- // Infer the value type from the container type and position.
- // resultTypes = op.types (array to populate)
- // containerTypes = array of container types from parsing
- // positionArg = position attribute value (ArrayAttr or array)
- if (containerTypes && containerTypes.length > 0) {
- const containerType = containerTypes[0];
- // Use the position argument if provided, otherwise fall back to op.attributes
- const positionAttr = positionArg === undefined ? op.attributes.get('position') : positionArg;
- const position = positionAttr && positionAttr.value ? positionAttr.value : positionAttr;
- if (position && Array.isArray(position) && position.length > 0) {
- const valueType = this.getInsertExtractValueElementType(containerType, position);
- if (valueType && resultTypes.length === 0) {
- resultTypes.push(new _.Type(valueType));
- }
- }
- }
- }
- parseLLVMLinkage(parser, op /*, args */) {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const linkage = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- op.addAttribute('linkage', linkage);
- }
- }
- parseOpBundles(parser, op /*, args */) {
- // Parse operation bundles: [] or ["tag"()] or ["tag"(%0, %1 : i32, i32), ...]
- // Returns: null if not present, true if success, throws on failure
- // args[0] = $op_bundle_operands - operands for bundles
- // args[1] = type($op_bundle_operands) - types
- // args[2] = $op_bundle_tags - tags attribute
- if (!parser.parseOptionalLSquare()) {
- return null; // Not present (equivalent to std::nullopt)
- }
- if (parser.parseOptionalRSquare()) {
- return true; // Success
- }
- const opBundles = [];
- do {
- const tag = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseLParen();
- const bundleOperands = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- bundleOperands.push(parser.parseAttribute());
- } while (parser.parseOptionalComma());
- parser.parseColon();
- do {
- parser.parseType();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- opBundles.push({ tag, operands: bundleOperands });
- } while (parser.parseOptionalComma());
- parser.parseRSquare();
- if (opBundles.length > 0) {
- op.addAttribute('op_bundle_tags', opBundles);
- }
- return true; // Success
- }
- parseShuffleType(parser, op, v1Types, resTypes, maskAttr) {
- // custom<ShuffleType>(ref(type($v1)), type($res), ref($mask))
- // Computes the result type: same element type as input, length = mask.size()
- if (resTypes.length > 0) {
- return; // Result type already set
- }
- if (!v1Types || v1Types.length === 0 || !maskAttr) {
- return;
- }
- const v1Type = v1Types[0];
- const typeStr = v1Type.toString ? v1Type.toString() : String(v1Type);
- const vecMatch = typeStr.match(/^vector<(?:\[)?(\d+)(?:\])?\s*x\s*(.+)>$/);
- if (!vecMatch) {
- return;
- }
- const elemType = vecMatch[2];
- const isScalable = typeStr.includes('[') && typeStr.includes(']');
- // Get mask length
- let maskLen = 0;
- const mask = maskAttr.value === undefined ? maskAttr : maskAttr.value;
- if (Array.isArray(mask)) {
- maskLen = mask.length;
- } else if (typeof mask === 'string') {
- const stripped = mask.replace(/[[\]]/g, '').trim();
- if (stripped) {
- maskLen = stripped.split(',').length;
- }
- }
- if (maskLen > 0) {
- const resultTypeStr = isScalable ?
- `vector<[${maskLen}]x${elemType}>` :
- `vector<${maskLen}x${elemType}>`;
- op.addTypes([new _.Type(resultTypeStr)]);
- }
- }
- // Parse switch operation cases
- // Where case: integer `:` bb-id (`(` ssa-use-and-type-list `)`)?
- parseSwitchOpCases(parser, op /*, args */) {
- // args[0] is ref(type($value)) - the flag type
- // args[1] is $case_values - attribute to populate
- // args[2] is $caseDestinations - successors array
- // args[3] is $caseOperands - operands for each case
- // args[4] is type($caseOperands) - types for case operands
- if (!parser.parseOptionalLSquare()) {
- return;
- }
- if (parser.parseOptionalRSquare()) {
- return;
- }
- const caseValues = [];
- const caseDestinations = [];
- const caseOperands = [];
- while (parser.getToken().isNot(_.Token.r_square) && parser.getToken().isNot(_.Token.eof)) {
- // Handle negative case values: -1, -2, etc.
- let sign = 1;
- if (parser.consumeIf(_.Token.minus)) {
- sign = -1;
- }
- if (parser.getToken().isNot(_.Token.integer) && parser.getToken().isNot('number')) {
- throw new mlir.Error(`Expected integer case value at ${parser.location()}`);
- }
- const value = sign * parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken();
- caseValues.push(value);
- parser.parseColon();
- const successor = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- caseDestinations.push(successor);
- if (parser.parseOptionalLParen()) {
- const operands = [];
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.eof)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- operands.push({ name: operand });
- if (!parser.parseOptionalComma()) {
- break;
- }
- } else {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- let idx = 0;
- while (parser.getToken().isNot(_.Token.r_paren) && idx < operands.length) {
- const type = parser.parseType();
- if (operands[idx]) {
- operands[idx].type = type;
- }
- idx++;
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- caseOperands.push(operands);
- } else {
- caseOperands.push([]);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- if (caseValues.length > 0) {
- op.addAttribute('case_values', caseValues);
- }
- if (caseDestinations.length > 0) {
- if (!op.successors) {
- op.successors = [];
- }
- for (const dest of caseDestinations) {
- op.successors.push({ name: dest });
- }
- }
- // Note: caseOperands handling would require more complex logic
- // to properly associate operands with their successors
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (result.op === 'llvm.func') {
- return this.parseLLVMFuncOp(parser, result);
- }
- if (result.op === 'llvm.mlir.global') {
- return this.parseLLVMGlobalOp(parser, result);
- }
- if (result.op === 'llvm.mlir.alias') {
- return this.parseLLVMAliasOp(parser, result);
- }
- if (result.op === 'llvm.alloca') {
- return this.parseLLVMAllocaOp(parser, result);
- }
- if (result.op === 'llvm.call') {
- return this.parseLLVMCallOp(parser, result);
- }
- if (result.op === 'llvm.call_intrinsic') {
- return this.parseLLVMCallIntrinsicOp(parser, result);
- }
- if (result.op === 'llvm.invoke') {
- return this.parseLLVMInvokeOp(parser, result);
- }
- if (result.op === 'llvm.landingpad') {
- return this.parseLLVMLandingpadOp(parser, result);
- }
- if (result.op === 'llvm.icmp' || result.op === 'llvm.fcmp') {
- return this.parseLLVMCmpOp(parser, result);
- }
- if (result.op.startsWith('llvm.intr.')) {
- if (opInfo.metadata.assemblyFormat) {
- return super.parseOperation(parser, result);
- }
- return this.parseLLVMIntrinsicOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseLLVMGlobalOp(parser, result) {
- const linkageKeywords = ['external', 'available_externally', 'linkonce', 'linkonce_odr', 'weak', 'weak_odr', 'appending', 'internal', 'private', 'extern_weak', 'common'];
- if (parser.getToken().is(_.Token.bare_identifier) && linkageKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('linkage', __spelling);
- }
- const visibilityKeywords = ['default', 'hidden', 'protected'];
- if (parser.getToken().is(_.Token.bare_identifier) && visibilityKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('visibility_', __spelling);
- }
- if (parser.parseOptionalKeyword('thread_local')) {
- result.addAttribute('thread_local_', true);
- }
- const unnamedAddrKeywords = ['unnamed_addr', 'local_unnamed_addr'];
- if (parser.getToken().is(_.Token.bare_identifier) && unnamedAddrKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('unnamed_addr', __spelling);
- }
- if (parser.parseOptionalKeyword('constant')) {
- result.addAttribute('constant', true);
- }
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- result.addAttribute('value', value);
- }
- parser.parseRParen();
- if (parser.parseOptionalKeyword('comdat')) {
- parser.parseLParen();
- const comdat = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- parser.parseRParen();
- result.addAttribute('comdat', comdat);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.types = [type];
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- parseLLVMAliasOp(parser, result) {
- const linkageKeywords = ['external', 'available_externally', 'linkonce', 'linkonce_odr', 'weak', 'weak_odr', 'internal', 'private'];
- if (parser.getToken().is(_.Token.bare_identifier) && linkageKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('linkage', __spelling);
- }
- const visibilityKeywords = ['default', 'hidden', 'protected'];
- if (parser.getToken().is(_.Token.bare_identifier) && visibilityKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('visibility_', __spelling);
- }
- if (parser.parseOptionalKeyword('thread_local')) {
- result.addAttribute('thread_local_', true);
- }
- const unnamedAddrKeywords = ['unnamed_addr', 'local_unnamed_addr'];
- if (parser.getToken().is(_.Token.bare_identifier) && unnamedAddrKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('unnamed_addr', __spelling);
- }
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('alias_type', type);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- // custom<GEPIndices>($dynamicIndices, $rawConstantIndices)
- parseGEPIndices(parser, op, operands, attrName) {
- // Note: the reference uses 'none' delimiter with TableGen handling brackets,
- // but mlir.js expects '[' already consumed and needs to handle ']' terminator
- const rawConstantIndices = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- const constIndex = parser.parseOptionalInteger();
- if (constIndex === null) {
- const operand = parser.parseOperand();
- operands.push(operand);
- rawConstantIndices.push(-2147483648);
- } else {
- rawConstantIndices.push(constIndex);
- }
- parser.parseOptionalComma();
- }
- if (rawConstantIndices.length > 0) {
- op.addAttribute(attrName || 'rawConstantIndices', rawConstantIndices);
- }
- }
- parseIndirectBrOpSucessors(parser, op /*, args */) {
- // All operands listed first, then colon, then all types
- parser.parseLSquare();
- const segmentSizes = [];
- if (parser.getToken().isNot(_.Token.r_square)) {
- do {
- const successor = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- if (!op.successors) {
- op.successors = [];
- }
- op.successors.push({ name: successor });
- const unresolvedOperands = [];
- const types = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.colon)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- types.push(type);
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- }
- for (let i = 0; i < unresolvedOperands.length; i++) {
- const type = i < types.length ? types[i] : null;
- parser.resolveOperand(unresolvedOperands[i], type, op.operands);
- }
- segmentSizes.push(unresolvedOperands.length);
- } while (parser.parseOptionalComma());
- }
- parser.parseRSquare();
- if (segmentSizes.length > 0) {
- op.addAttribute('indbr_operand_segments', segmentSizes);
- }
- }
- parseLLVMAllocaOp(parser, result) {
- // llvm.alloca [inalloca] %arraySize x !elemType : (i64) -> !llvm.ptr
- if (parser.parseOptionalKeyword('inalloca')) {
- result.addAttribute('inalloca', true);
- }
- const arraySize = parser.parseOperand();
- parser.parseKeyword('x');
- const elemType = parser.parseType();
- result.addAttribute('elem_type', elemType);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- const fnType = parser.parseFunctionType();
- if (fnType instanceof _.FunctionType) {
- parser.resolveOperands([arraySize], fnType.inputs, result.operands);
- result.addTypes(fnType.results);
- }
- return true;
- }
- parseLLVMCallOp(parser, result) {
- // llvm.call [cconv] [tailcall] @callee|%ptr (args) [vararg(type)] : func_type
- const cconvKeywords = ['ccc', 'fastcc', 'coldcc', 'cc', 'webkit_jscc', 'anyregcc', 'preserve_mostcc', 'preserve_allcc', 'preserve_nonecc', 'cxx_fast_tlscc', 'tailcc', 'swiftcc', 'swifttailcc', 'cfguard_checkcc', 'ghccc', 'arm_apcscc', 'arm_aapcscc', 'arm_aapcs_vfpcc', 'aarch64_vector_pcs', 'aarch64_sve_vector_pcs', 'aarch64_sme_preservemost_from_x0', 'aarch64_sme_preservemost_from_x2', 'msp430_intrcc', 'avr_intrcc', 'avr_signalcc', 'ptx_kernelcc', 'ptx_devicecc', 'spir_funccc', 'spir_kernelcc', 'intel_ocl_bicc', 'x86_64_sysvcc', 'win64cc', 'x86_fastcallcc', 'x86_stdcallcc', 'x86_thiscallcc', 'x86_vectorcallcc', 'x86_intrcc', 'amdgpu_vs', 'amdgpu_gs', 'amdgpu_ps', 'amdgpu_cs', 'amdgpu_kernel', 'amdgpu_kernelcc', 'x86_regcallcc', 'amdgpu_hs', 'msp430_builtincc', 'amdgpu_ls', 'amdgpu_es', 'aarch64_vfpcc', 'aarch64_sve_vfpcc', 'wasm_emscripten_invokecc', 'amdgpu_gfx', 'm68k_intrcc'];
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const value = parser.getTokenSpelling().str();
- if (cconvKeywords.includes(value) || /^cc_\d+$/.test(value)) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('CConv', __spelling);
- }
- }
- const tailcallKeywords = ['none', 'tail', 'musttail', 'notail'];
- if (parser.getToken().is(_.Token.bare_identifier) && tailcallKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('TailCallKind', __spelling);
- }
- let isDirect = false;
- let calleePtr = null;
- if (parser.getToken().is(_.Token.at_identifier)) {
- const callee = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('callee', callee);
- isDirect = true;
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- calleePtr = parser.parseOperand();
- }
- const unresolvedOperands = [];
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- unresolvedOperands.push(arg);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- if (parser.parseOptionalKeyword('vararg')) {
- parser.parseLParen();
- const varCalleeType = parser.parseType();
- result.addAttribute('var_callee_type', varCalleeType);
- parser.parseRParen();
- }
- if (parser.parseOptionalLSquare()) {
- if (!parser.parseOptionalRSquare()) {
- const opBundles = [];
- do {
- const tag = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseLParen();
- const bundleOperands = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- bundleOperands.push(parser.parseOperand());
- } while (parser.parseOptionalComma());
- parser.parseColon();
- do {
- parser.parseType();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- opBundles.push({ tag, operands: bundleOperands });
- } while (parser.parseOptionalComma());
- parser.parseRSquare();
- if (opBundles.length > 0) {
- result.addAttribute('op_bundle_tags', opBundles);
- }
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- let calleePtrType = null;
- if (!isDirect) {
- calleePtrType = parser.parseType();
- parser.parseComma();
- }
- const sig = parser.parseFunctionSignature();
- if (calleePtr) {
- parser.resolveOperand(calleePtr, calleePtrType, result.operands);
- }
- parser.resolveOperands(unresolvedOperands, sig.argTypes, result.operands);
- if (sig.resultTypes.length > 0) {
- result.types = sig.resultTypes.map((t) => t.toString());
- }
- return true;
- }
- parseLLVMCallIntrinsicOp(parser, result) {
- const intrinName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('intrin', intrinName);
- const unresolvedOperands = [];
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- unresolvedOperands.push(arg);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- // Parse operation bundles: [] or ["tag"()] or ["tag"(%0, %1 : i32, i32), ...]
- if (parser.parseOptionalLSquare()) {
- if (!parser.parseOptionalRSquare()) {
- const opBundles = [];
- do {
- const tag = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseLParen();
- const bundleOperands = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- bundleOperands.push(parser.parseOperand());
- } while (parser.parseOptionalComma());
- parser.parseColon();
- do {
- parser.parseType();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- opBundles.push({ tag, operands: bundleOperands });
- } while (parser.parseOptionalComma());
- parser.parseRSquare();
- if (opBundles.length > 0) {
- result.addAttribute('op_bundle_tags', opBundles);
- }
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- const sig = parser.parseFunctionSignature();
- parser.resolveOperands(unresolvedOperands, sig.argTypes, result.operands);
- if (sig.resultTypes.length > 0) {
- result.types = sig.resultTypes.map((t) => t.toString());
- }
- return true;
- }
- parseLLVMCmpOp(parser, result) {
- // llvm.icmp "eq" %lhs, %rhs : i32
- const predicate = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('predicate', predicate);
- const lhs = parser.parseOperand();
- parser.parseComma();
- const rhs = parser.parseOperand();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperands([lhs, rhs], [type, type], result.operands);
- // Result type is i1 (or vector/tensor of i1 for vector/tensor operands)
- let resultType = new _.IntegerType('i1');
- if (type instanceof _.VectorType) {
- resultType = new _.VectorType(type.shape, new _.IntegerType('i1'), type.scalableDims);
- } else if (type instanceof _.RankedTensorType) {
- resultType = new _.RankedTensorType(type.shape, new _.IntegerType('i1'), type.encoding);
- }
- result.addTypes([resultType]);
- return true;
- }
- parseLLVMIntrinsicOp(parser, result) {
- const unresolvedOperands = [];
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- const types = parser.parseColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.types = [resultType];
- }
- return true;
- }
- parseLLVMInvokeOp(parser, result) {
- const cconvKeywords = ['ccc', 'fastcc', 'coldcc', 'cc', 'webkit_jscc', 'anyregcc', 'preserve_mostcc', 'preserve_allcc', 'preserve_nonecc', 'cxx_fast_tlscc', 'tailcc', 'swiftcc', 'swifttailcc', 'cfguard_checkcc', 'ghccc'];
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const value = parser.getTokenSpelling().str();
- if (cconvKeywords.includes(value) || /^cc_\d+$/.test(value)) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('CConv', __spelling);
- }
- }
- let isDirect = false;
- let funcPtr = null;
- if (parser.getToken().is(_.Token.at_identifier)) {
- isDirect = true;
- const callee = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('callee', callee);
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- funcPtr = parser.parseOperand();
- }
- const unresolvedOperands = [];
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- parser.parseKeyword('to');
- const normalDest = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- result.successors = result.successors || [];
- const normalSucc = { label: normalDest };
- if (parser.parseOptionalLParen()) {
- normalSucc.operands = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- normalSucc.operands.push(operand);
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- }
- result.successors.push(normalSucc);
- parser.parseKeyword('unwind');
- const unwindDest = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const unwindSucc = { label: unwindDest };
- if (parser.parseOptionalLParen()) {
- unwindSucc.operands = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- unwindSucc.operands.push(operand);
- if (parser.parseOptionalColon()) {
- parser.parseType();
- }
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- }
- result.successors.push(unwindSucc);
- if (parser.parseOptionalKeyword('vararg')) {
- parser.parseLParen();
- const varargType = parser.parseType();
- result.addAttribute('var_callee_type', varargType);
- parser.parseRParen();
- }
- if (parser.parseOptionalLSquare()) {
- if (!parser.parseOptionalRSquare()) {
- const opBundles = [];
- do {
- const tag = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseLParen();
- const bundleOperands = [];
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- bundleOperands.push(parser.parseOperand());
- } while (parser.parseOptionalComma());
- parser.parseColon();
- do {
- parser.parseType();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- opBundles.push({ tag, operands: bundleOperands });
- } while (parser.parseOptionalComma());
- parser.parseRSquare();
- if (opBundles.length > 0) {
- result.addAttribute('op_bundle_tags', opBundles);
- }
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.parseColon();
- let calleePtrType = null;
- if (!isDirect) {
- calleePtrType = parser.parseType();
- parser.parseComma();
- }
- const sig = parser.parseFunctionSignature();
- if (funcPtr) {
- parser.resolveOperand(funcPtr, calleePtrType, result.operands);
- }
- parser.resolveOperands(unresolvedOperands, sig.argTypes, result.operands);
- if (sig.resultTypes.length > 0) {
- result.types = sig.resultTypes.map((t) => t.toString());
- }
- return true;
- }
- parseLLVMLandingpadOp(parser, result) {
- if (parser.parseOptionalKeyword('cleanup')) {
- result.addAttribute('cleanup', true);
- }
- while (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- parser.consumeToken(_.Token.bare_identifier); // 'catch' or 'filter'
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- parser.parseRParen();
- }
- parser.parseColon();
- const resultType = parser.parseType();
- result.types = [resultType];
- return true;
- }
- parseLLVMFuncOp(parser, result) {
- const linkageKeywords = ['external', 'available_externally', 'linkonce', 'linkonce_odr', 'weak', 'weak_odr', 'appending', 'internal', 'private', 'extern_weak', 'common'];
- if (parser.getToken().is(_.Token.bare_identifier) && linkageKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('linkage', __spelling);
- }
- const visibilityKeywords = ['default', 'hidden', 'protected'];
- if (parser.getToken().is(_.Token.bare_identifier) && visibilityKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('visibility_', __spelling);
- }
- const unnamedAddrKeywords = ['unnamed_addr', 'local_unnamed_addr'];
- if (parser.getToken().is(_.Token.bare_identifier) && unnamedAddrKeywords.includes(parser.getTokenSpelling().str())) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('unnamed_addr', __spelling);
- }
- const cconvKeywords = ['ccc', 'fastcc', 'coldcc', 'cc', 'webkit_jscc', 'anyregcc', 'preserve_mostcc', 'preserve_allcc', 'preserve_nonecc', 'cxx_fast_tlscc', 'tailcc', 'swiftcc', 'swifttailcc', 'cfguard_checkcc', 'ghccc', 'arm_apcscc', 'arm_aapcscc', 'arm_aapcs_vfpcc', 'aarch64_vector_pcs', 'aarch64_sve_vector_pcs', 'aarch64_sme_preservemost_from_x0', 'aarch64_sme_preservemost_from_x2', 'msp430_intrcc', 'avr_intrcc', 'avr_signalcc', 'ptx_kernelcc', 'ptx_devicecc', 'spir_funccc', 'spir_kernelcc', 'intel_ocl_bicc', 'x86_64_sysvcc', 'win64cc', 'x86_fastcallcc', 'x86_stdcallcc', 'x86_thiscallcc', 'x86_vectorcallcc', 'x86_intrcc', 'amdgpu_vs', 'amdgpu_gs', 'amdgpu_ps', 'amdgpu_cs', 'amdgpu_kernel', 'amdgpu_kernelcc', 'x86_regcallcc', 'amdgpu_hs', 'msp430_builtincc', 'amdgpu_ls', 'amdgpu_es', 'aarch64_vfpcc', 'aarch64_sve_vfpcc', 'wasm_emscripten_invokecc', 'amdgpu_gfx', 'm68k_intrcc'];
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const value = parser.getTokenSpelling().str();
- if (cconvKeywords.includes(value) || /^cc_\d+$/.test(value)) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('CConv', __spelling);
- }
- }
- parser.parseSymbolName('sym_name', result.attributes);
- const argResult = parser.parseFunctionArgumentList(true);
- const params = argResult.arguments.map((a) => a.type);
- const results = [];
- const resultAttrs = [];
- if (parser.parseOptionalArrow()) {
- parser.parseFunctionResultList(results, resultAttrs);
- }
- const returnType = results.length > 0 ? results[0] : null;
- const type = new _.LLVM.LLVMFunctionType(returnType, params, argResult.isVariadic);
- result.addAttribute('function_type', new _.TypeAttrOf(type));
- if (parser.parseOptionalKeyword('vscale_range')) {
- parser.parseLParen();
- const minRange = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseComma();
- const maxRange = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseRParen();
- result.addAttribute('vscale_range', `(${minRange}, ${maxRange})`);
- }
- if (parser.parseOptionalKeyword('comdat')) {
- parser.parseLParen();
- const comdat = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- parser.parseRParen();
- result.addAttribute('comdat', comdat);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- // llvm.func is IsolatedFromAbove
- parser.parseRegion(region, argResult.arguments, /* isIsolatedNameScope */ true);
- }
- return true;
- }
- };
- _.ROCDLDialect = class extends _.LLVM.LLVMDialect {
- constructor(operations) {
- super(operations, 'rocdl');
- }
- parseOperation(parser, result) {
- if (result.op === 'rocdl.raw.buffer.load') {
- return this.parseRawBufferLoadOp(parser, result);
- }
- if (result.op === 'rocdl.raw.buffer.store') {
- return this.parseRawBufferStoreOp(parser, result);
- }
- if (result.op === 'rocdl.raw.buffer.atomic.fadd') {
- return this.parseRawBufferAtomicOp(parser, result);
- }
- if (result.op === 'rocdl.raw.buffer.atomic.fmax') {
- return this.parseRawBufferAtomicOp(parser, result);
- }
- if (result.op === 'rocdl.raw.buffer.atomic.smax') {
- return this.parseRawBufferAtomicOp(parser, result);
- }
- if (result.op === 'rocdl.raw.buffer.atomic.umin') {
- return this.parseRawBufferAtomicOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseRawBufferLoadOp(parser, result) {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- parser.parseOptionalComma();
- }
- parser.parseColon();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- return true;
- }
- parseRawBufferStoreOp(parser, result) {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- parser.parseOptionalComma();
- }
- parser.parseColon();
- parser.parseType();
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- return true;
- }
- parseRawBufferAtomicOp(parser, result) {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- parser.parseOptionalComma();
- }
- parser.parseColon();
- parser.parseType();
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- return true;
- }
- };
- _.XSMMDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xsmm');
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (result.op === 'xsmm.unary.invoke') {
- return this.parseUnaryInvokeOp(parser, result);
- }
- if (result.op.startsWith('xsmm.') && result.op.includes('.invoke') && opInfo.metadata.hasCustomAssemblyFormat && !opInfo.metadata.assemblyFormat) {
- return this.parseGemmInvokeOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseUnaryInvokeOp(parser, result) {
- const unresolvedOperands = [];
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- parser.parseEqual();
- unresolvedOperands.push(parser.parseOperand());
- parser.parseLParen();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- parser.parseRParen();
- parser.parseColon();
- parser.parseType();
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- return true;
- }
- parseGemmInvokeOp(parser, result) {
- const unresolvedOperands = [];
- unresolvedOperands.push(parser.parseOperand());
- parser.parseComma();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- parser.parseEqual();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- parser.parseComma();
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- while (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseOptionalComma();
- }
- }
- } else if (parser.getToken().is(_.Token.bare_identifier)) {
- const keyword = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseEqual();
- const attrValue = parser.parseAttribute(new _.IntegerType('i64'));
- result.addAttribute(keyword, attrValue);
- } else {
- break;
- }
- }
- parser.parseColon();
- parser.parseType();
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- return true;
- }
- };
- _.StdxDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'stdx');
- }
- parseOperation(parser, result) {
- if (result.op === 'stdx.closure') {
- return this.parseClosureOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseClosureOp(parser, result) {
- const sig = parser.parseFunctionSignatureWithArguments(false);
- const argTypes = sig.arguments.map((a) => a.type);
- const type = { inputs: argTypes, results: sig.resultTypes };
- result.addAttribute('type', type);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, sig.arguments);
- }
- return true;
- }
- };
- _.VMDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'vm');
- this.registerCustomDirective('BranchTableCases', this.parseBranchTableCases.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'vm.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'vm.cond_fail') {
- // or: vm.cond_fail %status, "message"
- // or: vm.cond_fail %cond, %status
- // or: vm.cond_fail %status
- const unresolvedOperands = [];
- const firstOp = parser.parseOperand();
- unresolvedOperands.push(firstOp);
- if (parser.parseOptionalComma()) {
- // Could be second operand or message
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const secondOp = parser.parseOperand();
- unresolvedOperands.push(secondOp);
- // Optional message
- if (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.string)) {
- const msg = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('message', msg);
- }
- }
- } else if (parser.getToken().is(_.Token.string)) {
- const msg = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('message', msg);
- }
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, result.operands);
- }
- return true;
- }
- if (result.op === 'vm.import') {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- if (parser.parseOptionalKeyword('optional')) {
- result.addAttribute('is_optional', true);
- }
- parser.parseSymbolName('sym_name', result.attributes);
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.skip('(');
- }
- const inputs = [];
- const results = [];
- const resultAttrs = [];
- if (parser.parseOptionalArrow()) {
- parser.parseFunctionResultList(results, resultAttrs);
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- return true;
- }
- if (result.op === 'vm.export') {
- const functionRef = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('function_ref', functionRef);
- if (parser.parseOptionalKeyword('as')) {
- parser.parseLParen();
- const exportName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('export_name', exportName);
- parser.parseRParen();
- } else {
- result.addAttribute('export_name', functionRef);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- return true;
- }
- if (result.op.startsWith('vm.global.') && !result.op.startsWith('vm.global.store.') && !result.op.startsWith('vm.global.load.') && result.op !== 'vm.global.address') {
- result.compatibility = true;
- parser.parseOptionalVisibilityKeyword(result.attributes);
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'mutable') {
- const mutable = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('is_mutable', mutable);
- }
- parser.parseSymbolName('sym_name', result.attributes);
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addAttribute('type', type);
- }
- if (parser.parseOptionalEqual()) {
- const initialValue = parser.parseAttribute();
- result.addAttribute('initial_value', initialValue);
- }
- return true;
- }
- if (result.op === 'vm.initializer') {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'vm.rodata.inline') {
- result.compatibility = true;
- if (parser.getToken().is(_.Token.string)) {
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('name', name);
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addTypes(parser.parseOptionalColonTypeList());
- if (parser.parseOptionalEqual()) {
- const value = parser.parseAttribute();
- // Handle type annotation after the value (e.g., dense<...> : vector<21xi8>)
- if (parser.parseOptionalColon()) {
- const valueType = parser.parseType();
- value.type = valueType;
- }
- result.addAttribute('value', value);
- }
- return true;
- }
- if (result.op === 'vm.const.i32.zero') {
- result.compatibility = true;
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.floatliteral) || parser.getToken().is(_.Token.string)) {
- const value = parser.parseAttribute();
- result.addAttribute('value', value.value === undefined ? value : value.value);
- } else if (parser.getToken().is(_.Token.at_identifier)) {
- const symbol = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('rodata', symbol);
- }
- parser.parseOptionalAttrDict(result.attributes);
- const types = parser.parseOptionalColonTypeList();
- result.addTypes(types.length > 0 ? types : [new _.IntegerType('i32')]);
- return true;
- }
- // Handle vm.switch.ref operation
- if (result.op === 'vm.switch.ref') {
- const unresolvedOperands = [];
- const indexUnresolved = parser.parseOperand();
- unresolvedOperands.push(indexUnresolved);
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const value = parser.parseOperand();
- unresolvedOperands.push(value);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- parser.parseKeyword('else');
- const defaultValueUnresolved = parser.parseOperand();
- unresolvedOperands.push(defaultValueUnresolved);
- parser.parseOptionalAttrDict(result.attributes);
- let resultType = null;
- if (parser.parseOptionalColon()) {
- resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, resultType, result.operands);
- }
- return true;
- }
- // Handle vm.call and vm.call.variadic
- // Variadic has complex syntax like: @callee(op1, op2, [(tuple1), (tuple2)])
- if (result.op === 'vm.call' || result.op === 'vm.call.variadic') {
- result.compatibility = true;
- if (parser.getToken().is(_.Token.at_identifier)) {
- const callee = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('callee', callee);
- }
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.l_square)) {
- // Skip complex nested structures in variadic calls
- parser.skip('[');
- parser.parseOptionalComma(); // consume trailing comma if present
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- parser.parseOptionalComma(); // consume trailing comma if present
- } else {
- // Unexpected token, break to avoid infinite loop
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseOptionalAttrDict(result.attributes);
- // vm.call.variadic has special syntax with '...' ellipsis
- if (parser.parseOptionalColon()) {
- if (result.op === 'vm.call.variadic') {
- parser.skip('(');
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- } else {
- // Regular vm.call - Reference: uses functional-type(operands, results)
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(result.operands, type.inputs);
- result.addTypes(type.results);
- }
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseBranchTableCases(parser, op /*, args */) {
- if (parser.parseOptionalKeyword('default')) {
- parser.parseColon();
- const defaultDest = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- op.successors = op.successors || [];
- const succ = { dest: defaultDest };
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- const operands = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- operands.push(parser.parseOperand());
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.comma)) {
- parser.parseType();
- parser.parseOptionalComma();
- }
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- succ.operands = operands;
- }
- op.successors.push(succ);
- parser.parseOptionalComma();
- }
- const caseValues = [];
- while (parser.getToken().is(_.Token.integer)) {
- const caseValue = parser.parseInteger();
- caseValues.push(caseValue);
- parser.parseColon();
- const caseDest = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- const caseSucc = { dest: caseDest };
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- const operands = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- operands.push(parser.parseOperand());
- }
- if (parser.parseOptionalColon()) {
- while (parser.getToken().isNot(_.Token.r_paren) && parser.getToken().isNot(_.Token.comma)) {
- parser.parseType();
- parser.parseOptionalComma();
- }
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- caseSucc.operands = operands;
- }
- op.successors.push(caseSucc);
- parser.parseOptionalComma();
- }
- if (caseValues.length > 0) {
- op.addAttribute('case_values', caseValues);
- }
- }
- };
- _.MathDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'math');
- this.registerCustomAttribute('Arith_FastMathAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- }
- };
- _.TMTensorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tm_tensor');
- }
- };
- _.MLProgramDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'ml_program');
- this.registerCustomDirective('TypedInitialValue', this.parseTypedInitialValue.bind(this));
- this.registerCustomDirective('TokenOrdering', this.parseTokenOrdering.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'ml_program.func' || result.op === 'ml_program.subgraph') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseTokenOrdering(parser, result) {
- if (!parser.parseOptionalKeyword('ordering')) {
- return;
- }
- parser.parseLParen();
- if (parser.parseOptionalLParen()) {
- parser.parseRParen();
- } else {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const tok = parser.parseOperand();
- parser.resolveOperand(tok, null, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- parser.parseArrow();
- const produceType = parser.parseType();
- result.addAttribute('produceTokenType', { value: produceType, hidden: true });
- parser.parseRParen();
- }
- parseTypedInitialValue(parser, op, typeAttr, valueAttr) {
- if (parser.parseOptionalLParen()) {
- const attr = parser.parseAttribute();
- if (parser.parseOptionalColon()) {
- attr.type = parser.parseType();
- }
- parser.parseRParen();
- op.addAttribute(valueAttr, attr.value === undefined ? attr : attr.value);
- }
- parser.parseColon();
- const type = parser.parseType();
- op.addAttribute(typeAttr, type);
- }
- };
- _.IREEGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'iree_gpu');
- }
- };
- _.TFDeviceDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tf_device');
- }
- parseOperation(parser, result) {
- if (result.op === 'tf_device.replicate') {
- return this.parseReplicateOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseReplicateOp(parser, result) {
- let n = 1;
- if (!parser.parseOptionalLParen()) {
- parser.parseOptionalAttrDict(result.attributes);
- } else if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseRParen();
- parser.parseOptionalAttrDict(result.attributes);
- } else {
- do {
- if (parser.getToken().is(_.Token.l_square)) {
- const unresolvedInputs = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- unresolvedInputs.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- parser.parseRSquare();
- break;
- }
- }
- parser.parseKeyword('as');
- parser.parseOperand(); // block arg
- parser.parseColon();
- const type = parser.parseType();
- for (const input of unresolvedInputs) {
- parser.resolveOperand(input, type, result.operands);
- }
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedValue = parser.parseOperand();
- parser.parseKeyword('as');
- parser.parseOperand(); // block arg
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(unresolvedValue, type, result.operands);
- } else {
- break;
- }
- } while (parser.parseOptionalComma());
- parser.parseRParen();
- parser.parseOptionalAttrDict(result.attributes);
- }
- n = result.attributes.get('n').value;
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- if (region.blocks.length > 0) {
- const block = region.blocks[0];
- if (block.operations.length > 0) {
- const terminator = block.operations[block.operations.length - 1];
- if (terminator.operands) {
- for (const operand of terminator.operands) {
- if (operand.type) {
- for (let i = 0; i < n; i++) {
- result.addTypes([operand.type]);
- }
- }
- }
- }
- }
- }
- }
- return true;
- }
- };
- _.TFGDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfg');
- }
- getOperation(opName) {
- let op = super.getOperation(opName);
- if (!op) {
- this.registerOperandName(opName, {});
- op = super.getOperation(opName);
- }
- return op;
- }
- parseOperation(parser, result) {
- if (result.op === 'tfg.func') {
- if (parser.parseOptionalKeyword('generic')) {
- result.addAttribute('generic', true);
- }
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'tfg.return') {
- let dataOperands = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- dataOperands = parser.parseOperandList('paren');
- }
- const controlOperands = [];
- const controlRetAttrs = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const ctlDep = parser.parseOperand();
- controlOperands.push(ctlDep);
- if (parser.getToken().is(_.Token.l_brace)) {
- const attrs = new Map();
- parser.parseAttributeDict(attrs);
- controlRetAttrs.push(Object.fromEntries(attrs));
- } else {
- controlRetAttrs.push({});
- }
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- }
- if (controlRetAttrs.length > 0) {
- result.addAttribute('control_ret_attrs', controlRetAttrs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const types = parser.parseTypeListNoParens();
- parser.resolveOperands(dataOperands, types, result.operands);
- } else {
- parser.resolveOperands(dataOperands, dataOperands.map(() => null), result.operands);
- }
- parser.resolveOperands(controlOperands, controlOperands.map(() => null), result.operands);
- return true;
- }
- const opInfo = result.name.getRegisteredInfo();
- if (!opInfo.metadata.assemblyFormat) {
- this.parseTFGOperation(parser, result);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseTFGOperation(parser, result) {
- let unresolvedArgs = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- unresolvedArgs = parser.parseOperandList('paren');
- }
- const unresolvedCtls = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedCtls.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- }
- if (parser.parseOptionalKeyword('device')) {
- parser.parseLParen();
- const device = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseRParen();
- result.addAttribute('device', device);
- }
- if (parser.parseOptionalKeyword('name')) {
- parser.parseLParen();
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseRParen();
- result.addAttribute('_mlir_name', name);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedArgs, type.inputs, result.operands);
- parser.resolveOperands(unresolvedCtls, unresolvedCtls.map(() => new _.Type('!tfg.control')), result.operands);
- result.addTypes(type.results);
- } else {
- const types = [type];
- while (parser.parseOptionalComma()) {
- types.push(parser.parseType());
- }
- parser.resolveOperands(unresolvedArgs, types, result.operands);
- parser.resolveOperands(unresolvedCtls, unresolvedCtls.map(() => new _.Type('!tfg.control')), result.operands);
- }
- } else {
- parser.resolveOperands(unresolvedArgs, unresolvedArgs.map(() => null), result.operands);
- parser.resolveOperands(unresolvedCtls, unresolvedCtls.map(() => new _.Type('!tfg.control')), result.operands);
- }
- result.addTypes([new _.Type('!tfg.control')]);
- }
- };
- _.TFExecutorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tf_executor');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- const type = `!${dialect}.${typeName}`;
- if (typeName === 'control' || typeName === 'token') {
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op === 'tf_executor.graph') {
- return this.parseGraphOp(parser, result);
- }
- if (result.op === 'tf_executor.island') {
- return this.parseIslandOp(parser, result);
- }
- if (result.op === 'tf_executor.Enter') {
- return this.parseEnterOp(parser, result);
- }
- if (result.op === 'tf_executor._SwitchN') {
- const unresolvedData = parser.parseOperand();
- parser.parseComma();
- const unresolvedIndex = parser.parseOperand();
- parser.parseKeyword('of');
- const numOuts = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- result.addAttribute('num_outs', numOuts);
- let unresolvedControlInputs = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- unresolvedControlInputs = parser.parseOperandList('paren');
- }
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(unresolvedData, type, result.operands);
- parser.resolveOperand(unresolvedIndex, new _.RankedTensorType([], new _.IntegerType('i32'), null), result.operands);
- parser.resolveOperands(unresolvedControlInputs, unresolvedControlInputs.map(() => new _.Type('!tf_executor.control')), result.operands);
- for (let i = 0; i < numOuts; i++) {
- result.addTypes([type]);
- }
- result.addTypes([new _.Type('!tf_executor.control')]);
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (result.op === 'tf_executor.Switch' || result.op === 'tf_executor.Merge' ||
- result.op === 'tf_executor.LoopCond' || result.op === 'tf_executor.Exit') {
- // These ops have hasCustomAssemblyFormat: true but no assemblyFormat in metadata
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedOperands, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => type), result.operands);
- result.addTypes([type]);
- if (result.op === 'tf_executor.Switch') {
- result.addTypes([type]);
- }
- if (result.op === 'tf_executor.Merge') {
- result.addTypes([new _.RankedTensorType([], new _.IntegerType('i32'), null)]);
- }
- result.addTypes([new _.Type('!tf_executor.control')]);
- }
- } else {
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseEnterOp(parser, result) {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseKeyword('frame');
- const frameName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('frame_name', frameName);
- if (parser.parseOptionalKeyword('parallel_iterations')) {
- const parallelIterations = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('parallel_iterations', parseInt(parallelIterations, 10));
- } else {
- result.addAttribute('parallel_iterations', 10);
- }
- const isConstant = parser.parseOptionalKeyword('constant');
- result.addAttribute('is_constant', isConstant);
- parser.parseColon();
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedOperands, type.inputs, result.operands);
- result.addTypes(type.results);
- } else {
- const resolveTypes = unresolvedOperands.map((v, i) => i === 0 ? type : new _.Type('!tf_executor.control'));
- parser.resolveOperands(unresolvedOperands, resolveTypes, result.operands);
- result.addTypes([type]);
- result.addTypes([new _.Type('!tf_executor.control')]);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseGraphOp(parser, result) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- if (region.blocks && region.blocks.length > 0) {
- const [block] = region.blocks;
- if (block.operations && block.operations.length > 0) {
- const lastOp = block.operations[block.operations.length - 1];
- if (lastOp.name.getStringRef() === 'tf_executor.fetch' && lastOp.operands) {
- for (const operand of lastOp.operands) {
- const typeStr = operand.type ? operand.type.toString() : '';
- if (operand.type && typeStr !== '!tf_executor.control') {
- result.addTypes([operand.type]);
- }
- }
- }
- }
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- parseIslandOp(parser, result) {
- // or: tf_executor.island {...}
- // or: tf_executor.island(%control_inputs) {...}
- if (parser.getToken().is(_.Token.l_paren)) {
- const unresolvedOperands = parser.parseOperandList('paren');
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- const region = result.addRegion();
- if (parser.parseOptionalKeyword('wraps')) {
- const wrappedOp = parser.parseGenericOperation();
- result.addAttribute('wrappedOp', wrappedOp);
- for (const opResult of wrappedOp.results) {
- result.addTypes([opResult.type]);
- }
- } else if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseRegion(region);
- if (region.blocks.length > 0) {
- const block = region.blocks[region.blocks.length - 1];
- if (block.operations.length > 0) {
- const terminator = block.operations[block.operations.length - 1];
- if (terminator.name.getStringRef() === 'tf_executor.yield') {
- for (const operand of terminator.operands) {
- result.addTypes([operand.type]);
- }
- }
- }
- }
- }
- result.addTypes([new _.Type('!tf_executor.control')]);
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- };
- _.TFFrameworkDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tf_framework');
- }
- };
- _.TFRDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfr');
- }
- parseOperation(parser, result) {
- if (result.op === 'tfr.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.CoreRTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'corert');
- }
- parseOperation(parser, result) {
- if (result.op === 'corert.executeop' || result.op === 'corert.executeop.seq') {
- const isSeq = result.op === 'corert.executeop.seq';
- const opHandlerOperands = parser.parseOperandList('paren');
- for (const operand of opHandlerOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- const opNameAttr = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('op_name', opNameAttr);
- const operandOperands = parser.parseOperandList('paren');
- for (const operand of operandOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const funcAttrs = new Map();
- parser.parseAttributeDict(funcAttrs);
- result.addAttribute('op_func_attrs', Object.fromEntries(funcAttrs));
- }
- if (parser.parseOptionalColon()) {
- const resultCount = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- if (isSeq) {
- result.addTypes([new _.Type('!tfrt.chain')]);
- }
- const tensorHandleType = new _.Type('!corert.tensorhandle');
- for (let i = 0; i < resultCount; i++) {
- result.addTypes([tensorHandleType]);
- }
- } else if (isSeq) {
- result.addTypes([new _.Type('!tfrt.chain')]);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TFRTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfrt');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- const simpleTypes = ['chain', 'string', 'dist_context', 'device', 'tensor_type'];
- if (simpleTypes.includes(typeName)) {
- return new _.Type(type);
- }
- if (typeName === 'tensor') {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- // Fallback for unknown tfrt types
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (!opInfo) {
- return false;
- }
- if (opInfo.metadata.assemblyFormat === 'operands attr-dict') {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, result.operands);
- }
- this.inferResultTypes(result, new Map());
- return true;
- }
- if (result.op === 'tfrt.call') {
- parser.parseSymbolName('callee', result.attributes);
- const unresolvedOperands = parser.parseOperandList('paren');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseFunctionType();
- if (type) {
- if (type.inputs) {
- parser.resolveOperands(unresolvedOperands, type.inputs, result.operands);
- }
- if (type.results) {
- type.results.forEach((resultType) => {
- result.addTypes([resultType]);
- });
- }
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- if (result.op === 'tfrt.return') {
- if (parser.getToken().isNot(_.Token.kw_loc) && parser.getToken().isNot(_.Token.eof)) {
- const unresolvedOperands = parser.parseOperandList();
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- }
- return true;
- }
- if (result.op === 'tfrt.repeat.i32') {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseOptionalAttrDict(result.attributes);
- }
- const types = parser.parseOptionalColonTypeList();
- result.addTypes(types);
- if (unresolvedOperands.length > 0) {
- const i32Type = new _.IntegerType('i32');
- parser.resolveOperand(unresolvedOperands[0], i32Type, result.operands);
- const loopOperands = unresolvedOperands.slice(1);
- parser.resolveOperands(loopOperands, types, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'tfrt.if') {
- const unresolvedOperands = parser.parseOperandList();
- if (parser.parseOptionalKeyword('attributes')) {
- parser.parseOptionalAttrDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType) {
- if (funcType.inputs) {
- parser.resolveOperands(unresolvedOperands, funcType.inputs, result.operands);
- }
- if (funcType.results) {
- for (const resultType of funcType.results) {
- result.addTypes([resultType]);
- }
- }
- }
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const thenRegion = {};
- parser.parseRegion(thenRegion);
- result.regions.push(thenRegion);
- }
- if (parser.parseOptionalKeyword('else')) {
- if (parser.getToken().is(_.Token.l_brace)) {
- const elseRegion = {};
- parser.parseRegion(elseRegion);
- result.regions.push(elseRegion);
- }
- }
- return true;
- }
- if (result.op === 'tfrt.parallel_for.i32') {
- const startUnresolved = parser.parseOperand();
- parser.parseKeyword('to');
- const endUnresolved = parser.parseOperand();
- parser.parseKeyword('fixed');
- const blockSizeUnresolved = parser.parseOperand();
- let additionalArgs = [];
- if (parser.parseOptionalComma()) {
- additionalArgs = parser.parseOperandList();
- }
- const types = parser.parseOptionalColonTypeList();
- const i32Type = new _.IntegerType('i32');
- parser.resolveOperand(startUnresolved, i32Type, result.operands);
- parser.resolveOperand(endUnresolved, i32Type, result.operands);
- parser.resolveOperand(blockSizeUnresolved, i32Type, result.operands);
- parser.resolveOperands(additionalArgs, types, result.operands);
- result.addTypes([new _.Type('!tfrt.chain')]);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'tfrt.parallel_call.i32') {
- const startUnresolved = parser.parseOperand();
- parser.parseKeyword('to');
- const endUnresolved = parser.parseOperand();
- parser.parseKeyword('fixed');
- const blockSizeUnresolved = parser.parseOperand();
- const callee = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('callee', callee);
- const additionalArgs = parser.parseOperandList('paren');
- const types = parser.parseOptionalColonTypeList();
- const i32Type = new _.IntegerType('i32');
- parser.resolveOperand(startUnresolved, i32Type, result.operands);
- parser.resolveOperand(endUnresolved, i32Type, result.operands);
- parser.resolveOperand(blockSizeUnresolved, i32Type, result.operands);
- parser.resolveOperands(additionalArgs, types, result.operands);
- result.addTypes([new _.Type('!tfrt.chain')]);
- return true;
- }
- if (result.op === 'tfrt.while') {
- const condUnresolved = parser.parseOperand();
- const bodyFn = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('body_fn', bodyFn);
- parser.parseLParen();
- const argsUnresolved = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const arg = parser.parseOperand();
- argsUnresolved.push(arg);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalKeyword('parallel_iterations')) {
- parser.parseLParen();
- const parallelIterations = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('parallel_iterations', parseInt(parallelIterations, 10));
- parser.parseRParen();
- }
- parser.parseColon();
- const inputTypes = parser.parseTypeListParens();
- parser.parseArrow();
- const resultTypes = parser.parseTypeListParens();
- parser.resolveOperand(condUnresolved, new _.IntegerType('i1'), result.operands);
- parser.resolveOperands(argsUnresolved, inputTypes, result.operands);
- for (const resultType of resultTypes) {
- result.addTypes([resultType]);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TFRTFallbackAsyncDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfrt_fallback_async');
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (!opInfo) {
- return false;
- }
- if (result.op === 'tfrt_fallback_async.batch_function') {
- parser.parseKeyword('device');
- parser.parseLParen();
- const device = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- parser.parseRParen();
- result.addAttribute('device', device);
- const funcName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('f', funcName);
- const unresolvedOperands = parser.parseOperandList('paren');
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (parser.parseOptionalColon()) {
- const resultCount = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken();
- for (let i = 0; i < resultCount; i++) {
- result.addTypes([new _.Type('!tfrt_fallback.tf_tensor')]);
- }
- }
- return true;
- }
- if (result.op === 'tfrt_fallback_async.createop' || result.op.startsWith('tfrt_fallback_async.executeop')) {
- const isCreateOp = result.op === 'tfrt_fallback_async.createop';
- const hasChain = isCreateOp || result.op.includes('.seq');
- const hasAllocator = result.op.includes('.allocator');
- if ((hasChain || hasAllocator) && parser.getToken().is(_.Token.l_paren)) {
- const chainOperands = parser.parseOperandList('paren');
- for (const operand of chainOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- while (parser.getToken().isNot(_.Token.colon) && parser.getToken().isNot(_.Token.l_brace)) {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const key = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (parser.parseOptionalLParen()) {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseRParen();
- result.addAttribute(key, value);
- }
- } else if (parser.getToken().is(_.Token.string)) {
- const opNameAttr = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('op_name', opNameAttr);
- if (parser.getToken().is(_.Token.l_paren)) {
- const unresolvedOperands = parser.parseOperandList('paren');
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- break;
- } else {
- break;
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- }
- if (isCreateOp) {
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'num_args') {
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseLParen();
- const numArgs = parser.getToken().getSpelling().str();
- parser.consumeToken();
- parser.parseRParen();
- result.addAttribute('num_args', parseInt(numArgs, 10));
- }
- // createop always returns !tfrt.chain
- result.addTypes([new _.Type('!tfrt.chain')]);
- } else if (parser.parseOptionalColon()) {
- const resultCount = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken();
- result.addAttribute('result_count', resultCount);
- if (hasChain) {
- result.addTypes([new _.Type('!tfrt.chain')]);
- }
- for (let i = 0; i < resultCount; i++) {
- result.addTypes([new _.Type('!tfrt_fallback.tf_tensor')]);
- }
- } else if (hasChain) {
- result.addTypes([new _.Type('!tfrt.chain')]);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TileDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tile');
- }
- parseOperation(parser, result) {
- // tile.contract has format: tile.contract agg, combo, operands... attributes : types -> result
- // Example: %1 = tile.contract add, mul, %0, %arg0, %arg1 {sink = #map0, srcs = [#map1, #map2]} : tensor<f32>, tensor<1x256xf32>, tensor<256x512xf32> -> tensor<1x512xf32>
- if (result.op === 'tile.contract') {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const agg = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('agg', agg);
- }
- parser.parseOptionalComma();
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const combo = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('combo', combo);
- }
- parser.parseOptionalComma();
- const unresolvedOperands = parser.parseOperandList();
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- parser.resolveOperands(unresolvedOperands, parser.parseOptionalColonTypeList(), result.operands);
- if (parser.parseOptionalArrow()) {
- const resultTypes = parser.parseFunctionResultTypes();
- result.addTypes(resultTypes);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.PXADialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'pxa');
- }
- parseOperation(parser, result) {
- if (result.op === 'pxa.reduce' || result.op === 'pxa.vector_reduce') {
- const agg = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('agg', agg);
- const unresolvedVal = parser.parseOperand();
- parser.parseOptionalComma();
- const unresolvedMemref = parser.parseOperand();
- parser.skip('[');
- parser.parseOptionalAttrDict(result.attributes);
- let memrefType = null;
- let valType = null;
- if (parser.parseOptionalColon()) {
- memrefType = parser.parseType();
- result.addTypes([memrefType]);
- if (result.op === 'pxa.vector_reduce' && parser.parseOptionalComma()) {
- valType = parser.parseType();
- }
- }
- parser.resolveOperand(unresolvedVal, valType, result.operands);
- parser.resolveOperand(unresolvedMemref, memrefType, result.operands);
- return true;
- }
- if (result.op === 'pxa.load' || result.op === 'pxa.vector_load') {
- const unresolvedMemref = parser.parseOperand();
- parser.skip('[');
- parser.parseOptionalAttrDict(result.attributes);
- let memrefType = null;
- if (parser.parseOptionalColon()) {
- memrefType = parser.parseType();
- if (result.op === 'pxa.vector_load' && parser.parseOptionalComma()) {
- const vectorType = parser.parseType();
- result.addTypes([vectorType]);
- } else if (result.op === 'pxa.load' && memrefType) {
- // Result type is element type of memref
- const elementType = memrefType.elementType || memrefType;
- result.addTypes([elementType]);
- }
- }
- parser.resolveOperand(unresolvedMemref, memrefType, result.operands);
- return true;
- }
- if (result.op === 'pxa.generic') {
- const unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- if (parser.parseOptionalColon()) {
- parser.consumeToken(_.Token.hash_identifier); // Skip affine map reference
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalLess()) {
- const reduction = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- result.addAttribute('reduction', reduction);
- parser.parseGreater();
- }
- if (parser.getToken().is(_.Token.at_identifier)) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('kernel', __spelling);
- }
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- unresolvedOperands.push(operand);
- if (parser.getToken().is(_.Token.l_square)) {
- parser.skip('[');
- }
- if (parser.parseOptionalColon()) {
- parser.consumeToken(_.Token.hash_identifier); // Skip affine map reference
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalKeyword('tile')) {
- parser.parseColon();
- const tile = parser.skip('[');
- result.addAttribute('tile', tile);
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const funcType = parser.parseFunctionType();
- if (funcType && funcType.inputs) {
- parser.resolveOperands(unresolvedOperands, funcType.inputs, result.operands);
- }
- if (funcType && funcType.results) {
- for (const resultType of funcType.results) {
- result.addTypes([resultType]);
- }
- }
- } else {
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, result.operands);
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.ToyDialect = class extends _.HLODialect {
- constructor(operations) {
- super(operations, 'toy');
- }
- parseOperation(parser, result) {
- if (result.op === 'toy.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- // toy.constant: {attrs} dense<...> : type
- if (result.op === 'toy.constant') {
- parser.parseOptionalAttrDict(result.attributes);
- const value = parser.parseAttribute();
- result.addAttribute('value', value.value === undefined ? value : value.value);
- result.addTypes([value.type]);
- return true;
- }
- // toy.mul, toy.add: %lhs, %rhs : type
- if (result.op === 'toy.mul' || result.op === 'toy.add') {
- result.operands = parser.parseOperandList();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- const types = parser.parseOptionalColonTypeList();
- if (types.length > 0) {
- const [type] = types;
- for (const operand of result.operands) {
- operand.type = type;
- }
- result.addTypes(types);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.SdfgDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'sdfg');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (typeName === 'stream' && parser.getToken().is('_')) {
- parser.consumeToken('_');
- const suffix = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (suffix === 'array') {
- type += `_${suffix}`;
- }
- }
- if (typeName === 'array' || typeName === 'stream' || typeName === 'memlet' || type.endsWith('stream_array')) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op === 'sdfg.sdfg' || result.op === 'sdfg.nested_sdfg' || result.op === 'sdir.sdfg') {
- parser.parseOptionalAttrDict(result.attributes);
- const inputResult = parser.parseFunctionArgumentList();
- const inputs = inputResult.arguments.map((a) => a.type);
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let results = [];
- if (parser.parseOptionalArrow()) {
- const outputResult = parser.parseFunctionArgumentList();
- results = outputResult.arguments.map((a) => a.type);
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'sdfg.tasklet' || result.op === 'sdir.tasklet') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- const blockArgs = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const operand = parser.parseOperand();
- let blockArgName = operand;
- let type = null;
- if (parser.parseOptionalKeyword('as')) {
- blockArgName = parser.parseOperand();
- parser.parseColon();
- type = parser.parseType();
- } else {
- parser.parseColon();
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- blockArgs.push({ value: blockArgName, type });
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const type = parser.parseType();
- result.addTypes([type]);
- parser.parseOptionalComma();
- }
- } else {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, blockArgs);
- }
- return true;
- }
- if (result.op === 'sdfg.consume') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- // Parse named results: (pe: %p, elem: %e)
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.bare_identifier)) {
- parser.consumeToken(_.Token.bare_identifier); // name like 'pe' or 'elem'
- if (parser.parseOptionalColon()) {
- parser.parseOperand(); // Parse %p or %e but don't store
- result.types.push(null);
- }
- } else if (parser.getToken().is(_.Token.percent_identifier) || parser.getToken().is(_.Token.r_paren)) {
- break;
- } else {
- throw new mlir.Error(`Expected named result in sdfg.consume but got '${parser.getTokenSpelling().str()}' ${parser.location()}`);
- }
- parser.parseOptionalComma();
- }
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'sdfg.state' || result.op === 'sdir.state') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.getToken().is(_.Token.at_identifier)) {
- parser.parseSymbolName('sym_name', result.attributes);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- const region = result.addRegion();
- parser.parseRegion(region);
- return true;
- }
- if (result.op === 'sdfg.alloc' || result.op === 'sdir.alloc' || result.op === 'sdir.alloc_transient' || result.op === 'sdir.alloc_stream') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedOperands = [];
- if (parser.getToken().is(_.Token.l_paren)) {
- unresolvedOperands.push(...parser.parseOperandList('paren'));
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let allocType = null;
- if (parser.parseOptionalColon()) {
- allocType = parser.parseType();
- result.addTypes([allocType]);
- }
- const indexType = new _.IndexType();
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => indexType), result.operands);
- return true;
- }
- if (result.op === 'sdfg.store' || result.op === 'sdir.store') {
- parser.parseOptionalAttrDict(result.attributes);
- const valueOp = parser.parseOperand();
- parser.parseOptionalComma();
- const arrayOp = parser.parseOperand();
- const indices = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- indices.push(parser.parseOperand());
- } else {
- parser.consumeToken();
- }
- if (parser.getToken().is(_.Token.comma)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRSquare();
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.parseOptionalColon()) {
- const valueType = parser.parseType();
- parser.parseOptionalArrow();
- const arrayType = parser.parseType();
- parser.resolveOperand(valueOp, valueType, result.operands);
- parser.resolveOperand(arrayOp, arrayType, result.operands);
- const indexType = new _.IndexType();
- parser.resolveOperands(indices, indices.map(() => indexType), result.operands);
- } else {
- parser.resolveOperand(valueOp, null, result.operands);
- parser.resolveOperand(arrayOp, null, result.operands);
- parser.resolveOperands(indices, indices.map(() => null), result.operands);
- }
- return true;
- }
- if (result.op === 'sdfg.load' || result.op === 'sdir.load') {
- parser.parseOptionalAttrDict(result.attributes);
- const arrayOp = parser.parseOperand();
- const indices = [];
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- indices.push(parser.parseOperand());
- } else {
- parser.consumeToken();
- }
- if (parser.getToken().is(_.Token.comma)) {
- parser.parseOptionalComma();
- }
- }
- parser.parseOptionalRSquare();
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.parseOptionalColon()) {
- const arrayType = parser.parseType();
- parser.parseOptionalArrow();
- const resultType = parser.parseType();
- parser.resolveOperand(arrayOp, arrayType, result.operands);
- const indexType = new _.IndexType();
- parser.resolveOperands(indices, indices.map(() => indexType), result.operands);
- result.addTypes([resultType]);
- } else {
- parser.resolveOperand(arrayOp, null, result.operands);
- parser.resolveOperands(indices, indices.map(() => null), result.operands);
- }
- return true;
- }
- if (result.op === 'sdfg.map' || result.op === 'sdir.map') {
- parser.parseOptionalAttrDict(result.attributes);
- const params = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const param = parser.parseOperand();
- params.push(param);
- }
- if (parser.getToken().is(_.Token.comma)) {
- parser.parseOptionalComma();
- }
- }
- }
- if (parser.parseOptionalEqual()) {
- parser.skip('(');
- }
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'to') {
- parser.parseOptionalKeyword('to');
- parser.skip('(');
- }
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'step') {
- parser.parseOptionalKeyword('step');
- parser.skip('(');
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'sdir.consume') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- parser.consumeToken();
- }
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'sdfg.edge' || result.op === 'sdir.edge') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.percent_identifier)) {
- parser.consumeToken(_.Token.bare_identifier); // label like 'ref'
- parser.parseColon();
- }
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(operand, type, result.operands);
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.at_identifier)) {
- const src = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('src', src);
- }
- parser.parseOptionalArrow();
- if (parser.getToken().is(_.Token.at_identifier)) {
- const dst = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('dst', dst);
- }
- return true;
- }
- if (result.op === 'sdfg.sym' || result.op === 'sdir.sym') {
- if (parser.parseOptionalLParen()) {
- const expr = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('expr', expr);
- parser.parseOptionalRParen();
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- result.addTypes([type]);
- }
- return true;
- }
- if (result.op === 'sdfg.copy' || result.op === 'sdir.copy') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedSrc = parser.parseOperandList('none');
- const unresolvedDst = [];
- if (parser.parseOptionalArrow()) {
- unresolvedDst.push(...parser.parseOperandList('none'));
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- const allUnresolved = unresolvedSrc.concat(unresolvedDst);
- parser.resolveOperands(allUnresolved, allUnresolved.map(() => type), result.operands);
- } else {
- const allUnresolved = unresolvedSrc.concat(unresolvedDst);
- parser.resolveOperands(allUnresolved, allUnresolved.map(() => null), result.operands);
- }
- return true;
- }
- if (result.op === 'sdfg.libcall' || result.op === 'sdir.libcall') {
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.getToken().is(_.Token.string)) {
- const libname = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('libname', libname);
- }
- const unresolvedOperands = parser.parseOperandList('paren');
- const types = [];
- if (parser.parseOptionalColon()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- types.push(parser.parseType());
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- }
- parser.resolveOperands(unresolvedOperands, types.length > 0 ? types : unresolvedOperands.map(() => null), result.operands);
- return true;
- }
- if (result.op === 'sdfg.get_access' || result.op === 'sdir.get_access') {
- let unresolvedOperand = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperand = parser.parseOperand();
- }
- if (parser.parseOptionalColon()) {
- const inputType = parser.parseType();
- if (unresolvedOperand) {
- parser.resolveOperand(unresolvedOperand, inputType, result.operands);
- }
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- } else if (unresolvedOperand) {
- // No explicit type - try to resolve operand and infer type from its definition
- parser.resolveOperand(unresolvedOperand, null, result.operands);
- // If operand was resolved with a type, use it as result type
- if (result.operands.length > 0 && result.operands[0].type) {
- result.addTypes([result.operands[0].type]);
- } else {
- // Fallback: use a generic memref type
- result.addTypes([new _.Type('memref<*xf32>')]);
- }
- }
- return true;
- }
- if (result.op === 'sdir.call') {
- const callee = parser.parseOptionalSymbolName();
- if (callee) {
- result.addAttribute('callee', callee);
- }
- if (parser.getToken().is(_.Token.l_paren)) {
- const unresolvedOperands = parser.parseOperandList('paren');
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- if (parser.parseOptionalColon()) {
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- parser.parseType();
- parser.parseOptionalComma();
- }
- }
- if (parser.parseOptionalArrow()) {
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- }
- return true;
- }
- if (result.op === 'sdfg.alloc_symbol' || result.op === 'sdir.alloc_symbol') {
- if (parser.parseOptionalLParen()) {
- const sym = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('sym', sym);
- parser.parseOptionalRParen();
- }
- return true;
- }
- if (result.op === 'sdfg.return') {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedOperands = parser.parseOperandList('none');
- const types = parser.parseOptionalColonTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- }
- return true;
- }
- if (result.op === 'sdfg.stream_push' || result.op === 'sdir.stream_push') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedValue = parser.parseOperand();
- parser.parseOptionalComma();
- const unresolvedStream = parser.parseOperand();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let valueType = null;
- let streamType = null;
- if (parser.parseOptionalColon()) {
- valueType = parser.parseType();
- parser.parseOptionalArrow();
- streamType = parser.parseType();
- }
- parser.resolveOperand(unresolvedValue, valueType, result.operands);
- parser.resolveOperand(unresolvedStream, streamType, result.operands);
- return true;
- }
- if (result.op === 'sdfg.stream_pop' || result.op === 'sdir.stream_pop') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedStream = parser.parseOperand();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let streamType = null;
- if (parser.parseOptionalColon()) {
- streamType = parser.parseType();
- parser.parseOptionalArrow();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- parser.resolveOperand(unresolvedStream, streamType, result.operands);
- return true;
- }
- if (result.op === 'sdfg.stream_length' || result.op === 'sdir.stream_length') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedStream = parser.parseOperand();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let streamType = null;
- if (parser.parseOptionalColon()) {
- streamType = parser.parseType();
- parser.parseOptionalArrow();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- parser.resolveOperand(unresolvedStream, streamType, result.operands);
- return true;
- }
- if (result.op === 'sdfg.view_cast' || result.op === 'sdir.view_cast') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedInput = parser.parseOperand();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let inputType = null;
- if (parser.parseOptionalColon()) {
- inputType = parser.parseType();
- parser.parseOptionalArrow();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- parser.resolveOperand(unresolvedInput, inputType, result.operands);
- return true;
- }
- if (result.op === 'sdfg.subview' || result.op === 'sdir.subview') {
- parser.parseOptionalAttrDict(result.attributes);
- const unresolvedInput = parser.parseOperand();
- while (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.consumeToken();
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- let inputType = null;
- if (parser.parseOptionalColon()) {
- inputType = parser.parseType();
- parser.parseOptionalArrow();
- const resultType = parser.parseType();
- result.addTypes([resultType]);
- }
- parser.resolveOperand(unresolvedInput, inputType, result.operands);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TFLDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfl');
- // Operations that use parseOneResultSameOperandTypeOp in tfl_ops.cc
- this._binaryOps = new Set([
- 'add', 'sub', 'mul', 'div', 'floor_div', 'pow', 'squared_difference',
- 'less', 'less_equal', 'greater', 'greater_equal', 'not_equal',
- 'logical_and', 'logical_or'
- ]);
- }
- parseOperation(parser, result) {
- const opKind = result.op.substring('tfl.'.length);
- if (opKind === 'control_node') {
- if (parser.parseOptionalLParen()) {
- const unresolvedOperands = parser.parseOperandList('none');
- parser.parseRParen();
- // control_node operands don't have types parsed inline
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- if (parser.parseOptionalKeyword('controls')) {
- const region = { blocks: [{ operations: [] }] };
- const innerOp = parser.parseGenericOperation();
- region.blocks[0].operations.push(innerOp);
- result.regions.push(region);
- // Result types are inferred from inner op results + control type
- for (const opResult of innerOp.results) {
- result.addTypes([opResult.type]);
- }
- result.addTypes([new _.Type('!tfl.control')]);
- } else if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- // Result types from yield terminator + control type
- const block = region.blocks[region.blocks.length - 1];
- const yieldOp = block.operations[block.operations.length - 1];
- if (yieldOp && yieldOp.operands) {
- for (const operand of yieldOp.operands) {
- result.addTypes([operand.type]);
- }
- }
- result.addTypes([new _.Type('!tfl.control')]);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (this._binaryOps.has(opKind)) {
- // Or: (operands) <properties> : fn-type (generic form)
- if (parser.getToken().is(_.Token.l_paren)) {
- parser.parseLParen();
- const unresolvedOperands = parser.parseOperandList('none');
- parser.parseRParen();
- if (parser.parseOptionalLess()) {
- result.propertiesAttr = parser.parseAttribute();
- parser.parseGreater();
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const fnType = parser.parseType();
- if (fnType instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedOperands, fnType.inputs, result.operands);
- result.addTypes(fnType.results);
- } else {
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => fnType), result.operands);
- result.addTypes([fnType]);
- }
- } else {
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- return true;
- }
- // Compact form: %a, %b attr-dict : type
- const unresolvedOperands = parser.parseOperandList('none');
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => type), result.operands);
- result.addTypes([type]);
- } else {
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TFDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tf');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (typeName === 'resource' || typeName === 'variant') {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- if (typeName === 'string' || typeName === 'control') {
- return new _.Type(type);
- }
- return null;
- }
- };
- _.TFTypeDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tf_type');
- this.simpleTypes = new Set([
- 'string', 'qint8', 'qint16', 'qint32', 'quint8', 'quint16',
- 'f32ref', 'f64ref', 'uint4ref', 'int4ref', 'uint8ref', 'int8ref',
- 'uint16ref', 'int16ref', 'uint32ref', 'int32ref', 'uint64ref', 'int64ref',
- 'stringref', 'boolref', 'quint8ref', 'qint8ref', 'quint16ref', 'qint16ref',
- 'qint32ref', 'bfloat16ref', 'complex64ref', 'complex128ref', 'halfref',
- 'resourceref', 'variantref',
- 'float8e4m3fnref', 'float8e5m2ref', 'float8e4m3fnuzref',
- 'float8e4m3b11fnuzref', 'float8e5m2fnuzref'
- ]);
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- // Handle parametrized types like resource<>, variant<>, resource_handle<>
- if (typeName === 'resource' || typeName === 'variant' || typeName === 'resource_handle') {
- if (parser.parseOptionalLess()) {
- const subtypes = [];
- while (parser.getToken().isNot(_.Token.greater)) {
- subtypes.push(parser.parseType());
- parser.parseOptionalComma();
- }
- parser.parseGreater();
- return new _.Type(`${type}<${subtypes.join(', ')}>`);
- }
- return new _.Type(type);
- }
- if (this.simpleTypes.has(typeName)) {
- return new _.Type(type);
- }
- // Fallback for unknown tf_type types
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- };
- _.TransformDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'transform');
- this.registerCustomDirective('PackedOrDynamicIndexList', this.parsePackedOrDynamicIndexList.bind(this));
- this.registerCustomDirective('SemiFunctionType', this.parseSemiFunctionType.bind(this));
- this.registerCustomDirective('SequenceOpOperands', this.parseSequenceOpOperands.bind(this));
- this.registerCustomDirective('ForeachMatchSymbols', this.parseForeachMatchSymbols.bind(this));
- this.registerCustomDirective('TransformMatchDims', this.parseTransformMatchDims.bind(this));
- this.registerCustomDirective('ApplyRegisteredPassOptions', this.parseApplyRegisteredPassOptions.bind(this));
- this.registerCustomDirective('AlternativesOpSelectedRegion', this.parseAlternativesOpSelectedRegion.bind(this));
- this.registerCustomDirective('ContinuousTileSizeTypes', this.parseContinuousTileSizeTypes.bind(this));
- this.registerCustomDirective('MultitileSizesTypes', this.parseMultitileSizesTypes.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'transform.named_sequence') {
- return this.parseNamedSequenceOp(parser, result);
- }
- // C++-only operation: transform.test_transform_op ["message"]
- // Defined in mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp
- if (result.op === 'transform.test_transform_op') {
- if (parser.getToken().is(_.Token.string)) {
- const message = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('message', message);
- }
- return true;
- }
- // LinalgTransformOps.cpp:3009 SplitOp::parse
- if (result.op === 'transform.structured.split') {
- const unresolvedTarget = parser.parseOperand();
- parser.parseKeyword('after');
- let unresolvedDynamicChunk = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedDynamicChunk = parser.parseOperand();
- } else {
- const staticChunkSizes = parser.parseInteger();
- result.addAttribute('static_chunk_sizes', staticChunkSizes);
- }
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const targetType = parser.parseType();
- parser.resolveOperand(unresolvedTarget, targetType, result.operands);
- result.addTypes([targetType]);
- if (unresolvedDynamicChunk && parser.parseOptionalComma()) {
- const chunkType = parser.parseType();
- parser.resolveOperand(unresolvedDynamicChunk, chunkType, result.operands);
- } else if (unresolvedDynamicChunk) {
- // Default to index type if chunk type not specified
- parser.resolveOperand(unresolvedDynamicChunk, null, result.operands);
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseSequenceOpOperands(parser, op /*, args */) {
- const unresolvedOperands = [];
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (parser.parseOptionalComma()) {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- }
- }
- if (parser.parseOptionalColon()) {
- parser.parseOptionalLParen();
- const types = parser.parseTypeListNoParens();
- parser.resolveOperands(unresolvedOperands, types, op.operands);
- parser.parseOptionalRParen();
- } else {
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, op.operands);
- }
- }
- }
- parseForeachMatchSymbols(parser, op, matchersAttr, actionsAttr) {
- const matchers = [];
- const actions = [];
- do {
- const matcher = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- parser.parseArrow();
- const action = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- matchers.push(matcher);
- actions.push(action);
- } while (parser.parseOptionalComma());
- op.addAttribute(matchersAttr, matchers);
- op.addAttribute(actionsAttr, actions);
- }
- parseTransformMatchDims(parser, op, dimsAttr, invertedAttr, allAttr) {
- if (parser.parseOptionalKeyword('all')) {
- op.addAttribute(allAttr, true);
- return;
- }
- const isInverted = parser.parseOptionalKeyword('except');
- if (isInverted) {
- parser.parseLParen();
- }
- const dims = [];
- do {
- const value = parser.parseOptionalInteger();
- if (value !== null) {
- dims.push(value);
- }
- } while (parser.parseOptionalComma());
- if (isInverted) {
- parser.parseRParen();
- op.addAttribute(invertedAttr, true);
- }
- op.addAttribute(dimsAttr, dims);
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (typeName === 'any' && parser.getToken().is('_')) {
- parser.consumeToken('_');
- const suffix = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- type += `_${suffix}`;
- }
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- parseNamedSequenceOp(parser, result) {
- parser.parseOptionalVisibilityKeyword(result.attributes);
- parser.parseSymbolName('sym_name', result.attributes);
- const argResult = parser.parseFunctionArgumentList();
- const inputs = argResult.arguments.map((a) => a.type);
- const results = [];
- const resultAttrs = [];
- if (parser.parseOptionalArrow()) {
- parser.parseFunctionResultList(results, resultAttrs);
- }
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType(inputs, results)));
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- parseSemiFunctionType(parser, op /* , args */) {
- const hasLParen = parser.parseOptionalLParen();
- const argType = parser.parseType();
- if (op.operands.length > 0) {
- op.operands[0].type = argType;
- }
- if (!hasLParen) {
- return;
- }
- parser.parseRParen();
- parser.parseArrow();
- if (parser.parseOptionalLParen()) {
- let idx = 0;
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const type = parser.parseType();
- if (idx < op.types.length) {
- op.types[idx] = type;
- } else {
- op.addTypes([type]);
- }
- idx++;
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- } else {
- const type = parser.parseType();
- if (op.types.length > 0) {
- op.types[0] = type;
- } else {
- op.addTypes([type]);
- }
- }
- }
- parsePackedOrDynamicIndexList(parser, op, packedName, dynamicName, staticAttrName) {
- const dynamicOperands = [];
- const dynamicTypes = [];
- const staticValues = [];
- let packedOperand = null;
- // Check for packed syntax: *(%operand)
- if (parser.consumeIf(_.Token.star)) {
- parser.parseLParen();
- if (parser.getToken().is(_.Token.percent_identifier)) {
- packedOperand = parser.parseOperand();
- }
- parser.parseRParen();
- } else if (parser.parseOptionalLSquare()) {
- // List syntax: [int, %operand, int, ...]
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const value = parser.parseOperand();
- dynamicOperands.push(value);
- staticValues.push(-9223372036854775808); // ShapedType::kDynamic
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- dynamicTypes.push(type);
- } else if (parser.getToken().is(_.Token.integer) || parser.getToken().is('number')) {
- const intVal = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken();
- staticValues.push(intVal);
- } else {
- break;
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- }
- if (packedOperand && packedName) {
- parser.resolveOperand(packedOperand, null, op.operands);
- }
- if (dynamicName) {
- for (let i = 0; i < dynamicOperands.length; i++) {
- const type = i < dynamicTypes.length ? dynamicTypes[i] : null;
- parser.resolveOperand(dynamicOperands[i], type, op.operands);
- }
- }
- if (staticAttrName && staticValues.length > 0) {
- op.addAttribute(staticAttrName, staticValues);
- }
- }
- parseContinuousTileSizeTypes(parser, result, targetTypes, tileSizesTypes, chunkSizesTypes) {
- const funcType = parser.parseType();
- if (funcType instanceof _.FunctionType) {
- if (funcType.inputs.length !== 1 || funcType.results.length !== 1) {
- throw new mlir.Error(`Expected a trailing functional type with one argument and one result.`);
- }
- targetTypes.push(funcType.inputs[0]);
- tileSizesTypes.push(funcType.results[0]);
- chunkSizesTypes.push(funcType.results[0]);
- }
- }
- parseMultitileSizesTypes(parser, result, targetTypes, lowSizeTypes, highSizeTypes, splitPointTypes) {
- const funcType = parser.parseType();
- if (funcType instanceof _.FunctionType) {
- if (funcType.inputs.length !== 1 || funcType.results.length !== 1) {
- throw new mlir.Error(`Expected a trailing functional type with one argument and one result.`);
- }
- targetTypes.push(funcType.inputs[0]);
- lowSizeTypes.push(funcType.results[0]);
- highSizeTypes.push(funcType.results[0]);
- splitPointTypes.push(funcType.results[0]);
- }
- }
- parseApplyRegisteredPassOptions(parser, result) {
- if (!parser.parseOptionalLBrace()) {
- return;
- }
- const options = {};
- while (parser.getToken().isNot(_.Token.r_brace)) {
- let key = null;
- if (parser.getToken().is(_.Token.string)) {
- key = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else {
- key = parser.parseOptionalKeyword();
- }
- parser.parseEqual();
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- options[key] = `#transform.param_operand<${result.operands.length - 1}>`;
- } else if (parser.getToken().is(_.Token.l_square)) {
- parser.parseOptionalLSquare();
- const arr = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- arr.push(`#transform.param_operand<${result.operands.length - 1}>`);
- } else {
- const val = parser.parseAttribute();
- arr.push(val);
- }
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- options[key] = arr;
- } else {
- const value = parser.parseAttribute();
- options[key] = value;
- }
- parser.parseOptionalComma();
- }
- parser.parseRBrace();
- result.addAttribute('options', options);
- }
- parseAlternativesOpSelectedRegion(parser, result) {
- if (parser.getToken().is(_.Token.integer)) {
- const value = parser.parseInteger();
- result.addAttribute('selected_region_attr', value);
- } else if (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- };
- _.TestDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'test');
- this.blobManager = new Map();
- this.registerCustomDirective('CustomOptionalOperand', this.parseCustomOptionalOperand.bind(this));
- this.registerCustomDirective('CustomDirectiveOperands', this.parseCustomDirectiveOperands.bind(this));
- this.registerCustomDirective('CustomDirectiveOperandsAndTypes', this.parseCustomDirectiveOperandsAndTypes.bind(this));
- this.registerCustomDirective('CustomDirectiveResults', this.parseCustomDirectiveResults.bind(this));
- this.registerCustomDirective('CustomDirectiveWithTypeRefs', this.parseCustomDirectiveWithTypeRefs.bind(this));
- this.registerCustomDirective('CustomDirectiveRegions', this.parseCustomDirectiveRegions.bind(this));
- this.registerCustomDirective('CustomDirectiveSuccessors', this.parseCustomDirectiveSuccessors.bind(this));
- this.registerCustomDirective('CustomDirectiveAttrDict', this.parseCustomDirectiveAttrDict.bind(this));
- this.registerCustomDirective('CustomDirectiveAttributes', this.parseCustomDirectiveAttributes.bind(this));
- this.registerCustomDirective('CustomDirectiveSpacing', this.parseCustomDirectiveSpacing.bind(this));
- this.registerCustomDirective('CustomDirectiveOptionalOperandRef', this.parseCustomDirectiveOptionalOperandRef.bind(this));
- this.registerCustomDirective('UsingPropertyInCustom', this.parseUsingPropertyInCustom.bind(this));
- this.registerCustomDirective('IntProperty', this.parseIntProperty.bind(this));
- this.registerCustomDirective('SumProperty', this.parseSumProperty.bind(this));
- this.registerCustomDirective('SwitchCases', this.parseSwitchCases.bind(this));
- this.registerCustomDirective('DimensionList', this.parseDimensionList.bind(this));
- this.registerCustomDirective('OptionalCustomParser', this.parseOptionalCustomParser.bind(this));
- this.registerCustomDirective('OptionalLoc', this.parseOptionalLoc.bind(this));
- this.registerCustomDirective('DummyRegionRef', this.parseDummyRegionRef.bind(this));
- this.registerCustomDirective('DummySuccessorRef', this.parseDummySuccessorRef.bind(this));
- this.registerCustomType('CompoundNestedOuterType', this.parseCompoundNestedOuterType.bind(this));
- this.registerCustomType('CompoundNestedInnerType', this.parseCompoundNestedInnerType.bind(this));
- this.registerCustomType('CompoundTypeA', this.parseCompoundTypeA.bind(this));
- this.registerCustomAttribute('TestBitEnumAttr', this.parseEnumFlagsAngleBracketComma.bind(this));
- this.registerCustomAttribute('TestBitEnumVerticalBarAttr', this.parseEnumFlagsAngleBracketPipe.bind(this));
- this.registerCustomAttribute('TestEnumAttr', this.parseTestEnumAttr.bind(this));
- this.registerCustomAttribute('TestEnumProp', this.parseTestEnumAttr.bind(this));
- this.registerCustomAttribute('TestEnumPropAttrForm', this.parseTestEnumPropAttrForm.bind(this));
- this.registerCustomAttribute('TestBitEnumProp', this.parseTestBitEnumProp.bind(this));
- this.registerCustomAttribute('TestBitEnumPropNamed', this.parseTestBitEnumPropNamed.bind(this));
- this.registerCustomAttribute('TestArrayOfUglyAttrs', this.parseTestArrayOfUglyAttrs.bind(this));
- this.registerCustomAttribute('TestArrayOfInts', this.parseTestArrayOfInts.bind(this));
- this.registerCustomAttribute('TestArrayOfEnums', this.parseTestArrayOfEnums.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'test.conversion_func_op') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- if (result.op === 'test.region_if') {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseColon();
- const inputTypes = parser.parseTypeList();
- parser.resolveOperands(unresolvedOperands, inputTypes, result.operands);
- parser.parseArrow();
- const outputTypes = parser.parseFunctionResultTypes();
- for (const t of outputTypes) {
- result.addTypes([t]);
- }
- parser.parseKeyword('then');
- const thenRegion = {};
- parser.parseRegion(thenRegion);
- result.regions.push(thenRegion);
- parser.parseKeyword('else');
- const elseRegion = {};
- parser.parseRegion(elseRegion);
- result.regions.push(elseRegion);
- parser.parseKeyword('join');
- const joinRegion = {};
- parser.parseRegion(joinRegion);
- result.regions.push(joinRegion);
- return true;
- }
- if (result.op === 'test.affine_scope' || result.op === 'test.single_no_terminator_custom_asm_op') {
- const region = result.addRegion();
- parser.parseRegion(region);
- return true;
- }
- if (result.op === 'test.with_nice_properties') {
- result.compatibility = true;
- let label = null;
- if (parser.getToken().is(_.Token.string)) {
- label = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else {
- label = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- }
- parser.parseKeyword('is');
- const negative = parser.consumeIf(_.Token.minus);
- const value = parser.parseInteger();
- result.addAttribute('prop', { label, value: negative ? -value : value });
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- // Test operation with default-valued properties and UnitProp
- if (result.op === 'test.with_default_valued_properties') {
- result.compatibility = true;
- if (parser.parseOptionalKeyword('na')) {
- // All defaults
- } else {
- const a = parser.parseInteger();
- result.addAttribute('a', a);
- if (parser.getToken().is(_.Token.string)) {
- const __spelling = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('b', __spelling);
- }
- if (parser.getToken().is(_.Token.integer) || parser.getToken().is(_.Token.minus)) {
- const neg = parser.consumeIf(_.Token.minus);
- const c = parser.parseInteger();
- result.addAttribute('c', neg ? -c : c);
- }
- if (parser.consumeIf(_.Token.kw_unit)) {
- result.addAttribute('unit', true);
- } else if (parser.parseOptionalKeyword('unit_absent')) {
- result.addAttribute('unit', false);
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- // Test operation with optional properties using some<...> syntax
- if (result.op === 'test.with_optional_properties') {
- result.compatibility = true;
- const parseOptionalValue = () => {
- if (parser.parseOptionalKeyword('some')) {
- parser.parseLess();
- let value = null;
- if (parser.parseOptionalKeyword('none')) {
- value = null;
- } else if (parser.consumeIf(_.Token.kw_unit)) {
- value = true;
- } else if (parser.getToken().is(_.Token.string)) {
- value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else {
- const neg = parser.consumeIf(_.Token.minus);
- value = parser.parseInteger();
- if (neg) {
- value = -value;
- }
- }
- parser.parseGreater();
- return { some: value };
- }
- if (parser.getToken().is(_.Token.string)) {
- const __retval = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- return __retval;
- }
- const neg = parser.consumeIf(_.Token.minus);
- const value = parser.parseInteger();
- return neg ? -value : value;
- };
- const knownAttrs = new Set(['anAttr', 'simple', 'simplei8', 'simpleui8', 'nonTrivialStorage', 'hasDefault', 'nested', 'longSyntax', 'hasUnit', 'maybeUnit']);
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() !== 'loc') {
- const tokenValue = parser.getTokenSpelling().str();
- // Stop if this looks like an operation name (not a known attribute)
- if (!knownAttrs.has(tokenValue)) {
- break;
- }
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- if (name === 'hasUnit') {
- result.addAttribute(name, true);
- } else if (parser.parseOptionalEqual()) {
- result.addAttribute(name, parseOptionalValue());
- } else {
- break;
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- if (result.op === 'test.wrapping_region') {
- parser.parseKeyword('wraps');
- const region = result.addRegion();
- const block = { operations: [] };
- region.blocks = [block];
- const wrappedOp = parser.parseGenericOperation();
- block.operations.push(wrappedOp);
- if (wrappedOp.results) {
- for (const opResult of wrappedOp.results) {
- result.addTypes([opResult.type]);
- }
- }
- return true;
- }
- if (result.op === 'test.pretty_printed_region') {
- result.operands = parser.parseOperandList();
- if (parser.parseOptionalKeyword('start')) {
- const innerOpName = parser.parseCustomOperationName();
- result.addAttribute('inner_op', innerOpName);
- parser.parseKeyword('end');
- parser.parseColon();
- const fnType = parser.parseFunctionType();
- if (fnType && fnType.inputs) {
- parser.resolveOperands(result.operands, fnType.inputs);
- }
- if (fnType && fnType.results) {
- for (let i = 0; i < fnType.results.length; i++) {
- result.addTypes([fnType.results[i]]);
- }
- }
- parser.parseOptionalLocationSpecifier();
- } else {
- parser.parseLParen();
- const region = result.addRegion();
- parser.parseRegion(region);
- parser.parseRParen();
- parser.parseColon();
- const fnType = parser.parseFunctionType();
- if (fnType && fnType.inputs) {
- parser.resolveOperands(result.operands, fnType.inputs);
- }
- if (fnType && fnType.results) {
- for (let i = 0; i < fnType.results.length; i++) {
- result.addTypes([fnType.results[i]]);
- }
- }
- }
- return true;
- }
- if (result.op === 'test.isolated_region') {
- const operand = parser.parseOperand();
- const indexType = new _.IndexType();
- parser.resolveOperand(operand, indexType, result.operands);
- const region = result.addRegion();
- parser.parseRegion(region, [{ value: operand.toString(), type: new _.IndexType() }]);
- return true;
- }
- if (result.op === 'test.string_attr_pretty_name') {
- const numResults = parser.getNumResults();
- for (let i = 0; i < numResults; i++) {
- result.addTypes([new _.IntegerType('i32')]);
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (!result.attributes.has('names') && numResults > 0) {
- const names = [];
- for (let i = 0; i < numResults; i++) {
- const resultName = parser.getResultName(i);
- // Only use names that don't start with a digit (ref impl line 572)
- if (resultName && resultName.length > 0 && !/^\d/.test(resultName)) {
- names.push(new _.StringAttr(resultName));
- } else {
- names.push(new _.StringAttr(''));
- }
- }
- result.attributes.set('names', new _.ArrayAttr(names));
- }
- return true;
- }
- if (result.op === 'test.with_bounds_region') {
- parser.parseOptionalAttrDict(result.attributes);
- const argName = parser.parseOperand();
- parser.parseColon();
- const argType = parser.parseType();
- const region = result.addRegion();
- const arg = { value: argName, type: argType };
- parser.parseRegion(region, [arg]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseTestBitEnumProp(parser, type) {
- if (type.values.includes(parser.getTokenSpelling().str())) {
- return this.parseEnumFlags(parser, type, ',');
- }
- return null;
- }
- parseTestEnumAttr(parser, type) {
- const token = parser.getToken();
- const spelling = token.getSpelling().str();
- if (token && type.values && type.values.includes(spelling)) {
- parser.consumeToken();
- return new _.TypedAttr(spelling, null);
- }
- return null;
- }
- parseTestEnumPropAttrForm(parser) {
- return parser.parseOptionalAttribute();
- }
- parseTestBitEnumPropNamed(parser) {
- if (parser.parseOptionalKeyword('bit_enum')) {
- if (parser.parseOptionalLess()) {
- const flags = [];
- while (parser.getToken().isNot(_.Token.greater)) {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- flags.push(value);
- parser.parseOptionalComma();
- }
- parser.parseGreater();
- return new _.TypedAttr(`bit_enum<${flags.join(', ')}>`, null);
- }
- }
- return null;
- }
- // Parse CompoundNestedOuterType: assemblyFormat = "`<` `i` $inner `>`"
- // Full form: !test.cmpnd_nested_outer<i !test.cmpnd_inner<...>>
- // Elided form: <i <...>>
- parseCompoundNestedOuterType(parser) {
- parser.parseLess();
- parser.parseKeyword('i');
- // Parse $inner - could be full (!test.cmpnd_inner<...>) or elided (<...>)
- const inner = parser.getToken().is(_.Token.exclamation_identifier) ? parser.parseType() : this.parseCompoundNestedInnerType(parser);
- parser.parseGreater();
- return new _.Type(`!test.cmpnd_nested_outer<i ${inner}>`);
- }
- // Parse CompoundNestedInnerType: assemblyFormat = "`<` $some_int $cmpdA `>`"
- // Full form: !test.cmpnd_inner<42 !test.cmpnd_a<...>>
- // Elided form: <42 <...>>
- parseCompoundNestedInnerType(parser) {
- parser.parseLess();
- const someInt = parser.parseInteger();
- // Parse $cmpdA - could be full (!test.cmpnd_a<...>) or elided (<...>)
- const cmpdA = parser.getToken().is(_.Token.exclamation_identifier) ? parser.parseType() : this.parseCompoundTypeA(parser);
- parser.parseGreater();
- return new _.Type(`!test.cmpnd_inner<${someInt} ${cmpdA}>`);
- }
- // Parse CompoundTypeA: hasCustomAssemblyFormat = 1
- // Example: <1, !test.smpla, [5, 6]>
- parseCompoundTypeA(parser) {
- parser.parseLess();
- const width = parser.parseInteger();
- parser.parseComma();
- const oneType = parser.parseType();
- parser.parseComma();
- parser.parseLSquare();
- const arrayOfInts = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- arrayOfInts.push(parser.parseInteger());
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- parser.parseGreater();
- return new _.Type(`!test.cmpnd_a<${width}, ${oneType}, [${arrayOfInts.join(', ')}]>`);
- }
- parseOptionalLoc(parser, op, attrName = 'loc') {
- const loc = parser.parseOptionalLocationSpecifier();
- if (loc) {
- op.addAttribute(attrName, loc);
- } else {
- op.addAttribute(attrName, parser.location());
- }
- }
- parseDummyRegionRef() {
- }
- parseDummySuccessorRef() {
- }
- parseTestAttrUgly(parser) {
- parser.parseKeyword('begin');
- const attr = parser.parseAttribute();
- parser.parseKeyword('end');
- return attr;
- }
- // Parse TestArrayOfUglyAttrs: assemblyFormat = "`[` (`]`) : ($value^ ` ` `]`)?"
- parseTestArrayOfUglyAttrs(parser) {
- parser.parseLSquare();
- const elements = [];
- parser.parseCommaSeparatedListUntil(_.Token.r_square, () => {
- elements.push(this.parseTestAttrUgly(parser));
- });
- return new _.ArrayAttr(elements);
- }
- // Parse TestArrayOfInts: assemblyFormat = "`[` (`]`) : ($value^ `]`)?"
- parseTestArrayOfInts(parser) {
- parser.parseLSquare();
- const elements = [];
- parser.parseCommaSeparatedListUntil(_.Token.r_square, () => {
- const isNegative = parser.consumeIf(_.Token.minus);
- const value = parser.parseInteger();
- elements.push(new _.IntegerAttr(null, isNegative ? -value : value));
- });
- return new _.ArrayAttr(elements);
- }
- // Parse TestArrayOfEnums: elements are SimpleEnumAttr values (a, b, etc.)
- parseTestArrayOfEnums(parser) {
- parser.parseLSquare();
- const elements = [];
- parser.parseCommaSeparatedListUntil(_.Token.r_square, () => {
- if (parser.getToken().is(_.Token.string)) {
- elements.push(new _.TypedAttr(parser.parseString(), null));
- } else {
- const value = parser.parseKeyword();
- elements.push(new _.TypedAttr(value, null));
- }
- });
- return new _.ArrayAttr(elements);
- }
- parseOptionalCustomParser(parser, op, attrName = 'attr') {
- if (!parser.parseOptionalKeyword('foo')) {
- return null; // Optional group not taken
- }
- const attr = parser.parseAttribute();
- op.addAttribute(attrName, attr.value);
- return true;
- }
- parseDimensionList(parser, op, attrName = 'dimension_list') {
- const dims = [];
- if (parser.parseOptionalLSquare()) {
- parser.parseOptionalRSquare();
- op.addAttribute(attrName, []);
- return;
- }
- for (;;) {
- if (parser.consumeIf(_.Token.question)) {
- dims.push(-1);
- } else if (parser.getToken().is(_.Token.integer)) {
- dims.push(parser.parseInteger());
- } else {
- break;
- }
- const token = parser.getToken();
- const spelling = token.getSpelling().str();
- if (token && token.kind === _.Token.bare_identifier && spelling.startsWith('x')) {
- const rest = spelling.slice(1);
- if (rest === '') {
- parser.consumeToken();
- } else if (/^\d+$/.test(rest)) {
- parser.consumeToken();
- dims.push(parseInt(rest, 10));
- } else if (rest === '?') {
- parser.consumeToken();
- dims.push(-1);
- } else {
- break;
- }
- } else {
- break;
- }
- }
- op.addAttribute(attrName, dims);
- }
- parseCustomOptionalOperand(parser, result) {
- if (parser.parseOptionalLParen()) {
- const unresolvedOperand = parser.parseOperand();
- parser.resolveOperand(unresolvedOperand, null, result.operands);
- parser.parseRParen();
- }
- }
- // Custom directive: operand [, optOperand] -> (varOperands)
- parseCustomDirectiveOperands(parser, result) {
- const unresolvedRequired = parser.parseOperand();
- parser.resolveOperand(unresolvedRequired, null, result.operands);
- if (parser.parseOptionalComma()) {
- const unresolvedOptional = parser.parseOperand();
- parser.resolveOperand(unresolvedOptional, null, result.operands);
- }
- parser.parseArrow();
- parser.parseLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedVar = parser.parseOperand();
- parser.resolveOperand(unresolvedVar, null, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- // Custom directive: operands and types together
- parseCustomDirectiveOperandsAndTypes(parser, result) {
- this.parseCustomDirectiveOperands(parser, result);
- this.parseCustomDirectiveResults(parser, result);
- }
- // Custom directive: : type [, optType] -> (varTypes)
- parseCustomDirectiveResults(parser, result) {
- parser.parseColon();
- const type = parser.parseType();
- if (result.operands.length > 0) {
- result.operands[0].type = type;
- }
- if (parser.parseOptionalComma()) {
- const optType = parser.parseType();
- if (result.operands.length > 1) {
- result.operands[1].type = optType;
- }
- }
- parser.parseArrow();
- parser.parseLParen();
- let idx = 2; // Start after first two operands
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const varType = parser.parseType();
- if (result.operands.length > idx) {
- result.operands[idx].type = varType;
- }
- idx++;
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parseCustomDirectiveWithTypeRefs(parser, result) {
- // Parses: type_refs_capture : type [, type] -> (types)
- parser.parseKeyword('type_refs_capture');
- this.parseCustomDirectiveResults(parser, result);
- }
- parseCustomDirectiveRegions(parser, result) {
- const region = result.addRegion();
- parser.parseRegion(region);
- while (parser.parseOptionalComma()) {
- const varRegion = result.addRegion();
- parser.parseRegion(varRegion);
- }
- }
- parseCustomDirectiveSuccessors(parser, result) {
- if (!result.successors) {
- result.successors = [];
- }
- const successor = {};
- successor.label = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- result.successors.push(successor);
- while (parser.parseOptionalComma()) {
- const varSuccessor = {};
- varSuccessor.label = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.caret_identifier);
- result.successors.push(varSuccessor);
- }
- }
- parseCustomDirectiveAttrDict(parser, result) {
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- }
- parseCustomDirectiveAttributes(parser, result) {
- const attr = parser.parseAttribute();
- result.addAttribute('attr', attr);
- if (parser.parseOptionalComma()) {
- const optAttr = parser.parseAttribute();
- result.addAttribute('optAttr', optAttr);
- }
- }
- parseCustomDirectiveSpacing(parser, op, attrName) {
- if (attrName) {
- const name = attrName.name || attrName;
- const attr = parser.parseAttribute();
- op.addAttribute(name, attr);
- }
- }
- parseCustomDirectiveOptionalOperandRef(parser) {
- parser.parseInteger();
- }
- parseSwitchCases(parser, result) {
- const caseValues = [];
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === 'case') {
- parser.parseKeyword('case');
- const value = parser.parseInteger();
- caseValues.push(value);
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- result.addAttribute('cases', `array<i64: ${caseValues.join(', ')}>`);
- }
- parseUsingPropertyInCustom(parser, op, propArg) {
- const values = [];
- parser.parseLSquare();
- while (parser.getToken().isNot(_.Token.r_square)) {
- const value = parser.parseInteger();
- values.push(value);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- if (propArg) {
- const propName = typeof propArg === 'string' ? propArg : propArg.name;
- op.addAttribute(propName, `array<i64: ${values.join(', ')}>`);
- }
- }
- parseIntProperty(parser, op, propArg) {
- const value = parser.parseInteger();
- if (propArg) {
- const propName = typeof propArg === 'string' ? propArg : propArg.name;
- op.addAttribute(propName, value);
- }
- }
- parseSumProperty(parser, op, propArg) {
- const second = parser.parseInteger();
- parser.parseEqual();
- parser.parseInteger(); // sum value (validation skipped)
- if (propArg) {
- const propName = typeof propArg === 'string' ? propArg : propArg.name;
- op.addAttribute(propName, second);
- }
- }
- inferResultTypes(op, vars) {
- if (op.op === 'test.format_infer_type' || op.op === 'test.format_infer_type2') {
- op.addTypes([new _.IntegerType('i16')]);
- return;
- }
- if (op.op === 'test.format_types_match_attr') {
- const valueAttr = op.attributes.get('value');
- if (valueAttr && valueAttr.type) {
- op.addTypes([valueAttr.type]);
- return;
- }
- }
- if (op.op === 'test.format_infer_type_all_operands_and_types' ||
- op.op === 'test.format_infer_type_all_types_one_operand' ||
- op.op === 'test.format_infer_type_all_types_two_operands' ||
- op.op === 'test.format_infer_type_all_types' ||
- op.op === 'test.format_infer_type_variadic_operands') {
- for (const operand of op.operands) {
- if (operand.type) {
- op.addTypes([operand.type]);
- }
- }
- return;
- }
- if (op.op === 'test.format_infer_type_regions') {
- const region = op.regions[0];
- const block = region?.blocks?.[0];
- if (block?.arguments) {
- for (const arg of block.arguments) {
- if (arg.type) {
- op.addTypes([arg.type]);
- }
- }
- }
- return;
- }
- if (op.op === 'test.with_properties_and_inferred_type') {
- const lhsAttr = op.attributes.get('lhs');
- const rhsAttr = op.getAttr('rhs');
- const lhs = lhsAttr && lhsAttr.value !== undefined ? lhsAttr.value : (lhsAttr || 0);
- const rhs = rhsAttr !== undefined && rhsAttr !== null ? rhsAttr : 0;
- const width = (typeof lhs === 'number' ? lhs : 0) + (typeof rhs === 'number' ? rhs : 0);
- op.addTypes([new _.IntegerType(`i${width}`)]);
- return;
- }
- super.inferResultTypes(op, vars);
- }
- declareResource(key) {
- if (!this.blobManager.has(key)) {
- this.blobManager.set(key, new _.DenseResourceElementsHandle(key));
- }
- return this.blobManager.get(key);
- }
- getResourceKey(handle) {
- return handle.key;
- }
- parseResource(entry) {
- const blob = entry.parseAsBlob();
- this.blobManager.get(entry.key).blob = blob;
- }
- };
- _.triton = {};
- _.triton.PointerType = class extends _.Type {
- constructor(pointeeType, addressSpace) {
- super(null);
- this.pointeeType = pointeeType;
- this.addressSpace = addressSpace || 1;
- }
- static parse(parser) {
- parser.parseLess();
- const pointeeType = parser.parseType();
- let addressSpace = 1;
- if (parser.parseOptionalComma()) {
- addressSpace = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- }
- parser.parseGreater();
- return new _.triton.PointerType(pointeeType, addressSpace);
- }
- toString() {
- const pointeeStr = this.pointeeType?.toString ? this.pointeeType.toString() : this.pointeeType;
- if (this.addressSpace && this.addressSpace !== 1) {
- return `!tt.ptr<${pointeeStr}, ${this.addressSpace}>`;
- }
- return `!tt.ptr<${pointeeStr}>`;
- }
- };
- _.triton.TritonDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tt');
- this.registerCustomType('TT_Ptr', (parser) => _.triton.PointerType.parse(parser));
- this.registerCustomType('TT_TensorDescType', this.parseTensorDescType.bind(this));
- this.registerCustomType('TT_TensorPtr', (parser) => _.triton.PointerType.parse(parser));
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- // Handle ptr type specifically to properly parse pointee type
- if (typeName === 'ptr' && parser.getToken().is(_.Token.less)) {
- return _.triton.PointerType.parse(parser);
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- parseOperation(parser, result) {
- if (result.op === 'tt.func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseTensorDescType(parser) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- return new _.Type(`!tt.tensor_desc${content}`);
- }
- return null;
- }
- inferResultTypes(op, vars) {
- if (op.op === 'tt.load' && op.operands.length > 0) {
- const ptrType = op.operands[0].type;
- if (ptrType instanceof _.triton.PointerType) {
- op.addTypes([ptrType.pointeeType]);
- } else if (ptrType instanceof _.RankedTensorType) {
- const elementType = ptrType.elementType;
- if (elementType instanceof _.triton.PointerType) {
- const pointeeType = elementType.pointeeType;
- op.addTypes([new _.RankedTensorType(ptrType.shape, pointeeType, ptrType.encoding)]);
- return;
- }
- }
- }
- if (op.op === 'tt.unsplat') {
- const srcType = op.operands[0]?.type;
- if (srcType instanceof _.RankedTensorType) {
- op.addTypes([srcType.elementType]);
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.triton.gpu = {};
- _.triton.gpu.TritonGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'ttg');
- this.registerCustomType('TTG_MemDescType', this.parseMemDescType.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'ttg.warp_specialize') {
- const unresolvedOperands = [];
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- unresolvedOperands.push(parser.parseOperand());
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseComma();
- }
- }
- parser.parseRParen();
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- parser.parseKeyword('default');
- const defaultRegion = {};
- parser.parseRegion(defaultRegion);
- result.regions.push(defaultRegion);
- const partitionNumWarps = [];
- let partitionIndex = 0;
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getTokenSpelling().str() === `partition${partitionIndex}`) {
- parser.consumeToken(_.Token.bare_identifier);
- const argResult = parser.parseFunctionArgumentList();
- parser.parseKeyword('num_warps');
- parser.parseLParen();
- const numWarps = parser.getToken().getSpelling().str();
- parser.consumeToken();
- partitionNumWarps.push(parseInt(numWarps, 10));
- parser.parseRParen();
- const partitionRegion = {};
- partitionRegion.arguments = argResult.arguments;
- parser.parseRegion(partitionRegion);
- if (!result.regions[1]) {
- result.regions[1] = { blocks: [{ operations: [] }] };
- }
- partitionIndex++;
- }
- parser.parseColon();
- const fnType = parser.parseType();
- if (fnType instanceof _.FunctionType) {
- result.addAttribute('function_type', new _.TypeAttrOf(fnType));
- parser.resolveOperands(unresolvedOperands, fnType.inputs, result.operands);
- result.addTypes(fnType.results);
- } else {
- result.addAttribute('function_type', new _.TypeAttrOf(new _.FunctionType([], [fnType])));
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- result.addTypes([fnType]);
- }
- if (partitionNumWarps.length > 0) {
- result.addAttribute('partitionNumWarps', { type: 'array', element_type: 'i32', value: partitionNumWarps });
- }
- return true;
- }
- if (result.op === 'ttg.barrier') {
- const flags = [];
- flags.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.bare_identifier);
- while (parser.parseOptionalVerticalBar()) {
- flags.push(parser.getToken().getSpelling().str());
- parser.consumeToken(_.Token.bare_identifier);
- }
- result.addAttribute('addrSpace', new _.TypedAttr(flags.join('|')));
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- parseMemDescType(parser) {
- // Handle shorthand MemDescType notation: <dims x elementType, attributes...>
- // Full notation would be: !ttg.memdesc<dims x elementType, attributes...>
- if (parser.getToken().isNot(_.Token.less)) {
- return null;
- }
- const content = parser.skip('<');
- return new _.Type(`!ttg.memdesc<${content}>`);
- }
- };
- _.GluonDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'gluon');
- }
- };
- _.TritonNvidiaGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'ttng');
- this.registerCustomDirective('Token', this.parseToken.bind(this));
- this.registerCustomDirective('BarriersAndPreds', this.parseBarriersAndPreds.bind(this));
- }
- parseToken(parser, op, depOperands, tokenTypeArr) {
- // custom<Token>($acc_dep, type($token))
- // depOperands = operands array, tokenTypeArr = types array
- if (!parser.parseOptionalLSquare()) {
- return;
- }
- // Push token type to the types array
- if (Array.isArray(tokenTypeArr)) {
- tokenTypeArr.push(new _.Type('!ttng.async.token'));
- }
- if (parser.getToken().is(_.Token.r_square)) {
- parser.parseRSquare();
- return;
- }
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const dep = parser.parseOperand();
- if (!Array.isArray(depOperands)) {
- throw new mlir.Error(`Expected depOperands to be an array ${parser.location()}`);
- }
- depOperands.push(dep);
- }
- parser.parseRSquare();
- }
- parseBarriersAndPreds(parser, op, barrierOperands, predOperands) {
- while (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const barrier = parser.parseOperand();
- if (!Array.isArray(barrierOperands)) {
- throw new mlir.Error(`Expected barrierOperands to be an array ${parser.location()}`);
- }
- barrierOperands.push(barrier);
- if (parser.parseOptionalLSquare()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const pred = parser.parseOperand();
- if (!Array.isArray(predOperands)) {
- throw new mlir.Error(`Expected predOperands to be an array ${parser.location()}`);
- }
- predOperands.push(pred);
- }
- parser.parseRSquare();
- }
- }
- }
- }
- parseType(parser, dialect) {
- const mnemonic = parser.parseOptionalKeyword();
- if (mnemonic) {
- let type = `!${dialect}.${mnemonic}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- };
- _.TritonAMDGPUDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'amdg');
- this.registerCustomType('TT_Ptr', 'tt.ptr');
- this.registerCustomType('TT_TensorPtr', 'tt.ptr');
- this.registerCustomType('TTG_MemDescType', 'ttg.memdesc');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- };
- _.ProtonDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'proton');
- }
- };
- _.MichelsonDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'michelson');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- if ((typeName === 'big' || typeName === 'chain' || typeName === 'key') && parser.getToken().is('_')) {
- parser.consumeToken('_');
- const suffix = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- type += `_${suffix}`;
- }
- const simpleTypes = ['int', 'bytes', 'operation', 'nat', 'string', 'unit', 'bool', 'mutez', 'timestamp', 'address', 'key', 'signature', 'chain_id', 'key_hash'];
- if (simpleTypes.includes(type.substring(11))) { // Remove "!michelson." prefix
- return new _.Type(type);
- }
- const typesWithParams = ['pair', 'list', 'option', 'or', 'map', 'big_map', 'set', 'contract', 'lambda'];
- if (typesWithParams.includes(type.substring(11))) {
- if (parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- return null;
- }
- };
- _.PlanDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'plan');
- this.registerCustomDirective('WithValuesTypes', this.parseWithValuesTypes.bind(this));
- }
- parseWithValuesTypes(parser, op, elementsRef, resultTypes, elementsTypes) {
- const resultType = parser.parseType();
- resultTypes.push(resultType);
- if (op.types.length === 0) {
- op.addTypes([resultType]);
- } else {
- op.types[0] = resultType;
- }
- // Derive element types from tensor's element type (matching C++ getElementTypeOrSelf)
- const elementType = resultType.elementType || resultType;
- // Populate element types for each element operand
- if (elementsRef && elementsRef.length > 0) {
- for (let i = 0; i < elementsRef.length; i++) {
- elementsTypes.push(elementType);
- }
- }
- }
- };
- _.KernelDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'kernel');
- this.registerCustomDirective('KernelFunctionalType', this.parseKernelFunctionalType.bind(this));
- }
- parseKernelFunctionalType(parser, op /*, args */) {
- parser.parseLParen();
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- parser.parseType();
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- parser.parseArrow();
- const resultTypes = [];
- if (parser.parseOptionalLParen()) {
- if (parser.getToken().isNot(_.Token.r_paren)) {
- do {
- resultTypes.push(parser.parseType());
- } while (parser.parseOptionalComma());
- }
- parser.parseRParen();
- } else {
- resultTypes.push(parser.parseType());
- }
- if (resultTypes.length > 0) {
- op.addTypes(resultTypes);
- }
- }
- inferResultTypes(op, vars) {
- // kernel.combiner: result types are the first half of the input types
- if (op.op === 'kernel.combiner') {
- const inputsEntry = vars.get('inputs');
- if (inputsEntry?.types?.length > 0) {
- const numPairs = Math.floor(inputsEntry.types.length / 2);
- const resultTypes = inputsEntry.types.slice(0, numPairs);
- op.addTypes(resultTypes);
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.TensorRTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tensorrt');
- this.registerCustomAttribute('TensorRT_TopKOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ScatterModeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ResizeSelectorAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ResizeRoundModeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ResizeModeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ResizeCoordinateTransformationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ReduceOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_PaddingModeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_MatrixOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_LoopOutputAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_GatherModeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_FillOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ElementWiseOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_ActivationTypeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_UnaryOperationAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_TripLimitAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomAttribute('TensorRT_PoolingTypeAttr', this.parseEnumAttrBracket.bind(this));
- this.registerCustomDirective('StaticIndexI64Array', this.parseStaticIndexI64Array.bind(this));
- this.registerCustomDirective('StaticIndexI32Array', this.parseStaticIndexI32Array.bind(this));
- }
- parseEnumAttrBracket(parser) {
- if (parser.getToken().is(_.Token.less)) {
- parser.parseLess();
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseGreater();
- return { value };
- }
- return null;
- }
- parseStaticIndexI64Array(parser, op, attrName = 'broadcast_dims') {
- const values = [];
- do {
- const value = parser.parseOptionalInteger();
- if (value === null) {
- break;
- }
- values.push(value);
- } while (parser.parseOptionalComma());
- op.addAttribute(attrName, values);
- }
- parseStaticIndexI32Array(parser, op, attrName = 'static_values') {
- const values = [];
- do {
- const value = parser.parseOptionalInteger();
- if (value === null) {
- break;
- }
- values.push(value);
- } while (parser.parseOptionalComma());
- op.addAttribute(attrName, values);
- }
- parseOperation(parser, result) {
- if (result.op === 'tensorrt.for') {
- return this.parseForOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseForOp(parser, result) {
- const inductionVar = parser.parseOperand();
- parser.parseEqual();
- const unresolvedLb = parser.parseOperand();
- parser.resolveOperand(unresolvedLb, null, result.operands);
- parser.parseKeyword('to');
- const unresolvedUb = parser.parseOperand();
- parser.resolveOperand(unresolvedUb, null, result.operands);
- parser.parseKeyword('step');
- const unresolvedStep = parser.parseOperand();
- parser.resolveOperand(unresolvedStep, null, result.operands);
- parser.parseKeyword('init');
- const regionArgs = [{ name: inductionVar.name, type: null }];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- let iterArgName = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const iterArg = parser.parseOperand();
- iterArgName = iterArg.name;
- }
- if (parser.parseOptionalEqual()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedInit = parser.parseOperand();
- parser.resolveOperand(unresolvedInit, null, result.operands);
- if (iterArgName) {
- regionArgs.push({ name: iterArgName, type: null });
- }
- }
- }
- parser.parseOptionalComma();
- }
- }
- result.addTypes(parser.parseOptionalArrowTypeList());
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, regionArgs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- };
- _.executor = {};
- _.executor.TableType = class extends _.Type {
- constructor(body) {
- super(null);
- this.body = body;
- }
- toString() {
- const bodyStr = this.body.map((t) => t.toString()).join(', ');
- return `!executor.table<${bodyStr}>`;
- }
- };
- _.executor.ExecutorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'executor');
- this.registerCustomType('Executor_Table', this.parseTable.bind(this));
- this.registerCustomDirective('ExecutorMixedIndices', this.parseExecutorMixedIndices.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'executor.func') {
- parser.parseFunctionOp(result, true);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- inferResultTypes(op, vars) {
- if (op.op === 'executor.table.get') {
- // ExtractTableValueOp::inferReturnTypes
- const tableType = op.operands[0]?.type;
- const index = op.attributes.get('index');
- if (tableType instanceof _.executor.TableType && index !== undefined) {
- const idx = typeof index === 'number' ? index : parseInt(index, 10);
- if (idx >= 0 && idx < tableType.body.length) {
- op.addTypes([tableType.body[idx]]);
- return;
- }
- }
- }
- super.inferResultTypes(op, vars);
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- if (typeName === 'table') {
- if (parser.getToken().is(_.Token.less)) {
- parser.parseLess();
- const body = [];
- if (parser.getToken().isNot(_.Token.greater)) {
- do {
- const elementType = parser.parseType();
- body.push(elementType);
- } while (parser.parseOptionalComma());
- }
- parser.parseGreater();
- return new _.executor.TableType(body);
- }
- return new _.executor.TableType([]);
- }
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- parseTable(parser) {
- if (parser.getToken().is(_.Token.less)) {
- parser.parseLess();
- const body = [];
- if (parser.getToken().isNot(_.Token.greater)) {
- do {
- const elementType = parser.parseType();
- body.push(elementType);
- } while (parser.parseOptionalComma());
- }
- parser.parseGreater();
- return new _.executor.TableType(body);
- }
- return null;
- }
- parseExecutorMixedIndices(parser, op /*, args */) {
- const unresolvedOperands = [];
- do {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- } else {
- parser.parseAttribute();
- }
- } while (parser.parseOptionalComma());
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, op.operands);
- }
- }
- };
- _.TFRTTestDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfrt_test');
- this.registerCustomDirective('OptionalLoc', this.parseOptionalLoc.bind(this));
- this.registerCustomDirective('DummyRegionRef', this.parseDummyRegionRef.bind(this));
- this.registerCustomDirective('DummySuccessorRef', this.parseDummySuccessorRef.bind(this));
- }
- parseOptionalLoc(parser, op, attrName = 'loc') {
- const loc = parser.parseOptionalLocationSpecifier();
- if (loc) {
- op.addAttribute(attrName, loc);
- } else {
- op.addAttribute(attrName, parser.location());
- }
- }
- parseDummyRegionRef() {
- }
- parseDummySuccessorRef() {
- }
- parseOperation(parser, result) {
- const opInfo = result.name.getRegisteredInfo();
- if (!opInfo) {
- return false;
- }
- if (opInfo.metadata.assemblyFormat === 'operands attr-dict') {
- const unresolvedOperands = [];
- while (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- for (const unresolved of unresolvedOperands) {
- parser.resolveOperand(unresolved, null, result.operands);
- }
- this.inferResultTypes(result, new Map());
- return true;
- }
- if (result.op === 'tfrt_test.do.async') {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolvedOperands = parser.parseOperandList('none');
- parser.resolveOperands(unresolvedOperands, unresolvedOperands.map(() => null), result.operands);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseFunctionType();
- if (type && type.results) {
- type.results.forEach((resultType) => {
- result.addTypes([resultType]);
- });
- }
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- return true;
- }
- if (result.op === 'tfrt_test.benchmark') {
- if (parser.getToken().is(_.Token.string)) {
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- result.addAttribute('name', name);
- }
- parser.parseLParen();
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolved = parser.parseOperand();
- let type = null;
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- parser.resolveOperand(unresolved, type, result.operands);
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- while (parser.getToken().is(_.Token.bare_identifier) && parser.getToken().isNot(_.Token.l_brace)) {
- const name = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseEqual();
- let value = null;
- if (parser.getToken().is(_.Token.integer)) {
- value = parser.parseInteger();
- } else if (parser.getToken().is(_.Token.string)) {
- value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.string);
- } else {
- value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- }
- result.addAttribute(name, value);
- parser.parseOptionalComma();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region);
- }
- result.addTypes([new _.Type('!tfrt.chain')]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.XeVMDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xevm');
- }
- };
- _.VMVXDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'vmvx');
- }
- };
- _.MLRTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'mlrt');
- }
- };
- _.TFRTTensorDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfrt_tensor');
- }
- };
- _.TFRTDHTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfrt_dht');
- }
- };
- _.TFDDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'tfd');
- }
- };
- _.ACCDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'acc');
- this.registerCustomDirective('Var', this.parseVar.bind(this));
- this.registerCustomDirective('AccVar', this.parseAccVar.bind(this));
- this.registerCustomDirective('VarPtrType', this.parseVarPtrType.bind(this));
- this.registerCustomDirective('DeviceTypeOperandsWithKeywordOnly', this.parseDeviceTypeOperandsWithKeywordOnly.bind(this));
- this.registerCustomDirective('LoopControl', this.parseLoopControl.bind(this));
- this.registerCustomDirective('WaitClause', this.parseWaitClause.bind(this));
- this.registerCustomDirective('NumGangs', this.parseNumGangs.bind(this));
- this.registerCustomDirective('DeviceTypeOperands', this.parseDeviceTypeOperands.bind(this));
- this.registerCustomDirective('GangClause', this.parseGangClause.bind(this));
- this.registerCustomDirective('CombinedConstructsLoop', this.parseCombinedConstructsLoop.bind(this));
- this.registerCustomDirective('RecipeSym', this.parseRecipeSym.bind(this));
- this.registerCustomDirective('OperandWithKeywordOnly', this.parseOperandWithKeywordOnly.bind(this));
- this.registerCustomDirective('OperandsWithKeywordOnly', this.parseOperandsWithKeywordOnly.bind(this));
- this.registerCustomDirective('DeviceTypeOperandsWithSegment', this.parseDeviceTypeOperandsWithSegment.bind(this));
- this.registerCustomDirective('BindName', this.parseBindName.bind(this));
- this.registerCustomDirective('RoutineGangClause', this.parseRoutineGangClause.bind(this));
- this.registerCustomDirective('DeviceTypeArrayAttr', this.parseDeviceTypeArrayAttr.bind(this));
- }
- // custom<Var>($var) - receives ctx.get('var').operands
- parseVar(parser, op, operands) {
- if (!parser.parseOptionalKeyword('varPtr')) {
- parser.parseKeyword('var');
- }
- parser.parseLParen();
- operands.push(parser.parseOperand());
- }
- // custom<AccVar>($accVar, type($accVar)) - self-contained with type inline
- parseAccVar(parser, op, operands, types) {
- if (!parser.parseOptionalKeyword('accPtr')) {
- parser.parseKeyword('accVar');
- }
- parser.parseLParen();
- operands.push(parser.parseOperand());
- parser.parseColon();
- types.push(parser.parseType());
- parser.parseRParen();
- }
- // custom<VarPtrType>(type($var), $varType) - receives ctx.get('var').types
- parseVarPtrType(parser, op, types, varTypeAttrName) {
- const type = parser.parseType();
- types.push(type);
- parser.parseRParen();
- if (parser.parseOptionalKeyword('varType')) {
- parser.parseLParen();
- const varType = parser.parseType();
- op.addAttribute(varTypeAttrName || 'varType', varType);
- parser.parseRParen();
- }
- }
- parseDeviceTypeOperandsWithKeywordOnly(parser, op, operands, types, deviceTypesVar, keywordOnlyVar) {
- if (!parser.parseOptionalLParen()) {
- op.addAttribute(keywordOnlyVar, [{ value: 'none' }]);
- return;
- }
- const keywordOnlyAttrs = [];
- let needComma = false;
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- const attr = parser.parseAttribute();
- keywordOnlyAttrs.push(attr);
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- needComma = true;
- }
- if (keywordOnlyAttrs.length > 0) {
- op.addAttribute(keywordOnlyVar, keywordOnlyAttrs);
- }
- if (needComma) {
- parser.parseOptionalComma();
- }
- const unresolvedOperands = [];
- const operandTypes = [];
- const deviceTypes = [];
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- unresolvedOperands.push(operand);
- operandTypes.push(type);
- if (parser.parseOptionalLSquare()) {
- const deviceType = parser.parseAttribute();
- deviceTypes.push(deviceType);
- parser.parseRSquare();
- } else {
- deviceTypes.push({ value: 'none' });
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.resolveOperands(unresolvedOperands, operandTypes, op.operands);
- if (deviceTypes.length > 0) {
- op.addAttribute(deviceTypesVar, deviceTypes);
- }
- }
- parseLoopControl(parser, result) {
- const inductionVars = [];
- if (parser.parseOptionalKeyword('control')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const value = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- inductionVars.push({ value, type });
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- parser.parseEqual();
- parser.parseLParen();
- const lowerbound = parser.parseOperandList();
- const lowerboundTypes = parser.parseColonTypeList();
- parser.resolveOperands(lowerbound, lowerboundTypes, result.operands);
- parser.parseRParen();
- parser.parseKeyword('to');
- parser.parseLParen();
- const upperbound = parser.parseOperandList();
- const upperboundTypes = parser.parseColonTypeList();
- parser.resolveOperands(upperbound, upperboundTypes, result.operands);
- parser.parseRParen();
- parser.parseKeyword('step');
- parser.parseLParen();
- const step = parser.parseOperandList();
- const stepTypes = parser.parseColonTypeList();
- parser.resolveOperands(step, stepTypes, result.operands);
- parser.parseRParen();
- }
- const region = result.addRegion();
- parser.parseRegion(region, inductionVars);
- }
- parseWaitClause(parser, result) {
- if (parser.getToken().isNot(_.Token.l_paren)) {
- return;
- }
- parser.parseLParen();
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.parseAttribute();
- parser.parseOptionalComma();
- }
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOptionalComma();
- }
- }
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.parseOptionalLBrace()) {
- parser.parseOptionalKeyword('devnum');
- parser.parseOptionalColon();
- while (!parser.parseOptionalRBrace()) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- parser.parseOptionalComma();
- }
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- }
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- }
- parseNumGangs(parser, result) {
- while (parser.parseOptionalLBrace()) {
- while (!parser.parseOptionalRBrace()) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- parser.parseOptionalComma();
- }
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseOptionalComma();
- }
- }
- parseDeviceTypeOperands(parser, result) {
- while (parser.getToken().is(_.Token.percent_identifier)) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseOptionalComma();
- }
- }
- parseGangClause(parser, op, gangOperands, gangTypes, gangArgTypeVar, gangDeviceTypeVar, gangSegmentsVar, gangOnlyVar) {
- if (!parser.parseOptionalLParen()) {
- op.addAttribute(gangOnlyVar, [{ value: 'none' }]);
- return;
- }
- const gangOnlyAttrs = [];
- let needComma = false;
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- const attr = parser.parseAttribute();
- gangOnlyAttrs.push(attr);
- parser.parseOptionalComma();
- }
- parser.parseRSquare();
- needComma = true;
- }
- if (gangOnlyAttrs.length > 0) {
- op.addAttribute(gangOnlyVar, gangOnlyAttrs);
- }
- if (needComma) {
- parser.parseOptionalComma();
- }
- const gangArgTypes = [];
- const deviceTypes = [];
- const segments = [];
- while (parser.parseOptionalLBrace()) {
- let segmentCount = 0;
- while (parser.getToken().isNot(_.Token.r_brace)) {
- let argType = 'Num';
- if (parser.parseOptionalKeyword('num')) {
- parser.parseEqual();
- argType = 'Num';
- } else if (parser.parseOptionalKeyword('dim')) {
- parser.parseEqual();
- argType = 'Dim';
- } else if (parser.parseOptionalKeyword('static')) {
- parser.parseEqual();
- argType = 'Static';
- }
- gangArgTypes.push({ value: argType });
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, op.operands);
- segmentCount++;
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRBrace();
- segments.push(segmentCount);
- if (parser.parseOptionalLSquare()) {
- const deviceType = parser.parseAttribute();
- deviceTypes.push(deviceType);
- parser.parseRSquare();
- } else {
- deviceTypes.push({ value: 'none' });
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- if (gangArgTypes.length > 0) {
- op.addAttribute('gangOperandsArgType', gangArgTypes);
- }
- if (deviceTypes.length > 0) {
- op.addAttribute('gangOperandsDeviceType', deviceTypes);
- }
- if (segments.length > 0) {
- op.addAttribute('gangOperandsSegments', segments);
- }
- }
- parseCombinedConstructsLoop(parser, result) {
- const attr = parser.parseAttribute();
- result.addAttribute('combined', attr);
- }
- parseRecipeSym(parser, result) {
- const attr = parser.parseAttribute();
- result.addAttribute('recipe', attr);
- }
- parseOperandWithKeywordOnly(parser, result) {
- if (parser.getToken().isNot(_.Token.l_paren)) {
- return;
- }
- parser.parseLParen();
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.parseAttribute();
- parser.parseOptionalComma();
- }
- if (parser.getToken().isNot(_.Token.r_paren)) {
- parser.parseOptionalComma();
- }
- }
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const operand = parser.parseOperand();
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperand(operand, type, result.operands);
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- }
- parseOperandsWithKeywordOnly(parser, result) {
- // Handles format: (%v1, %v2 : t1, t2) where all operands are listed before colon
- // and all types are listed after colon
- if (parser.getToken().isNot(_.Token.l_paren)) {
- return;
- }
- parser.parseLParen();
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseRParen();
- return;
- }
- const unresolvedOperands = [];
- do {
- unresolvedOperands.push(parser.parseOperand());
- } while (parser.parseOptionalComma() && parser.getToken().isNot(_.Token.colon));
- parser.parseColon();
- const types = [];
- for (let i = 0; i < unresolvedOperands.length; i++) {
- if (i > 0) {
- parser.parseComma();
- }
- types.push(parser.parseType());
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- parser.parseRParen();
- }
- parseDeviceTypeOperandsWithSegment(parser, result) {
- this.parseNumGangs(parser, result);
- }
- parseBindName(parser, result) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- const attr = parser.parseAttribute();
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- result.addAttribute('bind', attr);
- parser.parseOptionalComma();
- }
- }
- parseRoutineGangClause(parser, result) {
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.parseOptionalKeyword('dim')) {
- parser.parseColon();
- }
- const value = parser.parseAttribute();
- if (parser.parseOptionalLSquare()) {
- parser.parseAttribute();
- parser.parseRSquare();
- }
- result.addAttribute('gangDim', value);
- parser.parseOptionalComma();
- }
- parser.parseRParen();
- } else if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.parseAttribute();
- parser.parseOptionalComma();
- }
- }
- }
- parseDeviceTypeArrayAttr(parser) {
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- parser.parseAttribute();
- parser.parseOptionalComma();
- }
- }
- }
- };
- _.smt = {};
- _.smt.BitVectorType = class extends _.Type {
- constructor(width) {
- super(null);
- this.width = width;
- }
- toString() {
- return `!smt.bv<${this.width}>`;
- }
- };
- _.smt.SMTDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'smt');
- }
- parseOperation(parser, result) {
- if (result.op === 'smt.eq' || result.op === 'smt.distinct') {
- const unresolvedOperands = parser.parseOperandList();
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const type = parser.parseType();
- const types = unresolvedOperands.map(() => type);
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- result.addTypes([new _.Type('!smt.bool')]);
- return true;
- }
- if (result.op === 'smt.bv.repeat') {
- const count = parser.parseInteger();
- result.addAttribute('count', count);
- parser.parseKeyword('times');
- const unresolvedOperand = parser.parseOperand();
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const inputType = parser.parseType();
- parser.resolveOperand(unresolvedOperand, inputType, result.operands);
- // RepeatOp::parse - result width = input width * count
- const inputWidth = inputType instanceof _.smt.BitVectorType ? inputType.width : 0;
- const resultWidth = inputWidth * count;
- result.addTypes([new _.smt.BitVectorType(resultWidth)]);
- return true;
- }
- if (result.op === 'smt.int.constant') {
- const value = parser.parseInteger();
- result.addAttribute('value', value);
- parser.parseOptionalAttrDict(result.attributes);
- result.addTypes([new _.Type('!smt.int')]);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseType(parser) {
- const typeName = parser.parseOptionalKeyword();
- if (typeName === 'bv' && parser.getToken().is(_.Token.less)) {
- parser.parseLess();
- const width = parseInt(parser.getToken().getSpelling().str(), 10);
- parser.consumeToken(_.Token.integer);
- parser.parseGreater();
- return new _.smt.BitVectorType(width);
- }
- if (typeName) {
- let type = `!smt.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- return null;
- }
- inferResultTypes(op, vars) {
- if (op.op === 'smt.bv.concat') {
- // ConcatOp::inferReturnTypes - result width = lhs width + rhs width
- if (op.operands.length >= 2) {
- const lhsType = op.operands[0].type;
- const rhsType = op.operands[1].type;
- const lhsWidth = lhsType instanceof _.smt.BitVectorType ? lhsType.width : 0;
- const rhsWidth = rhsType instanceof _.smt.BitVectorType ? rhsType.width : 0;
- const resultWidth = lhsWidth + rhsWidth;
- op.addTypes([new _.smt.BitVectorType(resultWidth)]);
- return;
- }
- }
- super.inferResultTypes(op, vars);
- }
- };
- _.MPMDDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'mpmd');
- this.registerCustomType('mesh_tensor', this.parseMeshTensorType.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'mpmd.named_computation') {
- return this.parseNamedComputationOp(parser, result);
- }
- if (result.op === 'mpmd.fragment') {
- return this.parseFragmentOp(parser, result);
- }
- if (result.op === 'mpmd.fragment_call') {
- return this.parseFragmentCallOp(parser, result);
- }
- if (result.op === 'mpmd.for') {
- return this.parseForOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseNamedComputationOp(parser, result) {
- // mpmd.named_computation<"name"(count)> (%inputs) (%block_args) { region } : (types) -> types
- parser.parseLess();
- // Parse single UserOriginAttr in short format: "name"(count)
- const origin = this.parseUserOriginAttr(parser);
- result.addAttribute('origin', origin);
- parser.parseGreater();
- const unresolvedInputs = parser.parseOperandList('paren');
- const entryArguments = this.parseBlockArguments(parser);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, entryArguments);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedInputs, type.inputs, result.operands);
- result.addTypes(type.results);
- }
- }
- return true;
- }
- parseBlockArguments(parser) {
- return parser.parseArgumentList('optionalParen', true);
- }
- parseFragmentOp(parser, result) {
- // mpmd.fragment<mesh="m1", origin=["f1"], stage_id=N> (%inputs) {attrs} (%block_args) { region } : (types) -> type
- parser.parseLess();
- while (parser.getToken().isNot(_.Token.greater)) {
- const attrName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseEqual();
- // Use custom parser for origin attribute (array of UserOriginAttr)
- const attrValue = attrName === 'origin' ? this.parseOriginArray(parser) : parser.parseAttribute();
- result.addAttribute(attrName, attrValue);
- parser.parseOptionalComma();
- }
- parser.parseGreater();
- const unresolvedInputs = parser.parseOperandList('paren');
- parser.parseOptionalAttrDict(result.attributes);
- const entryArguments = this.parseBlockArguments(parser);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, entryArguments);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedInputs, type.inputs, result.operands);
- result.addTypes(type.results);
- }
- }
- return true;
- }
- parseFragmentCallOp(parser, result) {
- // mpmd.fragment_call<mesh="m1", origin=["f1"]> @callee(%args) {attrs} : (types) -> type
- parser.parseLess();
- while (parser.getToken().isNot(_.Token.greater)) {
- const attrName = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.bare_identifier);
- parser.parseEqual();
- // Use custom parser for origin attribute (array of UserOriginAttr)
- const attrValue = attrName === 'origin' ? this.parseOriginArray(parser) : parser.parseAttribute();
- result.addAttribute(attrName, attrValue);
- parser.parseOptionalComma();
- }
- parser.parseGreater();
- const callee = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('callee', callee);
- const unresolvedArgs = parser.parseOperandList('paren');
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(result.attributes);
- }
- if (parser.parseOptionalColon()) {
- const type = parser.parseType();
- if (type instanceof _.FunctionType) {
- parser.resolveOperands(unresolvedArgs, type.inputs, result.operands);
- result.addTypes(type.results);
- }
- }
- return true;
- }
- parseForOp(parser, result) {
- // Use parseOperandList for simple operand list (no types in syntax)
- const inputs = parser.parseOperandList('paren');
- const types = inputs.map(() => null);
- parser.resolveOperands(inputs, types, result.operands);
- parser.parseOptionalAttrDict(result.attributes);
- const entryArguments = this.parseBlockArguments(parser);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, entryArguments);
- }
- if (parser.parseOptionalColon()) {
- const resultTypes = [];
- do {
- resultTypes.push(parser.parseType());
- } while (parser.parseOptionalComma());
- result.addTypes(resultTypes);
- }
- return true;
- }
- parseMeshTensorType(parser, prefix) {
- // Parse !mpmd.mesh_tensor<"mesh_name", tensor<shape>, sharding=<...>>
- parser.parseLess();
- const meshName = parser.parseString();
- parser.parseComma();
- const tensorType = parser.parseType();
- const result = { name: prefix, meshName, tensorType };
- if (parser.parseOptionalComma()) {
- if (parser.parseOptionalKeyword('sharding')) {
- parser.parseEqual();
- result.sharding = parser.parseAttribute();
- }
- }
- parser.parseGreater();
- return result;
- }
- // Parse UserOriginAttr in short format: "name"(count) where (count) is optional
- parseUserOriginAttr(parser) {
- const name = parser.parseString();
- let transposeCount = 0;
- if (parser.parseOptionalLParen()) {
- const count = parser.parseInteger();
- transposeCount = count;
- parser.parseRParen();
- }
- return { name, transposeCount };
- }
- parseOriginArray(parser) {
- const origins = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const origin = this.parseUserOriginAttr(parser);
- origins.push(origin);
- parser.parseOptionalComma();
- }
- return origins;
- }
- };
- _.SdyDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'sdy');
- this.registerCustomDirective('StrippedTensorShardingPerValueAttr', this.parseStrippedTensorShardingPerValueAttr.bind(this));
- this.registerCustomDirective('SingleBlockRegionNoBlockId', this.parseSingleBlockRegionNoBlockId.bind(this));
- this.registerCustomAttribute('Sdy_ListOfAxisRefLists', this.parseListOfAxisRefLists.bind(this));
- this.registerCustomAttribute('Sdy_ManualAxes', this.parseManualAxes.bind(this));
- this.registerCustomAttribute('Sdy_AllToAllParamList', this.parseAllToAllParamList.bind(this));
- this.registerCustomAttribute('Sdy_TensorSharding', this.parseTensorShardingAttrWrap.bind(this));
- this.registerCustomAttribute('Sdy_AxisRefList', this.parseAxisRefListWrap.bind(this));
- }
- parseOperation(parser, result) {
- if (result.op === 'sdy.constant') {
- return this.parseConstantOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- parseStrippedTensorShardingPerValueAttr(parser, op, attrName) {
- const shardings = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const sharding = this.parseTensorShardingAttr(parser);
- shardings.push(sharding);
- parser.parseOptionalComma();
- }
- if (attrName) {
- op.addAttribute(attrName, shardings);
- }
- return shardings;
- }
- parseSingleBlockRegionNoBlockId(parser, op /*, args */) {
- const entryArguments = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- const value = parser.parseOperand();
- let type = null;
- const attrs = [];
- if (parser.parseOptionalColon()) {
- type = parser.parseType();
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- parser.parseAttributeDict(attrs);
- }
- entryArguments.push({ value, type, attributes: attrs.length > 0 ? attrs : undefined });
- parser.parseOptionalComma();
- }
- }
- const region = op.addRegion();
- parser.parseRegion(region, entryArguments);
- return region;
- }
- parseTensorShardingAttr(parser) {
- parser.parseLess();
- let meshOrRef = null;
- if (parser.getToken().is(_.Token.at_identifier)) {
- meshOrRef = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- } else if (parser.parseOptionalKeyword('mesh')) {
- meshOrRef = this.parseMeshAttr(parser);
- } else {
- throw new mlir.Error(`Expected '@' or 'mesh', but got '${parser.getTokenSpelling().str()}' ${parser.location()}`);
- }
- parser.parseComma();
- const dimShardings = this.parseDimensionShardings(parser);
- let replicatedAxes = null;
- let unreducedAxes = null;
- while (parser.parseOptionalComma()) {
- if (parser.parseOptionalKeyword('replicated')) {
- parser.parseEqual();
- replicatedAxes = this.parseAxisRefList(parser);
- } else if (parser.parseOptionalKeyword('unreduced')) {
- parser.parseEqual();
- unreducedAxes = this.parseAxisRefList(parser);
- }
- }
- parser.parseGreater();
- return { meshOrRef, dimShardings, replicatedAxes, unreducedAxes };
- }
- parseMeshAttr(parser) {
- parser.parseLess();
- const axes = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const name = parser.parseString();
- parser.parseEqual();
- const size = parser.parseInteger();
- axes.push({ name, size });
- parser.parseOptionalComma();
- }
- let deviceIds = null;
- if (parser.parseOptionalComma()) {
- if (parser.parseOptionalKeyword('device_ids')) {
- parser.parseEqual();
- deviceIds = parser.parseAttribute().value;
- }
- }
- parser.parseGreater();
- return { axes, deviceIds };
- }
- parseDimensionShardings(parser) {
- const dimShardings = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const dimSharding = this.parseDimensionShardingAttr(parser);
- dimShardings.push(dimSharding);
- parser.parseOptionalComma();
- }
- return dimShardings;
- }
- parseDimensionShardingAttr(parser) {
- const axes = [];
- parser.parseLBrace();
- let isClosed = true;
- while (!parser.parseOptionalRBrace()) {
- if (parser.getToken().is(_.Token.question)) {
- parser.consumeToken(_.Token.question);
- isClosed = false;
- parser.parseRBrace();
- break;
- }
- const axis = this.parseAxisRefAttr(parser);
- axes.push(axis);
- if (parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.question)) {
- parser.consumeToken(_.Token.question);
- isClosed = false;
- parser.parseRBrace();
- break;
- }
- }
- }
- let priority = null;
- if (parser.getToken().is(_.Token.bare_identifier)) {
- const tokenValue = parser.getTokenSpelling().str();
- if (typeof tokenValue === 'string' && tokenValue.startsWith('p') && /^p\d+$/.test(tokenValue)) {
- parser.consumeToken(_.Token.bare_identifier);
- priority = parseInt(tokenValue.substring(1), 10);
- }
- }
- return { axes, isClosed, priority };
- }
- parseAxisRefAttr(parser) {
- const name = parser.parseString();
- let subAxisInfo = null;
- if (parser.parseOptionalColon()) {
- subAxisInfo = this.parseSubAxisInfo(parser);
- }
- return { name, subAxisInfo };
- }
- parseSubAxisInfo(parser) {
- parser.parseLParen();
- const preSize = parser.parseInteger();
- parser.parseRParen();
- const size = parser.parseInteger();
- return { preSize, size };
- }
- parseAxisRefList(parser) {
- const axes = [];
- parser.parseLBrace();
- while (!parser.parseOptionalRBrace()) {
- const axis = this.parseAxisRefAttr(parser);
- axes.push(axis);
- parser.parseOptionalComma();
- }
- return axes;
- }
- parseListOfAxisRefLists(parser) {
- const lists = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const list = this.parseAxisRefList(parser);
- lists.push(list);
- parser.parseOptionalComma();
- }
- return { value: lists };
- }
- parseManualAxes(parser) {
- const axes = [];
- parser.parseLBrace();
- while (!parser.parseOptionalRBrace()) {
- const axis = parser.parseString();
- axes.push(axis);
- parser.parseOptionalComma();
- }
- return { value: axes };
- }
- parseAllToAllParamList(parser) {
- const params = [];
- parser.parseLSquare();
- while (!parser.parseOptionalRSquare()) {
- const param = this.parseAllToAllParam(parser);
- params.push(param);
- parser.parseOptionalComma();
- }
- return { value: params };
- }
- parseAllToAllParam(parser) {
- const axes = this.parseAxisRefList(parser);
- parser.parseColon();
- const splitDim = parser.parseInteger();
- parser.parseArrow();
- const concatDim = parser.parseInteger();
- return { axes, splitDim, concatDim };
- }
- parseTensorShardingAttrWrap(parser) {
- return { value: this.parseTensorShardingAttr(parser) };
- }
- parseAxisRefListWrap(parser) {
- return { value: this.parseAxisRefList(parser) };
- }
- parseConstantOp(parser, result) {
- parser.parseOptionalAttrDict(result.attributes);
- const attr = parser.parseDenseElementsAttr();
- result.addAttribute('value', attr.value || attr);
- if (attr.type) {
- if (result.types.length === 0) {
- result.addTypes([attr.type]);
- } else {
- result.types[0] = attr.type;
- }
- }
- return true;
- }
- };
- _.XlaDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xla');
- }
- parseOperation(parser, result) {
- if (result.op === 'xla.apply_indexing') {
- const indexingMapAttr = parser.parseAttribute();
- result.addAttribute('indexing_map_attr', indexingMapAttr);
- const operands = [];
- if (parser.parseOptionalLParen()) {
- parser.parseOperands(operands);
- parser.parseRParen();
- }
- if (parser.parseOptionalLSquare()) {
- parser.parseOperands(operands);
- parser.parseRSquare();
- }
- const indexType = new _.IndexType();
- for (const operand of operands) {
- parser.resolveOperand(operand, indexType, result.operands);
- }
- parser.parseOptionalAttrDict(result.attributes);
- const map = indexingMapAttr.getIndexingMap().GetAffineMap();
- const numResults = map.getNumResults();
- for (let i = 0; i < numResults; i++) {
- result.addTypes([indexType]);
- }
- return true;
- }
- // xla.loop (%dims)[%ivs] -> (%map_results) in #map iter_args(%args = %inits) -> (types) { body }
- if (result.op === 'xla.loop') {
- const unresolvedDims = [];
- const unresolvedInits = [];
- const regionArgs = [];
- const indexType = new _.IndexType();
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedDims.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- // Parse [%ivs] -> (%map_results) - these are block arguments
- if (parser.parseOptionalLSquare()) {
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const iv = parser.parseOperand();
- regionArgs.push({ name: iv.name, type: indexType });
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- }
- if (parser.parseOptionalArrow()) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const mapResult = parser.parseOperand();
- regionArgs.push({ name: mapResult.name, type: indexType });
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseKeyword('in');
- const map = parser.parseAttribute();
- result.addAttribute('indexing_map_attr', map);
- const iterArgNames = [];
- if (parser.parseOptionalKeyword('iter_args')) {
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const iterArg = parser.parseOperand();
- iterArgNames.push(iterArg.name);
- }
- if (parser.parseOptionalEqual()) {
- unresolvedInits.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- const resultTypes = [];
- if (parser.parseOptionalArrow()) {
- if (parser.parseOptionalLParen()) {
- const types = parser.parseTypeListNoParens();
- parser.parseRParen();
- for (const t of types) {
- result.addTypes([t]);
- resultTypes.push(t);
- }
- } else {
- const type = parser.parseType();
- result.addTypes([type]);
- resultTypes.push(type);
- }
- }
- for (let i = 0; i < iterArgNames.length; i++) {
- const argType = resultTypes[i] || null;
- regionArgs.push({ name: iterArgNames[i], type: argType });
- }
- for (const dim of unresolvedDims) {
- parser.resolveOperand(dim, indexType, result.operands);
- }
- for (let i = 0; i < unresolvedInits.length; i++) {
- const initType = resultTypes[i] || null;
- parser.resolveOperand(unresolvedInits[i], initType, result.operands);
- }
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, regionArgs);
- }
- parser.parseOptionalAttrDict(result.attributes);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- parseAttribute(parser, type) {
- const keyword = parser.parseOptionalKeyword();
- if (keyword === 'indexing_map') {
- return _.IndexingMapAttr.parse(parser, type);
- }
- return null;
- }
- };
- _.XlaGpuDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xla_gpu');
- }
- parseOperation(parser, result) {
- // xla_gpu.shuffle_reduce(%ops) to N combiner=@func {attrs} : types
- if (result.op === 'xla_gpu.shuffle_reduce') {
- const unresolvedOperands = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseKeyword('to');
- const maxDistance = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- result.addAttribute('max_distance', parseInt(maxDistance, 10));
- parser.parseKeyword('combiner');
- parser.parseEqual();
- const combiner = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('combiner', new _.SymbolRefAttr(`@${combiner}`));
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const types = parser.parseTypeList();
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- result.addTypes(types);
- } else {
- for (const operand of unresolvedOperands) {
- parser.resolveOperand(operand, null, result.operands);
- }
- }
- return true;
- }
- // xla_gpu.reduce (%inputs) inits(%inits) dimensions=[...] combiner=@func {attrs} : in_types to out_types
- if (result.op === 'xla_gpu.reduce') {
- const unresolvedInputs = [];
- const unresolvedInits = [];
- if (parser.parseOptionalLParen()) {
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedInputs.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- }
- parser.parseKeyword('inits');
- parser.parseLParen();
- while (parser.getToken().isNot(_.Token.r_paren)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedInits.push(parser.parseOperand());
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRParen();
- parser.parseKeyword('dimensions');
- parser.parseEqual();
- const dimensions = parser.parseAttribute();
- result.addAttribute('dimensions', dimensions);
- parser.parseKeyword('combiner');
- parser.parseEqual();
- const combiner = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.at_identifier);
- result.addAttribute('combiner', new _.SymbolRefAttr(`@${combiner}`));
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalColon()) {
- const inputTypes = parser.parseTypeList();
- parser.resolveOperands(unresolvedInputs, inputTypes, result.operands);
- parser.parseKeyword('to');
- const outputTypes = parser.parseTypeList();
- result.addTypes(outputTypes);
- parser.resolveOperands(unresolvedInits, outputTypes, result.operands);
- } else {
- for (const input of unresolvedInputs) {
- parser.resolveOperand(input, null, result.operands);
- }
- for (const init of unresolvedInits) {
- parser.resolveOperand(init, null, result.operands);
- }
- }
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.XTileDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xtile');
- }
- parseOperation(parser, result) {
- // xtile.entry_func - function-like op with custom format
- if (result.op === 'xtile.entry_func') {
- parser.parseFunctionOp(result, false);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- };
- _.TritonXlaDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'triton_xla');
- // Mark operations with custom directives as having custom assembly format
- this.registerOperandName('triton_xla.extract', { name: 'triton_xla.extract', hasCustomAssemblyFormat: true });
- this.registerOperandName('triton_xla.insert', { name: 'triton_xla.insert', hasCustomAssemblyFormat: true });
- }
- parseOperation(parser, result) {
- // triton_xla.extract from $src as memref<shape, layout> [offsets] [sizes] [strides] : result_type
- if (result.op === 'triton_xla.extract') {
- parser.parseKeyword('from');
- const unresolvedSrc = parser.parseOperand();
- parser.parseKeyword('as');
- // Parse AsMemRefType: memref<shape, layout> - the memref type describes the pointer's layout
- const memrefType = parser.parseType();
- // Extract shape and layout from the memref type
- this._extractMemRefInfo(memrefType, result, 'src_shape', 'src_layout');
- this.parseDynamicIndexList(parser, result, 'offsets', 'static_offsets');
- this.parseDynamicIndexList(parser, result, 'sizes', 'static_sizes');
- this.parseDynamicIndexList(parser, result, 'strides', 'static_strides');
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const resultType = parser.parseType();
- parser.resolveOperand(unresolvedSrc, memrefType, result.operands);
- result.addTypes([resultType]);
- return true;
- }
- // triton_xla.insert $src into $dst as memref<shape, layout> [offsets] [sizes] [strides] : src_type
- if (result.op === 'triton_xla.insert') {
- const unresolvedSrc = parser.parseOperand();
- parser.parseKeyword('into');
- const unresolvedDst = parser.parseOperand();
- parser.parseKeyword('as');
- // Parse AsMemRefType: memref<shape, layout> - the memref type describes the pointer's layout
- const memrefType = parser.parseType();
- // Extract shape and layout from the memref type
- this._extractMemRefInfo(memrefType, result, 'dst_shape', 'dst_layout');
- this.parseDynamicIndexList(parser, result, 'offsets', 'static_offsets');
- this.parseDynamicIndexList(parser, result, 'sizes', 'static_sizes');
- this.parseDynamicIndexList(parser, result, 'strides', 'static_strides');
- parser.parseOptionalAttrDict(result.attributes);
- parser.parseColon();
- const srcType = parser.parseType();
- parser.resolveOperand(unresolvedSrc, srcType, result.operands);
- parser.resolveOperand(unresolvedDst, memrefType, result.operands);
- return true;
- }
- return super.parseOperation(parser, result);
- }
- _extractMemRefInfo(memrefType, op, shapeAttrName, layoutAttrName) {
- // Extract shape and layout from memref type like memref<512x1x128xbf16, #xtile.layout<[2, 1, 0]>>
- const typeStr = memrefType.toString();
- const shapeMatch = typeStr.match(/memref<([\d?x]+)/);
- if (shapeMatch) {
- const dims = shapeMatch[1].split('x').filter((d) => d).map((d) => d === '?' ? -1 : parseInt(d, 10));
- op.addAttribute(shapeAttrName, dims);
- }
- const layoutMatch = typeStr.match(/#[a-z_.]+\.<\[([^\]]+)\]>/);
- if (layoutMatch) {
- const layout = layoutMatch[1].split(',').map((s) => parseInt(s.trim(), 10));
- op.addAttribute(layoutAttrName, layout);
- }
- }
- parseDynamicIndexList(parser, op, dynamicName, staticName) {
- // Parse [val1, val2, ...] where vals can be %ssa or integer constants
- parser.parseLSquare();
- const staticValues = [];
- while (parser.getToken().isNot(_.Token.r_square)) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const unresolved = parser.parseOperand();
- parser.resolveOperand(unresolved, null, op.operands);
- staticValues.push(-9223372036854775808n); // ShapedType::kDynamic
- } else if (parser.getToken().is(_.Token.integer)) {
- const value = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.integer);
- staticValues.push(BigInt(value));
- }
- if (!parser.parseOptionalComma()) {
- break;
- }
- }
- parser.parseRSquare();
- op.addAttribute(staticName, staticValues);
- }
- };
- _.EnsembleDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'ensemble');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- const type = `!${dialect}.${typeName}`;
- const simpleTypes = ['physical_qubit', 'virtual_qubit', 'cbit', 'gate', 'gate_distribution', 'connectivity_graph'];
- if (simpleTypes.includes(typeName)) {
- return new _.Type(type);
- }
- return null;
- }
- };
- _.PolyDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'poly');
- }
- parseType(parser, dialect) {
- const typeName = parser.parseOptionalKeyword();
- if (!typeName) {
- return null;
- }
- let type = `!${dialect}.${typeName}`;
- // poly.poly<N> type has a degree bound parameter
- if (typeName === 'poly' && parser.getToken().is(_.Token.less)) {
- const content = parser.skip('<');
- type += content;
- }
- return new _.Type(type);
- }
- };
- _.NoisyDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'noisy');
- }
- parseType(parser, dialect) {
- if (parser.getToken().is(_.Token.inttype)) {
- const inttype = parser.getToken().getSpelling().str();
- parser.consumeToken(_.Token.inttype);
- return new _.Type(`!${dialect}.${inttype}`);
- }
- const typeName = parser.parseOptionalKeyword();
- if (typeName) {
- let type = `!${dialect}.${typeName}`;
- if (parser.getToken().is(_.Token.less)) {
- type += parser.skip('<');
- }
- return new _.Type(type);
- }
- return null;
- }
- };
- _.XtenNNDialect = class extends _.Dialect {
- constructor(operations) {
- super(operations, 'xten_nn');
- }
- parseAttribute(parser) {
- const keyword = parser.parseOptionalKeyword();
- if (keyword === 'dict_loc') {
- parser.parseLParen();
- const content = parser.skip('(');
- return new _.XtenNNDictLoc(content);
- }
- return null;
- }
- parseOperation(parser, result) {
- if (result.op === 'xten_nn.subgraph') {
- return this.parseSubgraphOp(parser, result);
- }
- if (result.op === 'xten_nn.quantize' || result.op === 'xten_nn.dequantize') {
- return this.parseQuantizeOp(parser, result);
- }
- if (result.op === 'xten_nn.kernel') {
- return this.parseKernelOp(parser, result);
- }
- return super.parseOperation(parser, result);
- }
- // Parse: (%name = %value : type, ...) { body } -> result_types
- // or: (%arg : type, ...) -> result_types (no body)
- parseSubgraphOp(parser, result) {
- const entryArgs = [];
- if (parser.parseOptionalLParen()) {
- while (!parser.parseOptionalRParen()) {
- // Parse either (%name = %value : type) or (%value : type)
- let blockArgName = null;
- let operandRef = null;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- const firstOperand = parser.parseOperand();
- if (parser.parseOptionalEqual()) {
- // (%name = %value : type) form
- blockArgName = firstOperand;
- if (parser.getToken().is(_.Token.percent_identifier)) {
- operandRef = parser.parseOperand();
- }
- } else {
- // (%value : type) form - value is both operand and block arg name
- operandRef = firstOperand;
- blockArgName = firstOperand;
- }
- }
- parser.parseColon();
- const type = parser.parseType();
- if (operandRef) {
- parser.resolveOperands([operandRef], [type], result.operands);
- }
- entryArgs.push({ value: blockArgName, type });
- if (!parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseOptionalRParen();
- }
- break;
- }
- }
- }
- parser.parseOptionalAttrDictWithKeyword(result.attributes);
- if (parser.getToken().is(_.Token.l_brace)) {
- const region = result.addRegion();
- parser.parseRegion(region, entryArgs, /* enableNameShadowing */ true);
- }
- if (parser.parseOptionalArrow()) {
- const types = parser.getToken().is(_.Token.l_paren) ? parser.parseTypeListParens() : parser.parseTypeListNoParens();
- result.addTypes(types);
- }
- return true;
- }
- // Parse: (%input: type) {attrs} -> type
- parseQuantizeOp(parser, result) {
- if (!parser.parseOptionalLParen()) {
- return false;
- }
- const unresolvedOperands = [];
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseColon();
- const type = parser.parseType();
- parser.resolveOperands(unresolvedOperands, [type], result.operands);
- unresolvedOperands.length = 0;
- if (!parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseOptionalRParen();
- }
- break;
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- result.addTypes(parser.parseArrowTypeList());
- return true;
- }
- // Parse: "name" (%arg : type, ...) instantiation_args [...] {attrs} -> result_types
- parseKernelOp(parser, result) {
- const nameAttr = parser.parseAttribute();
- result.addAttribute('name', nameAttr);
- if (parser.parseOptionalLParen()) {
- const unresolvedOperands = [];
- const types = [];
- while (!parser.parseOptionalRParen()) {
- if (parser.getToken().is(_.Token.percent_identifier)) {
- unresolvedOperands.push(parser.parseOperand());
- }
- parser.parseColon();
- types.push(parser.parseType());
- if (!parser.parseOptionalComma()) {
- if (parser.getToken().is(_.Token.r_paren)) {
- parser.parseOptionalRParen();
- }
- break;
- }
- }
- parser.resolveOperands(unresolvedOperands, types, result.operands);
- }
- if (parser.parseOptionalKeyword('instantiation_args')) {
- const instArgs = [];
- const instArgNames = [];
- if (parser.parseOptionalLSquare()) {
- while (!parser.parseOptionalRSquare()) {
- // Parse either "name" = value or just value
- if (parser.getToken().is(_.Token.string)) {
- const nameOrValue = parser.parseAttribute();
- if (parser.parseOptionalEqual()) {
- // "name" = value form
- instArgNames.push(nameOrValue);
- instArgs.push(parser.parseAttribute());
- } else {
- // Just a string value
- instArgs.push(nameOrValue);
- }
- } else {
- instArgs.push(parser.parseAttribute());
- }
- parser.parseOptionalComma();
- }
- }
- if (instArgs.length > 0) {
- result.addAttribute('instantiation_args', instArgs);
- }
- if (instArgNames.length > 0) {
- result.addAttribute('instantiation_arg_names', instArgNames);
- }
- }
- parser.parseOptionalAttrDict(result.attributes);
- if (parser.parseOptionalArrow()) {
- const types = parser.getToken().is(_.Token.l_paren) ? parser.parseTypeListParens() : parser.parseTypeListNoParens();
- result.addTypes(types);
- }
- return true;
- }
- };
- mlir.Metadata = class {
- static async open(context) {
- if (!mlir.Metadata._metadata) {
- const data = await context.request('mlir-metadata.json');
- mlir.Metadata._metadata = new mlir.Metadata(data);
- }
- return mlir.Metadata._metadata;
- }
- constructor(data) {
- this.operations = new Map();
- if (data) {
- const operations = JSON.parse(data);
- for (const op of operations) {
- const [dialectName] = op.name.split('.');
- if (!this.operations.has(dialectName)) {
- this.operations.set(dialectName, []);
- }
- this.operations.get(dialectName).push(op);
- }
- }
- }
- type(name) {
- const [dialectName] = name.split('.');
- const operations = this.operations.get(dialectName);
- if (operations) {
- const op = operations.find((op) => op.name === name);
- if (op) {
- return op;
- }
- }
- return { name };
- }
- };
- mlir.Error = class extends Error {
- constructor(message) {
- super(message);
- this.name = 'Error loading MLIR model.';
- }
- };
- export const ModelFactory = mlir.ModelFactory;
|