Kan endast allokera 2GB av 8GB på GeForce GTX 1080

Permalänk
Medlem

Kan endast allokera 2GB av 8GB på GeForce GTX 1080

Hej

Har precis stött på en begränsning i OpenCL som gör att jag endast kan allokera 2GB på minnet som finns på grafikkortet.

CL_DEVICE_MAX_MEM_ALLOC_SIZE: 2048 MByte
CL_DEVICE_GLOBAL_MEM_SIZE: 8192 MByte

Är detta en begränsning i specifikationen eller vet någon hur jag kan nyttja alla 8GB?

När jag använder

clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, Sizeof.cl_double * n, dst, null);

Så får jag:
Allocated global memory for input data: 4119.873 MB
Allocated global memory for result data: 4119.873 MB

GPU init was 294 ms
Exception in thread "main" org.jocl.CLException: CL_OUT_OF_HOST_MEMORY
at org.jocl.CL.checkResult(CL.java:787)
at org.jocl.CL.clCreateBuffer(CL.java:6013)
at DriverDP.main(DriverDP.java:155)

Permalänk
Medlem
Permalänk
Medlem

Tack för det.

Synd med en sådan begränsning, med mitt förra kort (AMD R9 380) så kunde jag allokera 100%.

Hoppas Kronos ändrar detta så snart som möjligt.

Tänk att köpa en kort med 12GB och sedan bara kunna använda kanske 3GB, inte helt ok.

Permalänk
Medlem

Får man fråga vad det är för projekt som du jobbar på?

Permalänk
Medlem
Skrivet av Alotiat:

Får man fråga vad det är för projekt som du jobbar på?

Visst

Jag håller på med ett redigerings program för videos (4K upplösning).

Istället för att utföra renderingen på CPU så görs det via GPUn (målet med programmet)

Tanken är även att det ska finnas stöd för multi GPU. Dvs att en video då kan redigeras på tex. 4st GTX 1080 parallellt.

Idag använder jag bara 1st GTX 1080 och det klarar att rendera 10 min 4K video på ca 4min.
Detta tar ca 23 min när jag använder 16st trådar på min Ryzen 1700X.

Men tyvärr så har jag stött på lite problem så att jag bara kan använda 2GB av minnet på grafikkortet.
Upplösningen 4K tar en hel del utrymme

.

Permalänk
Datavetare

Nvidia verkar initialt ha tolkat detta från specifikationen på fel sätt

CL_DEVICE_MAX_MEM_ALLOC_SIZE Max size of memory object allocation in bytes. The minimum value is max (1/4th of CL_DEVICE_GLOBAL_MEM_SIZE, 128*1024*1024)

Står ju faktiskt att CL_DEVICE_MAX_MEM_ALLOC_SIZE minst ska vara det största av CL_DEVICE_GLOBAL_MEM_SIZE/4 och 128 MB, står inget om att CL_DEVICE_MAX_MEM_ALLOC_SIZE inte kan vara större.

Det verkar finnas någon teknisk anledning för varför man inte vill sätta CL_DEVICE_MAX_MEM_ALLOC_SIZE==CL_DEVICE_GLOBAL_MEM_SIZE då inte heller AMD eller Intel gör detta. I AMDs fall verkar standardvärdet vara

CL_DEVICE_MAX_MEM_ALLOC_SIZE=min(4 GB, CL_DEVICE_GLOBAL_MEM_SIZE/2)

men går tydligen att ändra konfigurationen för AMD så man kan få upp till 90 % av CL_DEVICE_GLOBAL_MEM_SIZE (men har aldrig sett det gå över 4 GB).

I alla implementationer kan man allokera allt VRAM, men som @Snorren påpekar kan det kräva flera allokeringar.

Då du kör på GTX1080, varför kör du OpenCL och inte CUDA? CUDA har inte denna begränsning. Vidare är CUDA klart populärare jämfört med OpenCL av en anledning: är enklare och därmed snabbare att jobba i CUDA.

Detta är något Khronos måste fokusera på, CUDA är låst till Nvidia men på det stora hela är CUDA så pass mycket enklare att jobba i att det rätt snabbt blir värt kostnaden att låsa sig till Nvidia. SyCL ser ut att kunna vara precis det öppna svar på CUDA som vi väntat på

