5 Replies Latest reply on Jan 22, 2010 7:31 AM by Fr4nz

    Getting too short chrono times from OpenCL profiling...

    Fr4nz

      Hi,

      I'm getting way too short chrono times when I profile the execution of my kernel (in the order of tenths of a second, when it should report seconds of execution).

      The profiling code is report below here...did I make any error in the code?

      I'm working under Ubuntu 9.04 32-bit, ATI 5770 and hotfix 9.12 drivers.

      cl_ulong startTime, endTime; cl_event chronoEvent; errCode = clEnqueueNDRangeKernel(this->coda, this->kernel, 1, NULL, &totWorkItems, &dimWorkGroup, 0, NULL, &chronoEvent); if(errCode != CL_SUCCESS) printf("\nFallimento esecuzione kernel! %d", errCode); // Punto di sincronizzazione: prima di procedere con la lettura dei risultati, aspetto che // i vari work-item abbiano finito di lavorare; clFinish(this->coda); // Computo del tempo di esecuzione su GPU; // TODO: Computa tempi troppo brevi, da fixare! errCode = clGetEventProfilingInfo(chronoEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime, NULL); if(errCode != CL_SUCCESS) printf("\nFallimento profiling! %d", errCode); errCode = clGetEventProfilingInfo(chronoEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime, NULL); if(errCode != CL_SUCCESS) printf("\nFallimento profiling! %d", errCode); double elapsedTime = (endTime-startTime); printf("\nTempo trascorso: %.3f", elapsedTime/1000000000); clReleaseEvent(chronoEvent);

        • Getting too short chrono times from OpenCL profiling...
          n0thing

          Your profiling code looks ok, how many global threads you are using?

            • Getting too short chrono times from OpenCL profiling...
              Fr4nz

               

              Originally posted by: n0thing Your profiling code looks ok, how many global threads you are using?

               

              Work-group size is 128 (but I'm experimenting various values), total number of work-items vary (we are in the order of thousands of work items). No suggestions from AMD? :-(

              I'm starting to think that OpenCL profiling is broken under Ubuntu 32-bit...

                • Getting too short chrono times from OpenCL profiling...
                  genaganna

                   

                  Originally posted by: Fr4nz

                   

                  Work-group size is 128 (but I'm experimenting various values), total number of work-items vary (we are in the order of thousands of work items). No suggestions from AMD? :-(

                   

                  I'm starting to think that OpenCL profiling is broken under Ubuntu 32-bit...

                   

                  Fr4nz,

                           Could you please post test case to reproduce this issue?

                    • Getting too short chrono times from OpenCL profiling...
                      Fr4nz

                       

                      Originally posted by: genaganna 

                      Fr4nz,

                      Could you please post test case to reproduce this issue?

                      The class which manages the GPU and executes the kernel is part of a large project...it would be useless to report all the code here. I'll past the code of the aforementioned class (and the kernel code), maybe you can spot if there's something wrong there. In any case measured execution times are always wrong.

                      More info about the code attached: the method in the host code which creates the queue with profiling enabled is "parseSpecsGPUAllocQueue()", while the one which exectues the kernel and measure its execution time is "executeGPUIntersect()".

                      My specs are reported in first post. Thank you for your help.

                       

                      // **** HOST CODE (header) ********* #ifndef GPUINTERFACE_H_ #define GPUINTERFACE_H_ // Inclusione necessaria della libreria OpenCL... #include "CL/cl.h" class GPUInterface { private: // Campi privati classe; /** * Struct raffigurante il contesto associato alla GPU; */ cl_context contesto; /** * Questo campo rappresenta la coda dei comandi associata al contesto; */ cl_command_queue coda; /** * Campo che rappresenta il codice sorgente del kernel da compilare; */ cl_program programma; /** * Questo campo rappresenta un kernel compilato; */ cl_kernel kernel; /** * Campo contenente la memoria massima allocabile per un oggetto sul dispositivo OpenCL; */ long maxGlobMem; /** * Campo contenente la dimensione delle memorie locali; ricordiamo che una memoria locale * e' condivisa da tutti i work-item di un work-group. */ long maxLocMem; /** * 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 k, int prefixCkSize, int numItems, int dimRow); 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; inoltre alloca la coda dei comandi associata alla GPU; */ void parseSpecsGPUAllocQueue(); /** * 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_ */ // ***** HOST CODE (CPP) ******** GPUInterface.cpp // 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 OpenCL operante * su UNA GPU OpenCL-compatibile. */ GPUInterface::GPUInterface() { // Inizializzo un contesto OpenCL; initContext(); // Recupero le caratteristiche salienti del dispositivo OpenCL da usare poi durante // l'esecuzione dei thread. Inoltre, alloco la CODA DEI COMANDI; parseSpecsGPUAllocQueue(); // Compilo il kernel da eseguire successivamente sul dispositivo OpenCL e creo l'oggetto // "kernel" che gestirà, durante l'esecuzione dei kernel, la configurazione dei parametri // della funzione che il kernel rappresenta; compileKernel(); } /** * Distruttore principale della classe; */ GPUInterface::~GPUInterface() { // Libero le risorse occupate dalle varie struct OpenCL; clReleaseKernel(this->kernel); clReleaseProgram(this->programma); clReleaseCommandQueue(this->coda); clReleaseContext(this->contesto); } // 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 k, int prefixCkSize, int numItems, int dimRow) { // Codice di errore tornato dalle varie funzioni; cl_int errCode; // **** Creazione dei BUFFER __global usati dai kernel in INPUT/OUTPUT **** // 1) Creazione buffer di input del VD; unsigned int dimVD = numItems*dimRow*sizeof(unsigned int); cl_mem inputVDBufferCL = clCreateBuffer(this->contesto, CL_MEM_READ_ONLY, (size_t) dimVD, NULL, &errCode); if(errCode != CL_SUCCESS) { printf("\nFallimento inizializzazione buffer VD! %d", errCode); exit(1); } // 2) Creazione buffer di input della tabella dei candidati; unsigned int dimPrefix = k*prefixCkSize*sizeof(unsigned int); cl_mem inputPrefixBufferCL = clCreateBuffer(this->contesto, CL_MEM_READ_ONLY, (size_t) dimPrefix, NULL, &errCode); if(errCode != CL_SUCCESS) { printf("\nFallimento inizializzazione buffer intersezioni! %d", errCode); exit(1); } // 3) Creazione buffer di output; unsigned int dimSupports = prefixCkSize*sizeof(unsigned int); cl_mem outputBufferCL = clCreateBuffer(this->contesto, CL_MEM_WRITE_ONLY, (size_t) dimSupports, NULL, &errCode); if(errCode != CL_SUCCESS) { printf("\nFallimento inizializzazione buffer dei supporti! %d", errCode); exit(1); } // **** FINE creazione buffer globali **** // **** INIZIO setup INDEX SPACE **** // E' estremamente importante utilizzare dimensioni consone per i work-group! // Cio' e' legato al modo in cui i thread vengono impacchettati implicitamente in blocchi // dalla GPU (warps); di conseguenza bisogna calcolare le dimensioni dei work-group tenendo // conto di questo e del numero totale di intersezioni che bisogna effettuare; // Numero totale delle intersezioni da effettuare; size_t numIntersections = prefixCkSize; // Dimensione di un work-group: e' importante usare almeno 64 thread x work-group, dato che // su HW Nvidia i warp hanno dim. 32, su ATI 64 ed e' quindi consigliabile usare un // multiplo di 64!); size_t dimWorkGroup = 128; // Numero di work-items in totale; dev'essere maggiore/uguale al numero di intersezioni; // NOTA: il valore impostato dev'essere divisibile per la dimensione di un workgroup! size_t totWorkItems; if(dimWorkGroup > numIntersections) totWorkItems = numIntersections + (dimWorkGroup-numIntersections); else { totWorkItems = numIntersections + (dimWorkGroup-(numIntersections%dimWorkGroup)); } // FINE setup INDEX space; // **** Inizializzazione dei parametri del kernel **** // 1) Input buffer VD; errCode = clSetKernelArg(this->kernel, 0, sizeof(cl_mem), (void *)&inputVDBufferCL); if(errCode != CL_SUCCESS) printf("Argomento kernel inizializzato male! %d", errCode); // 2) Input buffer candidati; errCode = clSetKernelArg(this->kernel, 1, sizeof(cl_mem), (void *)&inputPrefixBufferCL); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male! %d", errCode); // 3) Inizializzazione buffer temporaneo slice VD; il buffer serve a // memorizzare una "slice" del dataset; verra' allocata, pertanto, // una quantita' di memoria pari al numero di item contenuti nel VD per 16 bytes, // ovvero 32*4 transazioni per ogni item; questo in maniera tale da poter usare // tutte le unita' funzionali degli engine durante le AND-intersezioni; errCode = clSetKernelArg(this->kernel, 2, numItems*sizeof(cl_uint4), NULL); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male dada! %d", errCode); // 4) Inizializzazione buffer temporaneo locale per le intersezioni; ogni thread del // work-group dovra' avere un buffer per questo scopo di 16 bytes (uint4); errCode = clSetKernelArg(this->kernel, 3, dimWorkGroup*sizeof(cl_uint4), NULL); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male dada! %d", errCode); // 5) Inizializzazione buffer temporaneo locale per il contenimento degli item // di un candidato assegnato ad un work-item errCode = clSetKernelArg(this->kernel, 4, k*dimWorkGroup*sizeof(cl_uint), NULL); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male dada! %d", errCode); // 6) Inizializzazione buffer temporaneo locale per i conteggi; ogni thread del // work-group dovra' avere un buffer per questo scopo di 16 bytes (uint4); errCode = clSetKernelArg(this->kernel, 5, dimWorkGroup*sizeof(cl_uint), NULL); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male dada! %d", errCode); // 7) Output buffer supporti; errCode = clSetKernelArg(this->kernel, 6, sizeof(cl_mem), (void *)&outputBufferCL); if(errCode != CL_SUCCESS) printf("Argomento kernel inizializzato male! %d", errCode); // 8) Indice dell'iterazione corrente; errCode = clSetKernelArg(this->kernel, 7, sizeof(int), (void *)&k); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male! %d", errCode); // 9) Numero di item contenuti nel VD; errCode = clSetKernelArg(this->kernel, 8, sizeof(int), (void *)&numItems); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male! %d", errCode); // 10) Dimensione di una riga di un item nell'array del vertical dataset; errCode = clSetKernelArg(this->kernel, 9, sizeof(int), (void *)&dimRow); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male! %d", errCode); // 11) Numero di intersezioni da eseguire; errCode = clSetKernelArg(this->kernel, 10, sizeof(int), (void *)&prefixCkSize); if(errCode != CL_SUCCESS) printf("\nArgomento kernel inizializzato male! %d", errCode); // **** Fine inizializzazione dei parametri del kernel **** // INIZIO lettura dei buffer di input dalla memoria host alla memoria globale del dispositivo; // 1) Lettura del VD; errCode = clEnqueueWriteBuffer(this->coda, inputVDBufferCL, CL_TRUE, 0, (size_t) dimVD, VD, 0, NULL, NULL); if(errCode != CL_SUCCESS) printf("\nErrore scrittura buffer dispositivo OpenCL! %d", errCode); // 2) Lettura del vettore contenente i candidati; errCode = clEnqueueWriteBuffer(this->coda, inputPrefixBufferCL, CL_TRUE, 0, (size_t) dimPrefix, prefixCkGPU, 0, NULL, NULL); if(errCode != CL_SUCCESS) printf("\nErrore scrittura buffer dispositivo OpenCL! %d", errCode); // **** Fine inizializzazione dei parametri del kernel **** // FINE lettura dei buffer di input; // Punto di sincronizzazione: attendo che la lettura sia stata effettuata per intero; clFinish(this->coda); // ESEGUO i threads; printf("\n\n**** Esecuzione %d_a iterazione su GPU ****", k); printf("\nNumero di itnersezioni da eseguire: %d", numIntersections); printf("\nWork items totali allocati: %d", totWorkItems); cl_event chronoEvent; errCode = clEnqueueNDRangeKernel(this->coda, this->kernel, 1, NULL, &totWorkItems, &dimWorkGroup, 0, NULL, &chronoEvent); if(errCode != CL_SUCCESS) printf("\nFallimento esecuzione kernel! %d", errCode); // Punto di sincronizzazione: prima di procedere con la lettura dei risultati, aspetto che // i vari work-item abbiano finito di lavorare; clFinish(this->coda); // Computo del tempo di esecuzione su GPU; cl_ulong startTime, endTime; errCode = clGetEventProfilingInfo(chronoEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime, NULL); if(errCode != CL_SUCCESS) printf("\nFallimento profiling! %d", errCode); errCode = clGetEventProfilingInfo(chronoEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime, NULL); if(errCode != CL_SUCCESS) printf("\nFallimento profiling! %d", errCode); // printf("\nDEBUG TIME: %ld - %ld", (unsigned long) startTime, (unsigned long) endTime); double elapsedTime = (endTime-startTime); printf("\nTempo trascorso: %.3f sec.", (elapsedTime/1000000000)); // Rilascio l'evento usato per cronometrare il tempo di esecuzione; clReleaseEvent(chronoEvent); // Leggo i risultati dall'output buffer; errCode = clEnqueueReadBuffer(this->coda, outputBufferCL, CL_TRUE, 0, (size_t) prefixCkSize*sizeof(cl_uint), support, 0, NULL, NULL); if(errCode != CL_SUCCESS) printf("\nErrore durante la lettura dei risultati! %d", errCode); // Punto di sincronizzazione: aspetto che la coda finisca tutte le operazioni ancora // eventualmente contenute al suo interno; clFinish(this->coda); // FINE comandi inviati alla GPU // Stampa di debug; // printf("\n\nDEBUG KERNEL %da iterazione: %d - %u", k, prefixCkSize, support[0]); // Release dei buffer input/output creati in precedenza; clReleaseMemObject(inputVDBufferCL); clReleaseMemObject(inputPrefixBufferCL); clReleaseMemObject(outputBufferCL); } // 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; // Numero di entry (piattaforme/dispositivi OpenCL) massime da memorizzare; cl_uint numMaxEntries = 1; // Struttura contenente la lista delle piattaforme disponibili; cl_platform_id platforms[numMaxEntries]; // Numero di piattaforme disponibili; cl_uint numPlatforms; // Recupero informazioni sulla piattaforma in cui viene eseguito il programma... errorCode = clGetPlatformIDs(numMaxEntries, platforms, &numPlatforms); if (errorCode != CL_SUCCESS) { printf("Non sono riuscito a ritrovare piattaforme..."); exit(1); } // printf("\nNumero di piattaforme: %d", numPlatforms); // Lista che conterra' i dispositivi ritrovati sulla piattaforma; cl_device_id devices[numMaxEntries]; // Numero di cl_uint numDevices; // Recupero il primo dispositivo GPU OpenCL-compatibile presente nella piattaforma; clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, numMaxEntries, devices, &numDevices); // Se sono su una piattaforma con GPU che non supportano OpenCL... if(numDevices == 0) { printf("\nNessuna GPU OpenCL-capable trovata, esco..."); exit(1); } printf("\nGPU OpenCL-compatibile trovata!"); // 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, // Flag; (cl_context_properties) platforms[0], // Specifica la piattaforma da utilizzare; // Zero finale; 0 }; // Creo un contesto sul dispositivo OpenCL-capable per l'esecuzione dei kernel this->contesto = clCreateContext(cps,numMaxEntries,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; // Prima passata: recupero la dimensione della lista dei dispositivi associata al contesto aperto; errCode = clGetContextInfo(contesto, CL_CONTEXT_DEVICES, 0, NULL, &numDevices); if(errCode != CL_SUCCESS) exit(1); // printf("\nNumero di dispositivi associati al contesto: %d", numDevices/sizeof(size_t)); // Seconda passata: recupero la lista dei dispositivi; cl_device_id listaDevices[1*(sizeof(size_t))]; // Mi faccio passare solo il primo dispositivo... errCode = clGetContextInfo(contesto, CL_CONTEXT_DEVICES, 1*sizeof(size_t), 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; // Display delle caratteristiche della GPU; printf("\n\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 L'esecuzione dei comandi avviene IN-ORDER (per cui verrà rispettato // l'ordine in cui vengono inviati i comandi alla coda!); this->coda = clCreateCommandQueue(this->contesto, listaDevices[0], CL_QUEUE_PROFILING_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"; const char* nomeFuncKernel = "intersect"; 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; 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); int numEntries = 1; size_t numDevices = numEntries*sizeof(size_t); // Prima recupero la dimensione di allocazione della lista di dispositivi; errCode = clGetContextInfo(this->contesto, CL_CONTEXT_DEVICES, 0, NULL, &numDevices); cl_device_id listaDevices[1]; // Poi recupero la lista; errCode = clGetContextInfo(this->contesto, CL_CONTEXT_DEVICES, numDevices, listaDevices, NULL); if(errCode != CL_SUCCESS) { printf("nFallimento associazione sorgente->contesto! Esco..."); exit(1); } // Compilazione del programma creato prima; errCode = clBuildProgram(this->programma, numEntries, 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 durante le iterazioni di DCI... this->kernel = clCreateKernel(this->programma, nomeFuncKernel, &errCode); if(errCode != CL_SUCCESS) { printf("\nFallimento creazione oggetto kernel! Esco..."); exit(1); } } // KERNEL CODE /** * Kernel OpenCL che esegue le intersezioni relative ad un insieme di candidati all'intersezione * k-esima. * * INPUT: * 1 - Puntatore al vettore contenente il vertical dataset (input buffer); * 2 - Puntatore al vettore contenente la lista di candidati (input buffer); * 3 - Puntatore alla porzione di memoria locale (shared) dedicata alla memorizzazione dello slice corrente del VD; * 4 - Puntatore alla porzione di memoria locale dedicata alla memorizzazione delle intersezioni temporanee relativa alla slice corrente; * 5 - Puntatore alla porzione di memoria locale contenente gli item costituenti di un candidato; * 6 - Puntatore alla porzione di memoria locale dedicata alla memorizzazione del conteggio di un candidato delegato ad uno specifico work-item; * 7 - Puntatore al vettore che conterrà i risultati finali dell'iterazione; * 8 - Indice dell'iterazione corrente; * 9 - Numero di item costituenti il VD; * 10 - Dimensione di una riga di item; * 11 - Numero totale di intersezioni da effettuare; * * EFFETTI: * Il vettore outputCounts verra' inizializzato con i supporti calcolati dei candidati; */ #define WAVEFRONT 64 // Dimensione di un warp sull'hardware ATI; #define BANKS 32 // Numero di banchi costituenti un'unita' di memoria locale sull'hw ATI; #define SIZEVECINT 4 // Banchi occupati da un uint4; __kernel void intersect(__global uint4 *inputVD, __global uint *inputPrefixCk, __local uint4 *tmpBuffer, __local uint4 *tmpIntersect, __local uint *tmpCandItems, __local uint *tmpCOUNTS, __global uint *outputCounts, const int k, const int numItems, const int dimRow, const int numIntersections) { // Info sulle proprieta' del work-group locale e dei work-item singoli; const int gid = get_global_id(0); const int lid = get_local_id(0); const int lsize = get_local_size(0); // Dimensione di una riga risistemata per uint4; const int dimRowVec4 = dimRow/4; // Maschera usata per contare il numero di bit "1" nel risultato dell'intersezione; const uint4 MASK = 1; // Inizializzazione variabile locale usata per i conteggi; tmpCOUNTS[lid] = 0; // Numero di thread che possono scrivere degli uint4 contemporaneamente su un dato numero di banchi // di memoria shared; const int modulo = BANKS/SIZEVECINT; // Copia in memoria locale degli item costituenti il candidato assegnato al thread corrente; int init = (gid < numIntersections) ? 0 : k; for(init; init<k; init++) // lid*lsize: serve affinche' un thread scriva sempre sul suo banco di riferimento senza // andare a scrivere in contemporanea su banchi acceduti da altri thread; tmpCandItems[lid + init*lsize] = inputPrefixCk[gid*k + init]; // Variabile private che serve a recuperare l'item correntemente considerato da prefixCk; int item; for(int i=0; i<dimRowVec4; i++) { // Inizializzazione slice corrente considerata dal VD; // Ogni thread scrive 16 bytes su 4 blocchi consecutivi: in totale abbiamo 8 thread per work-group che // scrivono, pertanto, tutti e 32 i blocchi della shared memory contemporaneamente senza conflitti; init = (lid < modulo) ? lid : numItems; for(int init=lid; init<numItems; init+=modulo) tmpBuffer[init] = inputVD[init*dimRowVec4 + i]; // Barriera: mi assicuro che la memoria locale sia stata inizializzata correttamente; barrier(CLK_LOCAL_MEM_FENCE); // Inizio computo k-intersezione sulla slice corrente; if(gid < numIntersections) { // Inizializzo il buffer delle intersezioni temporanee col primo item; tmpIntersect[lid] = tmpBuffer[tmpCandItems[lid]]; // Computo la k-intersezione sullo slice corrente; for(int j=1; j<k; j++) tmpIntersect[lid] &= tmpBuffer[tmpCandItems[lid + j*lsize]]; // Conteggio del supporto dello slice corrente; for(int j=0; j<32; j++) { tmpCOUNTS[lid] += (tmpIntersect[lid].x & MASK.x) + (tmpIntersect[lid].y & MASK.y) + (tmpIntersect[lid].z & MASK.z) + (tmpIntersect[lid].w & MASK.w); // Attenzione: lo shift va scritto così qui dentro... tmpIntersect[lid] = tmpIntersect[lid] >> MASK; } } // Fine computo k-intersezione slice; } // Fine ciclo slice corrente; // Barriera comune a tutti i thread che dovrebbe assicurare la corretta scrittura del buffer di output // (alcuni thread potrebbero rimanere indietro a causa dell'esecuzione in "warps"); // barrier(CLK_LOCAL_MEM_FENCE); if(gid < numIntersections) outputCounts[gid] = tmpCOUNTS[lid]; }

                        • Getting too short chrono times from OpenCL profiling...
                          Fr4nz

                          Yesterday I've replaced my Ubunut 9.04 32bit with the 64bit version: no luck, profiling times are still wrong (too short, it always gives 0.1-0.2 secs.).

                          Anyway, I've noticed (with my eyes, obviously) a good speedup in the kernel execution times, and this confirms that the 64-bit OpenCL ATI implementation is better than the 32-bit one.

                          Could someone help me with this profiling problem? It seems that the timer resolution (at least on 5770) is not 1ns...

                           

                          EDIT: Okay, actually profiling timing is correct, the GPU is so fast that outshines another part of the algorithm which I wronlgy excluded from the timing. Really amazing.