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
  • LingMan/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
  • olivia/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
  • iv-m/mesa
  • leftmostcat/mesa
  • gaoshunli/mesa
  • Nunwan/mesa
  • forbiddenlake/mesa
  • taras.pisetskyi/mesa
  • andrei.matraguna/mesa
  • mdziuban/mesa
  • zlice1/mesa
  • cfoch/mesa
  • kpouget/mesa
  • nikolaszimmermann/mesa
  • kennylevinsen/mesa
  • itistotalbotnet/mesa
  • bartek1516/mesa
  • kusma/mesa-old
  • ludvig.lindau/mesa
  • jpegxguy/mesa
  • Jie1zhang/mesa
  • daniellang/mesa
  • FalsePhilosopher/mesa-ps-4
  • marge-bot-mesa/mesa
1231 results
Show changes
Commits on Source (58)
Showing
with 294 additions and 201 deletions
......@@ -36,12 +36,12 @@ include:
- local: 'src/gallium/drivers/iris/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/lima/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/llvmpipe/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/panfrost/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/radeonsi/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/softpipe/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/virgl/ci/gitlab-ci.yml'
- local: 'src/gallium/drivers/zink/ci/gitlab-ci.yml'
- local: 'src/gallium/frontends/lavapipe/ci/gitlab-ci.yml'
- local: 'src/panfrost/ci/gitlab-ci.yml'
stages:
- sanity
......
......@@ -57,6 +57,7 @@ for var in \
MESA_VK_IGNORE_CONFORMANCE_WARNING \
MINIO_HOST \
NIR_VALIDATE \
PAN_I_WANT_A_BROKEN_VULKAN_DRIVER \
PAN_MESA_DEBUG \
PIGLIT_FRACTION \
PIGLIT_JUNIT_RESULTS \
......
......@@ -239,6 +239,7 @@
- src/gallium/winsys/panfrost/**/*
when: on_success
- changes: &panfrost_common_file_list
- src/panfrost/ci/*
- src/panfrost/include/*
- src/panfrost/lib/*
- src/panfrost/shared/*
......
......@@ -295,7 +295,7 @@ GLES3.2, GLSL ES 3.2 -- all DONE: i965/gen9+, radeonsi, virgl, llvmpipe, zink
Khronos, ARB, and OES extensions that are not part of any OpenGL or OpenGL ES version:
GL_ARB_bindless_texture DONE (nvc0, radeonsi)
GL_ARB_bindless_texture DONE (nvc0, radeonsi, zink)
GL_ARB_cl_event not started
GL_ARB_compute_variable_group_size DONE (i965/gen7+, nvc0, radeonsi, zink)
GL_ARB_ES3_2_compatibility DONE (i965/gen8+, radeonsi, virgl, zink)
......
......@@ -1423,12 +1423,21 @@ ASSERTED static bool is_dcc_supported_by_L2(const struct radeon_info *info,
/* 128B is recommended, but 64B can be set too if needed for 4K by DCN.
* Since there is no reason to ever disable 128B, require it.
* DCC image stores are always supported.
* If 64B is used, DCC image stores are unsupported.
*/
return surf->u.gfx9.color.dcc.independent_128B_blocks &&
surf->u.gfx9.color.dcc.max_compressed_block_size <= V_028C78_MAX_BLOCK_SIZE_128B;
}
static bool gfx10_DCN_requires_independent_64B_blocks(const struct radeon_info *info,
const struct ac_surf_config *config)
{
assert(info->chip_class >= GFX10);
/* For 4K, DCN requires INDEPENDENT_64B_BLOCKS = 1 and MAX_COMPRESSED_BLOCK_SIZE = 64B. */
return config->info.width > 2560 || config->info.height > 2560;
}
static bool is_dcc_supported_by_DCN(const struct radeon_info *info,
const struct ac_surf_config *config,
const struct radeon_surf *surf, bool rb_aligned,
......@@ -1460,8 +1469,7 @@ static bool is_dcc_supported_by_DCN(const struct radeon_info *info,
if (info->chip_class == GFX10 && surf->u.gfx9.color.dcc.independent_128B_blocks)
return false;
/* For 4K, DCN requires INDEPENDENT_64B_BLOCKS = 1. */
return ((config->info.width <= 2560 && config->info.height <= 2560) ||
return (!gfx10_DCN_requires_independent_64B_blocks(info, config) ||
(surf->u.gfx9.color.dcc.independent_64B_blocks &&
surf->u.gfx9.color.dcc.max_compressed_block_size == V_028C78_MAX_BLOCK_SIZE_64B));
default:
......@@ -2083,7 +2091,8 @@ static int gfx9_compute_surface(struct ac_addrlib *addrlib, const struct radeon_
surf->u.gfx9.color.dcc.max_compressed_block_size = V_028C78_MAX_BLOCK_SIZE_64B;
}
if (info->chip_class >= GFX10_3) {
if (info->chip_class >= GFX10_3 &&
gfx10_DCN_requires_independent_64B_blocks(info, config)) {
surf->u.gfx9.color.dcc.independent_64B_blocks = 1;
surf->u.gfx9.color.dcc.independent_128B_blocks = 1;
surf->u.gfx9.color.dcc.max_compressed_block_size = V_028C78_MAX_BLOCK_SIZE_64B;
......
......@@ -5022,7 +5022,7 @@ visit_load_input(isel_context* ctx, nir_intrinsic_instr* instr)
uint32_t attrib_stride = ctx->options->key.vs.vertex_attribute_strides[location];
unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[location];
unsigned binding_align = ctx->options->key.vs.vertex_binding_align[attrib_binding];
enum ac_fetch_format alpha_adjust = ctx->options->key.vs.alpha_adjust[location];
enum ac_fetch_format alpha_adjust = ctx->options->key.vs.vertex_alpha_adjust[location];
unsigned dfmt = attrib_format & 0xf;
unsigned nfmt = (attrib_format >> 4) & 0x7;
......@@ -5030,7 +5030,7 @@ visit_load_input(isel_context* ctx, nir_intrinsic_instr* instr)
unsigned mask = nir_ssa_def_components_read(&instr->dest.ssa) << component;
unsigned num_channels = MIN2(util_last_bit(mask), vtx_info->num_channels);
bool post_shuffle = ctx->options->key.vs.post_shuffle & (1 << location);
bool post_shuffle = ctx->options->key.vs.vertex_post_shuffle & (1 << location);
if (post_shuffle)
num_channels = MAX2(num_channels, 3);
......@@ -7388,9 +7388,9 @@ visit_load_sample_mask_in(isel_context* ctx, nir_intrinsic_instr* instr)
{
uint8_t log2_ps_iter_samples;
if (ctx->program->info->ps.uses_sample_shading) {
log2_ps_iter_samples = util_logbase2(ctx->options->key.fs.num_samples);
log2_ps_iter_samples = util_logbase2(ctx->options->key.ps.num_samples);
} else {
log2_ps_iter_samples = ctx->options->key.fs.log2_ps_iter_samples;
log2_ps_iter_samples = ctx->options->key.ps.log2_ps_iter_samples;
}
Builder bld(ctx->program, ctx->block);
......@@ -8000,7 +8000,7 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
}
case nir_intrinsic_load_barycentric_at_sample: {
uint32_t sample_pos_offset = RING_PS_SAMPLE_POSITIONS * 16;
switch (ctx->options->key.fs.num_samples) {
switch (ctx->options->key.ps.num_samples) {
case 2: sample_pos_offset += 1 << 3; break;
case 4: sample_pos_offset += 3 << 3; break;
case 8: sample_pos_offset += 7 << 3; break;
......@@ -8899,7 +8899,7 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
ctx->shader->info.stage == MESA_SHADER_TESS_EVAL);
Temp dst = get_ssa_temp(ctx, &instr->dest.ssa);
bld.copy(Definition(dst), Operand::c32(ctx->args->options->key.tcs.input_vertices));
bld.copy(Definition(dst), Operand::c32(ctx->args->options->key.tcs.tess_input_vertices));
break;
}
case nir_intrinsic_emit_vertex_with_counter: {
......@@ -10987,10 +10987,10 @@ export_fs_mrt_color(isel_context* ctx, int slot)
slot -= FRAG_RESULT_DATA0;
target = V_008DFC_SQ_EXP_MRT + slot;
col_format = (ctx->options->key.fs.col_format >> (4 * slot)) & 0xf;
col_format = (ctx->options->key.ps.col_format >> (4 * slot)) & 0xf;
bool is_int8 = (ctx->options->key.fs.is_int8 >> slot) & 1;
bool is_int10 = (ctx->options->key.fs.is_int10 >> slot) & 1;
bool is_int8 = (ctx->options->key.ps.is_int8 >> slot) & 1;
bool is_int10 = (ctx->options->key.ps.is_int10 >> slot) & 1;
bool is_16bit = values[0].regClass() == v2b;
/* Replace NaN by zero (only 32-bit) to fix game bugs if requested. */
......
......@@ -468,9 +468,9 @@ init_context(isel_context* ctx, nir_shader* shader)
ctx->range_ht = _mesa_pointer_hash_table_create(NULL);
ctx->ub_config.min_subgroup_size = 64;
ctx->ub_config.max_subgroup_size = 64;
if (ctx->shader->info.stage == MESA_SHADER_COMPUTE && ctx->options->key.cs.subgroup_size) {
ctx->ub_config.min_subgroup_size = ctx->options->key.cs.subgroup_size;
ctx->ub_config.max_subgroup_size = ctx->options->key.cs.subgroup_size;
if (ctx->shader->info.stage == MESA_SHADER_COMPUTE && ctx->args->shader_info->cs.subgroup_size) {
ctx->ub_config.min_subgroup_size = ctx->args->shader_info->cs.subgroup_size;
ctx->ub_config.max_subgroup_size = ctx->args->shader_info->cs.subgroup_size;
}
ctx->ub_config.max_workgroup_invocations = 2048;
ctx->ub_config.max_workgroup_count[0] = 65535;
......
......@@ -1667,6 +1667,7 @@ get_reg_create_vector(ra_ctx& ctx, RegisterFile& reg_file, Temp temp,
/* count variables to be moved and check "avoid" */
bool avoid = false;
bool linear_vgpr = false;
for (PhysReg j : reg_win) {
if (reg_file[j] != 0) {
if (reg_file[j] == 0xF0000000) {
......@@ -1677,17 +1678,20 @@ get_reg_create_vector(ra_ctx& ctx, RegisterFile& reg_file, Temp temp,
k += reg_file.test(reg, 1);
} else {
k += 4;
/* we cannot split live ranges of linear vgprs inside control flow */
if (ctx.assignments[reg_file[j]].rc.is_linear_vgpr()) {
if (ctx.block->kind & block_kind_top_level)
avoid = true;
else
break;
}
linear_vgpr |= ctx.assignments[reg_file[j]].rc.is_linear_vgpr();
}
}
avoid |= ctx.war_hint[j];
}
if (linear_vgpr) {
/* we cannot split live ranges of linear vgprs inside control flow */
if (ctx.block->kind & block_kind_top_level)
avoid = true;
else
continue;
}
if (avoid && !best_avoid)
continue;
......
......@@ -37,7 +37,7 @@
#define SMEM_MAX_MOVES (64 - ctx.num_waves * 4)
#define VMEM_MAX_MOVES (256 - ctx.num_waves * 16)
/* creating clauses decreases def-use distances, so make it less aggressive the lower num_waves is */
#define VMEM_CLAUSE_MAX_GRAB_DIST (ctx.num_waves * 8)
#define VMEM_CLAUSE_MAX_GRAB_DIST (ctx.num_waves * 2)
#define POS_EXP_MAX_MOVES 512
namespace aco {
......@@ -788,6 +788,7 @@ schedule_VMEM(sched_ctx& ctx, Block* block, std::vector<RegisterDemand>& registe
int window_size = VMEM_WINDOW_SIZE;
int max_moves = VMEM_MAX_MOVES;
int clause_max_grab_dist = VMEM_CLAUSE_MAX_GRAB_DIST;
bool only_clauses = false;
int16_t k = 0;
/* first, check if we have instructions before current to move down */
......@@ -822,12 +823,28 @@ schedule_VMEM(sched_ctx& ctx, Block* block, std::vector<RegisterDemand>& registe
/* We can't easily tell how much this will decrease the def-to-use
* distances, so just use how far it will be moved as a heuristic. */
part_of_clause =
grab_dist < clause_max_grab_dist && should_form_clause(current, candidate.get());
grab_dist < clause_max_grab_dist + k && should_form_clause(current, candidate.get());
}
/* if current depends on candidate, add additional dependencies and continue */
bool can_move_down = !is_vmem || part_of_clause || candidate->definitions.empty();
if (only_clauses) {
/* In case of high register pressure, only try to form clauses,
* and only if the previous clause is not larger
* than the current one will be.
*/
if (part_of_clause) {
int clause_size = cursor.insert_idx - cursor.insert_idx_clause;
int prev_clause_size = 1;
while (should_form_clause(current,
block->instructions[candidate_idx - prev_clause_size].get()))
prev_clause_size++;
if (prev_clause_size > clause_size + 1)
break;
} else {
can_move_down = false;
}
}
HazardResult haz =
perform_hazard_query(part_of_clause ? &clause_hq : &indep_hq, candidate.get(), false);
if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill ||
......@@ -838,6 +855,8 @@ schedule_VMEM(sched_ctx& ctx, Block* block, std::vector<RegisterDemand>& registe
break;
if (!can_move_down) {
if (part_of_clause)
break;
add_to_hazard_query(&indep_hq, candidate.get());
add_to_hazard_query(&clause_hq, candidate.get());
ctx.mv.downwards_skip(cursor);
......@@ -847,12 +866,20 @@ schedule_VMEM(sched_ctx& ctx, Block* block, std::vector<RegisterDemand>& registe
Instruction* candidate_ptr = candidate.get();
MoveResult res = ctx.mv.downwards_move(cursor, part_of_clause);
if (res == move_fail_ssa || res == move_fail_rar) {
if (part_of_clause)
break;
add_to_hazard_query(&indep_hq, candidate.get());
add_to_hazard_query(&clause_hq, candidate.get());
ctx.mv.downwards_skip(cursor);
continue;
} else if (res == move_fail_pressure) {
break;
only_clauses = true;
if (part_of_clause)
break;
add_to_hazard_query(&indep_hq, candidate.get());
add_to_hazard_query(&clause_hq, candidate.get());
ctx.mv.downwards_skip(cursor);
continue;
}
if (part_of_clause)
add_to_hazard_query(&indep_hq, candidate_ptr);
......
......@@ -259,7 +259,7 @@ load_sample_position(struct ac_shader_abi *abi, LLVMValueRef sample_id)
ptr = LLVMBuildBitCast(ctx->ac.builder, ptr, ac_array_in_const_addr_space(ctx->ac.v2f32), "");
uint32_t sample_pos_offset = radv_get_sample_pos_offset(ctx->args->options->key.fs.num_samples);
uint32_t sample_pos_offset = radv_get_sample_pos_offset(ctx->args->options->key.ps.num_samples);
sample_id = LLVMBuildAdd(ctx->ac.builder, sample_id,
LLVMConstInt(ctx->ac.i32, sample_pos_offset, false), "");
......@@ -275,9 +275,9 @@ load_sample_mask_in(struct ac_shader_abi *abi)
uint8_t log2_ps_iter_samples;
if (ctx->args->shader_info->ps.uses_sample_shading) {
log2_ps_iter_samples = util_logbase2(ctx->args->options->key.fs.num_samples);
log2_ps_iter_samples = util_logbase2(ctx->args->options->key.ps.num_samples);
} else {
log2_ps_iter_samples = ctx->args->options->key.fs.log2_ps_iter_samples;
log2_ps_iter_samples = ctx->args->options->key.ps.log2_ps_iter_samples;
}
LLVMValueRef result, sample_id;
......@@ -716,9 +716,9 @@ load_vs_input(struct radv_shader_context *ctx, unsigned driver_location, LLVMTyp
unsigned attrib_binding = ctx->args->options->key.vs.vertex_attribute_bindings[attrib_index];
unsigned attrib_offset = ctx->args->options->key.vs.vertex_attribute_offsets[attrib_index];
unsigned attrib_stride = ctx->args->options->key.vs.vertex_attribute_strides[attrib_index];
unsigned alpha_adjust = ctx->args->options->key.vs.alpha_adjust[attrib_index];
unsigned alpha_adjust = ctx->args->options->key.vs.vertex_alpha_adjust[attrib_index];
if (ctx->args->options->key.vs.post_shuffle & (1 << attrib_index)) {
if (ctx->args->options->key.vs.vertex_post_shuffle & (1 << attrib_index)) {
/* Always load, at least, 3 channels for formats that need to be shuffled because X<->Z. */
num_channels = MAX2(num_channels, 3);
}
......@@ -775,7 +775,7 @@ load_vs_input(struct radv_shader_context *ctx, unsigned driver_location, LLVMTyp
ctx->ac.i32_0, ctx->ac.i32_0, num_channels, data_format, num_format, 0, true);
}
if (ctx->args->options->key.vs.post_shuffle & (1 << attrib_index)) {
if (ctx->args->options->key.vs.vertex_post_shuffle & (1 << attrib_index)) {
LLVMValueRef c[4];
c[0] = ac_llvm_extract_elem(&ctx->ac, input, 2);
c[1] = ac_llvm_extract_elem(&ctx->ac, input, 1);
......@@ -899,9 +899,9 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, LLVMValueRef *values,
bool is_16bit = ac_get_type_size(LLVMTypeOf(values[0])) == 2;
if (ctx->stage == MESA_SHADER_FRAGMENT) {
unsigned index = target - V_008DFC_SQ_EXP_MRT;
unsigned col_format = (ctx->args->options->key.fs.col_format >> (4 * index)) & 0xf;
bool is_int8 = (ctx->args->options->key.fs.is_int8 >> index) & 1;
bool is_int10 = (ctx->args->options->key.fs.is_int10 >> index) & 1;
unsigned col_format = (ctx->args->options->key.ps.col_format >> (4 * index)) & 0xf;
bool is_int8 = (ctx->args->options->key.ps.is_int8 >> index) & 1;
bool is_int10 = (ctx->args->options->key.ps.is_int10 >> index) & 1;
LLVMValueRef (*packf)(struct ac_llvm_context * ctx, LLVMValueRef args[2]) = NULL;
LLVMValueRef (*packi)(struct ac_llvm_context * ctx, LLVMValueRef args[2], unsigned bits,
......@@ -1617,8 +1617,10 @@ handle_ngg_outputs_post_2(struct radv_shader_context *ctx)
LLVMValueRef provoking_vtx_in_prim = LLVMConstInt(ctx->ac.i32, 0, false);
/* For provoking vertex last mode, use num_vtx_in_prim - 1. */
if (ctx->args->options->key.vs.provoking_vtx_last)
provoking_vtx_in_prim = LLVMConstInt(ctx->ac.i32, ctx->args->options->key.vs.outprim, false);
if (ctx->args->options->key.vs.provoking_vtx_last) {
uint8_t outprim = si_conv_prim_to_gs_out(ctx->args->options->key.vs.topology);
provoking_vtx_in_prim = LLVMConstInt(ctx->ac.i32, outprim, false);
}
/* provoking_vtx_index = vtxindex[provoking_vtx_in_prim]; */
LLVMValueRef indices = ac_build_gather_values(&ctx->ac, vtxindex, 3);
......
......@@ -1237,30 +1237,6 @@ si_conv_gl_prim_to_gs_out(unsigned gl_prim)
}
}
static uint32_t
si_conv_prim_to_gs_out(enum VkPrimitiveTopology topology)
{
switch (topology) {
case VK_PRIMITIVE_TOPOLOGY_POINT_LIST:
case VK_PRIMITIVE_TOPOLOGY_PATCH_LIST:
return V_028A6C_POINTLIST;
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP:
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP_WITH_ADJACENCY:
return V_028A6C_LINESTRIP;
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP_WITH_ADJACENCY:
return V_028A6C_TRISTRIP;
default:
assert(0);
return 0;
}
}
static uint64_t
radv_dynamic_state_mask(VkDynamicState state)
{
......@@ -1859,7 +1835,7 @@ gfx9_get_gs_info(const struct radv_pipeline_key *key, const struct radv_pipeline
unsigned gs_num_invocations = MAX2(gs_info->gs.invocations, 1);
bool uses_adjacency;
switch (key->topology) {
switch (key->vs.topology) {
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST_WITH_ADJACENCY:
......@@ -2020,7 +1996,7 @@ gfx10_get_ngg_info(const struct radv_pipeline_key *key, struct radv_pipeline *pi
unsigned min_verts_per_prim = gs_type == MESA_SHADER_GEOMETRY ? max_verts_per_prim : 1;
unsigned gs_num_invocations = nir[MESA_SHADER_GEOMETRY] ? MAX2(gs_info->gs.invocations, 1) : 1;
bool uses_adjacency;
switch (key->topology) {
switch (key->vs.topology) {
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST_WITH_ADJACENCY:
......@@ -2409,7 +2385,7 @@ radv_link_shaders(struct radv_pipeline *pipeline,
ordered_shaders[i - 1]->info.inputs_read & VARYING_BIT_PSIZ;
bool topology_uses_psiz =
info->stage == pipeline->graphics.last_vgt_api_stage &&
((info->stage == MESA_SHADER_VERTEX && pipeline_key->topology == VK_PRIMITIVE_TOPOLOGY_POINT_LIST) ||
((info->stage == MESA_SHADER_VERTEX && pipeline_key->vs.topology == VK_PRIMITIVE_TOPOLOGY_POINT_LIST) ||
(info->stage == MESA_SHADER_TESS_EVAL && info->tess.point_mode) ||
(info->stage == MESA_SHADER_GEOMETRY && info->gs.output_primitive == GL_POINTS));
......@@ -2621,8 +2597,8 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
int first_non_void;
if (binding_input_rate & (1u << binding)) {
key.instance_rate_inputs |= 1u << location;
key.instance_rate_divisors[location] = instance_rate_divisors[binding];
key.vs.instance_rate_inputs |= 1u << location;
key.vs.instance_rate_divisors[location] = instance_rate_divisors[binding];
}
format_desc = vk_format_description(desc->format);
......@@ -2631,9 +2607,9 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
num_format = radv_translate_buffer_numformat(format_desc, first_non_void);
data_format = radv_translate_buffer_dataformat(format_desc, first_non_void);
key.vertex_attribute_formats[location] = data_format | (num_format << 4);
key.vertex_attribute_bindings[location] = desc->binding;
key.vertex_attribute_offsets[location] = desc->offset;
key.vs.vertex_attribute_formats[location] = data_format | (num_format << 4);
key.vs.vertex_attribute_bindings[location] = desc->binding;
key.vs.vertex_attribute_offsets[location] = desc->offset;
const struct ac_data_format_info *dfmt_info = ac_get_data_format_info(data_format);
unsigned attrib_align =
......@@ -2643,8 +2619,8 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
* skip updating vertex_binding_align in this case.
*/
if (desc->offset % attrib_align == 0)
key.vertex_binding_align[desc->binding] =
MAX2(key.vertex_binding_align[desc->binding], attrib_align);
key.vs.vertex_binding_align[desc->binding] =
MAX2(key.vs.vertex_binding_align[desc->binding], attrib_align);
if (!uses_dynamic_stride) {
/* From the Vulkan spec 1.2.157:
......@@ -2662,7 +2638,7 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
* avoid computing a wrong offset if it's initialized
* to something else than zero.
*/
key.vertex_attribute_strides[location] =
key.vs.vertex_attribute_strides[location] =
radv_get_attrib_stride(input_state, desc->binding);
}
......@@ -2687,7 +2663,7 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
break;
}
}
key.vertex_alpha_adjust[location] = adjust;
key.vs.vertex_alpha_adjust[location] = adjust;
switch (desc->format) {
case VK_FORMAT_B8G8R8A8_UNORM:
......@@ -2703,7 +2679,7 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
case VK_FORMAT_A2R10G10B10_SSCALED_PACK32:
case VK_FORMAT_A2R10G10B10_UINT_PACK32:
case VK_FORMAT_A2R10G10B10_SINT_PACK32:
key.vertex_post_shuffle |= 1 << location;
key.vs.vertex_post_shuffle |= 1 << location;
break;
default:
break;
......@@ -2713,25 +2689,25 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
const VkPipelineTessellationStateCreateInfo *tess =
radv_pipeline_get_tessellation_state(pCreateInfo);
if (tess)
key.tess_input_vertices = tess->patchControlPoints;
key.tcs.tess_input_vertices = tess->patchControlPoints;
const VkPipelineMultisampleStateCreateInfo *vkms =
radv_pipeline_get_multisample_state(pCreateInfo);
if (vkms && vkms->rasterizationSamples > 1) {
uint32_t num_samples = vkms->rasterizationSamples;
uint32_t ps_iter_samples = radv_pipeline_get_ps_iter_samples(pCreateInfo);
key.num_samples = num_samples;
key.log2_ps_iter_samples = util_logbase2(ps_iter_samples);
key.ps.num_samples = num_samples;
key.ps.log2_ps_iter_samples = util_logbase2(ps_iter_samples);
}
key.col_format = blend->spi_shader_col_format;
key.ps.col_format = blend->spi_shader_col_format;
if (pipeline->device->physical_device->rad_info.chip_class < GFX8) {
key.is_int8 = blend->col_format_is_int8;
key.is_int10 = blend->col_format_is_int10;
key.ps.is_int8 = blend->col_format_is_int8;
key.ps.is_int10 = blend->col_format_is_int10;
}
if (pipeline->device->physical_device->rad_info.chip_class >= GFX10) {
key.topology = pCreateInfo->pInputAssemblyState->topology;
key.vs.topology = pCreateInfo->pInputAssemblyState->topology;
const VkPipelineRasterizationStateCreateInfo *raster_info = pCreateInfo->pRasterizationState;
const VkPipelineRasterizationProvokingVertexStateCreateInfoEXT *provoking_vtx_info =
......@@ -2739,7 +2715,7 @@ radv_generate_graphics_pipeline_key(const struct radv_pipeline *pipeline,
PIPELINE_RASTERIZATION_PROVOKING_VERTEX_STATE_CREATE_INFO_EXT);
if (provoking_vtx_info &&
provoking_vtx_info->provokingVertexMode == VK_PROVOKING_VERTEX_MODE_LAST_VERTEX_EXT) {
key.provoking_vtx_last = true;
key.vs.provoking_vtx_last = true;
}
}
return key;
......@@ -2759,24 +2735,24 @@ static void
radv_fill_shader_keys(struct radv_device *device, struct radv_shader_variant_key *keys,
const struct radv_pipeline_key *key, nir_shader **nir)
{
keys[MESA_SHADER_VERTEX].vs.instance_rate_inputs = key->instance_rate_inputs;
keys[MESA_SHADER_VERTEX].vs.post_shuffle = key->vertex_post_shuffle;
keys[MESA_SHADER_VERTEX].vs.instance_rate_inputs = key->vs.instance_rate_inputs;
keys[MESA_SHADER_VERTEX].vs.post_shuffle = key->vs.vertex_post_shuffle;
for (unsigned i = 0; i < MAX_VERTEX_ATTRIBS; ++i) {
keys[MESA_SHADER_VERTEX].vs.instance_rate_divisors[i] = key->instance_rate_divisors[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_formats[i] = key->vertex_attribute_formats[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_bindings[i] = key->vertex_attribute_bindings[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_offsets[i] = key->vertex_attribute_offsets[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_strides[i] = key->vertex_attribute_strides[i];
keys[MESA_SHADER_VERTEX].vs.alpha_adjust[i] = key->vertex_alpha_adjust[i];
keys[MESA_SHADER_VERTEX].vs.instance_rate_divisors[i] = key->vs.instance_rate_divisors[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_formats[i] = key->vs.vertex_attribute_formats[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_bindings[i] = key->vs.vertex_attribute_bindings[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_offsets[i] = key->vs.vertex_attribute_offsets[i];
keys[MESA_SHADER_VERTEX].vs.vertex_attribute_strides[i] = key->vs.vertex_attribute_strides[i];
keys[MESA_SHADER_VERTEX].vs.alpha_adjust[i] = key->vs.vertex_alpha_adjust[i];
}
for (unsigned i = 0; i < MAX_VBS; ++i)
keys[MESA_SHADER_VERTEX].vs.vertex_binding_align[i] = key->vertex_binding_align[i];
keys[MESA_SHADER_VERTEX].vs.outprim = si_conv_prim_to_gs_out(key->topology);
keys[MESA_SHADER_VERTEX].vs.provoking_vtx_last = key->provoking_vtx_last;
keys[MESA_SHADER_VERTEX].vs.vertex_binding_align[i] = key->vs.vertex_binding_align[i];
keys[MESA_SHADER_VERTEX].vs.provoking_vtx_last = key->vs.provoking_vtx_last;
keys[MESA_SHADER_VERTEX].vs.topology = key->vs.topology;
if (nir[MESA_SHADER_TESS_CTRL]) {
keys[MESA_SHADER_VERTEX].vs_common_out.as_ls = true;
keys[MESA_SHADER_TESS_CTRL].tcs.input_vertices = key->tess_input_vertices;
keys[MESA_SHADER_TESS_CTRL].tcs.input_vertices = key->tcs.tess_input_vertices;
}
if (nir[MESA_SHADER_GEOMETRY]) {
......@@ -2838,50 +2814,21 @@ radv_fill_shader_keys(struct radv_device *device, struct radv_shader_variant_key
for (int i = 0; i < MESA_SHADER_STAGES; ++i)
keys[i].has_multiview_view_index = key->has_multiview_view_index;
keys[MESA_SHADER_FRAGMENT].fs.col_format = key->col_format;
keys[MESA_SHADER_FRAGMENT].fs.is_int8 = key->is_int8;
keys[MESA_SHADER_FRAGMENT].fs.is_int10 = key->is_int10;
keys[MESA_SHADER_FRAGMENT].fs.log2_ps_iter_samples = key->log2_ps_iter_samples;
keys[MESA_SHADER_FRAGMENT].fs.num_samples = key->num_samples;
if (nir[MESA_SHADER_COMPUTE]) {
unsigned subgroup_size = key->compute_subgroup_size;
unsigned req_subgroup_size = subgroup_size;
bool require_full_subgroups = key->require_full_subgroups;
if (!subgroup_size)
subgroup_size = device->physical_device->cs_wave_size;
unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.workgroup_size[0] *
nir[MESA_SHADER_COMPUTE]->info.workgroup_size[1] *
nir[MESA_SHADER_COMPUTE]->info.workgroup_size[2];
/* Games don't always request full subgroups when they should,
* which can cause bugs if cswave32 is enabled.
*/
if (device->physical_device->cs_wave_size == 32 &&
nir[MESA_SHADER_COMPUTE]->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size &&
local_size % RADV_SUBGROUP_SIZE == 0)
require_full_subgroups = true;
if (require_full_subgroups && !req_subgroup_size) {
/* don't use wave32 pretending to be wave64 */
subgroup_size = RADV_SUBGROUP_SIZE;
}
keys[MESA_SHADER_COMPUTE].cs.subgroup_size = subgroup_size;
}
keys[MESA_SHADER_FRAGMENT].fs.col_format = key->ps.col_format;
keys[MESA_SHADER_FRAGMENT].fs.is_int8 = key->ps.is_int8;
keys[MESA_SHADER_FRAGMENT].fs.is_int10 = key->ps.is_int10;
keys[MESA_SHADER_FRAGMENT].fs.log2_ps_iter_samples = key->ps.log2_ps_iter_samples;
keys[MESA_SHADER_FRAGMENT].fs.num_samples = key->ps.num_samples;
}
static uint8_t
radv_get_wave_size(struct radv_device *device, const VkPipelineShaderStageCreateInfo *pStage,
gl_shader_stage stage, const struct radv_shader_variant_key *key,
const struct radv_shader_info *info)
gl_shader_stage stage, const struct radv_shader_info *info)
{
if (stage == MESA_SHADER_GEOMETRY && !info->is_ngg)
return 64;
else if (stage == MESA_SHADER_COMPUTE) {
return key->cs.subgroup_size;
return info->cs.subgroup_size;
} else if (stage == MESA_SHADER_FRAGMENT)
return device->physical_device->ps_wave_size;
else
......@@ -2890,19 +2837,21 @@ radv_get_wave_size(struct radv_device *device, const VkPipelineShaderStageCreate
static uint8_t
radv_get_ballot_bit_size(struct radv_device *device, const VkPipelineShaderStageCreateInfo *pStage,
gl_shader_stage stage, const struct radv_shader_variant_key *key)
gl_shader_stage stage, const struct radv_shader_info *info)
{
if (stage == MESA_SHADER_COMPUTE && key->cs.subgroup_size)
return key->cs.subgroup_size;
if (stage == MESA_SHADER_COMPUTE && info->cs.subgroup_size)
return info->cs.subgroup_size;
return 64;
}
static void
radv_fill_shader_info(struct radv_pipeline *pipeline,
const VkPipelineShaderStageCreateInfo **pStages,
const struct radv_pipeline_key *pipeline_key,
struct radv_shader_variant_key *keys, struct radv_shader_info *infos,
nir_shader **nir)
{
struct radv_device *device = pipeline->device;
unsigned active_stages = 0;
unsigned filled_stages = 0;
......@@ -2987,11 +2936,39 @@ radv_fill_shader_info(struct radv_pipeline *pipeline,
radv_nir_shader_info_pass(pipeline->device, nir[i], pipeline->layout, &keys[i], &infos[i]);
}
if (nir[MESA_SHADER_COMPUTE]) {
unsigned subgroup_size = pipeline_key->cs.compute_subgroup_size;
unsigned req_subgroup_size = subgroup_size;
bool require_full_subgroups = pipeline_key->cs.require_full_subgroups;
if (!subgroup_size)
subgroup_size = device->physical_device->cs_wave_size;
unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.workgroup_size[0] *
nir[MESA_SHADER_COMPUTE]->info.workgroup_size[1] *
nir[MESA_SHADER_COMPUTE]->info.workgroup_size[2];
/* Games don't always request full subgroups when they should,
* which can cause bugs if cswave32 is enabled.
*/
if (device->physical_device->cs_wave_size == 32 &&
nir[MESA_SHADER_COMPUTE]->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size &&
local_size % RADV_SUBGROUP_SIZE == 0)
require_full_subgroups = true;
if (require_full_subgroups && !req_subgroup_size) {
/* don't use wave32 pretending to be wave64 */
subgroup_size = RADV_SUBGROUP_SIZE;
}
infos[MESA_SHADER_COMPUTE].cs.subgroup_size = subgroup_size;
}
for (int i = 0; i < MESA_SHADER_STAGES; i++) {
if (nir[i]) {
infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &keys[i], &infos[i]);
infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &infos[i]);
infos[i].ballot_bit_size =
radv_get_ballot_bit_size(pipeline->device, pStages[i], i, &keys[i]);
radv_get_ballot_bit_size(pipeline->device, pStages[i], i, &infos[i]);
}
}
......@@ -3058,7 +3035,7 @@ gather_tess_info(struct radv_device *device, nir_shader **nir, struct radv_shade
{
merge_tess_info(&nir[MESA_SHADER_TESS_EVAL]->info, &nir[MESA_SHADER_TESS_CTRL]->info);
unsigned tess_in_patch_size = pipeline_key->tess_input_vertices;
unsigned tess_in_patch_size = pipeline_key->tcs.tess_input_vertices;
unsigned tess_out_patch_size = nir[MESA_SHADER_TESS_CTRL]->info.tess.tcs_vertices_out;
/* Number of tessellation patches per workgroup processed by the current pipeline. */
......@@ -3468,7 +3445,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device,
}
radv_fill_shader_keys(device, keys, pipeline_key, nir);
radv_fill_shader_info(pipeline, pStages, keys, infos, nir);
radv_fill_shader_info(pipeline, pStages, pipeline_key, keys, infos, nir);
bool pipeline_has_ngg = (nir[MESA_SHADER_VERTEX] && keys[MESA_SHADER_VERTEX].vs_common_out.as_ngg) ||
(nir[MESA_SHADER_TESS_EVAL] && keys[MESA_SHADER_TESS_EVAL].vs_common_out.as_ngg);
......@@ -3641,7 +3618,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device,
pipeline->shaders[MESA_SHADER_FRAGMENT] = radv_shader_variant_compile(
device, modules[MESA_SHADER_FRAGMENT], &nir[MESA_SHADER_FRAGMENT], 1, pipeline->layout,
keys + MESA_SHADER_FRAGMENT, infos + MESA_SHADER_FRAGMENT, keep_executable_info,
pipeline_key, infos + MESA_SHADER_FRAGMENT, keep_executable_info,
keep_statistic_info, disable_optimizations, &binaries[MESA_SHADER_FRAGMENT]);
radv_stop_feedback(stage_feedbacks[MESA_SHADER_FRAGMENT], false);
......@@ -3657,7 +3634,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device,
radv_start_feedback(stage_feedbacks[MESA_SHADER_TESS_CTRL]);
pipeline->shaders[MESA_SHADER_TESS_CTRL] = radv_shader_variant_compile(
device, modules[MESA_SHADER_TESS_CTRL], combined_nir, 2, pipeline->layout, key,
device, modules[MESA_SHADER_TESS_CTRL], combined_nir, 2, pipeline->layout, pipeline_key,
&infos[MESA_SHADER_TESS_CTRL], keep_executable_info, keep_statistic_info,
disable_optimizations, &binaries[MESA_SHADER_TESS_CTRL]);
......@@ -3675,8 +3652,8 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device,
radv_start_feedback(stage_feedbacks[MESA_SHADER_GEOMETRY]);
pipeline->shaders[MESA_SHADER_GEOMETRY] = radv_shader_variant_compile(
device, modules[MESA_SHADER_GEOMETRY], combined_nir, 2, pipeline->layout,
&keys[pre_stage], &infos[MESA_SHADER_GEOMETRY], keep_executable_info,
device, modules[MESA_SHADER_GEOMETRY], combined_nir, 2, pipeline->layout, pipeline_key,
&infos[MESA_SHADER_GEOMETRY], keep_executable_info,
keep_statistic_info, disable_optimizations, &binaries[MESA_SHADER_GEOMETRY]);
radv_stop_feedback(stage_feedbacks[MESA_SHADER_GEOMETRY], false);
......@@ -3689,7 +3666,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device,
radv_start_feedback(stage_feedbacks[i]);
pipeline->shaders[i] = radv_shader_variant_compile(
device, modules[i], &nir[i], 1, pipeline->layout, keys + i, infos + i,
device, modules[i], &nir[i], 1, pipeline->layout, pipeline_key, infos + i,
keep_executable_info, keep_statistic_info, disable_optimizations, &binaries[i]);
radv_stop_feedback(stage_feedbacks[i], false);
......@@ -5728,9 +5705,9 @@ radv_generate_compute_pipeline_key(struct radv_pipeline *pipeline,
if (subgroup_size) {
assert(subgroup_size->requiredSubgroupSize == 32 ||
subgroup_size->requiredSubgroupSize == 64);
key.compute_subgroup_size = subgroup_size->requiredSubgroupSize;
key.cs.compute_subgroup_size = subgroup_size->requiredSubgroupSize;
} else if (stage->flags & VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT_EXT) {
key.require_full_subgroups = true;
key.cs.require_full_subgroups = true;
}
return key;
......
......@@ -347,34 +347,6 @@ struct radv_pipeline_cache {
VkAllocationCallbacks alloc;
};
struct radv_pipeline_key {
uint32_t instance_rate_inputs;
uint32_t instance_rate_divisors[MAX_VERTEX_ATTRIBS];
uint8_t vertex_attribute_formats[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_bindings[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_offsets[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_strides[MAX_VERTEX_ATTRIBS];
uint8_t vertex_binding_align[MAX_VBS];
enum ac_fetch_format vertex_alpha_adjust[MAX_VERTEX_ATTRIBS];
uint32_t vertex_post_shuffle;
unsigned tess_input_vertices;
uint32_t col_format;
uint32_t is_int8;
uint32_t is_int10;
uint8_t log2_ps_iter_samples;
uint8_t num_samples;
uint32_t has_multiview_view_index : 1;
uint32_t optimisations_disabled : 1;
uint32_t provoking_vtx_last : 1;
uint8_t topology;
/* Non-zero if a required subgroup size is specified via
* VK_EXT_subgroup_size_control.
*/
uint8_t compute_subgroup_size;
bool require_full_subgroups;
};
struct radv_shader_binary;
struct radv_shader_variant;
struct radv_pipeline_shader_stack_size;
......@@ -1687,6 +1659,8 @@ struct radv_event {
#define RADV_HASH_SHADER_ROBUST_BUFFER_ACCESS (1 << 14)
#define RADV_HASH_SHADER_ROBUST_BUFFER_ACCESS2 (1 << 15)
struct radv_pipeline_key;
void radv_hash_shaders(unsigned char *hash, const VkPipelineShaderStageCreateInfo **stages,
const struct radv_pipeline_layout *layout,
const struct radv_pipeline_key *key, uint32_t flags);
......@@ -2716,6 +2690,30 @@ si_conv_gl_prim_to_vertices(unsigned gl_prim)
}
}
static inline uint32_t
si_conv_prim_to_gs_out(enum VkPrimitiveTopology topology)
{
switch (topology) {
case VK_PRIMITIVE_TOPOLOGY_POINT_LIST:
case VK_PRIMITIVE_TOPOLOGY_PATCH_LIST:
return V_028A6C_POINTLIST;
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP:
case VK_PRIMITIVE_TOPOLOGY_LINE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_LINE_STRIP_WITH_ADJACENCY:
return V_028A6C_LINESTRIP;
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_FAN:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST_WITH_ADJACENCY:
case VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP_WITH_ADJACENCY:
return V_028A6C_TRISTRIP;
default:
assert(0);
return 0;
}
}
struct radv_extra_render_pass_begin_info {
bool disable_dcc;
};
......
......@@ -418,13 +418,13 @@ radv_shader_compile_to_nir(struct radv_device *device, struct vk_shader_module *
const struct radv_pipeline_key *key)
{
unsigned subgroup_size = 64, ballot_bit_size = 64;
if (key->compute_subgroup_size) {
if (key->cs.compute_subgroup_size) {
/* Only compute shaders currently support requiring a
* specific subgroup size.
*/
assert(stage == MESA_SHADER_COMPUTE);
subgroup_size = key->compute_subgroup_size;
ballot_bit_size = key->compute_subgroup_size;
subgroup_size = key->cs.compute_subgroup_size;
ballot_bit_size = key->cs.compute_subgroup_size;
}
nir_shader *nir;
......@@ -857,7 +857,7 @@ radv_lower_io_to_mem(struct radv_device *device, struct nir_shader *nir,
nir, device->physical_device->rad_info.chip_class, info->tcs.tes_reads_tess_factors,
info->tcs.tes_inputs_read, info->tcs.tes_patch_inputs_read, info->tcs.num_linked_inputs,
info->tcs.num_linked_outputs, info->tcs.num_linked_patch_outputs, true);
ac_nir_lower_tess_to_const(nir, pl_key->tess_input_vertices, info->num_tess_patches,
ac_nir_lower_tess_to_const(nir, pl_key->tcs.tess_input_vertices, info->num_tess_patches,
ac_nir_lower_patch_vtx_in | ac_nir_lower_num_patches);
return true;
......@@ -949,7 +949,7 @@ void radv_lower_ngg(struct radv_device *device, struct nir_shader *nir,
} else if (nir->info.stage == MESA_SHADER_VERTEX) {
/* Need to add 1, because: V_028A6C_POINTLIST=0, V_028A6C_LINESTRIP=1, V_028A6C_TRISTRIP=2, etc. */
num_vertices_per_prim = key->vs.outprim + 1;
num_vertices_per_prim = si_conv_prim_to_gs_out(pl_key->vs.topology) + 1;
/* Manually mark the instance ID used, so the shader can repack it. */
if (key->vs.instance_rate_inputs)
......@@ -1633,7 +1633,7 @@ struct radv_shader_variant *
radv_shader_variant_compile(struct radv_device *device, struct vk_shader_module *module,
struct nir_shader *const *shaders, int shader_count,
struct radv_pipeline_layout *layout,
const struct radv_shader_variant_key *key,
const struct radv_pipeline_key *key,
struct radv_shader_info *info, bool keep_shader_info,
bool keep_statistic_info, bool disable_optimizations,
struct radv_shader_binary **binary_out)
......
......@@ -76,8 +76,8 @@ struct radv_vs_variant_key {
/* For some formats the channels have to be shuffled. */
uint32_t post_shuffle;
/* Output primitive type. */
uint8_t outprim;
/* Topology. */
uint8_t topology;
/* Provoking vertex mode. */
bool provoking_vtx_last;
......@@ -100,17 +100,12 @@ struct radv_fs_variant_key {
uint32_t is_int10;
};
struct radv_cs_variant_key {
uint8_t subgroup_size;
};
struct radv_shader_variant_key {
union {
struct radv_vs_variant_key vs;
struct radv_fs_variant_key fs;
struct radv_tes_variant_key tes;
struct radv_tcs_variant_key tcs;
struct radv_cs_variant_key cs;
/* A common prefix of the vs and tes keys. */
struct radv_vs_out_key vs_common_out;
......@@ -118,6 +113,45 @@ struct radv_shader_variant_key {
bool has_multiview_view_index;
};
struct radv_pipeline_key {
uint32_t has_multiview_view_index : 1;
uint32_t optimisations_disabled : 1;
struct {
uint32_t instance_rate_inputs;
uint32_t instance_rate_divisors[MAX_VERTEX_ATTRIBS];
uint8_t vertex_attribute_formats[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_bindings[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_offsets[MAX_VERTEX_ATTRIBS];
uint32_t vertex_attribute_strides[MAX_VERTEX_ATTRIBS];
uint8_t vertex_binding_align[MAX_VBS];
enum ac_fetch_format vertex_alpha_adjust[MAX_VERTEX_ATTRIBS];
uint32_t vertex_post_shuffle;
uint32_t provoking_vtx_last : 1;
uint8_t topology;
} vs;
struct {
unsigned tess_input_vertices;
} tcs;
struct {
uint32_t col_format;
uint32_t is_int8;
uint32_t is_int10;
uint8_t log2_ps_iter_samples;
uint8_t num_samples;
} ps;
struct {
/* Non-zero if a required subgroup size is specified via
* VK_EXT_subgroup_size_control.
*/
uint8_t compute_subgroup_size;
bool require_full_subgroups;
} cs;
};
enum radv_compiler_debug_level {
RADV_COMPILER_DEBUG_LEVEL_PERFWARN,
RADV_COMPILER_DEBUG_LEVEL_ERROR,
......@@ -125,7 +159,7 @@ enum radv_compiler_debug_level {
struct radv_nir_compiler_options {
struct radv_pipeline_layout *layout;
struct radv_shader_variant_key key;
struct radv_pipeline_key key;
bool explicit_scratch_args;
bool clamp_shadow_reference;
bool robust_buffer_access;
......@@ -346,6 +380,8 @@ struct radv_shader_info {
bool uses_local_invocation_idx;
unsigned block_size[3];
uint8_t subgroup_size;
bool uses_sbt;
bool uses_ray_launch_size;
} cs;
......@@ -458,7 +494,7 @@ struct radv_shader_variant *radv_shader_variant_create(struct radv_device *devic
bool keep_shader_info);
struct radv_shader_variant *radv_shader_variant_compile(
struct radv_device *device, struct vk_shader_module *module, struct nir_shader *const *shaders,
int shader_count, struct radv_pipeline_layout *layout, const struct radv_shader_variant_key *key,
int shader_count, struct radv_pipeline_layout *layout, const struct radv_pipeline_key *key,
struct radv_shader_info *info, bool keep_shader_info, bool keep_statistic_info,
bool disable_optimizations, struct radv_shader_binary **binary_out);
......
......@@ -205,8 +205,6 @@ spec@ext_framebuffer_multisample@blit-mismatched-formats,Fail
spec@ext_framebuffer_multisample@interpolation 2 centroid-edges,Fail
spec@ext_framebuffer_multisample@interpolation 4 centroid-edges,Fail
spec@ext_framebuffer_object@fbo-blending-format-quirks,Fail
spec@ext_framebuffer_object@fbo-blending-formats,Fail
spec@ext_framebuffer_object@fbo-blending-formats@GL_RGB10,Fail
spec@ext_framebuffer_object@getteximage-formats init-by-clear-and-render,Fail
spec@ext_framebuffer_object@getteximage-formats init-by-rendering,Fail
spec@ext_gpu_shader4@execution@texelfetch@fs-texelfetch-isampler1darray,Fail
......
......@@ -4,6 +4,7 @@ glx@glx-multithread-texture
spec@arb_internalformat_query2@all internalformat_<x>_type pname checks
spec@!opengl 1.1@streaming-texture-leak
spec@!opengl 1.0@gl-1.0-blend-func
shaders@glsl-predication-on-large-array
# Extensions not supported
spec@arb_gpu_shader_fp64.*
......
......@@ -1964,17 +1964,20 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
spirv_op_to_string(opcode), elem_count, val->type->length);
nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
val->is_undef_constant = true;
for (unsigned i = 0; i < elem_count; i++) {
struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);
struct vtn_value *elem_val = vtn_untyped_value(b, w[i + 3]);
if (val->value_type == vtn_value_type_constant) {
elems[i] = val->constant;
if (elem_val->value_type == vtn_value_type_constant) {
elems[i] = elem_val->constant;
val->is_undef_constant = val->is_undef_constant &&
elem_val->is_undef_constant;
} else {
vtn_fail_if(val->value_type != vtn_value_type_undef,
vtn_fail_if(elem_val->value_type != vtn_value_type_undef,
"only constants or undefs allowed for "
"SpvOpConstantComposite");
/* to make it easier, just insert a NULL constant for now */
elems[i] = vtn_null_constant(b, val->type);
elems[i] = vtn_null_constant(b, elem_val->type);
}
}
......@@ -6039,6 +6042,29 @@ vtn_create_builder(const uint32_t *words, size_t word_count,
(b->generator_id == vtn_generator_glslang_reference_front_end &&
generator_version < 3);
/* Identifying the LLVM-SPIRV translator:
*
* The LLVM-SPIRV translator currently doesn't store any generator ID [1].
* Our use case involving the SPIRV-Tools linker also mean we want to check
* for that tool instead. Finally the SPIRV-Tools linker also stores its
* generator ID in the wrong location [2].
*
* [1] : https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/1223
* [2] : https://github.com/KhronosGroup/SPIRV-Tools/pull/4549
*/
const bool is_llvm_spirv_translator =
(b->generator_id == 0 &&
generator_version == vtn_generator_spirv_tools_linker) ||
b->generator_id == vtn_generator_spirv_tools_linker;
/* The LLVM-SPIRV translator generates Undef initializers for _local
* variables [1].
*
* [1] : https://github.com/KhronosGroup/SPIRV-LLVM-Translator/issues/1224
*/
b->wa_llvm_spirv_ignore_workgroup_initializer =
b->options->environment == NIR_SPIRV_OPENCL && is_llvm_spirv_translator;
/* words[2] == generator magic */
unsigned value_id_bound = words[3];
if (words[4] != 0) {
......
......@@ -604,6 +604,9 @@ struct vtn_value {
/* Valid for vtn_value_type_constant to indicate the value is OpConstantNull. */
bool is_null_constant:1;
/* Valid when all the members of the value are undef. */
bool is_undef_constant:1;
const char *name;
struct vtn_decoration *decoration;
struct vtn_type *type;
......@@ -696,6 +699,9 @@ struct vtn_builder {
/* True if we need to fix up CS OpControlBarrier */
bool wa_glslang_cs_barrier;
/* True if we need to ignore undef initializers */
bool wa_llvm_spirv_ignore_workgroup_initializer;
/* Workaround discard bugs in HLSL -> SPIR-V compilers */
bool uses_demote_to_helper_invocation;
bool convert_discard_to_demote;
......
......@@ -2015,7 +2015,16 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
unreachable("Should have been caught before");
}
if (initializer) {
/* Ignore incorrectly generated Undef initializers. */
if (b->wa_llvm_spirv_ignore_workgroup_initializer &&
initializer &&
storage_class == SpvStorageClassWorkgroup)
initializer = NULL;
/* Only initialize variable when there is an initializer and it's not
* undef.
*/
if (initializer && !initializer->is_undef_constant) {
switch (storage_class) {
case SpvStorageClassWorkgroup:
/* VK_KHR_zero_initialize_workgroup_memory. */
......@@ -2328,6 +2337,7 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
case SpvOpUndef: {
struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
val->type = vtn_get_type(b, w[1]);
val->is_undef_constant = true;
break;
}
......
......@@ -67,9 +67,6 @@ dEQP-VK.spirv_assembly.instruction.graphics.variable_pointers.graphics.writes_si
dEQP-VK.spirv_assembly.instruction.graphics.variable_pointers.graphics.writes_two_buffers_geom,Fail
dEQP-VK.spirv_assembly.instruction.graphics.variable_pointers.graphics.writes_two_buffers_vert,Fail
# Broken on all drivers: https://gitlab.freedesktop.org/mesa/mesa/-/issues/4582
dEQP-VK.wsi.display_control.register_device_event,Fail
# https://gitlab.khronos.org/Tracker/vk-gl-cts/-/issues/3052
# fixed by https://gerrit.khronos.org/c/vk-gl-cts/+/7837
bypass-dEQP-VK.renderpass.suballocation.subpass_dependencies.separate_channels.r8g8b8a8_unorm,Fail