Abstract-Det huvudsakliga syftet med detta dokument är att påskynda simuleringen prestanda SystemC konstruktioner på RTL abstraktionsnivå genom att utnyttja den höga graden av parallellism som erbjuds av dagens allmänna ändamål grafikprocessorer (GPGPU). Vårt tillvägagångssätt parallelliserar SystemC:s diskret-eventssimulering (DES) på GPGPU:er genom att omvandla modellen för beräkning av DES till en modell av samtidiga trådar som synkroniserar vid behov. Till skillnad från den kooperativa gängmodellen som används i SystemC:s referensimplementering, kan vår gängmodell utföra parallellt på det stora antalet enkla bearbetningsenheter som finns tillgängliga på GPU:er. Vår simuleringsinfrastruktur kallas SCGPSim 1 och den innehåller en källkodsöversättare (S2S) för att omvandla synthesizable SystemC-modeller till parallellt körbara program som riktar sig till en NVIDIA GPU. Översättaren behåller simuleringssemantiken i originaldesignen genom att tillämpa semantik som bevarar transformationer. De resulterande transformerade modeller som kartlagts på den massivt parallella arkitekturen hos GPU:er förbättrar simuleringseffektiviteten avsevärt. Preliminära experiment med olika stora exempel som AES, ALU och FIR har visat simuleringshastigheter från 30x till 100x. Med tanke på att våra transformationer ännu inte är optimerade tror vi att optimeringen kommer att förbättra simuleringsprestandan ytterligare. SystemC, ett modellerings- och simuleringsspråk, har använts för att göra tidig kompromissanalys och design-rymd prospektering. Det har visat sig att systemC-metoder förkortar konstruktionscykler som gör det möjligt för konstruktörer att möta de ständigt strama tid-till-marknad begränsningarna. En viktig bidragande faktor till SystemC:s framgång är dess förmåga att tillhandahålla mekanismer för snabb simulering. SystemC:s referensgenomförande gör dock inga försök att utnyttja de parallella arkitekturer som blir allt populärare. OSCI-implementationen av SystemC byggs med hjälp av trådbibliotek på applikationsnivå (kooperativa trådar) som resulterar i en oparalleliserbar simuleringskärna, eftersom kooperativa trådar inte kan schemaläggas av operativsystemet och därför inte kan sändas till olika bearbetningselement. Det är uppenbart att detta genomförande inte är utformat för att dra nytta av flerkärniga bearbetningssystem. En betydande insats har gått till att förbättra prestandan hos SystemC simuleringar [9, 7, 6 ], men bara en handfull försök att parallellisera dem [2, 8]. Detta dokument presenterar en källa-till-källa (S2S) översättare som omvandlar synthesizable SystemC designs till parallellt exe-1 Detta arbete finansierades delvis av en Air Force Labs, en AFOSR finansiering och NSF-CPA finansiering. cutable CUDA-program som körs på NVIDIA Graphics Processing Units (GPU). För att omvandla SystemC-konstruktionerna till CUDA-konstruktioner använder vi en programmeringsmodell där cooperativa trådar kan kartläggas till parallella trådar, och de kan simuleras effektivt vilket ger en storleksordningsminskning i simuleringstider under begränsningen att bevara befintliga SystemC-simuleringssemantik. Observera att vi använder semantiska-bevarande transformationer för att få exakt samma beteenden som när vi använder SystemC diskret-event schemaläggare. Vår översättare accepterar synthesizable SystemC-design; men vi kräver moduler i designen för att ha registrerade utgångar. Därför stöder vi för närvarande inte kombinatoriska vägar mellan moduler. Vi använder NVIDIA TESLA D870 GPU som exempel på målplattform. Detta dokument är organiserat enligt följande, Avsnitt 2 presenterar relaterade arbete. I Avsnitt 3 presenterar vi en kort introduktion på SystemC-simulatorn. I avsnitt 4 presenteras GPU-arkitekturen och CUDA-programmeringsmodellen. I Avsnitt 5 förklarar vi konverteringsförfarandet från SystemC till CUDA och den översättare som gör denna automatiska konvertering. Vi presenterar våra resultat i Avsnitt 6 och avslutar med framtida arbete i det sista avsnittet. Flera försök har gjorts i det förflutna att parallellisera en SystemC simulator, även om ingen av dem är riktade mot en massivt parallell plattform såsom en GPGPU. Denna förordning träder i kraft den tjugonde dagen efter det att den har offentliggjorts i Europeiska unionens officiella tidning. Visa en statisk omvandlingsalgoritm för SystemC-modeller. Detta arbete försöker förbättra simuleringen effektivitet genom att minska synkronisering overheads genom att ändra strukturen av modellen helt för att sammanfoga de olika modulerna och kooperativa trådar till en mycket högre granularitet gänga. År [2], Kaouane m.fl. har föreslagit ett konstruktionsflöde för en CELL-processorbaserad SystemC-simulering. Med begränsat antal synergistiska bearbetningselement, kräver det en hel del sammanhangsbrytare som minskar effektiviteten. Detta genomförande har också flera begränsningar, till exempel att processerna är känsliga för en enda signal, vilket är mycket sällsynt inom hårdvarumodellering. Dessutom delade de manuellt SystemC kod för att extrahera beräkna intensiva och kontrollera intensiva delar. Även om mycket arbete har gjorts i parallell diskret-event simulering, så vitt vi vet ingen verklig parallell implementering av SystemC existerar. I [8] Perumalla utvärderade diskret-event simulering på GPUs, men det gav ingen verklig hastighet upp. I själva verket var simuleringstiden mer än vad CPU skulle ta utan GPU:erna. Naguib m.fl. [3] visade att man genom processdelning kan påskynda SystemC-simuleringarna genom att automatiskt optimera modellen för simulering med resultat som visade 978-1-4244-5767-0/10/$26.00 2010 IEEE 2C-2 149 15% förbättring. Vårt tillvägagångssätt visar på en betydande förbättring av simuleringsprestandan jämfört med befintliga tillvägagångssätt. Vi implementerar en GPU-baserad cykelnoggrannhetssimulator genom att automatiskt utföra en käll-till-källa-omvandling på SystemC-konstruktioner till CUDA och utföra den effektivt på GPU. SystemC är ett designspråk på öppen källkod baserat på C++ som har sin egen simuleringskärna. SystemC-simuleringskärnan är en schemaläggare som hanterar både synkronisering och schemaläggning av samtidiga processer. Det implementerar en diskret-event schemaläggare som utför processer som svar på händelser. När en händelse inträffar, skjuts motsvarande process in i körbara processer kö för utförande. Schemaläggaren upprätthåller en kö av körbara processer, och det fortsätter att utföra dem tills de slutför utförandet eller nå en suspension vänta () uttalande. Kör alla processer slumpmässigt Bild 1 visar flödet av SystemC-simuleringskärnan. I initieringsfasen utförs alla processer i en ospecificerad ordning. I utvärderingsfasen väljs och utförs runnable-processer tills inga runnable-processer finns. Detta kan orsaka ytterligare processer för att bli redo att köras i denna fas på grund av omedelbara händelser anmälningar. När det inte finns några sådana färdiga processer går simulatorn in i uppdateringsfasen, där signalvärden uppdateras till de värden som beräknas i utvärderingsfasen. Vid denna tidpunkt, om det finns några väntande försenade anmälningar, går simulatorn in i den försenade anmälningsfasen som avgör vilka processer som är redo att köras på grund av de försenade anmälningarna och återgår till utvärderingsfasen. I annat fall, om det finns tidsinställda meddelanden, går simulatorn in i den tidsinställda anmälningsfasen där den för fram den aktuella simuleringstiden till den tidigaste tidsinställda anmälan och återinträder i utvärderingsfasen. Om det inte finns några tidsinställda meddelanden är simuleringen klar. Även om detta tillvägagångssätt fungerar perfekt för en enkärnig processor, det inte skala och dra nytta av multi-core/multi-processor miljö. IV CUDA PROGRAMING MODEL CUDA (Compute Unified Device Architecture) är en förlängning till C språk som utnyttjar bearbetningskraften hos GPU:er för att lösa komplexa beräkna-intensiva problem på ett effektivt sätt. Hög prestanda uppnås genom att lansera ett antal trådar och göra varje tråd köra en del av programmet parallellt. Dessa trådar är grupperade och ordnade som block. Varje block innehåller maximalt 512 trådar. Dessa block är ordnade i tvådimensionella matrisstrukturer för att bilda ett rutnät. Varje rutnät består av högst 65536 block. Figur 2 visar strukturen. Trådar från samma blockåtkomst snabbt delade on-chip-minnen och kan synkroniseras med inbyggda funktioner. Trådar från olika block kan kommunicera data med hjälp av det långsammare globala minnet. CUDA tråd exekvering skiljer sig från CPU-tråd exekvering i termer av schemaläggning av trådar. I CUDA är trådar grupperade för att bilda warps och dessa warps utförs i en-instruktion multipeltråd (SIMT) sätt. Alla trådar i en warp utför en instruktion åt gången. Om trådar som tillhör en warp avviker vid en databeroende villkorlig gren, utförs de sekventiellt och de alla konvergerar till samma körbana i senare instruktioner. Avdelningskontorets divergens gäller endast för trådar av en viss varp; trådar från olika varp utför självständigt på olika multiprocessorer. Vi måste skapa trådar som tillhör olika warps för att få upp maxhastigheten. HOST DEVICE Flaskhalsen i CUDA-programmeringsmodellen är dataöverföringshastigheten mellan den centrala bearbetningsenheten (CPU) och den grafiska bearbetningsenheten (GPU). För närvarande begränsas detta av PCIe-bussens maximala hastighet. Den övergripande prestandan hos applikationerna beror i hög grad på strategin för användning av GPU-minnet på chip. Frekvent minnesöverföring mellan CPU och GPU försämrar prestandan. Detta hindrar inte effektiviteten i vårt fall eftersom vi inte kommunicerar med GPU ofta. V SYSTEMC TILL CUDA En allmän SystemC-design består av tillverkare/stimulatorer, konsument/respons och kanal/beräkningsmoduler. Kanalen består av många moduler som modellerar beteendet hos hårdvaran. Det kan finnas en eller flera producenter och konsumenter. Figur 3 visar utförandemodellen för N-runnbara processer med en producent och en konsumentprocess. Simuleringen av konstruktionen fortsätter tills alla processer blir avstängda, d.v.s. Slutet på simuleringen är uppfyllt. Trådar skapade under SystemC simulering är kooperativa trådar och de utför effektivt sekventiellt även i en multiprocessormiljö. Figur 4 visar den föreslagna modellen 2C-2 150 för parallell körning av SystemC-konstruktionen. Enskilda trådar utförs parallellt och barriärer används för synkronisering mellan alla trådar. För att förhindra rastillståndet på delade variabler använder vi "Double buffring" på sådana variabler. I figur 4, trådarna utför i vertikal riktning och dataflöden i horisontell riktning. Trådar är färgade grå för viss del under utförandet cykel. I den gråfärgade regionen gör just den tråden meningsfulla beräkningar, annars väntar den bara vid barriären för de andra trådarna att slutföra sitt arbete. I början av simuleringscykeln börjar alla trådar köra parallellt, men endast Stimulus tråd gör meningsfullt arbete medan resten av trådarna träffar Barrier-1 och väntar på synkronisering med andra trådar. Det här är Barriärsynkronisering. Efter Barriär-1 väntar Stimulus tråd bara på Barriär-2, medan trådarna 2 till N-1 kör och når Barriär-2. Konsumenttråden som väntade på Barrier-2, gör nu meningsfullt arbete genom att spara den aktuella cykelutmatningen. Alla delade variabler har 2 kopior (Dubbel buffringsteknik) av sig själva. I slutet av Barrier-2 gör vi innehållet i båda kopiorna av variablerna detsamma, så att nästa nya data läses upp i nästa cykel. Barriär-1 säkerställer att inmatningen är tillgänglig innan kanaldriften startas, medan Barriär-2 säkerställer att utgångarna från nuvarande cykel lagras och att dubbla buffertar har samma innehåll i dem. Koncurrency extraheras genom att utföra trådarna 2 till N-1 parallellt. I detta avsnitt beskriver vi strukturen för CUDA-programmet som används i vårt genomförande. I CUDA-arkitekturen kan alla trådar se hela KERNEL-koden. Det enda sättet att få varje tråd att köra en del av hela koden är genom att särskilja den delen av koden med ett threadId. Listning-1 visar en CUDA KERNEL CODE struktur med fyra trådar som utför olika delar av hela koden för MAX CYCLES antal gånger. Den inbyggda funktionssynkroniseringen () används för synkronisering mellan de olika trådarna. Figur 5 visar det totala konstruktionsflödet för SCGPSim. Vi delar upp hela processen av SystemC simulering på CUDA i 5 olika faser som förklaras nedan i detalj.  / Code t h a t has t o has e x e c u t e d by t h r e a d 3 } __syncthreads ( ) ; / / B a r r i e r S y n c h r o n i z a t i o n } } C Design Flow SystemC design
Ett parallellt genomförande av SystemC-kärnan på en allmän grafisk bearbetningsenhet (GPGPU) införs i REF.
SCGPSim: A fast SystemC simulator on GPUs
{'venue': '2010 15th Asia and South Pacific Design Automation Conference (ASP-DAC)', 'journal': '2010 15th Asia and South Pacific Design Automation Conference (ASP-DAC)', 'mag_field_of_study': ['Computer Science']}