[CUDA] Globaal geheugen GPU *

[Matis]

Ik heb nu zelf een probleem met mn CUDA applicatie;

Misschien dat jullie me even kunnen helpen.

Ik wil 3 pointers welke vanaf mn GPU te benaderen zijn, globaal opslaan; zodat ze de volgende frame weer kan gebruiken

Hetvolgende is het probleem, ik wil gaan deinterlacen op mn GPU. Hiervoor heb ik 3 frames nodig.

En nu wil ik de drie pointers van deze frames ergens globaal opslaan, zodat wanneer mijn GPU weer in een in die lus komt, hij de laatste twee van de vorige keer kan gebruiken, zonder dat je die weer opnieuw hoeft te memcpy-en.

Ik heb in de CUDA Programming Guide 2.0 hetvolgende gevonden:

[b]4.2.2.1 __device__[/b]
The __device__ qualifier declares a variable that resides on the device.
At most one of the other type qualifiers defined in the next three sections may be
used together with __device__ to further specify which memory space the
variable belongs to. If none of them is present, the variable:
Resides in global memory space,
Has the lifetime of an application,
Is accessible from all the threads within the grid and from the host through the
runtime library.

Weet iemand hoe ik dit kan doen / oplossen?

Edit:
even een snippet-code erbij

C:

//Defines staan globaal!
__device__ int *d_pBuf_Y;
__device__ int *d_cBuf_Y;
__device__ int *d_nBuf_Y;
__device__ int *d_oBuf_Y;

main()
{
   cutilSafeCall(cudaMalloc((void**) &d_pBuf_Y, d_BufSize));
   cutilSafeCall(cudaMalloc((void**) &d_cBuf_Y, d_BufSize));
   cutilSafeCall(cudaMalloc((void**) &d_nBuf_Y, d_BufSize));
   cutilSafeCall(cudaMalloc((void**) &d_oBuf_Y, d_BufSize));
...
   while(1)
   {
      cutilSafeCall(cudaMemcpy(d_pBuf_Y, currentFrame, d_BufSize, cudaMemcpyHostToDevice));
      deinterlaceFrames<<< dimGrid, dimBlock >>>( d_pBuf_Y, d_oBuf_Y, startIndex);
      cutilSafeCall(cudaMemcpy(oBuf_Y, d_oBuf_Y, d_BufSize, cudaMemcpyDeviceToHost));
   }
}

__global__ void deinterlaceFrames(int *d_pBuf_Y, int *d_oBuf_Y, int startIndex)
{
   int *d_tBuf_Y = d_pBuf_Y;
   d_pBuf_Y = d_cBuf_Y;
   d_cBuf_Y = d_nBuf_Y;
   d_nBuf_Y = d_tBuf_Y;
   d_oBuf_Y = d_cBuf_Y;
   ...
}

Vergeten titel toe te voegen;

Moet zijn: [CUDA] Globaal geheugen GPU

[ Voor 3% gewijzigd door Matis op 26-02-2009 15:49 ]

[jvaneijk]

Is dit niet een functie iets en geen memcpy?

Wat wil je precies bereiken? En een global functie was dat niet een functie die uitgevoerd wordt op de CPU??

[Matis]

jvaneijk schreef op donderdag 26 februari 2009 @ 15:51:
Is dit niet een functie iets en geen memcpy?

Wat wil je precies bereiken? En een global functie was dat niet een functie die uitgevoerd wordt op de CPU??

Nee, dat is het niet; Een global-functie is een kernel functie, welke aan te roepen is vanuit de host en draait op het device.

Ik wil dus 3 pointers opslaan, zodat ik de volgende frame 2 van de 3 pointers kan hergebruiken.

[jvaneijk]

toaomatis schreef op donderdag 26 februari 2009 @ 15:53:
[...]


Nee, dat is het niet; Een global-functie is een kernel functie, welke aan te roepen is vanuit de host en draait op het device.

Ik wil dus 3 pointers opslaan, zodat ik de volgende frame 2 van de 3 pointers kan hergebruiken.

Worden die frames door verschillende threads berekend / gebruikt? Wat je kan in 1 run niet data ophalen van voorgaande threads.. De threads mogen niet afhankelijk van elkaar zijn dat is een beetje de basic knowledge achter CUDA architecture. Als ze niet afhankelijk van elkaar zijn zou ik het je zo snel even niet kunnen zeggen.

Wat zijn de foutmeldingen die je krijgt in deviceemu en is het op windows of linux iets meer info zou fijn zijn

[Matis]

jvaneijk schreef op donderdag 26 februari 2009 @ 15:57:
[...]