Om det finns orsaker till att köra just på OpenCL: har du testat att köra OpenCL på din CPU?
GPGPU på en dGPU är överlägset om flaskhalsen är minnesbandbredd eller FP32 kapacitet.

Bildbehandling kan ju mycket väl betyda heltalsoperationer och beroende på vad du gör kanske det inte alls är bandbreddskrävande (bara du kan svara på detta alt. beskriva mer i detalj var ditt program gör). I det läget kan faktiskt OpenCL på CPU prestera bättre jämfört med på GPU.

Det jag har noterat är att i alla fall Intels OpenCL implementation för CPU tenderar väldigt mycket bättre om man skickar in NULL som local_work_size argument till clEnqueueNDRangeKernel(). Detta kräver självklart att varje work-thread kan utföra sitt arbete oberoende från andra trådar i sin "work-group".

Om du inte handkodat SSE/AVX stöd i din CPU-version borde du se ca 4x prestandaboost om du jobbar med 32-bitars tal. Bildbehandling kan ju betyda att du jobbar med varje färgkanal separat och i så fall borde du se runt en tiopotens bättre prestanda.

Detta på Zen. På Intel borde du se ungefär den dubbla prestandabosten jämfört med Zen från att använda AVX2 och över tre gånger mer om det finns AVX512 stöd.

Visa signatur

Care About Your Craft: Why spend your life developing software unless you care about doing it well? - The Pragmatic Programmer

Permalänk
Medlem
Skrivet av Yoshman:

Nvidia verkar initialt ha tolkat detta från specifikationen på fel sätt

CL_DEVICE_MAX_MEM_ALLOC_SIZE Max size of memory object allocation in bytes. The minimum value is max (1/4th of CL_DEVICE_GLOBAL_MEM_SIZE, 128*1024*1024)

Står ju faktiskt att CL_DEVICE_MAX_MEM_ALLOC_SIZE minst ska vara det största av CL_DEVICE_GLOBAL_MEM_SIZE/4 och 128 MB, står inget om att CL_DEVICE_MAX_MEM_ALLOC_SIZE inte kan vara större.

Inte heller har jag sett att någon förändring av detta kommer ske inom nära framtid. Det är väldigt synd om det inte ska
gå att uttnyttja resurserna till max. På mina föregående AMD kort har det gått att använda i stort sett allt tillgängligt minne.

Skrivet av Yoshman:

Då du kör på GTX1080, varför kör du OpenCL och inte CUDA? CUDA har inte denna begränsning. Vidare är CUDA klart populärare jämfört med OpenCL av en anledning: är enklare och därmed snabbare att jobba i CUDA.

Detta är något Khronos måste fokusera på, CUDA är låst till Nvidia men på det stora hela är CUDA så pass mycket enklare att jobba i att det rätt snabbt blir värt kostnaden att låsa sig till Nvidia. SyCL ser ut att kunna vara precis det öppna svar på CUDA som vi väntat på

Började med OpenCL för många år sedan då jag endast hade AMD kort. Sedan har det rullat på. Skaffade mitt 1080 förra året.
Sedan delar jag inte idén med att stödja proprietär mjukvara om det finns andra alternativ

CUDA har jag aldrig provat men verkar klart intressant om denna minnesbegränsning inte finns där.

SyCL skulle man också läsa på lite kring och se hur det fungerar, har du provat SyCL själv?

Skrivet av Yoshman:

Om det finns orsaker till att köra just på OpenCL: har du testat att köra OpenCL på din CPU?
GPGPU på en dGPU är överlägset om flaskhalsen är minnesbandbredd eller FP32 kapacitet.

Bildbehandling kan ju mycket väl betyda heltalsoperationer och beroende på vad du gör kanske det inte alls är bandbreddskrävande (bara du kan svara på detta alt. beskriva mer i detalj var ditt program gör). I det läget kan faktiskt OpenCL på CPU prestera bättre jämfört med på GPU.

Jo då, har provat det också. I OpenCL är det snabbt gjort så att uträkningarna sker på CPU istället. Även växlingen mellan SP och DP går relativt snabbt att göra

Skillnaden blev stor mellan GPU och GPU, mycket stor till fördel GPU.

400-600% (kommer inte ihåg exakt) snabbare rendering när GPU används och att det inte är ännu större skillnad beror på att IO inte hänger med att mata GPUn med data. Kör jag skräddarsydda benchmarks där jag redan har läst in data till RAM så blir skillnaden större.

