Modo discussione | Modo lineare

Corralx · 06-05-2012 1:26

Salve, ho un problema con OpenCL e spero qualcuno possa aiutarmi perchè sto impazzendo (anche se non ho molte speranze vista la specificità della domanda).
Ringrazio in anticipo chi si prende la briga di leggersi il wall of text Asd

Sto portando il mio raytracer su OpenCL.
La struttura del programma è questa:
- Inizializzo un framebuffer con OpenGL a cui associo un renderbuffer come color attachment
- condivido il renderbuffer con OpenCL
- ogni frame scrivo il risultato pixel per pixel nel renderbuffer, e poi uso glBlitFramebuffer(...) per copiarne il contenuto nel back buffer

Usando un semplice kernel come il seguente, tutto funziona (come si suol dire so far so good):

Codice:

__kernel void rayTrace(__write_only image2d_t image)

{

    int2 coord = (int2)(get_global_id(0), get_global_id(1));

    float4 color = (float4)(1.f, .0f, .0f, 1.f);

    write_imagef(image, coord, color);

}

e mi ritrovo con la finestra tutta rossa (colore RGBA 1.0f,0.0f,0.0f,1.0f).

A questo punto vado a modificare il kernel per inserire il codice del raytracer, e la prima cosa che mi serve è reimplementare funzioni tipo gluUnProject per ottenere le coordinate in world space di un pixel (x,y) sul near plane della camera.
Qui cominciano i problemi.
Il seguente codice funziona:

Codice:

bool UnProject(const float3 win, const float16 modelMatrix, const float16 projMatrix, const int4 viewport, float3 obj)

{

    float finalMatrix[16];

    float in[4];

    float out[4];

    __MultMatrices(modelMatrix, projMatrix, finalMatrix);

    if (!__InvertMatrix(finalMatrix, finalMatrix))

        return false;

    in[0] = win[0];

    in[1] = win[1];

    in[2] = win[2];

    in[0] = (in[0] - viewport[0]) / viewport[2];

    in[1] = (in[1] - viewport[1]) / viewport[3];

    in[0] = in[0] * 2 - 1;

    in[1] = in[1] * 2 - 1;

    in[2] = in[2] * 2 - 1;

    __MultMatrixVec(finalMatrix, in, out);

    if (out[3] == 0.0)

        return false;

    out[0] /= out[3];

    out[1] /= out[3];

    out[2] /= out[3];

    obj[0] = out[0];

    obj[1] = out[1];

    obj[2] = out[2];

    return true;

}

Il problema è che tutti i parametri sono passati per valore, e in particolare, quando la funzione ritorna, dentro obj non trovo (ovviamente) il risultato voluto, ma quello passato come argomento.
Se passo ad una funzione definita come:

Codice:

bool UnProject(const float3* win, const float16* modelMatrix, const float16* projMatrix, const int4* viewport, float3* obj) {...}

in cui il codice è cambiato coerentemente per rispecchiare i cambiamenti, il kernel non viene eseguito.
Infatti la funzione clEnqueueNDRangeKernel(...) che lo manda in esecuzione ritorna con errore, nello specifico -54 ovvero CL_INVALID_WORK_GROUP_SIZE.

A questo punto non capisco più quello che succede.
Ovviamente cambiando semplicemente il codice il work group size non cambia.
Inoltre attualmente uso un work group size pari a 2, con global work size uguale a {800, 600} (la risoluzione del renderbuffer) e un local work size uguale a {20, 20} (che è quindi divisore del global work size).
Il numero massimo di local work size per la mia GPU è 512 che è maggiore di 20x20 = 400.
Questo è corretto è aderisce alle specifiche di OpenCL, oltre a funzionare perfettamente, finchè non passo i parametri delle funzioni per riferimento anzichè per valore dentro al kernel.

Qualcuno ha idea del motivo?
Sto mancando qualcosa?
E' un bug della mia implementazione?
Su google non ho trovato niente, e mi sembra un bug troppo grosso per non essere conosciuto.
Sto per spaccare il tavolo Asd

TheCrib · 06-05-2012 7:27

Non ho esperienza, pero' il discorso mi interessa e' ho fatto una search a volo.

Potrebbe essere una questione di 32/64 bit ?

Altrimenti non ti resta che fare varie prove.
Potresti provare a passare solo "obj" per riferimento.. etc etc.

Dimmi solo che nel codice finale non farai una inverse matrix per pixel 8)