Worden die frames door verschillende threads berekend / gebruikt? Wat je kan in 1 run niet data ophalen van voorgaande threads.. De threads mogen niet afhankelijk van elkaar zijn dat is een beetje de basic knowledge achter CUDA architecture. Als ze niet afhankelijk van elkaar zijn zou ik het je zo snel even niet kunnen zeggen.

Wat zijn de foutmeldingen die je krijgt in deviceemu en is het op windows of linux iets meer info zou fijn zijn

Ik krijg geen foutmeldingen, alleen krijg ik alleen het laatste frame te zien, in elke pointer die ik bekijk; Het maakt hierin dus geen klap uit of ik het vorige, huidige of laatste frame bekijk; ze zijn allemaal hetzelfde.

Er is maar 1 thread en ik heb netjes met syncs gewerkt; Het is dus een STA. Natuurlijk heb ik wel 16 threads op mn GPU gedefineerd, maar die werken allemaal op een ander stuk geheugen.

C++:

__constant__ int d_cBuf_Y[1280*720];

Dat mag ook niet, immers :

code:

make 
/tmp/tmpxft_0000188c_00000000-7_deinterlacer.cpp3.i(0): Error: Const space overflowed

[ Voor 14% gewijzigd door Matis op 26-02-2009 16:02 ]

[jvaneijk]

toaomatis schreef op donderdag 26 februari 2009 @ 16:00:
[...]


Ik krijg geen foutmeldingen, alleen krijg ik alleen het laatste frame te zien, in elke pointer die ik bekijk; Het maakt hierin dus geen klap uit of ik het vorige, huidige of laatste frame bekijk; ze zijn allemaal hetzelfde.

Er is maar 1 thread en ik heb netjes met syncs gewerkt; Het is dus een STA. Natuurlijk heb ik wel 16 threads op mn GPU gedefineerd, maar die werken allemaal op een ander stuk geheugen.

C++:

__constant__ int d_cBuf_Y[1280*720];

Dat mag ook niet, immers :

code:

make 
/tmp/tmpxft_0000188c_00000000-7_deinterlacer.cpp3.i(0): Error: Const space overflowed

Aan die melding te zien gebruik je teveel const memory. Daar zit een limiet aan namelijk. Weet even niet meer hoeveel want het is voor mij al bijna een half jaar dat ik voor het laatst wat met CUDA gedaan heb. maar ik zou dat even checken. Dat kan je doen door middel van het runnen van je project met een bepaalde flag -cubin geloof ik

Als je er niet uit komt kan je eventueel nog even decuda gebruiken is gemaakt door een PhD op utrecht geloof ik maar weet het niet meer zeker moet je even het cuda nvidia forum checken

[ Voor 8% gewijzigd door jvaneijk op 26-02-2009 16:27 ]

[Matis]

jvaneijk schreef op donderdag 26 februari 2009 @ 16:24:
[...]


Aan die melding te zien gebruik je teveel const memory. Daar zit een limiet aan namelijk. Weet even niet meer hoeveel want het is voor mij al bijna een half jaar dat ik voor het laatst wat met CUDA gedaan heb. maar ik zou dat even checken. Dat kan je doen door middel van het runnen van je project met een bepaalde flag -cubin geloof ik

Ja, dat klopt, maar ik gebruik om die reden ook geen __constant__ als type aanduiding.

Wel krijg ik devolgende melding als ik de gegevens uit de pointers wil gaan uitlezen:

code:

./deinterlacer.cu(46): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(47): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(48): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(49): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(50): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(51): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(52): Advisory: Cannot tell what pointer points to, assuming global memory space
./deinterlacer.cu(53): Advisory: Cannot tell what pointer points to, assuming global memory space

Misschien is dat juist het hele probleem!

[Erik Jan]

Ik heb verder geen specifieke kennis van CUDA, dus als ik er compleet naast zit, vergeet deze post maar

Maar je hoeft toch alleen maar een index bij te houden van welke van de 3 buffers de oudste frame bevat? Tenminste als ik de code een beetje goed interpreteer. Is die int startIndex daar niet voor bedoeld?

Even wat pseudocode toegevoegd aan jouw code om te verduidelijken wat ik bedoel.

C:

//Defines staan globaal!
__device__ int *buffer0;
__device__ int *buffer1;
__device__ int *buffer2;
__device__ int *output;