Har också skrivit ett standalone projekt utan OpenCL där 16st trådar används på en Ryzen 1700X och skillanden är stor milt sagt

Absoluta max prestanda blir förstås om data redan finns på minnet på grafikkortet då vi kan exkludera all data transfer mellan klient och host men det är såklart inget realistiskt test men snabbt går det om man verkligen vill maxa ur sin GPU

Skrivet av Yoshman:

Det jag har noterat är att i alla fall Intels OpenCL implementation för CPU tenderar väldigt mycket bättre om man skickar in NULL som local_work_size argument till clEnqueueNDRangeKernel(). Detta kräver självklart att varje work-thread kan utföra sitt arbete oberoende från andra trådar i sin "work-group".

Jodå, har provat det också men ser ingen prestandaskillnad alls. Blir dock lite enklare vilket alltid är bra.

Skrivet av Yoshman:

Om du inte handkodat SSE/AVX stöd i din CPU-version borde du se ca 4x prestandaboost om du jobbar med 32-bitars tal. Bildbehandling kan ju betyda att du jobbar med varje färgkanal separat och i så fall borde du se runt en tiopotens bättre prestanda.

Detta på Zen. På Intel borde du se ungefär den dubbla prestandabosten jämfört med Zen från att använda AVX2 och över tre gånger mer om det finns AVX512 stöd.

Intressant, 3ggr låter väldigt mycket om det nu uppnås i verkligheten också?

Detta behöver jag läsa mer om

Permalänk
Datavetare
Skrivet av Multithread:

SyCL skulle man också läsa på lite kring och se hur det fungerar, har du provat SyCL själv?

Har det på min rätt långa lista på saker att testa. Har inte hunnit dit än, mycket beroende på att jag har ett Nvidia Pascal kort i datorn jag primärt kör detta på (min stationära med Ryzen).

Har dock hunnit igenom ett gäng presentationer på YouTube kring SyCL, så har hyfsad koll vilket problem man försöker lösa med tekniken.

Skrivet av Multithread:

Jo då, har provat det också. I OpenCL är det snabbt gjort så att uträkningarna sker på CPU istället. Även växlingen mellan SP och DP går relativt snabbt att göra

Vilken ICD körde du för CPU? Lite ironiskt är Intels ICD faktiskt bättre även på Ryzen jämfört med AMDs (de verkar knapp bry sig om CPU-delen längre, de använder inget högre än SSE3).

Vet inte riktigt hur man skulle göra på Windows, installerade Intels OpenCL drivare på jobbdatorn (Skull Canyon NUC) och kopiera bara över OpenCL runtime till min stationära (båda system kör Ubuntu server 18.04).

Har två ICD installerade, Nvidias och Intel

OpenCL platforms * Intel(R) OpenCL NVIDIA CUDA CL_DEVICE_NAME : AMD Ryzen 7 2700X Eight-Core Processor CL_DEVICE_VERSION : OpenCL 1.2 (Build 475) CL_DRIVER_VERSION : 1.2.0.475

Enda problemet jag sett är att Intel stödjer OpenCL 2.0 både för GPU och CPU på Skylake, men samma ICD stödjer bara OpenCL 1.2 på Ryzen-datorn... AMDs CPU-stöd sträcker sig också bara till 1.2.

OpenCL platforms * Intel(R) OpenCL CL_DEVICE_NAME : Intel(R) Core(TM) i7-6770HQ CPU @ 2.60GHz CL_DEVICE_VERSION : OpenCL 2.0 (Build 475) CL_DRIVER_VERSION : 1.2.0.475

Tyvärr gäller 1.2 begränsningen även Nvidia på Linux. Har läst att man numera stödjer OpenCL 2.0, i alla fall i beta-form, under Windows. Men OpenCL är något som Nvidia behandlar rätt styvmoderligt, tycker det generellt sett är lättare att få förväntad prestanda på Nvidias GPUer med CUDA (om det beror på CUDA eller OpenCL är lite svårt att veta då CUDA inte fungerar på AMD eller Intel).

Skrivet av Multithread:

Skillnaden blev stor mellan GPU och GPU, mycket stor till fördel GPU.

