Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • panfrost/mesa
  • lima/mesa
  • anarsoul/mesa
  • mesa/mesa
  • kwg/mesa
  • curan/mesa
  • kusma/mesa
  • gfxstrand/mesa
  • ajax/mesa
  • FireBurn/mesa
  • bnieuwenhuizen/mesa
  • tarceri/mesa
  • chuckatkins/mesa
  • cmarcelo/mesa
  • anholt/mesa
  • rantogno/mesa
  • gerddie/mesa
  • keithp/mesa
  • chema/mesa
  • airlied/mesa
  • majanes/mesa
  • craftyguy/mesa
  • gurchetansingh/mesa
  • mvicomoya/mesa
  • kallisti5/mesa
  • bochecha/mesa
  • pauk.denis/mesa
  • dbaker/mesa
  • krh/mesa
  • jasuarez/mesa
  • sagarghuge/mesa
  • ofourdan/mesa
  • ickle/mesa
  • zzoon/mesa
  • Keenuts/mesa
  • tpohjola/mesa
  • velurimithun/mesa
  • cwabbott0/mesa
  • Richard_Yunchao/mesa
  • bkmgit/mesa
  • sima/mesa
  • GL/mesa
  • frkoenig/mesa
  • adelva1984/mesa
  • linyaa/mesa
  • slavslav/mesa
  • daniels/mesa
  • rellla/mesa
  • hakzsam/mesa
  • lyudess/panfrost-mesa
  • narmstrong/mesa-lima
  • robertfoss/mesa
  • nh/mesa
  • carlosg/mesa
  • jvesely/mesa
  • xlin16/mesa
  • narmstrong/panfrost-mesa
  • narmstrong/mesa
  • david.hanna11/mesa
  • idr/mesa
  • jljusten/mesa
  • antonovitch/mesa
  • arnomessiaen/mesa
  • mattst88/mesa
  • tpalli/mesa
  • mmha/mesa
  • llandwerlin/mesa
  • zumbi/mesa
  • rhyskidd/mesa
  • vlee/mesa
  • pmoreau/mesa
  • coypoop/mesa
  • jturney/mesa
  • strassek/mesa
  • olv/mesa
  • flto/mesa
  • iglosiggio/mesa
  • frohlich/mesa
  • karolherbst/mesa
  • elima/mesa
  • AtoningUnifex/mesa
  • tlonnber/mesa
  • samuelig/mesa
  • apinheiro/mesa
  • sjoerd/mesa
  • ahota/mesa
  • robh/mesa
  • daniel-schuermann/mesa
  • alistair23/mesa
  • linkmauve/mesa
  • elongbug/mesa
  • Haxk20/mesa
  • bartoldeman/mesa
  • daenzer/mesa
  • tjaalton/mesa
  • dhewg/mesa
  • Venemo/mesa
  • alyssa/mesa
  • Hi-Angel/mesa
  • carnaval/mesa
  • victure86/mesa
  • robclark/mesa
  • Vivek/mesa
  • tutankhamen/mesa
  • hopetech/mesa
  • marmeladema/mesa
  • ibriano/mesa
  • tanty/mesa
  • Oschowa/mesa
  • nroberts/mesa
  • zhanglei002/mesa
  • pal1000/mesa
  • xxxbxxx/mesa
  • lyudess/mesa
  • bentiss/mesa
  • afrantzis/mesa
  • yuq825/mesa
  • jamesxio/mesa
  • lkundrak/mesa
  • aphogat/mesa
  • ZeGentzy/mesa
  • aqxa1/mesa
  • mareko/mesa
  • austriancoder/mesa
  • jrtc27/mesa
  • skirk/mesa
  • grmat/mesa
  • heinrich.fink/mesa
  • el_christianito/mesa
  • davidriley/mesa
  • starnight/mesa
  • davidbepo/mesa
  • pendingchaos/mesa
  • icenowy/mesa
  • dominikd/mesa
  • zmike/mesa
  • tzimmermann/mesa
  • enunes/mesa
  • noblock/mesa
  • pzanoni/mesa
  • hygonsoc/mesa
  • lrusak/mesa
  • cap/mesa
  • pepp/mesa
  • axeldavy/mesa
  • domen55/mesa
  • john.stultz/mesa
  • jadahl/mesa
  • 10110111/mesa
  • alexvillacislasso/mesa
  • jzielins/mesa
  • pmanolova/mesa
  • nicholasbishop/mesa
  • david.boddie.puri.sm/mesa
  • agx/mesa
  • krzysztof.raszkowski/mesa
  • KhaledEmaraDev/mesa
  • itoral/mesa
  • shadeslayer/mesa
  • kszaq/mesa
  • haihao/mesa
  • wanderman.luck/mesa
  • Chunming-Zhou/mesa
  • emersion/mesa
  • agoldmints/mesa
  • marex/mesa
  • renchenglei/mesa
  • dbehr/mesa
  • zhen/mesa
  • ndufresne/mesa
  • abergmeier/mesa
  • grimkriegor/mesa
  • dlehman25/mesa
  • mmenzyns/mesa
  • jorgenatz/mesa
  • lostgoat/mesa
  • xexaxo/mesa
  • lynxeye/mesa
  • roman.stratiienko/mesa
  • brianp/mesa
  • Ericson2314/mesa
  • lfrb/mesa
  • dongwonk/mesa
  • pH5/mesa
  • q66/mesa
  • lepton/mesa
  • nia/mesa
  • AmanPatel599/mesa
  • currojerez/mesa
  • mol/mesa
  • danyspin97/mesa
  • zeising/mesa
  • mvlad/mesa
  • ssbertilson/mesa
  • Tofe/mesa
  • tomeu/mesa
  • kenmays/mesa
  • smbarber/mesa
  • dgstevens/mesa
  • YaLTeR/mesa
  • issor.oruam/mesa
  • cubanismo/mesa
  • bnf/mesa
  • libcg/mesa
  • bbrezillon/mesa
  • y2kenny/mesa
  • leandrohrb/mesa
  • psii/mesa
  • maos20008/mesa
  • MarijnS95/mesa
  • urjaman/mesa
  • banzr/mesa
  • rmader/mesa
  • daniels-test/mesa
  • thongthai/mesa
  • leoliu/mesa
  • baryluk/mesa
  • macieksolinski49/mesa
  • shikhar394/mesa
  • awatry/mesa
  • lordheavym/mesa
  • degasus/mesa
  • rg3igalia/mesa
  • miguelecasassanchez/mesa
  • italonicola/mesa
  • luis.mendes/mesa
  • bcrocker/mesa
  • pcercuei/mesa
  • Ikke/mesa
  • pelwell/mesa
  • John-Gee/mesa
  • Plagman/mesa
  • liyi42/mesa
  • Behem0th/mesa
  • kraj/mesa
  • abordado/mesa
  • romangg/mesa
  • sonny/mesa
  • ruanc/mesa
  • icecream95/mesa
  • tpyra/mesa
  • mariogrip/mesa
  • avdgrinten/mesa
  • Stephan/mesa
  • imbens/mesa
  • dh/mesa
  • stev47/mesa
  • haasn/mesa
  • RaspberryPiFan/mesa
  • sthibaul/mesa
  • zahraee.sm/mesa
  • brkho/mesa
  • lemon.py/mesa
  • sonicadvance1/mesa
  • Lasse/mesa
  • zhangjie/mesa
  • djogorchock/mesa
  • ShirishS/mesa
  • zhuravlev1134/mesa
  • castout/mesa
  • martin.fuzzey/mesa
  • idas/mesa
  • apesch/mesa
  • ngcortes/mesa
  • manu/mesa
  • JAYL/mesa
  • neochapay/mesa
  • w-flo/mesa
  • timkrause/mesa
  • ddavenport/mesa
  • thomash/mesa
  • andrzejuk.szymon/mesa
  • bkuhls/mesa
  • julianwi/mesa
  • aroth-fastprotect/mesa
  • DadSchoorse/mesa
  • berolinux/mesa
  • charmainel/mesa
  • shawn.p.huang/mesa
  • paska/mesa
  • hjl.tools/mesa
  • sroland/mesa
  • aacid/mesa
  • frog/mesa
  • scott-ph/mesa
  • awilfox/mesa
  • nyanmisaka/mesa
  • stalkerg/mesa
  • Cwiiis/mesa
  • imirkin/mesa
  • Bizjak/mesa
  • psreport/mesa
  • duncan.hopkins/mesa
  • maxice8/mesa
  • orbea/mesa
  • eschwartz/mesa
  • edmondo/mesa
  • jrfonseca/mesa
  • bbarladian/mesa
  • bhenden/mesa
  • alucowie/mesa
  • RAOF/mesa
  • leigh123linux/mesa
  • gtucker/mesa
  • LouisLi/mesa
  • sjnewbury/mesa
  • saavedra.pablo/mesa
  • ascent/mesa
  • wyqkp/mesa
  • Shao-Feng/mesa
  • jzhums/mesa
  • plevine457/mesa
  • tobiasjakobi/mesa
  • jpalus/mesa
  • yuzaipiaofei/mesa
  • jenatali/mesa
  • abhishek4/mesa
  • njha/mesa
  • konradybcio/mesa
  • mslusarz/mesa
  • fents290/mesa
  • edb/mesa
  • chewitt/mesa
  • ckolivas/mesa
  • blaztinn/mesa
  • cooperch/mesa
  • hikiko/mesa
  • adityaatluri/mesa
  • AniLeo/mesa
  • JibbityJobbity/mesa
  • rashedabdeltawab/mesa
  • ashafer/mesa
  • cme3000/mesa
  • farnoy/mesa
  • asheplyakov/mesa
  • kjliew/mesa
  • hanno/mesa
  • mikeroyal/mesa
  • ssingh/mesa
  • vitalyp/mesa
  • jpark37/mesa
  • mkasprza/mesa
  • WGH/mesa
  • mtmkls/mesa
  • vliaskov/mesa
  • kleinerm/mesa
  • hafixo/mesa
  • SureshG/mesa
  • mntmn/mesa
  • jbeich/mesa
  • yaongtime/mesa
  • xdontwox/mesa
  • jsg/mesa
  • EasyIP2023/mesa
  • Satyajit/mesa
  • jmunhoz/mesa
  • luigi.santivetti/mesa
  • italove/mesa
  • shawnguo/mesa
  • ricardoquesada/mesa
  • bcheng/mesa
  • Fahien/mesa
  • chrisf/mesa
  • alexhenrie/mesa
  • Yaong/mesa
  • gtskhadadze83/mesa
  • brightclark/mesa
  • gongmingqing/mesa
  • chturne/mesa
  • fjdegroo/mesa
  • dwg/mesa
  • hattrickcr7/mesa
  • Corngood/mesa
  • ccallawa/mesa
  • nh2/mesa
  • neobrain/mesa
  • yogeshmohan/mesa
  • felixonmars/mesa
  • hch12907/mesa
  • kode54/mesa
  • Sudland/mesa
  • woodychow/mesa
  • pichika/mesa
  • jschueller/mesa
  • tagr/mesa
  • Yang/mesa
  • ChaojiangLuo/mesa
  • vivekvpandya/mesa
  • yshui/mesa
  • av.linux.dev/mesa
  • DPA/mesa
  • m.tretter/mesa
  • mcoffin/mesa
  • johnkeeping/mesa
  • zhuyong/mesa
  • wareyang/mesa
  • jbates/mesa
  • boyzhang/mesa
  • padovan/mesa
  • xxmitsu/mesa
  • skarczew/mesa
  • alimon/mesa
  • RavenGSD/mesa
  • tguillem/mesa
  • mupuf/mesa
  • rui/mesa
  • ryanneph/mesa
  • Igortorrente/mesa
  • randrianasulu/mesa
  • dmrlawson/mesa
  • Jsn2win/mesa
  • krupatel/mesa
  • sahouston/mesa
  • tangm/mesa
  • qarmin/mesa
  • alex.kanavin/mesa
  • johnbradstreet/mesa
  • igor.v.kovalenko/mesa
  • OnikenX/mesa
  • martell/mesa
  • themaister/mesa
  • BranDougherty/mesa
  • MJDSys/mesa
  • segfaultxavi/mesa
  • pleath/mesa
  • tongwang-en/mesa
  • ptt-en/mesa
  • JoelLinn/mesa
  • Danil/mesa
  • theogen/mesa
  • luigifcruz/mesa
  • chenli/mesa
  • vjaquez/mesa
  • apopple/mesa
  • theozzhh79/mesa
  • BryanQuigley/mesa
  • rroohhh/mesa
  • caramelli/mesa
  • mforney/mesa
  • yannik/mesa
  • twoerner/mesa
  • zhangguoqing.kernel/mesa
  • swartzlib7/mesa
  • sergey042011/mesa
  • nicuborta123/mesa
  • lyintel/mesa
  • q4a/mesa-panfrost
  • q4a/mesa
  • alexander.kapshuk/mesa
  • stephane.marchesin/mesa
  • jeremyhu/mesa
  • bl4ckb0ne/mesa
  • geecandrey/mesa
  • wenxiaoming/mesa
  • billkris.ms/mesa
  • ella/mesa
  • luporl/mesa
  • danielzgtg/mesa
  • frediz/mesa
  • lubosz/mesa
  • gio/mesa
  • rsmith/mesa
  • neo.tzion/mesa
  • pundiramit/mesa
  • zongzi13545329/mesa
  • Mic92/mesa
  • mwnn/mesa
  • tamara-schmitz/mesa
  • luckyxxl/mesa
  • jcline/mesa
  • akihiko.odaki/mesa
  • yugang/mesa
  • lljgithub/mesa
  • leeball/mesa
  • dianders/mesa
  • tina/mesa
  • jschwartzentruber/mesa
  • maccraft/mesa
  • vsyrjala/mesa
  • timothee.chabat/mesa
  • adirat/mesa
  • egalli/mesa
  • zhenhanintel/mesa
  • ccrtnsp/mesa
  • tintou/mesa
  • mhenning/mesa
  • blaws05/mesa
  • Zamundaaa/mesa
  • retrooper/mesa
  • ernstp/mesa
  • xoddark/mesa
  • ericonr/mesa
  • axylp/mesa
  • Edgeworth/mesa
  • siyueyinghua/mesa
  • asahi/mesa
  • JN-Chen/mesa
  • alshi-msft/mesa
  • williamvds/mesa
  • kupper.pa/mesa
  • A-w-x/mesa
  • Lucretia/mesa
  • shannonliu011/mesa
  • bastianbeischer/mesa
  • CosineMath/mesa
  • zzyiwei/mesa
  • lucmann/mesa
  • spronovo/mesa
  • xgupta/mesa
  • oldherl/mesa
  • ezequielgarcia/mesa
  • mercuriete/mesa
  • lumag/mesa
  • msisov/mesa
  • jialina/mesa
  • Zakhrov/mesa
  • hexin.op/mesa
  • FrostKiwi/mesa
  • PaulKocialkowski/mesa
  • flightlessmango/mesa
  • arkamar/mesa
  • khfeng/mesa
  • mrisaacb/mesa
  • zhangxiaolin.a/mesa
  • chivay/mesa
  • Daft-Freak/mesa
  • iemaghni/mesa
  • amonakov/mesa
  • zhaofengli/mesa
  • mwalle/mesa
  • marcan/mesa
  • davidedmundson/mesa
  • qihangkong/mesa
  • philn/mesa
  • lfelipe/mesa
  • Mystro256/mesa
  • wicastC/mesa
  • Gofman/mesa
  • nurmukhametov/mesa
  • ishitatsuyuki/mesa
  • phomes/mesa
  • AlexZ/mesa
  • Ermine/mesa
  • mwen/mesa
  • zboszor/mesa
  • Tooniis/mesa
  • lorenz/mesa
  • tantan/mesa
  • liamwhite/mesa
  • dyang23/mesa
  • aswarup/mesa
  • shengyao/mesa
  • zhuyl/mesa
  • bvarner/mesa
  • jeff_shuai/mesa
  • jpewhacker/mesa
  • suijingfeng/mesa
  • svenny/mesa
  • xantares/mesa
  • nicuborta/mesa
  • rakko/mesa
  • WangChuan/mesa
  • thomas.wagner/mesa
  • janjanmostafa8/mesa
  • gallo/mesa
  • gh6h56j85ihg/mesa
  • quantum/mesa
  • liuyujun/mesa
  • edman007/mesa
  • mattvchandler/mesa
  • ebaker/mesa
  • mherrb/mesa
  • alxu/mesa
  • mikezackles/mesa
  • ArvindYadav/mesa-amd
  • satmandu/mesa
  • gdevi/mesa
  • gawin/mesa
  • rbrune/mesa
  • mwezdeck/mesa
  • derekf/mesa
  • amos/mesa
  • sadlerap/mesa
  • kbrenneman/mesa
  • augustin.zidek/mesa
  • arichardson/mesa
  • jasberc/mesa
  • Pokechu22/mesa
  • h0tc0d3/mesa
  • HayashiEsme/mesa
  • Lone_Wolf/mesa
  • maniraj87/mesa
  • bbeckett/mesa
  • scholzi1980/mesa
  • ssidhart/mesa
  • mstoeckl/mesa
  • shanshengwang/mesa
  • lukvec27/mesa
  • cristicc/mesa
  • autumnontape/mesa
  • smcv/mesa
  • elmarco/mesa
  • jchen10/mesa
  • rbernon/mesa
  • Cherser-s/mesa
  • mhillenbrand/mesa
  • t.clastres/mesa
  • jxzgithub/mesa
  • Daasin/mesa-bak
  • JoseExposito/mesa
  • stefandoesinger/mesa
  • lygstate/mesa
  • larumbe/mesa
  • renatopereyra/mesa
  • raininggibs/mesa
  • sin3point14/mesa
  • chris8136393/mesa
  • microlinux/mesa
  • Markus-included/mesa
  • Jason2013/mesa
  • xndcn/mesa
  • ryan.krattiger/mesa
  • badsector/mesa
  • hamarb123/mesa
  • frankbinns/mesa
  • yulang/mesa
  • mmp.dux/mesa
  • charles-lunarg/mesa
  • gt23263tq/mesa
  • KonstantinSeurer/mesa
  • ondracka/mesa
  • biju.das.au/mesa
  • apauk/mesa
  • zwuj/mesa
  • samael/mesa
  • croberts81/mesa
  • nanokatze/mesa
  • kwk/mesa
  • StefanBruens/mesa
  • parasyte/mesa
  • kvark/mesa
  • longxin2019/mesa
  • jorcrous/mesa
  • igor.torrente/mesa
  • AidoP/mesa
  • shmerl/mesa
  • preda/mesa
  • andrey-konovalov/mesa
  • peterh/mesa
  • otaviobp/mesa
  • zeekim/mesa
  • vkoul/mesa
  • bigon/mesa
  • beviu/mesa
  • alatiera/mesa
  • cvurdige/chandan-mesa
  • jesse.zhang/mesa
  • MTCoster/mesa
  • digetx/mesa
  • gbelgurr/mesa
  • akien/mesa
  • tales-aparecida/mesa
  • jocelyn/mesa
  • helen.fornazier/mesa
  • wingdeans/mesa
  • superm1/mesa
  • justsid/mesa
  • mallemsalam/mesa
  • rajnesh-kanwal/mesa
  • illiliti/mesa
  • ghishadow/mesa
  • YusufKhan-gamedev/mesa
  • jbillingsley/mesa
  • illwieckz/mesa
  • sergi/mesa
  • jevolk/mesa
  • arunpravin24/mesa
  • jeffk/mesa
  • bkylerussell/mesa
  • jjjighg/mesa
  • Ken2022/mesa
  • skywing/mesa
  • sharny/mesa
  • vitorhnn/mesa
  • a-wai/mesa
  • xperia64/mesa
  • yuesOctober/mesa
  • nouveau/mesa
  • zzag/mesa
  • yurikoles/mesa
  • dawnhan/mesa
  • garnet/mesa
  • andrescj-chromium/mesa
  • aleasto/mesa
  • zhangn1985/mesa
  • fangbaohui/mesa
  • tiago.shibata/mesa
  • flynnjiang/mesa
  • dikidiki26138/mesa
  • bnagappa/mesa
  • zhoubm/mesa
  • wuwh91/mesa
  • chandupokuru/mesa
  • JoniSt/mesa
  • pixelcluster/mesa
  • kuove/mesa
  • asahilina/mesa
  • evelez/mesa
  • skeller/mesa
  • limingchina1/mesa
  • panos-lunarg/mesa
  • zehortigoza/mesa
  • i509VCB/mesa
  • bigbear1385/mesa
  • tbettler/mesa
  • skyostil/mesa
  • Soroush/mesa
  • newbluemoon/mesa
  • jdtatz/mesa
  • chuansheng/mesa
  • jjun.9523/mesa
  • kai/mesa
  • jheaff1/mesa
  • TSnake41/mesa
  • flakylawww/mesa
  • alexfanqi/mesa
  • dakr/mesa
  • TheSpydog/mesa
  • ikshwaku.chauhan_amd/ikshwaku-mesa
  • justonli/mesa
  • Frogging101/mesa
  • sarojk/mesa
  • sjfricke/mesa
  • PixelyIon/mesa
  • adolfintel/mesa
  • vikramsingh0204/mesa
  • riteo/mesa
  • xphillyx/mesa
  • markboth/mesa
  • MaxKellermann/mesa
  • bitstreamout/mesa-fork
  • kg233/mesa-rkuixi-fork
  • zhililab/mesa
  • isinyaaa/mesa
  • swick/mesa
  • JeremyRand/mesa
  • thomas.devoogdt/mesa
  • Thaodan/mesa
  • delphij/mesa
  • zhanglianjie-163/mesa
  • d.ivo/mesa
  • os369510/mesa
  • shanep/mesa
  • alikates/mesa
  • Smokey365/mesa
  • nanotwerp/nanomesa
  • Beyley/mesa-glide
  • semjon00/mesa
  • andri/mesa
  • jadeja14890/mesa
  • bydavenet/mesa
  • mainiomano/mesa
  • flibitijibibo/mesa
  • endlesspring/mesa
  • kulikjak/mesa
  • lun/mesa
  • RandomShaper/mesa
  • georgeouzou/mesa
  • christophe-lunarg/mesa
  • Reiver/mesa
  • morphis/mesa
  • fvogt/mesa
  • Erdroy/mesa
  • antmonteiro/mesa
  • deathmist/mesa
  • michaelonchrome/mesa
  • sukhatri/mesa-sukhatri
  • tommydrum/mesa
  • jenneron/mesa
  • momo666750/mesa
  • antoniospg100/mesa
  • noah1510/mesa
  • mairacanal/mesa
  • rmckeever/mesa
  • jfoxwoosh/mesa
  • Triang3l/mesa
  • marysaka/mesa
  • morrisoncave56/mesa
  • MayeulC/mesa
  • himanshu.nayak/mesa-him
  • 8infy/mesa
  • ayaka/mesa
  • Chiitoo/mesa
  • jkqxz/mesa
  • zxs-un/mesa-riscv
  • zxs-un/mesa
  • pyuan/mesa
  • GermanAizek/mesa
  • StaticRocket/mesa
  • dwlsalmeida/mesa
  • heftig/mesa
  • mgorny/mesa
  • WebsterEndymion/mesa
  • qyliss/mesa
  • stilriv/mesa
  • zzcc231/mesa
  • davidwu2/mesa
  • iaguis/mesa
  • evan.quan/mesa
  • martty/mesa
  • nowrep/mesa
  • osy/mesa
  • Azedin1111/mesa
  • Quipyowert2/mesa
  • nihui/mesa
  • libyahhh519/mesa
  • filip.binkiewicz/mesa
  • Jing/mesa
  • ilikerackmounts/mesa
  • EmperorPenguin18/mesa
  • yixie/mesa
  • gustafullberg/mesa
  • etaash.mathamsetty/mesa
  • hongaoo/mesa
  • hlinander/mesa
  • crab2313/mesa
  • msahastr/mesa
  • driver1998/mesa
  • amber/mesa
  • CounterPillow/mesa
  • shangyatsen/mesa
  • rosefromthedead/mesa
  • ItzSwirlz/mesa
  • kwizart/mesa
  • sobkas/mesa
  • timvp/mesa
  • philipl/mesa
  • turol/mesa
  • antonino/mesa
  • valpackett/mesa
  • X547/mesa
  • alanc/mesa
  • soonray8989/mesa
  • xxlxng25/mesa
  • mivanchev/mesa
  • wangxiaoming321/mesa
  • yangyangdaji/mesa
  • atmitch75/mesa
  • Hazematman/mesa
  • gfx-ci-bot/mesa
  • zaps166/mesa
  • CFSworks/mesa
  • lynne/mesa
  • lorn10/mesa
  • MrPurple666/mesa
  • pino/mesa
  • ccawley2011/mesa
  • lukeyyyy/mesa
  • maririn312/mesa
  • mohamexiety/mesa
  • chenrui/mesa
  • gmitrano/mesa
  • ndeshpan/mesa
  • cheako/mesa
  • xzhan34/mesa
  • heitbaum/mesa
  • LunaFoxgirlVT/mesa
  • tengjinchung/mesa
  • wujiangGitHub/mesa
  • K900/mesa
  • lb90/mesa
  • DodoGTA/mesa-nvk
  • moiman/mesa
  • Rui511/mesa
  • andrealmeid/mesa
  • Macdu/mesa
  • raun/mesa
  • jannau/mesa
  • ouchuan/mesa
  • jdknight/mesa
  • weibinwu/mesa
  • davidre/mesa
  • fcui/mesa
  • rawoul/mesa
  • stolk/mesa
  • alan_previn_intel/mesa-alan-previn-features
  • Yogayaojia/mesa
  • lucaweiss/mesa
  • JiadongZhu/mesa
  • asriniva/mesa
  • nitin.reddy88/mesa
  • Sterophonick/mesa
  • Teddy-Kun/mesa
  • donaldrobson/mesa
  • jide/mesa
  • kkartaltepe/mesa
  • mbouron/mesa
  • pkubaj/mesa
  • nora/mesa
  • dvrogozh1/mesa
  • MrRobbin/mesa
  • Ristovski/mesa
  • veerabadhran/mesa
  • Kiskae/mesa
  • vaishali/mesa
  • gbeatty/mesa
  • doraskayo/mesa
  • benjaminl/mesa
  • Blisto/mesa
  • baikaishiuc/mesa
  • shanminchao/mesa
  • Samsuper/mesa
  • phodina/mesa
  • diederik/mesa
  • adamjoseph/mesa
  • xry111/mesa
  • dburkov05/mesa
  • penguin42/mesa
  • tpambor/mesa
  • stepri/mesa
  • vigneshraman/mesa
  • Julia/mesa
  • bishoman123/mesa
  • rosasco-wk/mesa
  • jani/mesa
  • smaeul/mesa
  • QwertyChouskie/mesa
  • K1ngst0m/mesa
  • conversy/mesa
  • WinLinux1028/mesa-fix
  • Nefsen402/mesa
  • tripzero/mesa
  • max8rr8/mesa
  • MaxVerevkin/mesa
  • BinHani/mesa
  • spencercw/mesa
  • jazzfool/mesa
  • renfeng/mesa
  • 246tnt/mesa
  • twisted89/mesa
  • iyes/mesa
  • krz/mesa
  • Conan_Kudo/mesa
  • kociap/mesa
  • huangrui666/mesa
  • mildsunrise/mesa
  • ids1024/mesa
  • VladSchillerimgtec/mesa
  • zwang20/mesa
  • LaserEyess/mesa
  • carsten.haitzler/mesa
  • doitsujin/mesa
  • fzatlouk/mesa
  • afd/mesa
  • ratatouillemorde/mesa
  • reillybrogan/mesa
  • vanvugt/mesa
  • ralphcampbell/mesa
  • llyyr/mesa
  • amdrexu/mesa
  • oreaus/mesa
  • liucong2/mesa
  • wangra/mesa
  • asuka-mio/mesa
  • slp/mesa
  • vimproved/mesa
  • antco/mesa
  • mtasaka/mesa
  • SoshyuCelia/mesa
  • cheyang/mesa
  • hmann/mesa
  • zfigura/mesa
  • haopiliu/upstream_vpelib
  • Sid127/mesa
  • jwillikers/mesa
  • 6by9/mesa
  • eerii/mesa
  • ElectrodeYT/mesa
  • rilian-la-te/mesa
  • yrlf/mesa
  • sunzhguy/mesa
  • andrewgazizov/mesa
  • nanonyme/mesa
  • lpy/mesa
  • parona/mesa
  • cgbowman/mesa
  • YukariChiba/mesa
  • jexposit/mesa
  • kiroma/mesa
  • theoparis/mesa
  • Tele42/mesa
  • Dudemanguy/mesa
  • blu/mesa
  • myelin/mesa
  • strongtz/mesa
  • detlev.c/mesa
  • MollySophia/mesa
  • dabrain34/mesa
  • jnoorman/mesa
  • leio/mesa
  • Reflux5231/mesa
  • lihongtao/mesa
  • anthony-linaro/mesa
  • tshikaboom/mesa
  • ramenguy99/mesa
  • vhorinek/mesa
  • GKraats/mesa_alu
  • russelltg/mesa
  • rgallaispou/mesa
  • pleasurefish/mesa
  • SupSuper/mesa
  • chaos_princess/mesa
  • chrisduerr/mesa
  • MouriNaruto/mesa
  • absurdistcode/mesa
  • daoxiang.gong/mesa
  • chyyran/mesa
  • Snowiiii/mesa
  • rooq/mesa
  • colinmarc/mesa
  • rpavlik/mesa
  • DarkFire01/mesa
  • VladimirTechMan/mesa
  • Anzel/mesa
  • catvinyl/mesa
  • h0lyalg0rithm/mesa
  • fweimer/mesa
  • Anorak/mesa
  • imagination/mesa
  • javierm/mesa
  • ericsmith/mesa
  • tranquillity-codes/mesa
  • dan.g.tob/mesa
  • ekurzinger/mesa
  • zdobersek/mesa-fork
  • dodger_mink/mesa
  • tursulin/mesa-llandwerlin
  • Vladimir-A/mesa
  • surafel911/mesa
  • thfr/mesa
  • City-busz/mesa
  • dtobolik/mesa
  • sarthakbhatt/mesa
  • Jorropo/mesa
  • rankinc/mesa
  • sushmave/mesa
  • IsaacMarovitz/mesa
  • vnovotny/mesa
  • smit-mayani/mesa
  • Yifan.Zhang/merge-gfx-patch
  • K0bin/mesa
  • qbojj/mesa
  • Valentine/mesa
  • SludgePhD/mesa
  • tong1wu/mesa
  • npopov/mesa
  • alexmrqt/mesa
  • chipitsine/mesa
  • Mstrodl/mesa
  • whot/mesa
  • cmsanta/mesa
  • kaiwenjon/mesa
  • ystreet/mesa
  • sixtyfourktec/mesa
  • yu-re-ka/mesa
  • vorporeal/mesa
  • MischaBaars/mesa
  • peytolee/mesa-radeonsi
  • zack/mesa
  • nadaouf/mesa
  • ahuillet/mesa
  • oscarbg/mesa
  • thsr/mesa
  • JCWasmx86/mesa
  • gonsolo/mesa
  • msizanoen1/mesa
  • subzeroxr3i/mesa
  • NaveenKumar/mesa
  • Civiloid/mesa
  • phreer/mesa
  • amazingfate/mesa
  • zsnow/mesa
  • aperezdc/mesa
  • cepsylon/mesa
  • Pipetto-crypto/mesa
  • yuanqingxiang520/mesa
  • jvutukur/mesa
  • zachary.battleman/mesa
  • ncopa/mesa
  • ao2/mesa
  • e.rosebrock/mesa
  • bbhtt/mesa
  • marvelrenju1/mesa
  • hmtheboy154/mesa
  • asurati/mesa
  • ahokananoevolution/mesa
  • FredFunk/mesa
  • zturtleman/mesa
  • MingcongBai/mesa
  • dbrouwer/mesa
  • wswsmao/mesa
  • dlundqvist/mesa
  • DDoSQc/mesa
  • mahkoh/mesa
  • TimHuang/mesa
  • jeremyg/mesa
  • n3rdopolis/mesa
  • CreativeCylon/mesa
  • rjodin/mesa
  • mastag/mesa
  • alihomafar/mesa
  • smartavionics/mesa
  • esdrastarsis/mesa
  • spottumu/mesa
  • ShenghuaLinINTEL/mesa
  • Calandracas/mesa
  • dougg3/mesa
  • jsimmons1/mesa
  • mcatanzaro/mesa
  • yinjiyao/mesa
  • Zhaojiale/mesa
  • agd5f/mesa
  • Jelgnum/add-RADV_UMR_EXTRA_ARGS-docs
  • jopadan/mesa
  • DeeptiPatil/mesa
  • alvinhochun/mesa
  • ShriramShastry/mesa
  • Jaakkonen/mesa
  • pmenzel/mesa
  • raki-huawei/mesa
  • xen0n/mesa
  • Zer0xFF/mesa
  • twitt-thornwaite/mesa
  • larsivsi/mesa
  • clemy/mesa
  • demarchi/mesa
  • calder/mesa
  • lemonflynn/mesa
  • asrivats/mesa
  • ketil.johnsen/mesa
  • Trigger.Huang/mesa
  • brad0/mesa
  • RussellLiu/mesa
  • quic_lkondred/mesa
  • mombasa/mesa
  • jmacnak/mesa
  • company/mesa
  • jna/mesa
  • utkuiseri-arm/mesa
  • ritalat/mesa
  • dsvensson/mesa
  • jules.blok/mesa
  • Drakulix/mesa
  • iforbes/mesa
  • zhanghe9702/mesa
  • slouken/mesa
  • chihchie/mesa
  • Guy1524/mesa
  • Hust-YL/hust-yl-mesa
  • ziyao233/mesa
  • MegWATTT/mesa
  • genhoayi/mesa
  • me-cafebabe/mesa
  • EBADBEEF/mesa
  • manueldun/mesa
  • xclaesse/mesa
  • chenyche_codegen/mesa
  • lcagustini/mesa
  • jkim/mesa
  • marc-hb/mesa
  • benh/mesa
  • shreeya/mesa
  • vladly/mesa
  • poweif/mesa
  • zeux/mesa
  • whao/mesa
  • Yanfeng-Mi/mesa
  • fuel-pcbox/mesa
  • sewn/mesa
  • JohannesKauffmann/mesa
  • barracuda156/mesa
  • provod/mesa
  • xytovl/mesa
  • dtgs1208/mesa
  • mbrost/mesa-bind-queues
  • fknfilewalker/mesa
  • benjarobin/mesa
  • jameshogan/mesa
  • caciottax86/mesa
  • Gobrosse/mesa
  • runderwo/mesa
  • orowith2os/oro-mesa
  • hnez/mesa
  • gizmo98/mesa-hwdb
  • tjk/mesa
  • sarbes/mesa
  • GabiAle97/mesa-termux
  • ashleysmithcol/mesa
  • SnowyCoder/mesa
  • aswolfers/mesa
  • InFamousBear/mesa
  • adalessandro/mesa
  • thesamesam/mesa
  • zzxyb/mesa
  • rpsingh.amd/mesa
  • lool/mesa
  • molinari/mesa
  • chrpil/mesa