main()
{
...
   int oudste = 0;
   while(1)
   {
      if (oudste == 0) {
         deinterlaceFrames<<< dimGrid, dimBlock >>>(buffer0, buffer1, buffer2);
         cutilSafeCall(cudaMemcpy(buffer0, currentFrame, d_BufSize, cudaMemcpyHostToDevice));
         oudste = 1;
      }
      elseif (oudste == 1) {
         deinterlaceFrames<<< dimGrid, dimBlock >>>(buffer1, buffer2, buffer0);
         cutilSafeCall(cudaMemcpy(buffer1, currentFrame, d_BufSize, cudaMemcpyHostToDevice));
         oudste = 2;
      }
      elseif (oudste == 2) {
         deinterlaceFrames<<< dimGrid, dimBlock >>>(buffer2, buffer0, buffer1);
         cutilSafeCall(cudaMemcpy(buffer2, currentFrame, d_BufSize, cudaMemcpyHostToDevice));
         oudste = 0;
      }
      cutilSafeCall(cudaMemcpy(oBuf_Y, output, d_BufSize, cudaMemcpyDeviceToHost));
   }
}

__global__ void deinterlaceFrames(int *frame0, int *frame1, int *frame2)
{
   ...
}

[Zoijar]

Ik snap niets van je kernel code... je assigned golbale pointers aan elkaar daar in een kernel. Dat gaat parallel gedraaid worden; kan nooit werken. Wat hierboven staat is gewoon een simpele oplossing. Je kan device pointers wel gewoon op je host aanpassen, alleen waar ze naar wijzen betekent niets op de host. (ik bedoel, je kan een pointer naar device geheugen gewoon op je host opslaan; dat hoeft niet per se ook in device geheugen) Maar je kan ze wel swappen. Dat is de basis ping-pong strategie: while (1) {swap(src,dst); execute kernel(src, dst);}

zie bv dit stukje uit een van mijn sources:

C:

// Device Global Memory
State *dev_Y1, *dev_Y2, *dev_Y3, *dev_Y4; // intermediate states
State *dev_k1, *dev_k2, *dev_k3;

...

CUDA_CALL(cudaMalloc((void**)&dev_Y1, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_Y2, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_Y3, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_Y4, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_k1, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_k2, state_size));
CUDA_CALL(cudaMalloc((void**)&dev_k3, state_size));

...
CUDA_CALL(cudaMemcpy(dev_Y1, host_Y1, sizeof(State)*g_const.nr_balls, cudaMemcpyHostToDevice));

RK4_step<<<dimGrid, dimBlock>>>(dev_Y1, dev_Y1, dev_Y2, dev_k1, 0.5f);
RK4_step<<<dimGrid, dimBlock>>>(dev_Y1, dev_Y2, dev_Y3, dev_k2, 0.5f);
RK4_step<<<dimGrid, dimBlock>>>(dev_Y1, dev_Y3, dev_Y4, dev_k3, 1.0f);
RK4_finalize<<<dimGrid, dimBlock>>>(dev_Y1, dev_Y4, dev_k1, dev_k2, dev_k3);

// Copy next state from device to host
CUDA_CALL(cudaMemcpy(host_Y1, dev_Y1, sizeof(State)*g_const.nr_balls, cudaMemcpyDeviceToHost));

...

// weights: 0.5, 0.5, 1.0
// Compute 1.0 ad-hoc solution to synchronize between blocks
__global__ void RK4_step(State* Y1, State* Yt, State* Yout, State* kout, float weight) {
    int global_id = blockIdx.x*blockDim.x + threadIdx.x;

    State initial = Y1[global_id];
    State k = get_k(Yt);

    kout[global_id] = k;
    Yout[global_id].velocity = initial.velocity + (k.velocity * weight);
    Yout[global_id].position = initial.position + (k.position * weight);
}

Daar staat ook geen __device__ voor de pointers.

[ Voor 72% gewijzigd door Zoijar op 26-02-2009 20:28 ]

[Matis]

Alvast bedankt voor jullie reactie's; Het feit dat je er niets van begrijpt zal wel aan mij liggen, vind het parralel programmeren nog een beetje onwennig.

Ik ga vandaag jullie ideeen eens proberen te implementeren.

Het werkt prima; ik krijg de gewenste pointers nu, met hun waarde. Ik doe nu elke lus 4 keer memcpy heen en 1 keer memcpy terug. Met een grootte van 921600 bytes.

Deze "dure" operatie kost nog geen 1 ms, dus niet echt noemenswaardig

[ Voor 36% gewijzigd door Matis op 27-02-2009 09:54 ]

[jvaneijk]

toaomatis schreef op vrijdag 27 februari 2009 @ 08:47:
Alvast bedankt voor jullie reactie's; Het feit dat je er niets van begrijpt zal wel aan mij liggen, vind het parralel programmeren nog een beetje onwennig.