För mig har det varit så att de saker där primära flaskhalsen är SP så är GPU överlägset, men om man jobbar med DP är det oftast snabbare på CPU. Intel är dock märkbart starkare, i7-6770HQ kommer ofta väldigt nära Ryzen-maskinen trots hälften så många kärnor. Är prestandamässigt i princip noll skillnad mellan SSE och AVX för Ryzen, medan det är väsentligt snabbare med AVX på Intel.

Skrivet av Multithread:

400-600% (kommer inte ihåg exakt) snabbare rendering när GPU används och att det inte är ännu större skillnad beror på att IO inte hänger med att mata GPUn med data. Kör jag skräddarsydda benchmarks där jag redan har läst in data till RAM så blir skillnaden större.

Låter rätt mycket vad man kan vänta sig om man gör något där GPU är stark, d.v.s där RAM-bandbredd eller SP beräkningar är flaskhals (numera även HP på de kort med bra stöd, d.v.s. inte konsumentversionen av Pascal...).

Lite lustiga här är att Intel iGPU är förvånansvärt stark på DP. DP är 1/4 av SP för Intel, 1/32 på konsument Pascal och 1/16 på Polaris och Vega. Så NUC:en på jobbet har en GPU som drar jämnt med Titan XP i DP (och ligger mellan Polaris 570 och 580)

Skrivet av Multithread:

Har också skrivit ett standalone projekt utan OpenCL där 16st trådar används på en Ryzen 1700X och skillanden är stor milt sagt

Absoluta max prestanda blir förstås om data redan finns på minnet på grafikkortet då vi kan exkludera all data transfer mellan klient och host men det är såklart inget realistiskt test men snabbt går det om man verkligen vill maxa ur sin GPU

DET är en av de få "killer-apps" som finns för starka iGPU. Med OpenCL 2.0 kan man då göra rätt coola saker med "shared virtual memory", samma adressrymd i host-kod som i OpenCL kernel -> man kan enkelt använda data med pekare ihop med OpenCL! Självklart tappar man massor med bandbredd när CPU och iGPU delar RAM, men finns vissa saker som det ovan när delat minne är en tillgång

Skrivet av Multithread:

Jodå, har provat det också men ser ingen prestandaskillnad alls. Blir dock lite enklare vilket alltid är bra.

Om du orkar kan du testa Intels ICD, är bara där jag sett detta ge en prestandaboost (en rätt rejäl sådan, var rätt besviken på OpenCL prestanda på CPU innan). AMD och Nvidia rekommenderar faktiskt emot att man sätter local_work_size argumentet till NULL i deras respektive ICD.

Skrivet av Multithread:

Intressant, 3ggr låter väldigt mycket om det nu uppnås i verkligheten ocskå?

Har ingen CPU med AVX512, bara sett andra posta resultat som visar att på klar boost över AVX(flyttal)/AVX2(heltal).

AVX512 har i teorin klar fördel över SSE/AVX om din kod innehåller villkor som tar olika vägar inom samma "work-group". Detta kräver lite extra instruktioner i SSE/AVX medan AVX512, likt GPUer, har HW-stöd för att bara maska bort resultatet för de trådar i en "work-group" där kroppen för ett villkor inte ska köras.

Visa signatur

Care About Your Craft: Why spend your life developing software unless you care about doing it well? - The Pragmatic Programmer

Permalänk
Medlem

Kan endast allokera 2GB av 8GB på GeForce GTX 1080

Skrivet av Yoshman:

Har det på min rätt långa lista på saker att testa. Har inte hunnit dit än, mycket beroende på att jag har ett Nvidia Pascal kort i datorn jag primärt kör detta på (min stationära med Ryzen).

Har dock hunnit igenom ett gäng presentationer på YouTube kring SyCL, så har hyfsad koll vilket problem man försöker lösa med tekniken.

Ska göra det samma, verkar intressant helt klart. Dags att lära sig något nytt

Skrivet av Yoshman:

Vilken ICD körde du för CPU? Lite ironiskt är Intels ICD faktiskt bättre även på Ryzen jämfört med AMDs (de verkar knapp bry sig om CPU-delen längre, de använder inget högre än SSE3).

Vet inte riktigt hur man skulle göra på Windows, installerade Intels OpenCL drivare på jobbdatorn (Skull Canyon NUC) och kopiera bara över OpenCL runtime till min stationära (båda system kör Ubuntu server 18.04).