Corralx · 06-05-2012 12:21

Avevo giá letto di questo thread ma:
- lavoro su osx, in cui sia size_t sia i pointer hanno dimensione 8 byte, e le mie variabili che contengono le dimensioni dei work group sono size_t, e alla clEnqueueNDRangeKernel le passo come (size_t*)&var
- non capisco perché senza pointer funziona ( in particolare senza pointer alla memoria __private del kernel, i pointer __global argomenti del __kernel non danno problemi)

Per quanto riguarda la matrice, spero di no Linguaccia

In futuro me la ricalcoleró quando serve e la terró nella memoria __global del kernel Sisi

EDIT: dimenticavo. Ho già provato a passare per riferimento solo un parametro, stesso risultato.

TheCrib · 06-05-2012 13:26

(06-05-2012 12:21)Corralx ha scritto: EDIT: dimenticavo. Ho già provato a passare per riferimento solo un parametro, stesso risultato.

La butto li.. e' possibile che quando passi tutto per valore, la routine viene ottimizzata in toto, quindi magari il problema c'e' ma e' latente.

Corralx · 06-05-2012 14:22

(06-05-2012 13:26)TheCrib ha scritto:
(06-05-2012 12:21)Corralx ha scritto: EDIT: dimenticavo. Ho già provato a passare per riferimento solo un parametro, stesso risultato.

La butto li.. e' possibile che quando passi tutto per valore, la routine viene ottimizzata in toto, quindi magari il problema c'e' ma e' latente.

Uhm potrei provare a passare al compilatore le flag per disabilitare le ottimizzazioni.
Mò provo Sisi

EDIT: No, niente.
Provato senza ottimizzazioni, con lo strict aliasing, ecc...
Sempre lo stesso risultato.

EDIT2: cercando un altro pò in giro ho trovato questo:
http://forums.nvidia.com/index.php?showtopic=223153
Sembra che nell'implementazioni di NVidia i puntatori siano tutti a 32 bit sul device.

Inoltre facendo qualche altra prova, il problema sembra essere nel deferenziare tali puntatori per scriverci dei valori (forse un qualche conflitto fra i tipi C come float, e i tipi vettoriali di OpenCL float4, float16 ecc).
Anche se ancora non capisco che ci azzecca CL_INVALID_WORK_GROUP_SIZE.
Probabilmente quell'errore mi stà mandando fuori strada.

Corralx · 07-05-2012 1:00

UPDATE:
Ho riscritto il kernel da zero, ristrutturando un pò le cose ed ora ho individuato e ristretto bene il problema.
Il problema è nella funzione che inverte la matrice:

Codice:

static bool __InvertMatrix(__private const float16* m, __private float16* invOut)