Ik ga vandaag jullie ideeen eens proberen te implementeren.

Het werkt prima; ik krijg de gewenste pointers nu, met hun waarde. Ik doe nu elke lus 4 keer memcpy heen en 1 keer memcpy terug. Met een grootte van 921600 bytes.

Deze "dure" operatie kost nog geen 1 ms, dus niet echt noemenswaardig

Pas op met het meten van je memcpy want die 1ms lijkt me niet goed. de overhead die er bestaat is immens, deze bedraagd ongeveer 150-200ms Ik heb in het begin ook altijd dit probleem gehad. Maar je moet zeker goed opletten met initialisatie van je geheugen.

Nogmaals: Geheugen kopieren van en naar GPU is ontzettend duur en bijna niet lonend als het zo ontzettend klein is dan krijg je die overhead echt nooit weg. Die ben je pas kwijt bij enkele MB's. Heb er onderzoek naar gedaan en schrok even toen ik die tijden waarnam.

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 10:41:
[...]


Pas op met het meten van je memcpy want die 1ms lijkt me niet goed. de overhead die er bestaat is immens, deze bedraagd ongeveer 150-200ms Ik heb in het begin ook altijd dit probleem gehad. Maar je moet zeker goed opletten met initialisatie van je geheugen.

Nogmaals: Geheugen kopieren van en naar GPU is ontzettend duur en bijna niet lonend als het zo ontzettend klein is dan krijg je die overhead echt nooit weg. Die ben je pas kwijt bij enkele MB's. Heb er onderzoek naar gedaan en schrok even toen ik die tijden waarnam.

Klopt, in de Programming Guide staat ook dat cudamemcpy heel duur is, maar ik heb het nu *netjes* werkend; nu flink gaan optimaliseren natuurlijk. Straks hoef ik immers nog maar 1 keer per frame naar de GPU te memcpyen en terug, immers heb ik de frames nodig om naar een file te schrijven. Het zal helemaal mooi zijn als ik ze meteen naar mn monitor zou kunnen streamen, maar dat is nu nog niet van toepassing.

Daarnaast wil ik graag nog weten wat jullie adviseren mbt die geheugen allocatie.

Ik heb nu dus 4 arrays van 920k groot met daarin int's.

In de Programming Guide hebben ze het over de volgende geheugenallocaties:

Linear memory is allocated using cudaMalloc() or cudaMallocPitch() and freed using cudaFree().

ook staat er:

cudaMallocPitch() is recommended for allocations of 2D arrays as it makes sure that the allocation is appropriately padded to meet the alignment requirements described in Section 5.1.2.1.

Ik gebruik zelf dus cudaMalloc mbv het cutilSafeCall(cudaMalloc((...)) commando.

Vertraagt de cutilSafeCall(...) de hele zwik? Of waar dient het letterlijk voor, er staat niets over in de programming guide en op het forum van NVidia wordt het wel gebruikt, maar de uitleg ontbreekt.

wel staat er hetvolgende in de broncode van de SDK

C++: cutil_inline.h

inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
    do {
        if( cudaSuccess != err) {
            fprintf(stderr, "cudaSafeCall() Runtime API error in file <%s>, line %i : %s.\n",
                    file, line, cudaGetErrorString( err) );
            exit(-1);
        }
    } while (0);
}

Zijn er nog workarrounds om die cudamemcpy sneller te maken? Bijvoorbeeld het defineren van het geheugen middels cudaMallocArray() oid?

Of kan ik het ook async versturen middels de cudaMemcpyAsync(...)?

