25 Replies Latest reply on Jan 15, 2010 9:45 PM by Fr4nz

    Cannot modify correctly __local uint4 vectors variables: serious compiler bug?

    Fr4nz
      If yes, it is very serious...

      Hello,

      after posting yesterday about my problems with __local uint4 vectors variables, I tried to understand better the problem: it seems that I can read these variables without problems but I can't modify them correctly in an "usual" way (see below what I mean).

      Consider this simple kernel (ignore the parameters that aren't used inside the kernel...):

       

      /**
       * Kernel: uint4 vector test;
       */
      __kernel
      void
      intersect(__global uint4 *inputVD,
                __global uint *inputPrefixCk,
            __global uint *outputCounts,
            __local uint4 *tmpBuffer,
                const int k,
                const int dimRow) {

          // Simple mask
          const uint4 MASK = (uint4) (1);

          // Recover global ID of the work item;
          const int gid = get_global_id(0);

          // Test: modify __local vector variable with vstore4;
          vstore4(4,0,tmpBuffer+gid); // OK here....

         // Test: modify __local vector variable with an operator (in this case "+");
          tmpBuffer[gid] += MASK; // NOT ok here: cannot modify tmpBuffer value correctly! Only "x" and "w" components are updated correctly...
          
          // Final output of the work-item;
          outputCounts[gid] = tmpBuffer[gid].x + tmpBuffer[gid].y + tmpBuffer[gid].z + tmpBuffer[gid].w;

      }



      As you can read from the source comments, when I modify a __local vector uint4 variable location with "vstore4" everything works, but if I try a stupid operation (in the example a sum with another uint4 vector variable) only components "x" and "w" (the first and the last) are updated.

      Is this normal? What is the explanation of this? A compiler bug maybe? The only workaround I've found in order to update correctly a __local vector variable is to use vstore4...in fact if I replace the non-working line with:

       

      // Test: store the result of the sum with vstore4 instead of using a simple sum...
          vstore4(tmpBuffer[gid] + MASK,0,tmpBuffer+gid); // Again ok here....


      Then everything works fine. So it really seems we have a problem when writing "implictly" (that is, without using vstore4 explicitly ) a vector, because "y" and "z" components aren't updated or they take strange values...

        • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
          genaganna

          Fr4nz,

                 Your posting code is some how missed. Please post again.

            • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
              Fr4nz

              Sorry for the inconvenient, first post is fixed now!

                • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                  omkaranathan

                  Fr4nz,

                  I'm not having any issues in modifying local variables with operators. Which version of OpenCL are you using? Also could you post the host code too?

                   

                    • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                      Fr4nz

                       

                      Originally posted by: omkaranathan Fr4nz,

                       

                      I'm not having any issues in modifying local variables with operators. Which version of OpenCL are you using? Also could you post the host code too?

                       

                       

                      Are you saying that the same kernel I've posted works correctly for you??

                      I'm using the latest 9.12 hotfix driver for linux under Ubuntu 9.04 32bit; also, I've an ATI 5770.

                      Anyway, below here is reported my host code. Note that the class functions are called in this order:

                      1. initContext (get platform, create context);
                      2. parseSpecsGPU (get the characteristics of the first device: the GPU);
                      3. compileKernel (compilation of the kernel source);
                      4. executeGPUIntersect (create a command queue, setup kernel args and issue the various threads to the videocard);

                       

                      Please, tell me if there's something wrong in the host code because this problem is making me crazy :-(

                      As a sidenote, in the meantime I will try to reinstall completely the video drivers in order to see if the problem disappear...

                      Thank you for the answer!

                      /* * Classe che rappresenta l'interfaccia per l'uso delle GPU OpenCL-compatibili. */ // Inclusione degli header utili per OpenCL; #include "GPUInterface.h" #define __NO_STD_VECTOR // Vengono usati i vector nel namespace cl; #define __NO_STD_STRING // Vengono usate le stringhe nel namespace cl; #include "CL/cl.hpp" // Inclusione degli header standard utilizzati; #include <utility> #include <cstdio> #include <cstdlib> #include <fstream> #include <iostream> // Costruttori/distruttore classe; /** * Costruttore principale della classe: questo costruttore inizializza un contesto operante su un * dispositivo GPU. Tutti i campi dell'istanza saranno inizializzati in base a cio'. */ GPUInterface::GPUInterface() { // Inizializzo un contesto OpenCL; initContext(); // Recupero le caratteristiche salienti del dispositivo OpenCL da usare poi durante l'esecuzione // della parte di DCI coinvolta; parseSpecsGPU(); // Compilo il kernel da eseguire successivamente sul dispositivo OpenCL; compileKernel(); } /** * Distruttore principale della classe; */ GPUInterface::~GPUInterface() { // Deallocazione del contesto; if(this->contesto != NULL) delete contesto; } // Metodi pubblici classe; /** * Questo metodo ritorna la quantita' di memoria globale presente sul dispositivo GPU * rilevato. */ long GPUInterface::getGPUGlobalMem() const { return(this->maxMem*4); } /** * Questo metodo ritorna la massima quantita' di memoria allocabile sul dispositivo GPU rilevato. */ long GPUInterface::getGPUMaxMemAllocable() const { return(this->maxMem); } /** * Questo metodo ritorna il numero di unita' funzionali contenute in ogni processing element * del corrente dispositivo OpenCL. */ int GPUInterface::getNumVectors() const { return(this->numVectors); } /** * Questo metodo si occupa di eseguire un'iterazione tramite intersezioni di vettori * di bit a partire dal dataset verticale sul dispositivo OpenCL rilevato nel sistema. */ void GPUInterface::executeGPUIntersect(unsigned int *VD, unsigned int *prefixCkGPU, unsigned int *support, int numItems, int dimRow, int prefixCkSize, int k) { // Codice di errore tornato dalle varie funzioni; cl_int err; // Creazione della coda dei comandi associata al dispositivo GPU-OpenCL; cl::vector<cl::Device> listaDevices = this->contesto->getInfo<CL_CONTEXT_DEVICES>(); cl::CommandQueue queue((*(this->contesto)), listaDevices[0]); queue.setProperty(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, CL_TRUE, NULL); // Creazione dell'input buffer relativo al vertical dataset; da notare che il flag // CL_MEM_COPY_HOST_PTR indica che il contenuto di VD dev'essere copiato nella memoria // del dispositivo; // NOTA: Non serve inizializzare come buffer singole variabili! Quelle le si possono passare // come semplice argomento al kernel... cl::Buffer inputVD = cl::Buffer(*(this->contesto), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(unsigned int)*numItems*dimRow, VD, &err); // Creazione dell'input buffer relativo al prefix array; cl::Buffer inputPrefixCkGPU = cl::Buffer(*(this->contesto), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(unsigned int)*k*prefixCkSize, prefixCkGPU, &err); // Creazione dell'output buffer relativo ai conteggi dei vari candidati; cl::Buffer outputSupport = cl::Buffer(*(this->contesto), CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(unsigned int)*prefixCkSize, support, &err); // Inizializzazione degli argomenti dei kernel: essenzialmente ogni argomento del kernel // viene legato contro un buffer creato in precedenza (a meno che un argomento non // sia un puntatore ma un riferimento semplice: in tal caso si puo' direttamente passare come // argomento la variabile stessa che contiene il valore interessato)! // Questo "legame" e' utile quando si vogliono usare comandi di scrittura/lettura appartenenti // all'istanza della coda dei comandi (vedi sotto); this->kernel->setArg(0, inputVD); this->kernel->setArg(1, inputPrefixCkGPU); this->kernel->setArg(2, outputSupport); this->kernel->setArg(3, (size_t) dimRow*sizeof(cl_uint), NULL); // Per le variabili locali si DEVE saltare la creazione del buffer, indicandone solo la dimensione da allocare sulla GPU! this->kernel->setArg(4, (size_t) prefixCkSize*sizeof(cl_uint4), NULL); // Per le variabili locali si DEVE saltare la creazione del buffer, indicandone solo la dimensione da allocare sulla GPU! this->kernel->setArg(5, k); this->kernel->setArg(6, dimRow); // Copia dei buffer di input verso la GPU... // NOTA: Il VD verra' visto internamente al kernel come un vettore uint4! queue.enqueueWriteBuffer(inputVD, CL_TRUE, 0, sizeof(unsigned int)*numItems*dimRow, VD, NULL, NULL); queue.enqueueWriteBuffer(inputPrefixCkGPU, CL_TRUE, 0, sizeof(unsigned int)*k*prefixCkSize, prefixCkGPU, NULL, NULL); // Handler di eventi: serve ad aspettare che un comando dato alla coda sia finito; cl::Event e; // Creazione index space e successiva esecuzione dell'insieme di thread; // Parametri importanti da ricordare: // Il terzo parametro setta il numero di work-item TOTALI da eseguire; // Il quarto parametro specifica, eventualmente, le dimensioni di un workgroup; err = queue.enqueueNDRangeKernel(*(this->kernel), cl::NullRange, cl::NDRange(prefixCkSize), cl::NullRange, 0, &e); e.wait(); // Recupero dal dispositivo i risultati ottenuti... err = queue.enqueueReadBuffer(outputSupport,CL_TRUE,0,sizeof(unsigned int)*prefixCkSize,support,NULL,&e); e.wait(); // Fine esecuzione dei kernel; le risorse occupate dalla coda dei comandi vengono liberate; err = queue.finish(); // printf("\n%d - %u", prefixCkSize, support[4]); } // Metodi private classe; /** * Questo metodo privato inizializza un contesto sulla piattaforma attuale. */ void GPUInterface::initContext() { // Rilevamento piattaforma; cl_int err; //create context containing the target devices cl::vector<cl::Platform> platforms; // Recupero delle piattaforme disponibili; err = cl::Platform::get(&platforms); if(err == CL_SUCCESS) printf("Piattaforma OpenCL rilevata..."); else { printf("\nPiattaforma OpenCL non rilevata... %d", err); // TODO: Da aggiungere l'uscita immediata dal programma in caso di // fallimento... } // Inizializza la prima piattaforma trovata; cl::Platform piattaforma(platforms[0]); // Recupero devices GPU OpenCL-compatibili dalla piattaforma rilevata; cl::vector<cl::Device> devices; err = platforms[0].getDevices(CL_DEVICE_TYPE_GPU, &devices); if(err == CL_SUCCESS) printf("\nGPU OpenCL-compatibile rilevata..."); else { printf("\nGPU OpenCL-compatibile non rilevata... %d", err); // TODO: Da aggiungere l'uscita immediata dal programma in caso di // fallimento... } // Creazione contesto sulla piattaforma ed il dispositivo GPU rilevato... cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[0])(), 0 }; this->contesto = new cl::Context(CL_DEVICE_TYPE_GPU, cps, NULL, NULL, &err); if (err == CL_SUCCESS) printf("contesto generato con successo!"); else { printf("\nContesto non generato... %d", err); // TODO: Da aggiungere l'uscita immediata dal programma in caso di // fallimento... } } /** * Questo metodo privato inizializza i campi dell'istanza relativi alle caratteristiche della * GPU trovata sull'attuale sistema. */ void GPUInterface::parseSpecsGPU() { // Recupero la lista di devices associata ad un contesto; in realta' usero solo // il primo dispositivo rilevato; cl::vector<cl::Device> devices = this->contesto->getInfo<CL_CONTEXT_DEVICES>(); // Interrogazione caratteristiche salienti del dispositivo GPU collegato al contesto; cl::string tmpString = devices[0].getInfo<CL_DEVICE_VENDOR>(); const char *vendor = tmpString.c_str(); // Vendor GPU tmpString = devices[0].getInfo<CL_DEVICE_NAME>(); const char *model = tmpString.c_str(); // Modello GPU int freq = devices[0].getInfo<CL_DEVICE_MAX_CLOCK_FREQUENCY>(); // Frequenza GPU; this->maxMem = devices[0].getInfo<CL_DEVICE_MAX_MEM_ALLOC_SIZE>(); // Memoria massima allocabile per un oggetto; this->localMem = devices[0].getInfo<CL_DEVICE_LOCAL_MEM_SIZE>(); // Dimensione memorie locali; int numSIMDEngines = devices[0].getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>(); // Numero di SIMD engines; this->numVectors = devices[0].getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT>(); // Numero di unita' vettoriali per processing element this->maxWGroup = devices[0].getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>(); // Dimensione di un workgroup // Stampa delle caratteristiche ritrovate; printf("\n\n**** Caratteristiche GPU rilevate ****"); printf("\nVendor: %s\nModel: %s\nOp.frequency: %d Mhz", vendor, model, freq); printf("\nQuantita' di memoria globale sul dispositivo: %ld bytes (%ld Mbytes)", 4*this->maxMem, 4*this->maxMem/1024/1024); printf("\nMemoria massima allocabile sul dispositivo per un oggetto: %ld bytes (%ld Mbytes) (1/4 del totale)", this->maxMem, this->maxMem/1024/1024); printf("\nDimensione delle memorie locali: %ld bytes (%ld Kbytes)", this->localMem, this->localMem/1024); printf("\nUnita' vettoriali per processing element: %d (consigliato l'uso del tipo \"%duint\")", this->numVectors, this->numVectors); printf("\nNumero di compute units: %d", numSIMDEngines); printf("\nDimensione massima di un workgroup: %d work-items", this->maxWGroup); printf("\n\n"); } /** * Questo metodo compila il kernel relativo alla parte di intersezioni eseguita sulla GPU. La * compilazione avviene all'inizio dell'esecuzione del programma una volta per tutte. */ void GPUInterface::compileKernel() { // NOTA: In eclipse il kernel deve trovarsi dentro la cartella root del progetto // affinche' sia vista; in condizioni di esecuzione "noramle", il kernel deve trovarsi // nella stessa cartella dell'eseguibile; const char* nomeKernel = "intersect.cl"; printf("**** COMPILAZIONE KERNEL OPENCL *****"); printf("\nNome kernel: %s", nomeKernel); // Apre il file del kernel in lettura; std::ifstream file(nomeKernel); // Parserizza il file nella stringa program; std::string textSrc(std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>())); // Copia il codice sorgente in un apposito contenitore che serve a gestirlo, successivamente, // per la compilazione; cl::Program::Sources source(1, std::make_pair(textSrc.c_str(), textSrc.length()+1)); // Associo il codice sorgente del kernel al contesto tramite la creazione // di un oggetto Program; cl::Program program((*(this->contesto)),source); // **** Compilazione del kernel per il dispositivo rilevato **** // Recupero la lista di devices associata ad un contesto; in realta' usero' solo // il primo dispositivo rilevato; cl::vector<cl::Device> devices = this->contesto->getInfo<CL_CONTEXT_DEVICES>(); cl_int err; err = program.build(devices); // A questo punto creo l'oggetto kernel vero e proprio. // Esso sara' collegato al programma sopra compilato. A questo punto posso // inoltre capire se il kernel e' stato compilato correttamente oppure no. this->kernel = new cl::Kernel(program, "intersect", &err); if(err == CL_SUCCESS) printf("\nCompilazione del kernel effettuata con successo!\n\n"); // L'esito con successo e' 0! else { printf("\nErrore durante la compilazione...\n\n"); exit(1); } }

                        • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                          omkaranathan

                          I have tried a test case which modifies local variable with operator. Even many of the samples are doing the same.

                          @code : Seems the header file is missing. The cpp file too seems to be incomplete.

                            • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                              Fr4nz

                               

                              Originally posted by: omkaranathan I have tried a test case which modifies local variable with operator. Even many of the samples are doing the same.

                               

                              @code : Seems the header file is missing.

                               

                              Then it's really strange that I only have problems with __local vectors...and it is very very strange that only x and w components are correctly vritten. I really can't see a reason. Anyway I'm going to reinstall drivers and SDK, I hope this will solve the situation. I will update you later.

                              Oh, below here is the header, I hope you can find something wrong inside my class...

                              /* * GPUInterface.h * * Created on: 11/gen/2010 * Author: fr4nz */ #ifndef GPUINTERFACE_H_ #define GPUINTERFACE_H_ // Forward declarations; da notare che le classi di cui faccio le f.d. appartengono al namespace cl, // per cui devo tenere conto di questo quando eseguo la f.d.! namespace cl { class Context; // Contesto per gestire il dispositivo OpenCL; class Kernel; } class GPUInterface { private: // Campi privati classe; /** * Puntatore che fa riferimento ad un oggetto contenente il contesto OpenCL in cui si * opera; */ cl::Context *contesto; /** * Oggetto kernel: rappresenta il kernel da eseguire sui processing elements * del dispositivo OpenCL. */ cl::Kernel *kernel; /** * Campo contenente la memoria massima allocabile per un oggetto sul dispositivo OpenCL; */ long maxMem; /** * Campo contenente la dimensione delle memorie locali; ricordiamo che una memoria locale * e' condivisa da tutti i work-item di un work-group. */ long localMem; /** * Campo contenente il numero di unita' funzionali per processing elements; e' sempre bene * vettorizzare i calcoli per sfruttare il piu' possibile la potenza computazionale del * dispositivo OpenCL. Tale informazione, pertanto, puo' dare un'idea su come bisogna * organizzare le strutture dati da passare al dispositivo. */ int numVectors; /** * Numero massimo di work-item per work-group. */ int maxWGroup; public: // Costruttori/distruttore classe; /** * Costruttore principale della classe; */ GPUInterface(); /** * Distruttore della classe; */ ~GPUInterface(); // Metodi pubblici classe; /** * Questo metodo ritorna la quantita' di memoria globale presente sul dispositivo GPU * rilevato. */ long getGPUGlobalMem() const; /** * Questo metodo ritorna la massima quantita' di memoria allocabile sul dispositivo GPU rilevato. */ long getGPUMaxMemAllocable() const; /** * Questo metodo ritorna il numero di unita' funzionali contenute in ogni processing element * del corrente dispositivo OpenCL. */ int getNumVectors() const; /** * Questo metodo si occupa di eseguire un'iterazione tramite intersezioni di vettori * di bit a partire dal dataset verticale sul dispositivo OpenCL rilevato nel sistema. * * INPUT: * - VD: Puntatore al vertical dataset su cui operare; * - prefixCkGPU: Puntatore alla struttura dati dell'array dei prefissi, opportunamente rielaborata * per l'uso su GPU; * - numItems: Numero di item raffigurati dal VD; * - dimRow: Dimensione di una riga d'item; * - prefixCkSize: Numero di candidati di cui calcolare il supporto; * - k: Indice dell'iterazione corrente; * * EFFETTI: * - La struttura dati supports verra' modificata in maniera tale da contenere i supporti * dei vari candidati; */ void executeGPUIntersect(unsigned int *VD, unsigned int *prefixCkGPU, unsigned int *support, int numItems, int dimRow, int prefixCkSize, int k); private: // Metodi privati classe; /** * Questo metodo inizializza un contesto sulla piattaforma attuale. */ void initContext(); /** * Questo metodo recupera le caratteristiche principali della GPU, utili al fine di * "tarare" l'algoritmo; */ void parseSpecsGPU(); /** * Questo metodo compila il kernel relativo alla parte di intersezioni eseguita sulla GPU. La * compilazione avviene all'inizio dell'esecuzione del programma una volta per tutte. */ void compileKernel(); }; #endif /* GPUINTERFACE_H_ */

                                • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                  Fr4nz

                                  UPDATE - No way: I've reinstalled drivers, SDK and follow the usual installation instructions: local vector variables still aren't updated in their y and z components when modified with operators. For example, consider this kernel:

                                   

                                  __kernel
                                  void
                                  intersect(__global uint *outputCounts,
                                        __local uint4 *tmpBuffer) {
                                     
                                      const int gid = get_global_id(0);

                                      vstore4(3,0,tmpBuffer+gid);

                                      tmpBuffer[gid] += tmpBuffer[gid];
                                     
                                      outputCounts[gid] = tmpBuffer[gid].y;
                                  }


                                  The value returned should be 6, but it gives 3 (the original value loaded with vstore4). Same thing with z component, while x and w are correct (6).

                                  Oh, tmpBuffer, as you can see from the host code, has size > 1: did you consider in your test cases trying to modify local vector variables with size > 1 ?

                                  PLEASE, test the above kernel and check if it gives correct result to you. The only thing that comes to my mind is that I've done something wrong in the host code (but sincerely I don't see where). I'll try to replace the use of cl.hpp with cl.h and see if something changes...

                                    • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                      omkaranathan

                                      The code provided(.cpp) is incomplete.

                                        • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                          Fr4nz

                                           

                                          Originally posted by: omkaranathan The code provided(.cpp) is incomplete.

                                           

                                          This class is part of a large project, but it is the ONLY class that manages the GPU. If there's a bug, it must be in this class, believe me. It would be impractical to post here 10 classes that doesn't have anything to do with the problem.

                                          Anyway, may you try my simple kernel posted above, with tmpBuffer having size > (1*cl_uint4), and tell me if it gives you correct results?

                                            • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                              nou

                                              try this three thing. separately.  use float instead of uint

                                              declare tmpBuffer inside of kernel

                                              __kernel void intersect(__global uint *outputCounts) {
                                                  __local uint4 *tmpBuffer[LOCAL_SIZE];

                                              use barrier(CLK_LOCAL_MEM_FENCE);

                                                • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                  Fr4nz

                                                   

                                                  Originally posted by: nou try this three thing. separately.  use float instead of uint

                                                   

                                                  declare tmpBuffer inside of kernel

                                                   

                                                  __kernel void intersect(__global uint *outputCounts) {     __local uint4 *tmpBuffer[LOCAL_SIZE];

                                                   

                                                  use barrier(CLK_LOCAL_MEM_FENCE);

                                                   

                                                   

                                                  Hi Nou,

                                                  now I'm rewriting the host code with cl.h library (instead of using c++ bindings), hoping that this strange "behaviour" disappears...if this will be unsuccesful, I will try your suggestions and report here the results. Thank you for your interest in my problem!

                                                  • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                    Fr4nz

                                                     

                                                    Originally posted by: nou try this three thing. separately.  use float instead of uint

                                                     

                                                    declare tmpBuffer inside of kernel

                                                     

                                                    __kernel void intersect(__global uint *outputCounts) {     __local uint4 *tmpBuffer[LOCAL_SIZE];

                                                     

                                                    use barrier(CLK_LOCAL_MEM_FENCE);

                                                     

                                                    Hi nou, I've made interesting findings with your tests I think!  First of all: I've replaced the old GPU init code, which used C++ bindings, with the new code that uses cl.h, but I had NO luck: same problems as before with __local vectors. Then, I've made the tests you've suggested and here are the results:

                                                    1. If I use float4 instead of uint4, only the "w" component is updated;
                                                    2. Here it is an important finding: If I use the barrier before the last line, then x and w components have always the right value, while "y" and "z" sometimes have the right value (6), sometimes have a random value; if I use the barrier after the last command, the "usual" behaviour is found (x and w OK, y and z NOT); this behaviour is strange and makes me think that the host code doesn't wait properly for the end of the executed instructions. I've posted my host code below, see "executeGPUIntersect()" function;
                                                    3. If I use a __local uint4 vector declared inside the kernel with a size given by get_local_size(0), only the "x" component is updated;

                                                    I hope that someone can point me to the solution, because I'm getting really tired of this strange problem

                                                    /* * Classe che rappresenta l'interfaccia per l'uso delle GPU OpenCL-compatibili. */ // Inclusione degli header utili per OpenCL; #include "GPUInterface.h" // Inclusione degli header standard utilizzati; #include <cstdio> #include <cstdlib> #include <fstream> #include <iostream> #include <string.h> // Costruttori/distruttore classe; /** * Costruttore principale della classe: questo costruttore inizializza un contesto operante su un * dispositivo GPU. Tutti i campi dell'istanza saranno inizializzati in base a cio'. */ GPUInterface::GPUInterface() { // Inizializzo un contesto OpenCL; initContext(); // Recupero le caratteristiche salienti del dispositivo OpenCL da usare poi durante // l'esecuzione della parte di DCI coinvolta. // Inoltre alloco la CODA DEI COMANDI; parseSpecsGPUAllocQueue(); // Compilo il kernel da eseguire successivamente sul dispositivo OpenCL e creo l'oggetto // "kernel" che lo gestisce; compileKernel(); } /** * Distruttore principale della classe; */ GPUInterface::~GPUInterface() { } // Metodi pubblici classe; /** * Questo metodo ritorna la quantita' di memoria globale presente sul dispositivo GPU * rilevato. */ long GPUInterface::getGPUGlobalMem() const { return((long)this->maxGlobMem*4); } /** * Questo metodo ritorna la massima quantita' di memoria allocabile sul dispositivo GPU rilevato. */ long GPUInterface::getGPUMaxMemAllocable() const { return(this->maxGlobMem); } /** * Questo metodo ritorna il numero di unita' funzionali contenute in ogni processing element * del corrente dispositivo OpenCL. */ int GPUInterface::getNumVectors() const { return(this->numVectors); } /** * Questo metodo si occupa di eseguire un'iterazione tramite intersezioni di vettori * di bit a partire dal dataset verticale sul dispositivo OpenCL rilevato nel sistema. */ void GPUInterface::executeGPUIntersect(unsigned int *VD, unsigned int *prefixCkGPU, unsigned int *support, int numItems, int dimRow, int prefixCkSize, int k) { // Codice di errore tornato dalle varie funzioni; cl_int errCode; // **** Creazione dei BUFFER __global usati dai kernel in INPUT/OUTPUT **** cl_mem outputBufferCL = clCreateBuffer(this->contesto, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, (size_t) prefixCkSize*sizeof(cl_uint), support, &errCode); if(errCode != CL_SUCCESS) printf("Buffer inizializzati male! %d", errCode); // **** FINE creazione buffer globali **** // **** Inizializzazione dei parametri del kernel **** errCode = clSetKernelArg(this->kernel, 0, sizeof(cl_mem), (void *)&outputBufferCL); if(errCode != CL_SUCCESS) printf("Argomento kernel inizializzato male! %d", errCode); errCode = clSetKernelArg(this->kernel, 1, prefixCkSize*sizeof(cl_uint4), NULL); if(errCode != CL_SUCCESS) printf("Argomento kernel inizializzato male! %d", errCode); // **** Fine inizializzazione dei parametri del kernel **** // INIZIO comandi inviati alla GPU; // Array di eventi usati per attendere che venga eseguito totalmente un determinato comando; cl_event events[2]; // Numero di thread totali; size_t globalThreads[1]= {prefixCkSize}; // Numero di thread locali per workgroup - NON DEFINITO (non serve con task parallelism...in teoria!); // ESEGUO i threads; errCode = clEnqueueNDRangeKernel(this->coda, this->kernel, 1, NULL, globalThreads, NULL, 0, NULL, &events[0]); if(errCode != CL_SUCCESS) printf("AIA! %d", errCode); // Aspetto che i thread finiscano l'esecuzione... errCode = clWaitForEvents(1, &events[0]); if(errCode != CL_SUCCESS) printf("AIA!"); // Leggo i risultati; errCode = clEnqueueReadBuffer(this->coda, outputBufferCL, CL_TRUE, 0, (size_t) prefixCkSize*sizeof(cl_uint), support, 0, NULL, &events[1]); // Aspetto che i thread finiscano l'esecuzione... errCode = clWaitForEvents(1, &events[1]); // Rilascio le risorse utilizzate; clReleaseEvent(events[0]); clReleaseEvent(events[1]); // FINE comandi inviati alla GPU printf("\nDEBUG KERNEL %d - %u", prefixCkSize, support[0]); } // Metodi private classe; /** * Questo metodo privato inizializza un contesto sulla piattaforma attuale. */ void GPUInterface::initContext() { printf("**** Inizializzazione GPU OpenCL-compatibile ****"); // Variabile per il recupero dei codici di errore cl_int errorCode; // Codice d'errore ritornato in seguito ad una operazione; // Numero di entries da memorizzare (serve ad allocare lo spazio necessario per i // i puntatori passati ad una funzione OpenCL); cl_uint numEntries = 1; // Variabili ausiliarie per recuperare informazioni sulla piattaforma corrente; // Lista delle piattaforme disponibili; cl_platform_id *platforms = (cl_platform_id*) malloc(numEntries*sizeof(cl_platform_id)); // Numero di piattaforme disponibili; cl_uint *numPlatforms = (cl_uint*) malloc(numEntries*sizeof(cl_uint)); // Recupero informazioni sulla piattaforma in cui viene eseguito il programma... // // Lista degli argomenti (INPUT/OUTPUT): // 1) Il primo parametro indica quante piattaforme possono essere aggiunte alla // lista rappresentata dal secondo parametro; se viene impostata a NULL, non ci // sarà un limite massimo di piattaforme aggiungibili; // 2) (puntatore, cl_platform_id) Il secondo parametro rappresenta la lista di piattaforme ritrovate; il valori // contenuto saranno gli identificativi delle varie piattaforme; // 3) (puntatore, cl_unit) Il terzo parametro rappresenta il numero di piattaforme ritrovate; // // OUTPUT: cl_int - L'esito dell'operazione; errorCode = clGetPlatformIDs(numEntries, platforms, numPlatforms); // printf("\nCodice restituito da clGetPlatformIDs: %d", errorCode); printf("\nNumero di piattaforme: %d", *numPlatforms); // Variabili relative ai dispositivi rilevati sulla piattaforma corrente cl_device_id *devices = (cl_device_id*) malloc(numEntries*sizeof(cl_device_id)); // Lista (vettore) dei dispositivi disponibili sulla piattaforma; cl_uint *numDevices = (cl_uint*) malloc(sizeof(cl_uint)); // Numero di dispositivi disponibili sulla piattaforma; // Recupero informazioni sui dispositivi presenti nella piattaforma; il risultato // sarà una lista di identificativi; // // Lista degli argomenti (INPUT/OUTPUT): // 1) Il primo parametro rappresenta la lista delle piattaforme; // 2) Il secondo parametro rappresenta il tipo di dispositivi da cui si vuole ricavare // l'ID (CPU, GPU, APU, etc.); // 3) Il terzo parametro indica quanti sono i dispositivi da aggiungere alla lista // ritornata da clGetDeviceIDs (se il quarto parametro è NULL, può assumere un valore // pari a zero: in tal caso il valore verrà ignorato); // 4) (puntatore, cl_device_id) Il quarto parametro rappresenta la lista di dispositivi // trovati; // 5) Il quinto parametro rappresenta il numero di dispositivi ritrovati sulla piattaforma // il cui tipo è quello indicato nel secondo parametro; // // OUTPUT: cl_int - L'esito dell'operazione; clGetDeviceIDs(*platforms,CL_DEVICE_TYPE_GPU,numEntries,devices,numDevices); // printf("\n\nCodice restituito da clGetDeviceIDs: %d", errorCode); printf("\nNumero di GPU OpenCL-compatibili trovate: %d", *numDevices); // Se sono su una piattaforma con GPU che non supportano OpenCL... if(numDevices == 0) { printf("\nNessuna GPU OpenCL-capable trovata, esco..."); exit(1); } // Creazione contesto sulla piattaforma ed il dispositivo GPU rilevato... // Inizializzo le variabile che caratterizza le proprietà di un contesto. Da notare // che è una lista costituita da coppie "(CL_CONTEXT_PLATFORM,id_platform)" e terminata // da uno zero finale; cl_context_properties cps[3] = { // Coppia; CL_CONTEXT_PLATFORM, (cl_context_properties) platforms[0], // DEBUG - Bisogna fare il cast? Va bene passare il primo argomento? Debuggare... // Zero finale; 0 }; // Creo un contesto sul dispositivo OpenCL-capable per l'esecuzione dei kernel // Lista degli argomenti: // 1) (puntatore) Il primo parametro rappresenta una lista di proprietà relative // al contesto; // 2) Il secondo parametro indica il numero di dispositivi da associare al contesto; // 3) Il terzo parametro è il vettore contenente gli identificativi dei dispositivi da associare; // 4) Il quarto parametro è un puntatore ad una eventuale funzione di callback (NULL); // 5) Il quinto parametro è un puntatore a dati associati alla funzione di callback (NULL); // 6) Il sesto parametro è un puntatore al codice da ritornare successivamente all'esecuzione della funzione; // dopo la creazione del contesto; this->contesto = clCreateContext(cps,*numDevices,devices,NULL,NULL,&errorCode); if(errorCode != CL_SUCCESS) { printf("Inizializzazione di un contesto fallita, esco..."); exit(1); } printf("\nContesto OpenCL inizializzato con successo!"); } /** * Questo metodo privato inizializza i campi dell'istanza relativi alle caratteristiche della * GPU trovata sull'attuale sistema. */ void GPUInterface::parseSpecsGPUAllocQueue() { // Codice d'errore restituito dalle funzioni della libreria CL; cl_int errCode; // Dimensione dell'oggetto ritornato dalla funzione (NOTA: ritorna la dimensione "grezza" in bytes, // NON il numero di elementi contenuti dentro il vettore!); size_t *numDevices = new size_t; // Prima passata: recupero la dimensione della lista dei dispositivi associata al contesto aperto; // // Lista degli argomenti: // 1) Il primo parametro è il contesto creato; // 2) Il secondo parametro è di tipo "enum" è serve a specificare il tipo di informazione // che si richiede alla funzione; // 3) Il terzo parametro indica la dimensione dei dati restituita dal quarto parametro // (vedi sotto); // 4) (puntatore) Il quarto parametro indica la locazione di memoria in cui immagazzinare // le informazioni restiuite da clGetContextInfo; se è NULL l'informazione non viene // immagazzinata; // 5) (puntatore) Il quinto parametro ritorna la dimensione effettiva dei dati ritornati; // in questo caso rappresenta il numero di dispositivi coinvolti dal contesto; errCode = clGetContextInfo(contesto, CL_CONTEXT_DEVICES, 0, NULL, numDevices); printf("\nNumero di dispositivi associati al contesto: %d", *numDevices/sizeof(size_t)); // Seconda passata: recupero la lista dei dispositivi; cl_device_id *listaDevices = new cl_device_id; errCode = clGetContextInfo(contesto, CL_CONTEXT_DEVICES, *numDevices, listaDevices, NULL); if(errCode != CL_SUCCESS) exit(1); // Recupero le caratteristiche salienti della GPU; // Variabili usate per il recupero; cl_uint gpuMhz; cl_ulong sizeMemGlob; cl_ulong sizeMemLoc; cl_uint computeUnits; cl_uint maxWGroup; cl_uint numVectUnits; errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(cl_ulong),&sizeMemGlob,NULL); errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_LOCAL_MEM_SIZE,sizeof(cl_ulong),&sizeMemLoc,NULL); errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_MAX_CLOCK_FREQUENCY,sizeof(cl_uint),&gpuMhz,NULL); errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(cl_uint),&computeUnits,NULL); errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_MAX_WORK_GROUP_SIZE,sizeof(cl_uint),&maxWGroup,NULL); errCode = clGetDeviceInfo(listaDevices[0],CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT,sizeof(cl_uint),&numVectUnits,NULL); // Inizializzazione dei campi della classe contenenti le caratteristiche d'interesse della GPU; this->maxGlobMem = (long) sizeMemGlob; this->maxLocMem = (long) sizeMemLoc; this->numVectors = (long) numVectUnits; // printf("\n\nCodice restituito da clGetDeviceIDs: %d", errorCode); printf("\nFrequenza operativa della GPU: %d Mhz", gpuMhz); printf("\nDimensione memoria principale: %ld bytes (%ld MBytes)", this->maxGlobMem*4, this->maxGlobMem*4/1024/1024); printf("\nDimensione memorie locali (una per ogni P.E.): %ld bytes (%ld Kbytes)", this->maxLocMem, this->maxLocMem/1024); printf("\nNumero di Processing Elements (alias \"SIMD engines\"): %d", computeUnits); printf("\nDimensione massima per un work-group: %d", maxWGroup); printf("\nNumero di unita' funzionali per processing element: %d", this->numVectors); // Inizializzazione della coda dei comandi: associo un dispositivo ed il contesto a cui // appartiene una coda dei comandi, da usare durante l'esecuzione per fornire comandi di // lettura/scrittura buffer e di esecuzione di un insieme di thread; this->coda = clCreateCommandQueue(this->contesto, listaDevices[0], CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &errCode); if(errCode != CL_SUCCESS) { printf("Fallimento nella creazione della coda dei comandi! Esco..."); exit(1); } printf("\nCoda dei comandi creata con successo!"); } /** * Questo metodo compila il kernel relativo alla parte di intersezioni eseguita sulla GPU. La * compilazione avviene all'inizio dell'esecuzione del programma una volta per tutte. */ void GPUInterface::compileKernel() { // Codice ritornato da una funzione CL; cl_int errCode; // NOTA: In eclipse il kernel deve trovarsi dentro la cartella root del progetto // affinche' sia vista; in condizioni di esecuzione "noramle", il kernel deve trovarsi // nella stessa cartella dell'eseguibile; const char* nomeKernel = "intersect.cl"; printf("\n\n**** COMPILAZIONE KERNEL OPENCL *****"); printf("\nNome kernel: %s", nomeKernel); // Apre il file del kernel in lettura; std::ifstream file(nomeKernel); // Parserizza il file nella stringa program; std::string textSrc(std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>())); // Converto il sorgente in un array di char utilizzabile dalla funzione CL; const char *source = textSrc.c_str(); size_t sizeSource[1] = {strlen(source)}; // Associazione di un codice sorgente al contesto; successivamente verra' compilato. // Argomenti: // 1) Contesto a cui associare il kernel; // 2) Non documentato, boh! // 3) Puntatore al vettore di caratteri che contiene il sorgente; // 4) Dimensione del sorgente; // 5) Puntatore al codice tornato dalla funzione; this->programma = clCreateProgramWithSource(this->contesto, 1, &source, sizeSource, &errCode); if(errCode != CL_SUCCESS) { printf("nFallimento associazione sorgente->contesto! Esco..."); exit(1); } printf("\nAssociazione sorgente->contesto creata correttamente!!"); // Recupero la lista dei dispositivi (conterra' solo il primo dispositivo, ovvero la GPU); size_t numDevices; // Prima recupero la dimensione di allocazione della lista di dispositivi; errCode = clGetContextInfo(this->contesto, CL_CONTEXT_DEVICES, 0, NULL, &numDevices); cl_device_id *listaDevices = new cl_device_id; // Poi recupero la lista; errCode = clGetContextInfo(this->contesto, CL_CONTEXT_DEVICES, numDevices, listaDevices, NULL); // Compilazione del programma creato prima; errCode = clBuildProgram(this->programma, 1, listaDevices, NULL, NULL, NULL); if(errCode != CL_SUCCESS) { // Recupero le informazioni di errore date dal compilatore; printf("\nFallimento compilazione kernel! Errori ritornati dal compilatore:\n"); size_t dimLog; clGetProgramBuildInfo(this->programma, listaDevices[0], CL_PROGRAM_BUILD_LOG, NULL, NULL, &dimLog); char logInfo[dimLog]; clGetProgramBuildInfo(this->programma, listaDevices[0], CL_PROGRAM_BUILD_LOG, dimLog, logInfo, NULL); printf("\n%s", logInfo); exit(1); } printf("\nCompilazione kernel effettuata correttamente!\n\n"); // Creazione oggetto kernel; da usare successivamente... this->kernel = clCreateKernel(this->programma, "intersect", &errCode); if(errCode != CL_SUCCESS) { printf("\nFallimento creazione oggetto kernel! Esco..."); exit(1); } }

                                                      • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                        nou

                                                        what is your global and local work size? for what size did you initialize tmpBuffer.

                                                        now i realize that your code will work only if you run only one workgroup. you can not address local arrays with global ID.

                                                        size_t lid = get_local_id(0);

                                                        tmpBuffer[lid] += tmpBuffer[lid];

                                                          • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                            Fr4nz

                                                             

                                                            Originally posted by: nou what is your global and local work size? for what size did you initialize tmpBuffer.


                                                            In order to illustrate why I made some choices with global/local work size I must illustrate you the problem I'm solving. The problem I want to solve with GPU is this: I have a set of vectors, let's say {1,2,3,4,5,6,etc}, and I want to make many AND-intersections on them in parallel, that is for example:

                                                            inters1 = 1 & 3 & 5;

                                                            inters2 = 2 & 5 & 6;

                                                            inters3 = 3 & 5 & 6;

                                                            ....

                                                            and so on. At the end, for every AND-intersection calculated threads count the "1" bits contained in their assignated intersection and and store the result in the output buffer (at the right position).

                                                            As you can see this is, potentially, a "task-parallel" problem with some possible optimizations; so, my "naive" idea is to use workgroups of size 1 and the total number of workgroups should be equal to the total number of the intersections I have to make. In fact, if you check my code I use this parameters in clEnqueueNDRangeKernel:

                                                             

                                                            size_t globalThreads[1]= {prefixCkSize};

                                                                errCode = clEnqueueNDRangeKernel(this->coda,
                                                                                                this->kernel,
                                                                                                1,
                                                                                                NULL,
                                                                                                globalThreads,
                                                                                                NULL,
                                                                                                0,
                                                                                                NULL,
                                                                                                &events[0]);


                                                            Taken into account what I've told you above, the size of tmpBuffer should be enough to contain at least "numTotalIntersections * sizeof(uint4)" bytes (remember that I have to return the number of "1" bits for every intersections, not the entire AND-intersections, so I can calculate an intersection taking into account only a slice at a time, with a lot of local memory saving).

                                                            And here we return to my initial problem: how the global/local size could influence the results I obtain with the simple kernels I've posted in this thread? This seems nonsensical to me.

                                                              • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                                nou

                                                                well local memory is when you need to work item share some data between.

                                                                local can be used as optimization if you read some sort of data with many work item then you can store this data into local memory because it is faste than global.

                                                                problem is addressing local arrays.see my previous post.

                                                                  • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                                    Fr4nz

                                                                     

                                                                    Originally posted by: nou well local memory is when you need to work item share some data between.

                                                                     

                                                                    local can be used as optimization if you read some sort of data with many work item then you can store this data into local memory because it is faste than global.

                                                                     

                                                                    problem is addressing local arrays.see my previous post.

                                                                     

                                                                    Nou you got the point: I want to use local memory because is faster (use it to store there temp intersections). I'm going to try your suggestion immediately, and see if this solves the problem.

                                                                    Anyway, get_global_id shouldn't work either? Every work-item has its own global id...

                                                                     

                                                                      • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                                        nou

                                                                        yes but in each workgroup each work item has own position. and each workgroup have unique local memory which is shared between work item in the same workgroup.

                                                                        when you pass NULL as local_workgroup_size it is on the driver decide optimal workgroup size. on ATi card it is multply of 64. global size must be divisible with local_size.

                                                                        //have 16 global work item and workgroup size 4 global_id 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 local_id 0 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3

                                                                          • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                                            Fr4nz

                                                                             

                                                                            Originally posted by: nou yes but in each workgroup each work item has own position. and each workgroup have unique local memory which is shared between work item in the same workgroup.

                                                                             

                                                                            when you pass NULL as local_workgroup_size it is on the driver decide optimal workgroup size. on ATi card it is multply of 64. global size must be divisible with local_size.

                                                                             

                                                                             

                                                                            RIght, I've modified the enqueueNDRange as here:

                                                                             

                                                                            // Numero di thread totali;
                                                                                size_t globalThreads[1]= {prefixCkSize};
                                                                                size_t localThreads[1]= {1};
                                                                                // Numero di thread locali per workgroup - NON DEFINITO (non serve con task parallelism...in teoria!);

                                                                                // ESEGUO i threads;
                                                                                errCode = clEnqueueNDRangeKernel(this->coda,
                                                                                                                this->kernel,
                                                                                                                1,
                                                                                                                NULL,
                                                                                                                globalThreads,
                                                                                                                localThreads,
                                                                                                                0,
                                                                                                                NULL,
                                                                                                                &events[0]);


                                                                            Still no luck.

                                                                    • Cannot modify correctly __local uint4 vectors variables: serious compiler bug?
                                                                      Fr4nz

                                                                      Unfortunately no way nou. Same problem also if I use get_local_id(): only x and w components have correct values after add. Test kernel was this:

                                                                       

                                                                      __kernel
                                                                      void
                                                                      intersect(__global uint *outputCounts,
                                                                            __local uint4 *tmpBuffer) {

                                                                      size_t lid = get_local_id(0);

                                                                          vstore4(4,0,tmpBuffer+lid);
                                                                          tmpBuffer[lid] += tmpBuffer[lid];
                                                                         

                                                                          outputCounts[lid] = tmpBuffer[lid].w;