1208 results
Show changes
Commits on Source (83)
Showing
with 189 additions and 101 deletions
......@@ -474,7 +474,7 @@ Vulkan 1.2 -- all DONE: anv, nvk, tu, vn
VK_EXT_descriptor_indexing DONE (anv, dzn, lvp, nvk, radv, tu, vn)
VK_EXT_host_query_reset DONE (anv, hasvk, lvp, nvk, panvk, pvr, radv, tu, v3dv, vn)
VK_EXT_sampler_filter_minmax DONE (anv, lvp, nvk, panvk/v10+, radv, tu, vn)
VK_EXT_scalar_block_layout DONE (anv, dzn, hasvk, lvp, nvk, pvr, radv/gfx7+, tu, vn, v3dv/vc7+)
VK_EXT_scalar_block_layout DONE (anv, dzn, hasvk, lvp, nvk, panvk, pvr, radv/gfx7+, tu, vn, v3dv/vc7+)
VK_EXT_separate_stencil_usage DONE (anv, dzn, hasvk, lvp, nvk, radv, tu, v3dv, vn)
VK_EXT_shader_viewport_index_layer DONE (anv, hasvk, lvp, nvk, radv, tu, vn)
......
......@@ -103,6 +103,7 @@
#define AMDGPU_GFX1150_RANGE 0x01, 0x40 //# 1 <= x < 64
#define AMDGPU_GFX1151_RANGE 0xC0, 0xFF //# 192 <= x < 255
#define AMDGPU_GFX1152_RANGE 0x40, 0x50 //# 64 <= x < 80
#define AMDGPU_GFX1153_RANGE 0x50, 0xC0 //# 80 <= x < 192
#define AMDGPU_PHOENIX1_RANGE 0x01, 0x80 //# 1 <= x < 128
#define AMDGPU_PHOENIX2_RANGE 0x80, 0xC0 //# 128 <= x < 192
#define AMDGPU_HAWK_POINT1_RANGE 0xC0, 0xF0 //# 192 <= x < 240
......@@ -181,6 +182,7 @@
#define ASICREV_IS_GFX1150(r) ASICREV_IS(r, GFX1150)
#define ASICREV_IS_GFX1151(r) ASICREV_IS(r, GFX1151)
#define ASICREV_IS_GFX1152(r) ASICREV_IS(r, GFX1152)
#define ASICREV_IS_GFX1153(r) ASICREV_IS(r, GFX1153)
#define ASICREV_IS_PHOENIX1(r) ASICREV_IS(r, PHOENIX1)
#define ASICREV_IS_PHOENIX2(r) ASICREV_IS(r, PHOENIX2)
......
......@@ -5,12 +5,17 @@ dEQP-VK.pipeline.shader_object_unlinked_spirv.sampler.view_type.2d_array.format.
dEQP-VK.pipeline.monolithic.creation_cache_control.compute_pipelines.duplicate_single_recreate_no_caching
dEQP-VK.pipeline.monolithic.sampler.view_type.1d_array.format.r8_snorm.min_reduce.comp_r_g_b_a.max
dEQP-VK.pipeline.monolithic.sampler.view_type.1d_array.format.r8_srgb.mag_reduce.comp_r_g_b_a.min
dEQP-VK.pipeline.monolithic.sampler.view_type.1d_array.format.r8_srgb.min_reduce.comp_r_g_b_a.max
dEQP-VK.pipeline.monolithic.sampler.view_type.1d_array.format.r8_srgb.min_reduce.comp_r_zero_zero_zero.max
dEQP-VK.pipeline.monolithic.sampler.view_type.2d_array.format.d16_unorm.min_reduce.comp_identity_zero_zero_zero.min
dEQP-VK.pipeline.monolithic.sampler.view_type.2d_array.format.d16_unorm.min_reduce.comp_r_zero_zero_zero.min
dEQP-VK.pipeline.monolithic.sampler.view_type.2d_array.format.r8_srgb.min_reduce.comp_r_g_b_a.min
dEQP-VK.pipeline.monolithic.sampler.view_type.3d.format.r8_unorm.min_reduce.comp_identity_zero_zero_zero.min
dEQP-VK.pipeline.monolithic.sampler.view_type.cube_array.format.r16_snorm.min_reduce.comp_identity_zero_zero_zero.max
dEQP-VK.pipeline.monolithic.sampler.view_type.cube.format.d32_sfloat.min_reduce.comp_identity_zero_zero_zero.max
dEQP-VK.pipeline.monolithic.sampler.view_type.cube.format.r8_snorm.min_reduce.comp_r_zero_zero_zero.max
dEQP-VK.pipeline.monolithic.sampler.view_type.cube.format.r8_snorm.min_reduce.comp_r_zero_zero_zero.min
dEQP-VK.pipeline.shader_object_unlinked_binary.interface_matching.vector_length.out_ivec2_in_ivec2_member_of_structure_in_block_vert_tesc_tese_geom_out_frag_in
dEQP-VK.pipeline.shader_object_unlinked_spirv.interface_matching.vector_length.out_ivec3_in_ivec2_member_of_structure_in_block_vert_geom_out_frag_in
......
......@@ -5,3 +5,5 @@ dEQP-VK.ray_tracing_pipeline.pipeline_no_null_shaders_flag.*
# Nightly run expectations update
dEQP-VK.pipeline.monolithic.image.suballocation.sampling_type.combined.view_type.1d_array.format.r8_uint.count_1.size.3x1_array_of_6
dEQP-VK.pipeline.shader_object_unlinked_spirv.render_to_image.core.1d_array.mipmap.a2b10g10r10_unorm_pack32_d32_sfloat_s8_uint
dEQP-VK.api.copy_and_blit.copy_commands2.image_to_buffer.2d_images.mip_copies_bc7_unorm_block_64x192_2_layers_transfer
......@@ -857,6 +857,7 @@ bool ac_query_gpu_info(int fd, void *dev_p, struct radeon_info *info,
identify_chip(GFX1150);
identify_chip(GFX1151);
identify_chip(GFX1152);
identify_chip(GFX1153);
break;
case FAMILY_GFX12:
identify_chip(GFX1200);
......
......@@ -351,10 +351,8 @@ emit_ps_color_clamp_and_alpha_test(nir_builder *b, lower_ps_state *s)
}
static void
emit_ps_mrtz_export(nir_builder *b, lower_ps_state *s)
emit_ps_mrtz_export(nir_builder *b, lower_ps_state *s, nir_def *mrtz_alpha)
{
nir_def *mrtz_alpha = s->options->alpha_to_coverage_via_mrtz ? s->color[0][3] : NULL;
/* skip mrtz export if no one has written to any of them */
if (!s->depth && !s->stencil && !s->sample_mask && !mrtz_alpha)
return;
......@@ -740,12 +738,18 @@ emit_ps_null_export(nir_builder *b, lower_ps_state *s)
static void
export_ps_outputs(nir_builder *b, lower_ps_state *s)
{
nir_def *mrtz_alpha = NULL;
b->cursor = nir_after_impl(b->impl);
/* Alpha-to-coverage should be before alpha-to-one. */
if (!s->options->no_depth_export && s->options->alpha_to_coverage_via_mrtz)
mrtz_alpha = s->color[0][3];
emit_ps_color_clamp_and_alpha_test(b, s);
if (!s->options->no_depth_export)
emit_ps_mrtz_export(b, s);
emit_ps_mrtz_export(b, s, mrtz_alpha);
/* When non-monolithic shader, RADV export mrtz in main part (except on
* RDNA3 for alpha to coverage) and export color in epilog.
......
......@@ -62,6 +62,7 @@ const char *ac_get_family_name(enum radeon_family family)
CASE(GFX1150);
CASE(GFX1151);
CASE(GFX1152);
CASE(GFX1153);
CASE(GFX1200);
CASE(GFX1201);
#undef CASE
......@@ -207,6 +208,8 @@ const char *ac_get_llvm_processor_name(enum radeon_family family)
return "gfx1151";
case CHIP_GFX1152:
return "gfx1152";
case CHIP_GFX1153:
return "gfx1153";
case CHIP_GFX1200:
return "gfx1200";
case CHIP_GFX1201:
......
......@@ -125,6 +125,7 @@ enum radeon_family
CHIP_GFX1150,
CHIP_GFX1151,
CHIP_GFX1152,
CHIP_GFX1153,
/* GFX12 (RDNA 4) */
CHIP_GFX1200,
CHIP_GFX1201,
......
......@@ -1617,7 +1617,7 @@ chain_branches(asm_context& ctx, std::vector<uint32_t>& out, branch_info& branch
}
const unsigned block_offset = insert_at + code.size();
branch_instr = bld.sopp(aco_opcode::s_branch, target);
branch_instr = bld.sopp(aco_opcode::s_branch, 0);
emit_sopp_instruction(ctx, code, branch_instr, true);
insert_code(ctx, out, insert_at, code.size(), code.data());
......
......@@ -319,9 +319,9 @@ perform_barrier(wait_ctx& ctx, wait_imm& imm, memory_sync_info sync, unsigned se
if (bar_scope_lds <= subgroup_scope)
events &= ~event_lds;
/* in non-WGP, the L1 (L0 on GFX10+) cache keeps all memory operations
/* Until GFX12, in non-WGP, the L1 (L0 on GFX10+) cache keeps all memory operations
* in-order for the same workgroup */
if (!ctx.program->wgp_mode && sync.scope <= scope_workgroup)
if (ctx.gfx_level < GFX12 && !ctx.program->wgp_mode && sync.scope <= scope_workgroup)
events &= ~(event_vmem | event_vmem_store | event_smem);
if (events)
......
......@@ -8232,13 +8232,19 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
Temp local_ids[3];
 
/* Thread IDs are packed in VGPR0, 10 bits per component. */
for (uint32_t i = 0; i < 3; i++) {
if (i == 0 && ctx->shader->info.workgroup_size[1] == 1 &&
ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable) {
local_ids[i] = get_arg(ctx, ctx->args->local_invocation_ids);
} else if (i == 2 || (i == 1 && ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable)) {
local_ids[0] = get_arg(ctx, ctx->args->local_invocation_ids);
if (ctx->shader->info.workgroup_size[1] > 1 || ctx->shader->info.workgroup_size[2] > 1 ||
ctx->shader->info.workgroup_size_variable) {
unsigned size_x = ctx->shader->info.workgroup_size_variable
? 1024
: util_next_power_of_two(ctx->shader->info.workgroup_size[0]);
Temp mask = bld.copy(bld.def(s1), Operand::c32(size_x - 1));
local_ids[0] = bld.vop2(aco_opcode::v_and_b32, bld.def(v1), mask, local_ids[0]);
}
for (uint32_t i = 1; i < 3; i++) {
if (i == 2 || (i == 1 && ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable)) {
local_ids[i] =
bld.vop2(aco_opcode::v_lshrrev_b32, bld.def(v1), Operand::c32(i * 10u),
get_arg(ctx, ctx->args->local_invocation_ids));
......@@ -11019,7 +11025,8 @@ export_fs_mrt_color(isel_context* ctx, const struct aco_ps_epilog_info* info, Te
}
 
static void
export_fs_mrtz(isel_context* ctx, Temp depth, Temp stencil, Temp samplemask, Temp alpha)
export_fs_mrtz(isel_context* ctx, const struct aco_ps_epilog_info* info, Temp depth, Temp stencil,
Temp samplemask, Temp alpha)
{
Builder bld(ctx->program, ctx->block);
unsigned enabled_channels = 0;
......@@ -11062,7 +11069,7 @@ export_fs_mrtz(isel_context* ctx, Temp depth, Temp stencil, Temp samplemask, Tem
}
 
if (alpha.id()) {
assert(ctx->program->gfx_level >= GFX11);
assert(ctx->program->gfx_level >= GFX11 || info->alpha_to_one);
values[3] = Operand(alpha);
enabled_channels |= 0x8;
}
......@@ -13507,6 +13514,9 @@ select_ps_epilog(Program* program, void* pinfo, ac_shader_config* config,
 
Builder bld(ctx.program, ctx.block);
 
bool has_mrtz_alpha = einfo->alpha_to_coverage_via_mrtz && einfo->colors[0].used;
Temp mrtz_alpha;
Temp colors[MAX_DRAW_BUFFERS][4];
for (unsigned i = 0; i < MAX_DRAW_BUFFERS; i++) {
if (!einfo->colors[i].used)
......@@ -13520,22 +13530,24 @@ select_ps_epilog(Program* program, void* pinfo, ac_shader_config* config,
colors[i][c] = emit_extract_vector(&ctx, color, c, col_types == ACO_TYPE_ANY32 ? v1 : v2b);
}
 
/* Store MRTZ.a before applying alpha-to-one if enabled. */
if (has_mrtz_alpha && i == 0)
mrtz_alpha = colors[0][3];
emit_clamp_alpha_test(&ctx, einfo, colors[i], i);
}
 
bool has_mrtz_depth = einfo->depth.used;
bool has_mrtz_stencil = einfo->stencil.used;
bool has_mrtz_samplemask = einfo->samplemask.used;
bool has_mrtz_alpha = einfo->alpha_to_coverage_via_mrtz && einfo->colors[0].used;
bool has_mrtz_export =
has_mrtz_depth || has_mrtz_stencil || has_mrtz_samplemask || has_mrtz_alpha;
if (has_mrtz_export) {
Temp depth = has_mrtz_depth ? get_arg(&ctx, einfo->depth) : Temp();
Temp stencil = has_mrtz_stencil ? get_arg(&ctx, einfo->stencil) : Temp();
Temp samplemask = has_mrtz_samplemask ? get_arg(&ctx, einfo->samplemask) : Temp();
Temp alpha = has_mrtz_alpha ? colors[0][3] : Temp();
 
export_fs_mrtz(&ctx, depth, stencil, samplemask, alpha);
export_fs_mrtz(&ctx, einfo, depth, stencil, samplemask, mrtz_alpha);
}
 
/* Export all color render targets */
......
......@@ -514,6 +514,17 @@ print_regs(ra_ctx& ctx, PhysRegInterval regs, const RegisterFile& reg_file)
}
}
bool
is_sgpr_writable_without_side_effects(amd_gfx_level gfx_level, PhysReg reg)
{
assert(reg < 256);
bool has_flat_scr_lo_gfx89 = gfx_level >= GFX8 && gfx_level <= GFX9;
bool has_flat_scr_lo_gfx7_or_xnack_mask = gfx_level <= GFX9;
return (reg <= vcc_hi || reg == m0) &&
(!has_flat_scr_lo_gfx89 || (reg != flat_scr_lo && reg != flat_scr_hi)) &&
(!has_flat_scr_lo_gfx7_or_xnack_mask || (reg != 104 || reg != 105));
}
unsigned
get_subdword_operand_stride(amd_gfx_level gfx_level, const aco_ptr<Instruction>& instr,
unsigned idx, RegClass rc)
......@@ -2080,6 +2091,14 @@ operand_can_use_reg(amd_gfx_level gfx_level, aco_ptr<Instruction>& instr, unsign
gfx_level >= GFX10); /* sdata can be vcc */
case Format::MUBUF:
case Format::MTBUF: return idx != 2 || gfx_level < GFX12 || reg != scc;
case Format::SOPK:
if (idx == 0 && reg == scc)
return false;
FALLTHROUGH;
case Format::SOP2:
case Format::SOP1:
return get_op_fixed_to_def(instr.get()) != (int)idx ||
is_sgpr_writable_without_side_effects(gfx_level, reg);
default:
// TODO: there are more instructions with restrictions on registers
return true;
......@@ -2906,7 +2925,8 @@ optimize_encoding_sopk(ra_ctx& ctx, RegisterFile& register_file, aco_ptr<Instruc
return;
unsigned literal_idx = instr->operands[1].isLiteral();
if (instr->operands[!literal_idx].physReg() >= 128)
PhysReg op_reg = instr->operands[!literal_idx].physReg();
if (!is_sgpr_writable_without_side_effects(ctx.program->gfx_level, op_reg))
return;
unsigned def_id = instr->definitions[0].tempId();
......
......@@ -447,6 +447,12 @@ radv_device_init_meta(struct radv_device *device)
bool loaded_cache = radv_load_meta_pipeline(device);
bool on_demand = !loaded_cache;
result = vk_meta_device_init(&device->vk, &device->meta_state.device);
if (result != VK_SUCCESS)
return result;
device->meta_state.device.pipeline_cache = device->meta_state.cache;
mtx_init(&device->meta_state.mtx, mtx_plain);
result = radv_device_init_meta_clear_state(device, on_demand);
......@@ -567,6 +573,9 @@ radv_device_finish_meta(struct radv_device *device)
radv_store_meta_pipeline(device);
vk_common_DestroyPipelineCache(radv_device_to_handle(device), device->meta_state.cache, NULL);
mtx_destroy(&device->meta_state.mtx);
if (device->meta_state.device.cache)
vk_meta_device_finish(&device->vk, &device->meta_state.device);
}
nir_builder PRINTFLIKE(3, 4)
......
......@@ -737,12 +737,6 @@ radv_device_init_accel_struct_build_state(struct radv_device *device)
device->meta_state.accel_struct_build.radix_sort = vk_create_radix_sort_u64(
radv_device_to_handle(device), &device->meta_state.alloc, device->meta_state.cache, radix_sort_config);
result = vk_meta_device_init(&device->vk, &device->meta_state.device);
if (result != VK_SUCCESS)
goto exit;
device->meta_state.device.pipeline_cache = device->meta_state.cache;
device->vk.as_build_ops = &build_ops;
device->vk.write_buffer_cp = radv_write_buffer_cp;
device->vk.flush_buffer_write_cp = radv_flush_buffer_write_cp;
......
......@@ -5304,10 +5304,10 @@ static struct radv_shader_part *
lookup_ps_epilog(struct radv_cmd_buffer *cmd_buffer)
{
struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
const struct radv_physical_device *pdev = radv_device_physical(device);
const struct radv_shader *ps = cmd_buffer->state.shaders[MESA_SHADER_FRAGMENT];
const struct radv_rendering_state *render = &cmd_buffer->state.render;
const struct radv_dynamic_state *d = &cmd_buffer->state.dynamic;
const struct radv_physical_device *pdev = radv_device_physical(device);
struct radv_ps_epilog_state state = {0};
uint8_t color_remap[MAX_RTS];
 
......@@ -5351,11 +5351,11 @@ lookup_ps_epilog(struct radv_cmd_buffer *cmd_buffer)
state.colors_written = ps->info.ps.colors_written;
 
if (ps->info.ps.exports_mrtz_via_epilog) {
assert(pdev->info.gfx_level >= GFX11);
state.export_depth = ps->info.ps.writes_z;
state.export_stencil = ps->info.ps.writes_stencil;
state.export_sample_mask = ps->info.ps.writes_sample_mask;
state.alpha_to_coverage_via_mrtz = d->vk.ms.alpha_to_coverage_enable;
state.alpha_to_coverage_via_mrtz =
d->vk.ms.alpha_to_coverage_enable && (pdev->info.gfx_level >= GFX11 || d->vk.ms.alpha_to_one_enable);
}
}
 
......@@ -10530,6 +10530,12 @@ radv_emit_db_shader_control(struct radv_cmd_buffer *cmd_buffer)
radeon_opt_set_context_reg(cmd_buffer, R_02806C_DB_SHADER_CONTROL, RADV_TRACKED_DB_SHADER_CONTROL,
db_shader_control);
} else {
/* Use the alpha value from MRTZ.a for alpha-to-coverage when alpha-to-one is also enabled.
* GFX11+ selects MRTZ.a by default if present.
*/
db_shader_control |= S_02880C_COVERAGE_TO_MASK_ENABLE(
pdev->info.gfx_level < GFX11 && d->vk.ms.alpha_to_coverage_enable && d->vk.ms.alpha_to_one_enable);
radeon_opt_set_context_reg(cmd_buffer, R_02880C_DB_SHADER_CONTROL, RADV_TRACKED_DB_SHADER_CONTROL,
db_shader_control);
}
......@@ -10870,7 +10876,8 @@ radv_emit_all_graphics_states(struct radv_cmd_buffer *cmd_buffer, const struct r
(cmd_buffer->state.dirty_dynamic &
(RADV_DYNAMIC_COLOR_WRITE_MASK | RADV_DYNAMIC_COLOR_BLEND_ENABLE | RADV_DYNAMIC_RASTERIZATION_SAMPLES |
RADV_DYNAMIC_LINE_RASTERIZATION_MODE | RADV_DYNAMIC_PRIMITIVE_TOPOLOGY | RADV_DYNAMIC_POLYGON_MODE |
RADV_DYNAMIC_ATTACHMENT_FEEDBACK_LOOP_ENABLE)))
RADV_DYNAMIC_ATTACHMENT_FEEDBACK_LOOP_ENABLE | RADV_DYNAMIC_ALPHA_TO_COVERAGE_ENABLE |
RADV_DYNAMIC_ALPHA_TO_ONE_ENABLE)))
radv_emit_db_shader_control(cmd_buffer);
 
if (info->indexed && info->indirect && cmd_buffer->state.dirty & RADV_CMD_DIRTY_INDEX_BUFFER)
......
......@@ -32,45 +32,44 @@ enum {
RADV_DEBUG_NOBINNING = 1ull << 17,
RADV_DEBUG_NO_NGG = 1ull << 18,
RADV_DEBUG_DUMP_META_SHADERS = 1ull << 19,
RADV_DEBUG_DISCARD_TO_DEMOTE = 1ull << 20,
RADV_DEBUG_LLVM = 1ull << 21,
RADV_DEBUG_FORCE_COMPRESS = 1ull << 22,
RADV_DEBUG_HANG = 1ull << 23,
RADV_DEBUG_IMG = 1ull << 24,
RADV_DEBUG_NO_UMR = 1ull << 25,
RADV_DEBUG_INVARIANT_GEOM = 1ull << 26,
RADV_DEBUG_NO_DISPLAY_DCC = 1ull << 27,
RADV_DEBUG_NO_TC_COMPAT_CMASK = 1ull << 28,
RADV_DEBUG_NO_VRS_FLAT_SHADING = 1ull << 29,
RADV_DEBUG_NO_ATOC_DITHERING = 1ull << 30,
RADV_DEBUG_NO_NGGC = 1ull << 31,
RADV_DEBUG_DUMP_PROLOGS = 1ull << 32,
RADV_DEBUG_NO_DMA_BLIT = 1ull << 33,
RADV_DEBUG_SPLIT_FMA = 1ull << 34,
RADV_DEBUG_DUMP_EPILOGS = 1ull << 35,
RADV_DEBUG_NO_FMASK = 1ull << 36,
RADV_DEBUG_SHADOW_REGS = 1ull << 37,
RADV_DEBUG_EXTRA_MD = 1ull << 38,
RADV_DEBUG_NO_GPL = 1ull << 39,
RADV_DEBUG_VIDEO_ARRAY_PATH = 1ull << 40,
RADV_DEBUG_NO_RT = 1ull << 41,
RADV_DEBUG_NO_MESH_SHADER = 1ull << 42,
RADV_DEBUG_NO_NGG_GS = 1ull << 43,
RADV_DEBUG_NO_ESO = 1ull << 44,
RADV_DEBUG_PSO_CACHE_STATS = 1ull << 45,
RADV_DEBUG_NIR_DEBUG_INFO = 1ull << 46,
RADV_DEBUG_DUMP_TRAP_HANDLER = 1ull << 47,
RADV_DEBUG_DUMP_VS = 1ull << 48,
RADV_DEBUG_DUMP_TCS = 1ull << 49,
RADV_DEBUG_DUMP_TES = 1ull << 50,
RADV_DEBUG_DUMP_GS = 1ull << 51,
RADV_DEBUG_DUMP_PS = 1ull << 52,
RADV_DEBUG_DUMP_TASK = 1ull << 53,
RADV_DEBUG_DUMP_MESH = 1ull << 54,
RADV_DEBUG_DUMP_CS = 1ull << 55,
RADV_DEBUG_DUMP_NIR = 1ull << 56,
RADV_DEBUG_DUMP_ASM = 1ull << 57,
RADV_DEBUG_DUMP_BACKEND_IR = 1ull << 58,
RADV_DEBUG_LLVM = 1ull << 20,
RADV_DEBUG_FORCE_COMPRESS = 1ull << 21,
RADV_DEBUG_HANG = 1ull << 22,
RADV_DEBUG_IMG = 1ull << 23,
RADV_DEBUG_NO_UMR = 1ull << 24,
RADV_DEBUG_INVARIANT_GEOM = 1ull << 25,
RADV_DEBUG_NO_DISPLAY_DCC = 1ull << 26,
RADV_DEBUG_NO_TC_COMPAT_CMASK = 1ull << 27,
RADV_DEBUG_NO_VRS_FLAT_SHADING = 1ull << 28,
RADV_DEBUG_NO_ATOC_DITHERING = 1ull << 29,
RADV_DEBUG_NO_NGGC = 1ull << 30,
RADV_DEBUG_DUMP_PROLOGS = 1ull << 31,
RADV_DEBUG_NO_DMA_BLIT = 1ull << 32,
RADV_DEBUG_SPLIT_FMA = 1ull << 33,
RADV_DEBUG_DUMP_EPILOGS = 1ull << 34,
RADV_DEBUG_NO_FMASK = 1ull << 35,
RADV_DEBUG_SHADOW_REGS = 1ull << 36,
RADV_DEBUG_EXTRA_MD = 1ull << 37,
RADV_DEBUG_NO_GPL = 1ull << 38,
RADV_DEBUG_VIDEO_ARRAY_PATH = 1ull << 39,
RADV_DEBUG_NO_RT = 1ull << 40,
RADV_DEBUG_NO_MESH_SHADER = 1ull << 41,
RADV_DEBUG_NO_NGG_GS = 1ull << 42,
RADV_DEBUG_NO_ESO = 1ull << 43,
RADV_DEBUG_PSO_CACHE_STATS = 1ull << 44,
RADV_DEBUG_NIR_DEBUG_INFO = 1ull << 45,
RADV_DEBUG_DUMP_TRAP_HANDLER = 1ull << 46,
RADV_DEBUG_DUMP_VS = 1ull << 47,
RADV_DEBUG_DUMP_TCS = 1ull << 48,
RADV_DEBUG_DUMP_TES = 1ull << 49,
RADV_DEBUG_DUMP_GS = 1ull << 50,
RADV_DEBUG_DUMP_PS = 1ull << 51,
RADV_DEBUG_DUMP_TASK = 1ull << 52,
RADV_DEBUG_DUMP_MESH = 1ull << 53,
RADV_DEBUG_DUMP_CS = 1ull << 54,
RADV_DEBUG_DUMP_NIR = 1ull << 55,
RADV_DEBUG_DUMP_ASM = 1ull << 56,
RADV_DEBUG_DUMP_BACKEND_IR = 1ull << 57,
RADV_DEBUG_DUMP_SHADERS = RADV_DEBUG_DUMP_VS | RADV_DEBUG_DUMP_TCS | RADV_DEBUG_DUMP_TES | RADV_DEBUG_DUMP_GS |
RADV_DEBUG_DUMP_PS | RADV_DEBUG_DUMP_TASK | RADV_DEBUG_DUMP_MESH | RADV_DEBUG_DUMP_CS |
RADV_DEBUG_DUMP_NIR | RADV_DEBUG_DUMP_ASM | RADV_DEBUG_DUMP_BACKEND_IR,
......
......@@ -222,7 +222,6 @@ radv_physical_device_init_cache_key(struct radv_physical_device *pdev)
key->emulate_rt = !!(instance->perftest_flags & RADV_PERFTEST_EMULATE_RT);
key->ge_wave32 = pdev->ge_wave_size == 32;
key->invariant_geom = !!(instance->debug_flags & RADV_DEBUG_INVARIANT_GEOM);
key->lower_discard_to_demote = !!(instance->debug_flags & RADV_DEBUG_DISCARD_TO_DEMOTE);
key->no_fmask = !!(instance->debug_flags & RADV_DEBUG_NO_FMASK);
key->no_ngg_gs = !!(instance->debug_flags & RADV_DEBUG_NO_NGG_GS);
key->no_rt = !!(instance->debug_flags & RADV_DEBUG_NO_RT);
......@@ -786,6 +785,7 @@ radv_physical_device_get_features(const struct radv_physical_device *pdev, struc
bool has_shader_image_float_minmax = pdev->info.gfx_level != GFX8 && pdev->info.gfx_level != GFX9 &&
pdev->info.gfx_level != GFX11 && pdev->info.gfx_level != GFX11_5;
bool has_fragment_shader_interlock = radv_has_pops(pdev);
const bool is_zink = instance->vk.app_info.engine_name && !strcmp(instance->vk.app_info.engine_name, "mesa zink");
*features = (struct vk_features){
/* Vulkan 1.0 */
......@@ -806,7 +806,11 @@ radv_physical_device_get_features(const struct radv_physical_device *pdev, struc
.depthBounds = true,
.wideLines = true,
.largePoints = true,
.alphaToOne = true,
/* FIXME: alpha-to-one is performed before alpha-to-coverage but it should be after.
* This is currently broken on RADV and the fixes aren't trivial to backport. Let's disable it
* for everything except Zink because it's a requirement and no application seems affected.
*/
.alphaToOne = is_zink,
.multiViewport = true,
.samplerAnisotropy = true,
.textureCompressionETC2 = pdev->info.has_etc_support || pdev->emulate_etc2,
......@@ -1190,7 +1194,7 @@ radv_physical_device_get_features(const struct radv_physical_device *pdev, struc
.extendedDynamicState3SampleLocationsEnable = pdev->info.gfx_level < GFX10,
.extendedDynamicState3LineRasterizationMode = true,
.extendedDynamicState3ExtraPrimitiveOverestimationSize = false,
.extendedDynamicState3AlphaToOneEnable = !pdev->use_llvm,
.extendedDynamicState3AlphaToOneEnable = is_zink && !pdev->use_llvm,
.extendedDynamicState3RasterizationStream = false,
.extendedDynamicState3ColorBlendAdvanced = false,
.extendedDynamicState3ViewportWScalingEnable = false,
......
......@@ -53,7 +53,6 @@ struct radv_physical_device_cache_key {
uint32_t emulate_rt : 1;
uint32_t ge_wave32 : 1;
uint32_t invariant_geom : 1;
uint32_t lower_discard_to_demote : 1;
uint32_t no_fmask : 1;
uint32_t no_ngg_gs : 1;
uint32_t no_rt : 1;
......
......@@ -1886,6 +1886,18 @@ radv_generate_graphics_state_key(const struct radv_device *device, const struct
if (state->ts)
key.ts.patch_control_points = state->ts->patch_control_points;
const bool alpha_to_coverage_unknown =
!state->ms || BITSET_TEST(state->dynamic, MESA_VK_DYNAMIC_MS_ALPHA_TO_COVERAGE_ENABLE);
const bool alpha_to_coverage_enabled = alpha_to_coverage_unknown || state->ms->alpha_to_coverage_enable;
const bool alpha_to_one_unknown = !state->ms || BITSET_TEST(state->dynamic, MESA_VK_DYNAMIC_MS_ALPHA_TO_ONE_ENABLE);
const bool alpha_to_one_enabled = alpha_to_one_unknown || state->ms->alpha_to_one_enable;
/* alpha-to-coverage is always exported via MRTZ on GFX11 but it's also using MRTZ when
* alpha-to-one is enabled (alpha to MRTZ.a and one to MRT0.a).
*/
key.ms.alpha_to_coverage_via_mrtz =
alpha_to_coverage_enabled && (pdev->info.gfx_level >= GFX11 || alpha_to_one_enabled);
if (state->ms) {
key.ms.sample_shading_enable = state->ms->sample_shading_enable;
if (!BITSET_TEST(state->dynamic, MESA_VK_DYNAMIC_MS_RASTERIZATION_SAMPLES) &&
......@@ -1894,10 +1906,6 @@ radv_generate_graphics_state_key(const struct radv_device *device, const struct
}
}
if (pdev->info.gfx_level >= GFX11 && state->ms) {
key.ms.alpha_to_coverage_via_mrtz = state->ms->alpha_to_coverage_enable;
}
if (state->ia) {
key.ia.topology = radv_translate_prim(state->ia->primitive_topology);
}
......@@ -1922,14 +1930,17 @@ radv_generate_graphics_state_key(const struct radv_device *device, const struct
key.ps.epilog = radv_pipeline_generate_ps_epilog_key(device, state);
if (pdev->info.gfx_level >= GFX11) {
/* On GFX11, alpha to coverage is exported via MRTZ when depth/stencil/samplemask are also
* exported. Though, when a PS epilog is needed and the MS state is NULL (with dynamic
* rendering), it's not possible to know the info at compile time and MRTZ needs to be
* exported in the epilog.
*/
key.ps.exports_mrtz_via_epilog =
key.ps.has_epilog && (!state->ms || BITSET_TEST(state->dynamic, MESA_VK_DYNAMIC_MS_ALPHA_TO_COVERAGE_ENABLE));
/* Alpha to coverage is exported via MRTZ when depth/stencil/samplemask are also exported.
* Though, when a PS epilog is needed and the MS state is NULL (with dynamic rendering), it's not
* possible to know the info at compile time and MRTZ needs to be exported in the epilog.
*/
if (key.ps.has_epilog) {
if (pdev->info.gfx_level >= GFX11) {
key.ps.exports_mrtz_via_epilog = alpha_to_coverage_unknown;
} else {
key.ps.exports_mrtz_via_epilog =
(alpha_to_coverage_unknown && alpha_to_one_enabled) || (alpha_to_one_unknown && alpha_to_coverage_enabled);
}
}
key.dynamic_rasterization_samples = BITSET_TEST(state->dynamic, MESA_VK_DYNAMIC_MS_RASTERIZATION_SAMPLES) ||
......@@ -2730,6 +2741,21 @@ radv_graphics_shaders_compile(struct radv_device *device, struct vk_pipeline_cac
/* Optimize varyings on lowered shader I/O (more efficient than optimizing I/O derefs). */
radv_graphics_shaders_link_varyings(stages);
/* Optimize constant clip/cull distance after linking to operate on scalar io in the last
* pre raster stage.
*/
radv_foreach_stage(i, active_nir_stages & (VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_TESSELLATION_EVALUATION_BIT))
{
if (stages[i].key.optimisations_disabled)
continue;
int64_t stage_start = os_time_get_nano();
NIR_PASS(_, stages[i].nir, nir_opt_clip_cull_const);
stages[i].feedback.duration += os_time_get_nano() - stage_start;
}
radv_fill_shader_info(device, RADV_PIPELINE_GRAPHICS, gfx_state, stages, active_nir_stages);
radv_declare_pipeline_args(device, stages, gfx_state, active_nir_stages);
......
......@@ -187,6 +187,13 @@ gather_intrinsic_store_output_info(const nir_shader *nir, const nir_intrinsic_in
const uint8_t gs_streams = nir_intrinsic_io_semantics(instr).gs_streams;
info->gs.output_streams[location] |= gs_streams << (component * 2);
}
if ((location == VARYING_SLOT_CLIP_DIST0 || location == VARYING_SLOT_CLIP_DIST1) && !io_sem.no_sysval_output) {
unsigned base = (location == VARYING_SLOT_CLIP_DIST1 ? 4 : 0) + component;
unsigned clip_array_mask = BITFIELD_MASK(nir->info.clip_distance_array_size);
info->outinfo.clip_dist_mask |= (write_mask << base) & clip_array_mask;
info->outinfo.cull_dist_mask |= (write_mask << base) & ~clip_array_mask;
}
}
static void
......@@ -303,6 +310,7 @@ gather_intrinsic_info(const nir_shader *nir, const nir_intrinsic_instr *instr, s
gather_intrinsic_load_input_info(nir, instr, info, gfx_state, stage_key);
break;
case nir_intrinsic_store_output:
case nir_intrinsic_store_per_vertex_output:
gather_intrinsic_store_output_info(nir, instr, info, consider_force_vrs);
break;
case nir_intrinsic_bvh64_intersect_ray_amd:
......@@ -1223,26 +1231,18 @@ radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *n
outinfo->writes_primitive_shading_rate_per_primitive = per_prim_mask & VARYING_BIT_PRIMITIVE_SHADING_RATE;
outinfo->export_prim_id_per_primitive = per_prim_mask & VARYING_BIT_PRIMITIVE_ID;
/* Clip/cull distances. */
outinfo->clip_dist_mask = (1 << nir->info.clip_distance_array_size) - 1;
outinfo->cull_dist_mask = (1 << nir->info.cull_distance_array_size) - 1;
outinfo->cull_dist_mask <<= nir->info.clip_distance_array_size;
int pos_written = 0x1;
outinfo->pos_exports = 1;
if (outinfo->writes_pointsize || outinfo->writes_viewport_index || outinfo->writes_layer ||
outinfo->writes_primitive_shading_rate)
pos_written |= 1 << 1;
unsigned num_clip_distances = util_bitcount(outinfo->clip_dist_mask);
unsigned num_cull_distances = util_bitcount(outinfo->cull_dist_mask);
outinfo->pos_exports++;
if (num_clip_distances + num_cull_distances > 0)
pos_written |= 1 << 2;
if (num_clip_distances + num_cull_distances > 4)
pos_written |= 1 << 3;
unsigned clip_cull_mask = outinfo->clip_dist_mask | outinfo->cull_dist_mask;
outinfo->pos_exports = util_bitcount(pos_written);
if (clip_cull_mask & 0x0f)
outinfo->pos_exports++;
if (clip_cull_mask & 0xf0)
outinfo->pos_exports++;
}
info->vs.needs_draw_id |= BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID);
......