{  

    __private float16 inv;

    __private float det;

    inv.s0 = m->s5*m->sA*m->sF - m->s5*m->sB*m->sE - m->s9*m->s6*m->sF

    + m->s9*m->s7*m->sE + m->sD*m->s6*m->sB - m->sD*m->s7*m->sA;

    inv.s4 = -m->s4*m->sA*m->sF + m->s4*m->sB*m->sE + m->s8*m->s6*m->sF

    - m->s8*m->s7*m->sE - m->sC*m->s6*m->sB + m->sC*m->s7*m->sA;

    inv.s8 = m->s4*m->s9*m->sF - m->s4*m->sB*m->sD - m->s8*m->s5*m->sF

    + m->s8*m->s7*m->sD + m->sC*m->s5*m->sB - m->sC*m->s7*m->s9;

    inv.sC = -m->s4*m->s9*m->sE + m->s4*m->sA*m->sD + m->s8*m->s5*m->sE

    - m->s8*m->s6*m->sD - m->sC*m->s5*m->sA + m->sC*m->s6*m->s9;

    inv.s1 = -m->s1*m->sA*m->sF + m->s1*m->sB*m->sE + m->s9*m->s2*m->sF

    - m->s9*m->s3*m->sE - m->sD*m->s2*m->sB + m->sD*m->s3*m->sA;

    inv.s5 = m->s0*m->sA*m->sF - m->s0*m->sB*m->sE - m->s8*m->s2*m->sF

    + m->s8*m->s3*m->sE + m->sC*m->s2*m->sB - m->sC*m->s3*m->sA;

    inv.s9 = -m->s0*m->s9*m->sF + m->s0*m->sB*m->sD + m->s8*m->s1*m->sF

    - m->s8*m->s3*m->sD - m->sC*m->s1*m->sB + m->sC*m->s3*m->s9;

    inv.sD = m->s0*m->s9*m->sE - m->s0*m->sA*m->sD - m->s8*m->s1*m->sE

    + m->s8*m->s2*m->sD + m->sC*m->s1*m->sA - m->sC*m->s2*m->s9;

    inv.s2 = m->s1*m->s6*m->sF - m->s1*m->s7*m->sE - m->s5*m->s2*m->sF

    + m->s5*m->s3*m->sE + m->sD*m->s2*m->s7 - m->sD*m->s3*m->s6;

    inv.s6 = -m->s0*m->s6*m->sF + m->s0*m->s7*m->sE + m->s4*m->s2*m->sF

    - m->s4*m->s3*m->sE - m->sC*m->s2*m->s7 + m->sC*m->s3*m->s6;

    inv.sA = m->s0*m->s5*m->sF - m->s0*m->s7*m->sD - m->s4*m->s1*m->SF

    + m->s4*m->s3*m->sD + m->sC*m->s1*m->s7 - m->sC*m->s3*m->s5;

    inv.sE = -m->s0*m->s5*m->sE + m->s0*m->s6*m->sD + m->s4*m->s1*m->sE

    - m->s4*m->s2*m->sD - m->sC*m->s1*m->s6 + m->sC*m->s2*m->s5;

    inv.s3 = -m->s1*m->s6*m->sB + m->s1*m->s7*m->sA + m->s5*m->s2*m->sB

    - m->s5*m->s3*m->sA - m->s9*m->s2*m->s7 + m->s9*m->s3*m->s6;

    inv.s7 = m->s0*m->s6*m->sB - m->s0*m->s7*m->sA - m->s4*m->s2*m->sB

    + m->s4*m->s3*m->sA + m->s8*m->s2*m->s7 - m->s8*m->s3*m->s6;

    inv.sB = -m->s0*m->s5*m->sB + m->s0*m->s7*m->s9 + m->s4*m->s1*m->sB

    - m->s4*m->s3*m->s9 - m->s8*m->s1*m->s7 + m->s8*m->s3*m->s5;

    inv.sF = m->s0*m->s5*m->sA - m->s0*m->s6*m->s9 - m->s4*m->s1*m->sA

    + m->s4*m->s2*m->s9 + m->s8*m->s1*m->s6 - m->s8*m->s2*m->s5;

    det = m->s0*inv.s0 + m->s1*inv.s4 + m->s2*inv.s8 + m->s3*inv.sC;

    /*

    if (det < .00001f)

        return false;

    */

    det = 1.f / det;

    /*

    invOut->s0 = inv.s0 * det;

    invOut->s1 = inv.s1 * det;

    invOut->s2 = inv.s2 * det;

    invOut->s3 = inv.s3 * det;

    invOut->s4 = inv.s4 * det;

    invOut->s5 = inv.s5 * det;

    invOut->s6 = inv.s6 * det;

    invOut->s7 = inv.s7 * det;

    invOut->s8 = inv.s8 * det;

    invOut->s9 = inv.s9 * det;

    invOut->sA = inv.sA * det;

    invOut->sB = inv.sB * det;

    invOut->sC = inv.sC * det;

    invOut->sD = inv.sD * det;

    invOut->sE = inv.sE * det;

    invOut->sF = inv.sF * det;

    */

    return true;

}

In particolare il problema sono i 2 pezzi commentati:
- l'if statement che controlla se det è 0
- l'assegnazione dei valori finali della matrice invertita al parametro invOut

Se decommento una qualsiasi di quelle righe, clEnqueueNDRangeKernel(...) ritorna CL_INVALID_WORK_GROUP_SIZE.
Questo non me lo spiego proprio, anche perchè:
- tra i 2 tipi di istruzioni non vedo nessun collegamento
- le stesse istruzioni le uso in altre funzioni senza nessun tipo di problema

In particolare l'if statement crasha qualsiasi condizione o variabile metta, tranne i casi banali di if(true) e if(false). Negli ultimi 2 casi funziona come dovrebbe.

L'assegnazione dei valori a invOut invece ritorna errore se assegno uno dei valori di inv, ovvero dell'altra matrice.
Se assegno a invOut un valore direttamente, o una variabile di tipo float funziona correttamente (a meno che la variabile di tipo float non sia inizializzata direttamente a partire dal valore di inv, in quel caso riporta comunque errore).