[ Voor 4% gewijzigd door Matis op 27-02-2009 11:08 . Reden: Te snel op verstuur gedrukt :( ]

[jvaneijk]

toaomatis schreef op vrijdag 27 februari 2009 @ 11:04:
[...]


Klopt, in de Programming Guide staat ook dat cudamemcpy heel duur is, maar ik heb het nu *netjes* werkend; nu flink gaan optimaliseren natuurlijk. Straks hoef ik immers nog maar 1 keer per frame naar de GPU te memcpyen en terug, immers heb ik de frames nodig om naar een file te schrijven. Het zal helemaal mooi zijn als ik ze meteen naar mn monitor zou kunnen streamen, maar dat is nu nog niet van toepassing.

Daarnaast wil ik graag nog weten wat jullie adviseren mbt die geheugen allocatie.

Ik heb nu dus 4 arrays van 920k groot met daarin int's.

In de Programming Guide hebben ze het over de volgende geheugenallocaties:

Linear memory is allocated using cudaMalloc() or cudaMallocPitch() and freed using cudaFree().

ook staat er:

cudaMallocPitch() is recommended for allocations of 2D arrays as it makes sure that the allocation is appropriately padded to meet the alignment requirements described in Section 5.1.2.1.

Ik gebruik zelf dus cudaMalloc mbv het cutilSafeCall(cudaMalloc((...)) commando.

Vertraagt de cutilSafeCall(...) de hele zwik? Of waar dient het letterlijk voor, er staat niets over in de programming guide en op het forum van NVidia wordt het wel gebruikt, maar de uitleg ontbreekt.

wel staat er hetvolgende in de broncode van de SDK

C++: cutil_inline.h

inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
    do {
        if( cudaSuccess != err) {
            fprintf(stderr, "cudaSafeCall() Runtime API error in file <%s>, line %i : %s.\n",
                    file, line, cudaGetErrorString( err) );
            exit(-1);
        }
    } while (0);
}

Zijn er nog workarrounds om die cudamemcpy sneller te maken? Bijvoorbeeld het defineren van het geheugen middels cudaMallocArray() oid?

Of kan ik het ook async versturen middels de cudaMemcpyAsync(...)?

cutil wordt gebruikt als utility tool. Het is een manier om cudacalls te voorzien van error informatie als het dus fout gaat. Ook is het makkelijk om via deze manier je cudafuncties aan te roepen vanuit C/C++ als je dus een groter programma krijgt en je een deel gebruikt in CUDA.

Wat betrecht het alloceren van geheugen is altijd een moeilijke geweest in CUDA. Ikzelf heb altijd 3D arrays gewerkt. Deze 3D arrays sloeg ik plat om er zo een 1D array van te maken. Het vergte wat denkwerk om mijn threadblocks enzo goed te krijgen maar uiteindelijk is he toch gelukt.

Wat ik ook nog even wil zeggen.. Probeer niet teveel dingen op te slaan want het opnieuw berekenen is vaak vele malen sneller dan het opslaan. In het begin was ik er niet huiverig om maar naarmate je de bottlenecks begint te doorzien van de grafische kaart is het duidelijk dat opslaan geen optie is, want dynamisch geheugen alloceren is ook niet mogelijk...

Voor de rest wens ik je heel erg veel succes met je project en zou graag zien wat het eindproduct geworden is.. Houdt me op de hoogte

[Zoijar]

Precies, een read uit globaal geheugen kost 400-600 cycles. Veel van die latency wordt wel verborgen achter meerdere parallelle reads, maar het blijft duur. Als je dat dan vergelijk met een multiply of add die 4 cycles kost, en een reciprocal sqrt die 16 cycles kost, dan zie je dat je al snel 35 sqrts kan uitvoeren in de tijd van 1 enkele geheugen read. Lokaal geheugen is wel snel, dat lees je in 4 cycles. De tactiek is dan ook om zo veel mogelijk het globale geheugen dat je gebruikt per benodigde block naar thread lokaal geheugen te laden in een keer. Dan hide je bijna alle latency en kan je daarna snel met de data werken.

[Matis]

Zoijar schreef op vrijdag 27 februari 2009 @ 12:44:
Precies, een read uit globaal geheugen kost 400-600 cycles. Veel van die latency wordt wel verborgen achter meerdere parallelle reads, maar het blijft duur. Als je dat dan vergelijk met een multiply of add die 4 cycles kost, en een reciprocal sqrt die 16 cycles kost, dan zie je dat je al snel 35 sqrts kan uitvoeren in de tijd van 1 enkele geheugen read. Lokaal geheugen is wel snel, dat lees je in 4 cycles. De tactiek is dan ook om zo veel mogelijk het globale geheugen dat je gebruikt per benodigde block naar thread lokaal geheugen te laden in een keer. Dan hide je bijna alle latency en kan je daarna snel met de data werken.

Precies; dat las ik net ook al:

ik heb nu hetvolgende:

C++: deinterlacer.cu

int *d_tBuf_Y = d_pBuf_Y;
d_pBuf_Y = d_cBuf_Y;
d_cBuf_Y = d_nBuf_Y;
d_nBuf_Y = d_tBuf_Y;
cudaMemcpy(d_nBuf_Y, yData, d_BufSize, cudaMemcpyHostToDevice);
d_oBuf_Y = d_cBuf_Y;       
deinterlaceFrames<<< dimGrid, dimBlock >>>( d_pBuf_Y, d_cBuf_Y, d_nBuf_Y, d_oBuf_Y, startIndex);
cudaMemcpy(oBuf_Y, d_oBuf_Y, d_BufSize, cudaMemcpyDeviceToHost);

Het probleem is nu dat mn CPU de frames niet snelgenoeg kan uitlezen van schijf, waardoor ik helaas niet echt lekker kan gaan optimaliseren.

[jvaneijk]

Ik heb de ballen verstand van wat je aan het doen bent maar is het niet mogelijk om je complete plaatje / filmpje in 1x op je videokaart te zetten in global mem? en daarop allerlei transformaties uitvoeren?

Nogmaals neem mij niet kwalijk maar ik weet niet wat het is of wat het doet sorry

[Zoijar]

Van disk naar host naar gpu naar host naar disk... dan is je computing niet de bottleneck Wat je misschien kan doen is asynchroon van disk lezen, zoveel je kan, terwijl CUDA bezig is met computing. Dan hide je alle computing time denk ik, en wordt het puur I/O bound.

Je moet idd kijken wat je precies wilt. Een video frame processen kost iha misschien een milliseconde (weet niet wat je precies doet), terwijl een enkele disk seek al ruim 10 ms kost. Geheugen uploaden naar GPU kan met ongeveer 3 GB/s (dat is mijn max met een pcie 2.0 bus en een GTX260). Downloaden van je GPU is meestal veel trager. Het ligt erg aan het buffer formaat, maar meestal is dat zo'n 800-1200MB/s oid. Een HD frame is zo'n 8MB. Dat betekent dat je over het downloaden van een HD frame zo'n 8ms doet. Dat is vaak al veel meer dan je computing tijd. Het zou kunnen dat het hierdoor sneller is om het op de host te computen met SSE.

Als je hier dus winst op wilt halen, dan moet je zorgen dat alles asynchroon gaat. Tijdens je disk reads en writes moeten er transfers van en naar de GPU lopen, en computing aan de gang zijn. Dat is denk ik wel haalbaar, en dan zou het met de snelheid van een file copy moeten lopen.

[ Voor 67% gewijzigd door Zoijar op 27-02-2009 13:43 ]

[Matis]

Zoijar schreef op vrijdag 27 februari 2009 @ 13:35:
Van disk naar host naar gpu naar host naar disk... dan is je computing niet de bottleneck Wat je misschien kan doen is asynchroon van disk lezen, zoveel je kan, terwijl CUDA bezig is met computing. Dan hide je alle computing time denk ik, en wordt het puur I/O bound.

Ja klopt, daar kwam ik net achter. Ik had een aantal frames in mn geheugen geladen (van mn C++ applicatie) en die naar de GPU gestuurd, maar ik kreeg geen performance winst; Immers duurt het schrijven van Host naar Device 600 klokslagen oid.

Ik moet zo maar eens async gaan proberen te implementeren.

Edit:
Heeft iemand toevallig iets van pseudo-code als voorbeeld. Ik heb denk ik maar 1 stream nodig, alleen toen ik dat implementeerde ging mn videokaart raar doen. Ik kreeg allemaal vreemde pixels over mn scherm; een reboot fixte het, maar ik ben toch een beetje bang geworden Kaart koste 200 ofzo, wil ik niet verstoken door een slecht programma

[ Voor 20% gewijzigd door Matis op 27-02-2009 14:15 . Reden: Extra tekst, Typo's ]

[jvaneijk]

toaomatis schreef op vrijdag 27 februari 2009 @ 13:53:
[...]

Ja klopt, daar kwam ik net achter. Ik had een aantal frames in mn geheugen geladen (van mn C++ applicatie) en die naar de GPU gestuurd, maar ik kreeg geen performance winst; Immers duurt het schrijven van Host naar Device 600 klokslagen oid.

Ik moet zo maar eens async gaan proberen te implementeren.

Dat wordt een plezier... Async calls geven weer hele andere problemen die je niet zomaar op lost.
En vaak is er niet eens zo heel veel winst mee te behalen. Wij hebben toen ook die afslag geprobeerd te maken in het AMC maar kwamen erachter dat het sync gewoon vele malen sneller bleek te gaan omdat onze bottleneck geheugen op de grafische kaart was..

[jvaneijk]

Je fried je kaart niet zomaar.. kan totaal geen kwaad... het is gewoon een probleem met je streams die je een beetje vol propt en je grafische driver er niets meer van snapt... Is vervelend soms maar heel normaal.. En onder windows probeer die watchdog te omzeilen... Maar volgens mij werk je met Unix systeem dus niet zo'n probleem.

[Zoijar]

Update ook je drivers even naar de nieuwste nvidia drivers. Drivers, en in sommige gevallen hardware, zijn vaak brak... daar merk je pas wat van als je de kaart wat stressed. Kan ook dat je kaart wat heet wordt, of dat je misschien overclocked. Ik moest m'n oude kaart zelfs iets down-clocken als ik demanding cuda applicaties wilde draaien, anders crashde hij na 30 seconden... op zich kan het geen kwaad.

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 14:20:
Je fried je kaart niet zomaar.. kan totaal geen kwaad... het is gewoon een probleem met je streams die je een beetje vol propt en je grafische driver er niets meer van snapt... Is vervelend soms maar heel normaal.. En onder windows probeer die watchdog te omzeilen... Maar volgens mij werk je met Unix systeem dus niet zo'n probleem.

Ja ubuntu 8.10 , maar ik vond het toch onplezierig om mn kaart zo te zien sterven (al was het maar voor even )

Zoijar schreef op vrijdag 27 februari 2009 @ 14:23:
Update ook je drivers even naar de nieuwste nvidia drivers. Drivers, en in sommige gevallen hardware, zijn vaak brak... daar merk je pas wat van als je de kaart wat stressed. Kan ook dat je kaart wat heet wordt, of dat je misschien overclocked. Ik moest m'n oude kaart zelfs iets down-clocken als ik demanding cuda applicaties wilde draaien, anders crashde hij na 30 seconden... op zich kan het geen kwaad.

Ik heb vlgns mij de laatste drivers; Volgens NVIDIA X Server draai ik op 180.29. Ja, de laatste dus!

Asynchroon mag dan wel problemen geven, maar dit gaat niet werken op deze synchrone manier. Ik *moet* realtime kunnen werken, lekker 24 a 25 fps. Dat haal ik nu nog niet.

[ Voor 2% gewijzigd door Matis op 27-02-2009 14:33 . Reden: +Linkje ]

[jvaneijk]

Beeld bewerkings programmatje aan het maken?

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 14:36:
Beeld bewerkings programmatje aan het maken?

Nope, deinterlacer; Nja soort van beeldbewerking

Ik ben eigenlijk niet helemaal eerlijk geweest. Ik ben met /voor Dirac bezig. Een opensource video coder en decoder, bedacht en opgezet door de BBC

http://diracvideo.org/
http://www.bbc.co.uk/rd/projects/dirac/

Edit:
Ik zie net dat ten tijde van het uitvoeren van mn programma mn GPU helemaal niet naar powermode gaat, wat dus inhoudt dat mn programma niet veel GPU kracht vraagt, maar alleen maar aan het wachten is op die memcalls.
De temp is ook prima gebleven, tussen de 60 en 64 grC.

[ Voor 62% gewijzigd door Matis op 27-02-2009 14:55 . Reden: Extra tekst, Typo's ]

[jvaneijk]

Mijn programma wat soms meer dan enkele miljarden berekeningen deed, ging op den duur richting de 95gr terwijl de treshold op 145gr stond ofzo en dat waren echt berekeningen die soms wel 20min konden duren en dan nog werdt hij niet super warm

en FDTD is nog in ontwikkeling maar dat neemt op de PC normaal al tussen de 5 en 80 uur in beslag dus op de GPU dachten we ong in een kwart te kunnen doen.. maar das een serieuse stress test voor hem

[ Voor 29% gewijzigd door jvaneijk op 27-02-2009 15:19 ]

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 15:18:
Mijn programma wat soms meer dan enkele miljarden berekeningen deed, ging op den duur richting de 95gr terwijl de treshold op 145gr stond ofzo en dat waren echt berekeningen die soms wel 20min konden duren en dan nog werdt hij niet super warm.

95graden dat is wel heel erg veel en een threshold van 145 mag hoop ik nooit gehaald worden. Dan smelt heel je PCB, maar het probleem is dus dat ik niet genoeg data aanbied aan de GPU, waardoor hij een beetje niets hoeft te doen en maar 1% gebruikt van de totale rekenkracht. Een beetje jammer, maar al wel een enorme boost t.o.v een reguliere CPU.

Ik ben bezig met de asynchrone communicatie, alleen dat was, zoals hierboven beschreven, lastiger dan ik dacht en gehoopt had.

Heeft een van jullie wel eens beelden naar het scherm geschreven mbv een CUDA applicatie?

[jvaneijk]

Nog nooit iets naar scherm hoeven schrijven maar dat is ook niet mogelijk vanuit CUDA. CUDA is puur en alleen computing en geen rendering. om te renderen moet je OpenGL of DX interop gebruiken. Check de voorbeelden in de SDK en probeer eens met die calculation tool te controleren hoeveel procent je gevuld heb van je GPU want 1% is wel heel erg weinig... miss. moet je dan je blokken aanpassen...

[Zoijar]

toaomatis schreef op vrijdag 27 februari 2009 @ 15:27:
Heeft een van jullie wel eens beelden naar het scherm geschreven mbv een CUDA applicatie?

Zie ook het andere cuda topic; daar heb je cudaglmpabufferobject voor, maar die is traag... maar er is dus een cuda/opengl koppeling aanwezig.

[Matis]

Zoijar schreef op vrijdag 27 februari 2009 @ 15:39:
[...]

Zie ook het andere cuda topic; daar heb je cudaglmpabufferobject voor, maar die is traag... maar er is dus een cuda/opengl koppeling aanwezig.

Ja, precies, maar als je bijvoorbeeld die nBody pakt, of de fluidsGL van de voorbeelden, dan projecteren ze de berekeningen naar het scherm, wat IMO precies is wat ik wil, daarmee sla ik de memcpDtoH over, wat me weer veel klokslagen scheelt.

Nja ik zal me daar eens in verdiepen.

[Zoijar]

Die kopieren ook de posities naar opengl buffers, en het opengl programma gebruikt die buffers tijdens het renderen. Het zou heel veel moeten schelen, maar wegens een brakke implementatie haal je geen 60+ GB/s die je zou verwachten.

[ Voor 33% gewijzigd door Zoijar op 27-02-2009 15:48 ]

[jvaneijk]

ik zal CUDA straks thuis weer eens even installeren.. is wel onder windows maar kan ik eens kijken hoe die mapbuffer in zijn werking gaat tis voor mij echt een enorme tijd geleden kom ik achter... Maar veel problemen die jij hebt ben ik toen in het begin tegenaan gelopen met versie 0.7 toen de tijd

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 15:51:
ik zal CUDA straks thuis weer eens even installeren.. is wel onder windows maar kan ik eens kijken hoe die mapbuffer in zijn werking gaat tis voor mij echt een enorme tijd geleden kom ik achter... Maar veel problemen die jij hebt ben ik toen in het begin tegenaan gelopen met versie 0.7 toen de tijd

Ik hoop niet dat ik je weer aangestoken heb met het hele CUDA gebeuren maar het zou wel fijn zijn voor wat pseudocode.

De async gegevensoverdracht heb ik bijna af.

[jvaneijk]

Ik heb geprobeerd het uit mijn geheugen te verbannen na iets meer dan een jaar ermee gewerkt te hebben... Het is zo ontzetten frustrerend als je een week lang met een fout zit en niemand maar dan ook echt niemand (zelfs de developers van CUDA niet) kan je helpen. En later merken ze bij Nvidia op dat het een bug is aan hun kant

[Matis]

jvaneijk schreef op vrijdag 27 februari 2009 @ 16:30:
Ik heb geprobeerd het uit mijn geheugen te verbannen na iets meer dan een jaar ermee gewerkt te hebben... Het is zo ontzetten frustrerend als je een week lang met een fout zit en niemand maar dan ook echt niemand (zelfs de developers van CUDA niet) kan je helpen. En later merken ze bij Nvidia op dat het een bug is aan hun kant

Ik begrijp je frustatie, ik vind dat ook altijd enorm vervelend .

Ik heb trouwens het idee dat ik die cudaMemcpyAsync verkeerd heb geinterpreteerd;
CMIIW. Nu ik het opnieuw lees, blijkt het dus de mogelijkheid te zijn, dat er een burst van arrays gestuurd wordt naar de GPU met die 600 klokslagen vertraging, zodat je die tijd kan opvangen, hiermee kun je ervoor zorgen dat er theoretisch nooit gewacht hoeft te worden.

Ik dacht dat je met dat Async zonder vertraging kon schrijven naar je GPU en lezen ervan, maar na het implementeren van de Async zoals ik dacht dat hij in eerste instantie zou werken, vertraagde mijn programma met ongeveer factor 150 .

Toen ging in wat code op het forum van Nvidia zoeken en toen viel het kwartje pas echt

[ Voor 0% gewijzigd door Matis op 27-02-2009 16:54 . Reden: Extra tekst, Typo's ]

[Matis]

Vandaag weer verder gegaan met het programmeren, maar loop nu tegen het volgende probleem op:

Mijn consumer (welke de afbeeldingen inlees) is in staat op sneller een frame van 3,5 MB in te lezen dan dat mijn consumer nodig heeft om binnen de while-lus naar het begin te springen

In assembly is dit 1 instructie, maar kennelijk denkt g++ daar anders over

Momenteel gefixed, het was een gedeeld geheugen issue.

[ Voor 7% gewijzigd door Matis op 02-03-2009 16:36 . Reden: Extra tekst, Typo's ]