Har två ICD installerade, Nvidias och Intel

OpenCL platforms * Intel(R) OpenCL NVIDIA CUDA CL_DEVICE_NAME : AMD Ryzen 7 2700X Eight-Core Processor CL_DEVICE_VERSION : OpenCL 1.2 (Build 475) CL_DRIVER_VERSION : 1.2.0.475

Jag vill minnas att det är NVIDIA CUDA men inte helt säkert. Jag har stöd för alla funktioner i OpenCL 2.0 så det borde vara NVIDIA CUDA. Var 1 år sedan jag installerade men allt jag laddade ned då var från Nvidia.

Skrivet av Yoshman:

Tyvärr gäller 1.2 begränsningen även Nvidia på Linux. Har läst att man numera stödjer OpenCL 2.0, i alla fall i beta-form, under Windows. Men OpenCL är något som Nvidia behandlar rätt styvmoderligt, tycker det generellt sett är lättare att få förväntad prestanda på Nvidias GPUer med CUDA (om det beror på CUDA eller OpenCL är lite svårt att veta då CUDA inte fungerar på AMD eller Intel).
För mig har det varit så att de saker där primära flaskhalsen är SP så är GPU överlägset, men om man jobbar med DP är det oftast snabbare på CPU. Intel är dock märkbart starkare, i7-6770HQ kommer ofta väldigt nära Ryzen-maskinen trots hälften så många kärnor. Är prestandamässigt i princip noll skillnad mellan SSE och AVX för Ryzen, medan det är väsentligt snabbare med AVX på Intel.

Ok, vad är det för beräkningar som utförts med DP där CPUn varit snabbare?
Var det en del vilkorssatser och få beräkningsmässigt uppgifter som utförts i en sådan kernel?

Även fast DP är tyngre så har mina kernels alltid varit mycket snabbare än respektive på min Ryzen 1700X.

Har gjort en del experiment för att se vart gränserna går. Exempelvis kernels som handlar om att söka i listor efter mönster eller data. I dessa fall har CPUn varit snabbare i alla tester jag gjort. För få beräkningar helt enkelt i mitt fall.

Visst att just en sådan implementation som sökning inte är optimal för en GPU är lätt uträknat. Syftet jag hade var mer åt hållet hur bra/dåligt resultatet skulle bli.

Skrivet av Yoshman:

DET är en av de få "killer-apps" som finns för starka iGPU. Med OpenCL 2.0 kan man då göra rätt coola saker med "shared virtual memory", samma adressrymd i host-kod som i OpenCL kernel -> man kan enkelt använda data med pekare ihop med OpenCL! Självklart tappar man massor med bandbredd när CPU och iGPU delar RAM, men finns vissa saker som det ovan när delat minne är en tillgång

Ja det konceptet är bra i många fall. Brukar använda det också

Refereras ibland som pinned memory i OpenCL doc och forum, även som SVM. Handlar om att man arbetar med att reservera minne
genom CL_MEM_USE_HOST_PTR.

Gav mina program kring 40-60% i total prestandavinst att använda SVM.
Det lägger dock till en del komplexitet men i många fall är det klart värt den tycket jag själv efter man lärt sig grunderna i OpenCL.

Skrivet av Yoshman:

Om du orkar kan du testa Intels ICD, är bara där jag sett detta ge en prestandaboost (en rätt rejäl sådan, var rätt besviken på OpenCL prestanda på CPU innan). AMD och Nvidia rekommenderar faktiskt emot att man sätter local_work_size argumentet till NULL i deras respektive ICD.

Låter som att det är värt att testa Intels också, ska göra det till helgen

Skrivet av Yoshman:

Har ingen CPU med AVX512, bara sett andra posta resultat som visar att på klar boost över AVX(flyttal)/AVX2(heltal).

AVX512 har i teorin klar fördel över SSE/AVX om din kod innehåller villkor som tar olika vägar inom samma "work-group". Detta kräver lite extra instruktioner i SSE/AVX medan AVX512, likt GPUer, har HW-stöd för att bara maska bort resultatet för de trådar i en "work-group" där kroppen för ett villkor inte ska köras.

Ok, det är jag inte alls så hemma på men skulle vilja lära mig mer om. En dag kanske man har en CPU med AVX512