Non so davvero come spiegarmelo questo :\

TheCrib · 07-05-2012 8:07

Lo so che l'ho gia' detto e che hai gia' provato senza ottimizzazioni.. ma secondo me internamente qualcosa viene comunque eliminata.

Quello che hanno in comune quei due blocchi di codice e' che quando sono commentati la funzione diventa inutile.

C'e' un qualche tipo di asm intermedio che si puo' vedere ? Magari qualcosa di specifico di NVidia, visto che non credo le specs di OpenCL si fermino al linguaggio pseudo-C.

Corralx · 07-05-2012 13:28

(07-05-2012 8:07)TheCrib ha scritto: Lo so che l'ho gia' detto e che hai gia' provato senza ottimizzazioni.. ma secondo me internamente qualcosa viene comunque eliminata.

Quello che hanno in comune quei due blocchi di codice e' che quando sono commentati la funzione diventa inutile.

C'e' un qualche tipo di asm intermedio che si puo' vedere ? Magari qualcosa di specifico di NVidia, visto che non credo le specs di OpenCL si fermino al linguaggio pseudo-C.

Non che io sappia.
C'era gDEbugger che debugga codice OpenCL, ma su OSX 10.7+ non funziona.

Comunque credo anche io la ragione sia quella.
In effetti il fatto che se metto true o false nella condizione booleana funzioni, mi fa pensare che il compilatore faccia dead code elimination e quindi "salti" il problema (del resto usa LLVM).

OpenCL usa C99 con qualche restrizione, nuove keyword e tipi di dati vettoriali.

EDIT: a quanto pare llvm 3.0 può compilare kernel OpenCL e emettere il suo bytecode o direttamente il binario da ispezionare.
Peccato che in llvm 3.1 (quello che ho io) il supporto sia stato rimosso per problemi Facepalm

Ricompilerò dal trunk di svn il 3.0....

Corralx · 08-05-2012 18:55

Per completezza (nella disgrazia) aggiungo gli ultimi sviluppi:
- Non riesco a chiedere il binario a OpenCL, a quanto pare crasha il programma se ci provo Asd

- Se compilo ed eseguo il codice su CPU funziona come dovrebbe, e chiaramente le funzioni di debug sono presenti solo li... Facepalm

Sempre più nella melma...

TheCrib · 08-05-2012 19:05

Great Pain Unit !

gso · 08-05-2012 21:04

Premesso che non ho mai usato OpenCL in vita mia e conto di non farlo mai...

(07-05-2012 1:00)Corralx ha scritto: Il problema è nella funzione che inverte la matrice:

Codice:

static bool __InvertMatrix(__private const float16* m, __private float16* invOut) { __private float16 inv; __private float det;

Queste strutture tipo float16 le hai scritte te?

Se le hai scritte te ti consiglio di importare "CL/cl_platform.h" e di sostituire le strutture/union con quelle ufficiali che sono cl_float16 , cl_float2 ecc... Prova anche a usare il tipo cl_float.

UPDATE: allora ho visto che questi float16 sono citati anche nella ref card OpenCL per cui, sebbene su linux non ci siano, da qualche altra parte esisteranno... Ma comunque rimpiazzerei tutto coi tipi cl_ .
Ho visto che per GPU Cuda va anche di moda usare i tipi delle librerie di cuda: (è un esempio senza l'uso di OpenCL)
http://rcabreral.blogspot.it/2010/04/inverting-complex-matrices-with-cuda.html

Per evitare problemi su problemi, fossi in te per il momento farei anche un rimpiazzo di tutti i bool con dei "cl_int" e toglierei più keyword possibili, soprattutto "static".

(07-05-2012 1:00)Corralx ha scritto: clEnqueueNDRangeKernel(...) ritorna CL_INVALID_WORK_GROUP_SIZE.
Questo non me lo spiego proprio, anche perchè:
- tra i 2 tipi di istruzioni non vedo nessun collegamento
- le stesse istruzioni le uso in altre funzioni senza nessun tipo di problema

L'errore dice proprio che i work items passati non sono della "grandezza giusta" che si aspettava... ergo mi sà che gli stai passando una lunghezza che nel caso della GPU è sbagliata.

Corralx · 08-05-2012 23:01

(08-05-2012 21:04)gso ha scritto: Premesso che non ho mai usato OpenCL in vita mia e conto di non farlo mai...

(07-05-2012 1:00)Corralx ha scritto: Il problema è nella funzione che inverte la matrice:

Codice:

static bool __InvertMatrix(__private const float16* m, __private float16* invOut) { __private float16 inv; __private float det;

Queste strutture tipo float16 le hai scritte te?

Se le hai scritte te ti consiglio di importare "CL/cl_platform.h" e di sostituire le strutture/union con quelle ufficiali che sono cl_float16 , cl_float2 ecc... Prova anche a usare il tipo cl_float.

UPDATE: allora ho visto che questi float16 sono citati anche nella ref card OpenCL per cui, sebbene su linux non ci siano, da qualche altra parte esisteranno... Ma comunque rimpiazzerei tutto coi tipi cl_ .
Ho visto che per GPU Cuda va anche di moda usare i tipi delle librerie di cuda: (è un esempio senza l'uso di OpenCL)
http://rcabreral.blogspot.it/2010/04/inverting-complex-matrices-with-cuda.html

Per evitare problemi su problemi, fossi in te per il momento farei anche un rimpiazzo di tutti i bool con dei "cl_int" e toglierei più keyword possibili, soprattutto "static".

(07-05-2012 1:00)Corralx ha scritto: clEnqueueNDRangeKernel(...) ritorna CL_INVALID_WORK_GROUP_SIZE.
Questo non me lo spiego proprio, anche perchè:
- tra i 2 tipi di istruzioni non vedo nessun collegamento
- le stesse istruzioni le uso in altre funzioni senza nessun tipo di problema

L'errore dice proprio che i work items passati non sono della "grandezza giusta" che si aspettava... ergo mi sà che gli stai passando una lunghezza che nel caso della GPU è sbagliata.

I tipi tipo float16 sono i tipi predefiniti di OpenCL.
Nel codice dell'applicazione vengono ridefiniti come cl_float16, mentre nel kernel di OpenCL, il C99 viene esteso col supporto nativo a quei tipi che si chiamano float16, float4 ecc...

Per quanto riguarda l'errore, il problema è che la grandezza è corretta, e non mi spiego come mai funzioni tutto a dovere, tranne quando una di quelle specifiche righe entra in gioco.
Non ha proprio senso che una condizione booleana o una assegnazione vadano a modificare il numero di thread in esecuzione, che peraltro viene deciso prima dell'esecuzione. Dead

gso · 08-05-2012 23:36

Boh, così però la vedo dura perchè a quanto pare non ci sono utenti con esperienze di OpenCL che abbiano avuto problemi di questo tipo...

Imo dovresti cercare di fare un campione piccolo concentrato di questo raytracer (una parola... Sorriso

) e postarlo qui, così uno se lo può scaricare e provare sulla sua postazione.

Io sarei portato per pensare a un casino coi tipi ma effettivamente può essere qualche altro problema...

Corralx · 09-05-2012 11:39

(08-05-2012 23:36)gso ha scritto: Boh, così però la vedo dura perchè a quanto pare non ci sono utenti con esperienze di OpenCL che abbiano avuto problemi di questo tipo...

Imo dovresti cercare di fare un campione piccolo concentrato di questo raytracer (una parola... ) e postarlo qui, così uno se lo può scaricare e provare sulla sua postazione.

Io sarei portato per pensare a un casino coi tipi ma effettivamente può essere qualche altro problema...

Difatti son giorni che cerco una spiegazione e nessuno ha incontrato tale problema.
Oggi se ho tempo creo un esempio minimale e lo carico su un repo git, così lo faccio provare ad un pò di gente.
Potrebbe essere colpa della mia implementazione ma mi sembra un bug troppo grosso per non essere stato trovato.

Cmq giusto come update, ho temporaneamente risolto il problema.
Dalla specifica ho dedotto che se passo NULL come local work size, lui cerca di fare il meglio che può (senza nessuna garanzia di ottimizzazione).
Così facendo funziona.
In particolare ho notato che riducendo il local work size a {10, 10} tutto funziona, ma chiaramente è estremamente più lento.

Una spiegazione sul perchè una assegnazione ad un puntatore dovrebbe far variare il numero di thread in esecuzione, e in particolare far superare il numero massimo, ancora non la ho Asd

Modo discussione \| Modo lineare Problema con OpenCL e i suoi kernel
Autore	Messaggio