{"id":2740,"date":"2018-01-25T16:53:59","date_gmt":"2018-01-25T16:53:59","guid":{"rendered":"http:\/\/www.colorwave.it\/blog\/?p=2740"},"modified":"2018-01-28T10:14:11","modified_gmt":"2018-01-28T10:14:11","slug":"gpu_kernel","status":"publish","type":"post","link":"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_kernel\/admin\/2018\/01\/","title":{"rendered":"gpu_kernel"},"content":{"rendered":"<p><strong>Scopo della classe<\/strong>: gpu_kernel \u00e8 stata progettata per facilitare la preparazione ed esecuzione dei kernel OpenCL, nascondendo tutti i dettagli\u00a0 delle operazioni necessarie, la cui scrittura complica notevolmente la stesura del codice: per eseguire una funzione kernel \u00e8 necessario innanzitutto individuare le piattaforme OpenCL presenti nell&#8217;hardware; quindi definire i context, richiamare un programma C OpenCL\u00a0 (sorgente o precompilato) contenente la funzione kernel, eventualmente compilarlo, quindi <a href=\"https:\/\/www.khronos.org\/registry\/OpenCL\/sdk\/1.0\/docs\/man\/xhtml\/clCreateKernel.html\" target=\"_blank\" rel=\"noopener\">creare un oggetto kernel<\/a>.<\/p>\n<p>Dopo questi passaggi sar\u00e0 necessario, per eseguire il kernel, definire gli argomenti della funzione, in genere oggetti allocati nel context OpenCL, infine Lanciare il kernel ed estrarre gli oggetti calcolati (oggetti di memoria e\/o immagini 2d o 3d.<\/p>\n<p>Di seguito \u00e8 riportato un diagramma della struttura di un&#8217;applicazione OpenCL:<img fetchpriority=\"high\" decoding=\"async\" data-attachment-id=\"2822\" data-permalink=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_kernel\/admin\/2018\/01\/attachment\/opencl_program_struct-2\/\" data-orig-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?fit=614%2C519&amp;ssl=1\" data-orig-size=\"614,519\" data-comments-opened=\"1\" data-image-meta=\"{&quot;aperture&quot;:&quot;0&quot;,&quot;credit&quot;:&quot;&quot;,&quot;camera&quot;:&quot;&quot;,&quot;caption&quot;:&quot;&quot;,&quot;created_timestamp&quot;:&quot;0&quot;,&quot;copyright&quot;:&quot;&quot;,&quot;focal_length&quot;:&quot;0&quot;,&quot;iso&quot;:&quot;0&quot;,&quot;shutter_speed&quot;:&quot;0&quot;,&quot;title&quot;:&quot;&quot;,&quot;orientation&quot;:&quot;0&quot;}\" data-image-title=\"opencl_program_struct\" data-image-description=\"\" data-image-caption=\"\" data-medium-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?fit=300%2C254&amp;ssl=1\" data-large-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?fit=614%2C519&amp;ssl=1\" tabindex=\"0\" role=\"button\" class=\"aligncenter wp-image-2822\" src=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?resize=463%2C392&#038;ssl=1\" alt=\"\" width=\"463\" height=\"392\" srcset=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?w=614&amp;ssl=1 614w, https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/opencl_program_struct-1.png?resize=300%2C254&amp;ssl=1 300w\" sizes=\"(max-width: 463px) 100vw, 463px\" data-recalc-dims=\"1\" \/><\/p>\n<p>Come si pu\u00f2 dedurre dall&#8217;analisi del diagramma UML,\u00a0 l&#8217;esecuzione di un kernel richiede la definizione di diversi oggetti, allocati utilizzando i comandi delle API OpenCL .<\/p>\n<p>La creazione ed esecuzione di un kernel richiede infatti la creazione di un programma e degli argomenti necessari; il diagramma che segue schematizza il metodo di creazione dello stesso<img decoding=\"async\" data-attachment-id=\"2823\" data-permalink=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_kernel\/admin\/2018\/01\/attachment\/creatingopenclprogram-2\/\" data-orig-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?fit=736%2C527&amp;ssl=1\" data-orig-size=\"736,527\" data-comments-opened=\"1\" data-image-meta=\"{&quot;aperture&quot;:&quot;0&quot;,&quot;credit&quot;:&quot;&quot;,&quot;camera&quot;:&quot;&quot;,&quot;caption&quot;:&quot;&quot;,&quot;created_timestamp&quot;:&quot;0&quot;,&quot;copyright&quot;:&quot;&quot;,&quot;focal_length&quot;:&quot;0&quot;,&quot;iso&quot;:&quot;0&quot;,&quot;shutter_speed&quot;:&quot;0&quot;,&quot;title&quot;:&quot;&quot;,&quot;orientation&quot;:&quot;0&quot;}\" data-image-title=\"CreatingOpenCLProgram\" data-image-description=\"\" data-image-caption=\"\" data-medium-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?fit=300%2C215&amp;ssl=1\" data-large-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?fit=640%2C458&amp;ssl=1\" tabindex=\"0\" role=\"button\" class=\"aligncenter wp-image-2823\" src=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?resize=468%2C335&#038;ssl=1\" alt=\"\" width=\"468\" height=\"335\" srcset=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?w=736&amp;ssl=1 736w, https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2018\/01\/CreatingOpenCLProgram-1.jpg?resize=300%2C215&amp;ssl=1 300w\" sizes=\"(max-width: 468px) 100vw, 468px\" data-recalc-dims=\"1\" \/><\/p>\n<!--nextpage-->\n<p>La lettura di questo post richiede una buona conoscenza dei paradigmi sottesi alle API di OpenCL;\u00a0 Per accedere a ulteriori informazioni si propongono questi articoli :<\/p>\n<ul>\n<li><a href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/siti-opencl-interessanti\/admin\/2017\/12\/\">una lista di siti OpenCL<\/a><\/li>\n<li><a href=\"https:\/\/www.colorwave.it\/blog\/uncategorized\/kernel-e-work-item\/admin\/2017\/12\/\">Kernel e WorkItem<\/a>;<\/li>\n<li><a href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/terminologia-opencl\/admin\/2017\/12\/\">Terminologia OpenCL\u00a0<\/a><\/li>\n<\/ul>\n<p>Descriveremo ora i metodi e le funzioni della classe :<\/p>\n<pre class=\"lang:default decode:true\" title=\"La classe gpu_kernel\">class AFX_EXT_CLASS gpu_kernel\r\n{\r\npublic:\r\n\tgpu_kernel();\r\n\t~gpu_kernel();\r\n\r\n\tchar *Name();\r\n\r\n\tvoid Set_id(cl_kernel  id);\r\n\tcl_kernel Get_id();\r\n\tgpu_program *Program();\r\n\tbool SetIndexSpace(size_t *index_spize_size, size_t *work_group_size = 0);\r\n\tbool Set_index_space_2d(size_t dx, size_t dy);\r\n\tbool FlushMemObjects();\r\n\tbool PushArg(gpu_mem_object *memobj);\r\n\tbool PushArg(cl_sampler sampler);\r\n\tbool PushArg(cl_int value);\r\n\tbool PushArg(cl_uint value);\r\n\tbool PushArg(cl_float value);\r\n\tbool PushArg(void *ptr, size_t size);\r\n\tbool PopArg(gpu_mem_object *buf);\r\n\tbool Exec();\r\n\tsize_t Get_preferred_multiple_workgroup_size();\r\n\tcl_ulong Get_local_mem_size();\r\n\tsize_t Get_workgroup_size();\r\n\tsize_t *Get_compile_workgroup_size(size_t compile_work_group_size[3]);\r\n\r\n\tfriend class gpu_program;\r\n\tfriend class gpu_buffer;\r\n\tfriend class GPU;\r\n\r\nprivate:\r\n\tbool SetName(char *name);\r\n\tbool SetProgram(gpu_program *prg);\r\n\tbool SetId(cl_kernel kid);\r\n\r\n\t__use_private_struct\r\n\r\n};\r\n<\/pre>\n<h4>Ottenere un&#8217;istanza gpu_kernel<\/h4>\n<p>Le istanze della classe la classe gpu_kernel sono create in due diversi modi;<\/p>\n<p>un modo, completamente automatico, \u00e8 gi\u00e0 stato presentato in <span style=\"color: #0000ff;\"><a style=\"color: #0000ff;\" href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_context\/admin\/2017\/11\/\" target=\"_blank\" rel=\"noopener\"> questo articolo di descrizione della classe gpu_context;\u00a0<\/a><\/span>Il comando\u00a0<span style=\"color: #339966;\"><em>gpu_kernel *kernel = OPCL::GPU::OpenKernel(prog_name, kernel_name)\u00a0<\/em><\/span>rimanda un kernel dati i nomi del programma e della funzione __kernel contenuta nel codice C OpenCl;<\/p>\n<p>un secondo modo \u00e8 quello di richiamare la funzione\u00a0OpenKernel\u00a0 di un&#8217;istanza <a href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_program\/admin\/2017\/12\/\" target=\"_blank\" rel=\"noopener\">gpu_program\u00a0<\/a>\u00a0(questo metodo \u00e8 comunque usato internamente dalla GPU::OpenKernel).<\/p>\n<p>Il primo metodo \u00e8 semplice da usare, come descritto nell&#8217;articolo citato; Il secondo metodo richiede un numero maggiore di passaggi ed \u00e8 utile solo in casi particolari dei quali scriveremo in seguito;<\/p>\n<!--nextpage-->\n<p>L&#8217;inserimento degli argomenti nel kernel \u00e8 fatta usando il gruppo di metodi gpu_kernel::PushArg(), ciascuno dei quali trasferisce nella coda degli argomenti tipologie di informazioni differenti : uno dei metodi usa le istanze derivate dalle classi\u00a0gpu_mem_object; un altro trasferisce una istanza di una struttura OpenCL, <a href=\"https:\/\/www.khronos.org\/registry\/OpenCL\/sdk\/1.0\/docs\/man\/xhtml\/sampler_t.html\" target=\"_blank\" rel=\"noopener\">sampler t<\/a>\u00a0, usata per definire le modalit\u00e0 di accesso ad una mappa immagine. Altri metodi PushArg permettodo di passare al kernel\u00a0 valori int, float,\u00a0 o buffer di byte;<\/p>\n<h4>I metodi PushArg()<\/h4>\n<p>Una volta ottenuta l&#8217;istanza gpu_kernel, per lanciarla occorrer\u00e0\u00a0 definire gli argomenti della funzione; allo scopo si usano i metodi PushArg() fornite dalla classe; Questo gruppo di funzioni <a href=\"https:\/\/www.khronos.org\/registry\/OpenCL\/sdk\/1.0\/docs\/man\/xhtml\/clSetKernelArg.html\" target=\"_blank\" rel=\"noopener\">inseriscono gli argomenti passati nella coda di comando assegnata al kernel<\/a>, per l&#8217;esecuzione successiva; I metodi PushArg incorporano un gruppo di chiamate alle API OpenCL piuttosto complesso: il loro uso semplifica notevolmente la\u00a0 vita del programmatore..<\/p>\n<p><img decoding=\"async\" data-attachment-id=\"283\" data-permalink=\"https:\/\/www.colorwave.it\/blog\/framework\/linterfaccia-operation_call\/admin\/2015\/05\/attachment\/fumetto\/\" data-orig-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" data-orig-size=\"24,24\" data-comments-opened=\"1\" data-image-meta=\"{&quot;aperture&quot;:&quot;0&quot;,&quot;credit&quot;:&quot;&quot;,&quot;camera&quot;:&quot;&quot;,&quot;caption&quot;:&quot;&quot;,&quot;created_timestamp&quot;:&quot;0&quot;,&quot;copyright&quot;:&quot;&quot;,&quot;focal_length&quot;:&quot;0&quot;,&quot;iso&quot;:&quot;0&quot;,&quot;shutter_speed&quot;:&quot;0&quot;,&quot;title&quot;:&quot;&quot;,&quot;orientation&quot;:&quot;0&quot;}\" data-image-title=\"fumetto\" data-image-description=\"\" data-image-caption=\"\" data-medium-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" data-large-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" tabindex=\"0\" role=\"button\" class=\"size-full wp-image-283 alignleft\" src=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?resize=24%2C24&#038;ssl=1\" alt=\"\" width=\"24\" height=\"24\" data-recalc-dims=\"1\" \/><\/p>\n<p>Questi metodi in genere non sono usati direttamente, ma\u00a0 gestiti internamente dalle classi derivate da gpu_mem_object, per costruire i metodi virtuali PushKernel(): leggere <a href=\"https:\/\/www.colorwave.it\/blog\/framework\/gpu_mem_object\/admin\/2017\/11\/\" target=\"_blank\" rel=\"noopener\">questo articolo per ulteriori chiarimenti ; <\/a>\u00a0Un esempio di questa modalit\u00e0 \u00e8 fornito nel frammento di codice che segue, e che \u00e8 stato usato anche <span style=\"color: #0000ff;\">in <a style=\"color: #0000ff;\" href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_context\/admin\/2017\/11\/\" target=\"_blank\" rel=\"noopener\">questo articolo di presentazione dei gpu_context<\/a><\/span><\/p>\n<pre class=\"lang:default decode:true\" title=\"Inserimento di argomenti del kernel\">\/\/ immagine sorgente\r\n\t\t\tgpu_image *src = new gpu_image(kernel, img_src, CL_MEM_READ_ONLY);\r\n\t\t\tsrc-&gt;Push_kernel();\r\n \r\n\t\t\t\/\/ immagine trasferita in forma di matrice..\r\n\t\t\tgpu_matrix_2d *dest = gpu_matrix_2d::New(kernel, img_dest, CL_MEM_WRITE_ONLY);\r\n\t\t\tdest-&gt;Push_kernel();\r\n \r\n\t\t\tgpu_screener *screener = gpu_screener::New(kernel, DTH);\r\n\t\t\tscreener-&gt;Push_kernel();\r\n \r\n\t\t\t\/\/ lut dei colori assegnati per dtv\r\n\t\t\tOPCL::gpu_rgb_lut *lut = OPCL::gpu_rgb_lut::New(kernel, MAXDTV + 1, CL_MEM_READ_ONLY, _$str-&gt;rgb);\r\n\t\t\tlut-&gt;Push_kernel();<\/pre>\n<p>Le classi<span style=\"color: #800000;\"><em> gpu_image, gpu_matrix_2d, gpu_screener e gpu_rgb_lut\u00a0<\/em><span style=\"color: #000000;\"><span style=\"color: #800000;\"> <span style=\"color: #000000;\">sono derivate da gpu_mem_object e come tali contengono la riscrittura della funzione virtuale gpu_mem_object::PushKernel. Questa funzione permette di scrivere sulla lista degli argomenti del kernel senza dovere indicarlo, poich\u00e8 il costrutture delle classi aveva gi\u00e0 richiesto questa informazione.\u00a0 In seguito mostreremo come costruire classe derivata da gpu_mem_object, chiarendo le modalit\u00e0 di riscrittura del metodo gpu_mem_object::PushKernel().<\/span><\/span><\/span><\/span><\/p>\n<p><!--nextpage-->Alcuni degli argomenti passati al kernel saranno elaborati durante l&#8217;esecuzione dello stesso; Queste informazioni potranno essere estratte dal programma host utilizzando il metodo gpu_kernel::PopArg. Il frammento di codice che segue mostra il codice per l&#8217;esecuzione del kernel e quello di estrazione dei risultati:<\/p>\n<pre class=\"lang:default decode:true\" title=\"Esecuzione del kernel ed estrazione dei dati calcolati\">if (kernel-&gt;Set_index_space_2d(ret_area-&gt;dx, ret_area-&gt;dy))\r\n\t\t\t{\r\n\t\t\t\tif (kernel-&gt;Exec())\r\n\t\t\t\t{\r\n\t\t\t\t\tbool res = kernel-&gt;PopArg(dest);\r\n\t\t\t\t\tMSGS name;\r\n\t\t\t\t\tSPRINTF(name, L\"%s_dth\", img_src-&gt;GetName());\r\n\t\t\t\t\timg_dest-&gt;InitDocInfo(name);\r\n\t\t\t\t\timg_dest-&gt;SetHtmlToolTip();\r\n\t\t\t\t}\r\n\t\t\t\telse\r\n\t\t\t\t{\r\n\t\t\t\t\tOPCL::GPU::Display_error_info();\r\n\t\t\t\t\t__delete (img_dest);\r\n\t\t\t\t}<\/pre>\n<p>Del metodo\u00a0Set_index_space_2d(..) contenuto nella prima riga scriveremo fra poco; per ora ricordiamo che definisce l&#8217;intervallo di indici usato dalla gpu per eseguire i singoli <a href=\"https:\/\/www.colorwave.it\/blog\/uncategorized\/kernel-e-work-item\/admin\/2017\/12\/\" target=\"_blank\" rel=\"noopener\">work_item<\/a>;<\/p>\n<p>Il metodo gpu_kernel::Exec() lancia in esecuzione il kernel e rimanda un valore <span style=\"color: #808080;\">true\u00a0<span style=\"color: #000000;\">se il codice \u00e8 stato eseguito correttamente; il passaggio successivo \u00e8 l&#8217;estrazione delle informazioni calcolata attraverso la gpu_kernel::PopArg():<\/span><\/span><\/p>\n<pre class=\"lang:default decode:true \" title=\"estrazione dati dal kernel\">bool res = kernel-&gt;PopArg(dest);<\/pre>\n<h4>PopArg()<\/h4>\n<p>La funzione riceve un puntatore ad una istanza di tipo gpu_mem_object e rimanda true se le informazioni sono state estratte correttamente; in questo esempio l&#8217;argomento <em>dest<\/em> passato \u00e8 lo stesso che in precedenza era stato trasferito come argomento al kernel:<\/p>\n<pre class=\"lang:default decode:true\" title=\"oggetto trasferito al kernel per essere calcolato\">gpu_matrix_2d *dest = gpu_matrix_2d::New(kernel, img_dest, CL_MEM_WRITE_ONLY);\r\n\t\t\tdest-&gt;Push_kernel();<\/pre>\n<p><img decoding=\"async\" data-attachment-id=\"283\" data-permalink=\"https:\/\/www.colorwave.it\/blog\/framework\/linterfaccia-operation_call\/admin\/2015\/05\/attachment\/fumetto\/\" data-orig-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" data-orig-size=\"24,24\" data-comments-opened=\"1\" data-image-meta=\"{&quot;aperture&quot;:&quot;0&quot;,&quot;credit&quot;:&quot;&quot;,&quot;camera&quot;:&quot;&quot;,&quot;caption&quot;:&quot;&quot;,&quot;created_timestamp&quot;:&quot;0&quot;,&quot;copyright&quot;:&quot;&quot;,&quot;focal_length&quot;:&quot;0&quot;,&quot;iso&quot;:&quot;0&quot;,&quot;shutter_speed&quot;:&quot;0&quot;,&quot;title&quot;:&quot;&quot;,&quot;orientation&quot;:&quot;0&quot;}\" data-image-title=\"fumetto\" data-image-description=\"\" data-image-caption=\"\" data-medium-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" data-large-file=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?fit=24%2C24&amp;ssl=1\" tabindex=\"0\" role=\"button\" class=\"size-full wp-image-283 alignleft\" src=\"https:\/\/i0.wp.com\/www.colorwave.it\/blog\/wp-content\/uploads\/2015\/05\/fumetto.png?resize=24%2C24&#038;ssl=1\" alt=\"\" width=\"24\" height=\"24\" data-recalc-dims=\"1\" \/><\/p>\n<p>Da quanto scritto si pu\u00f2 dedurre che le informazioni elaborate ed estratte dal kernel usando la PopArg sono necessariamente del tipo gpu_mem_object; inoltre va ricordato che queste istanze devono essere state create con\u00a0 il flag\u00a0<a href=\"https:\/\/www.colorwave.it\/blog\/framework\/gpu_mem_object\/admin\/2017\/11\/\" target=\"_blank\" rel=\"noopener\"><em><span style=\"color: #800000;\">CL_MEM_WRITE_ONLY<\/span> <\/em>o <\/a><span style=\"color: #800000;\"><em>CL_MEM_READ_WRITE\u00a0 <\/em><span style=\"color: #000000;\">definiti mediante le funzioni <a href=\"https:\/\/www.colorwave.it\/blog\/framework\/gpu_mem_object\/admin\/2017\/11\/\">gpu_mem_object::Set_mem_flags()<\/a><\/span><\/span><\/p>\n<p><!--nextpage-->Nell&#8217;esempio che stiamo trattando, il kernel ha il compito di ricevere due immagini: la prima una\u00a0gpu_image (sorgente) e la seconda, una\u00a0 gpu_matrix_2d, destinazione delle elaborazioni; La prima cosa da osservare (ma \u00e8 banale..) \u00e8 che l&#8217;immagine trasformata \u00e8 un&#8217;istanza di tipo diverso dal sorgente; ovviamente questa situazione dipende da ragioni interne al programma che si sta costruendo: la destinazione potrebbe essere un istogramma, una tabella dei colori presenti nel sorgente, o qualunque altra cosa. In questo caso la scelta \u00e8 stata quella di usare come destinazione una istanza di gpu_matrix_2d per calcolare un&#8217;immagine risultante, inserendola in una matrice bidimensionale; le classi gpu_image e gpu_matrix_2d hanno un costruttore che permette di inizializzare le stesse a partire da istanze della classe IMAGE di Colibri (le\u00a0 img_src e\u00a0 img_dest nel codice sopra riportato); entrambe le istanze src e des costruite hanno il compito di trasferire le mappe IMAGE fra CPU-&gt;GPU-&gt;CPU;<\/p>\n<p><em><span style=\"color: #800000;\">Poich\u00e8 la\u00a0<strong>gpu_image *src<\/strong> \u00e8 CL_MEM_READ_ONLY, la mappa associata sar\u00e0 automaticamente trasferita alla GPU; La\u00a0<strong>gpu_matrix_2d *dest\u00a0<\/strong> invece, essendo definita come\u00a0<\/span><span style=\"color: #800000;\"><em><a style=\"color: #800000;\" href=\"https:\/\/www.colorwave.it\/blog\/framework\/gpu_mem_object\/admin\/2017\/11\/\" target=\"_blank\" rel=\"noopener\">CL_MEM_WRITE_ONLY<\/a>,\u00a0 genera\u00a0 una mappa sulla GPU ma non trasferir\u00e0 informazioni<\/em>;<\/span><\/em><\/p>\n<p>la <span style=\"color: #800000;\">img_dest<\/span> passata al costruttore \u00e8 un&#8217;immagine bidimensionale definita dal tipo (RGB,Lab,BGRO..ecc) e dalle dimensioni; la\u00a0gpu_matrix_2d cos\u00ec inizializzata genera una mappa sulla GPU che ha lo stesso tipo di pixel, destinata ad accogliere il risultato dell&#8217;elaborazione; ad esecuzione effettuata, il comando kernel-&gt;PopArg(dest) effettua automaticamente il trasferimendo dei dati calcolati dalla GPU alla istanza <span style=\"color: #800000;\">img_dest<\/span> incapsulata nella\u00a0<span style=\"color: #800000;\">gpu_matrix_2d dest<\/span>;\u00a0 risultato finale: l&#8217;oggetto <span style=\"color: #993300;\">img_dest<\/span> conterr\u00e0 l&#8217;immagine trasformata della sorgente IMAGE <span style=\"color: #993300;\">src<\/span>;<\/p>\n<p>Questo non breve preambolo ha lo scopo di introdurre l&#8217;uso del concetto <a href=\"https:\/\/www.khronos.org\/registry\/OpenCL\/sdk\/2.1\/docs\/man\/xhtml\/ndrange.html\">NDRange\u00a0 di OpenCL<\/a>.<\/p>\n<p>Abbiamo gi\u00e0 spiegato nell&#8217;articolo <a href=\"https:\/\/www.colorwave.it\/blog\/uncategorized\/kernel-e-work-item\/admin\/2017\/12\/\">Kernel e Work-Item<\/a>\u00a0che, nel paradigma OpenCL, un\u00a0<strong><em>kernel<\/em>\u00a0<\/strong>\u00e8 una funzione in C-OpenCL (entry point di un programma) che \u00e8 eseguito su una\u00a0<em><a href=\"https:\/\/www.khronos.org\/assets\/uploads\/developers\/library\/2012-pan-pacific-road-show-June\/OpenCL-Details-Taiwan_June-2012.pdf\" target=\"_blank\" rel=\"noopener\">compute-unit<\/a>\u00a0<\/em>in pi\u00f9 istanze parallele. Ogni istanza del kernel \u00e8 definita\u00a0<em>work-item.\u00a0 <\/em>Nel nostro caso (l&#8217;esempio che stiamo analizzando) ogni istanza del kernel ricever\u00e0 le coordinate di un punto dell&#8217;immagine destinazione (quelle del pixel x,y) e, sulla base di esse estrarr\u00e0 le coordinate del pixel nel sorgente che sar\u00e0 trasformato ed inserito in posizione (x,y) destinazione; il numero totale di work-item usati sar\u00e0 nel nostro caso pari al numero di\u00a0 pixel nell&#8217;immagine destinazione; ovviamente non tutti i work-item saranno eseguiti in parallelo: a questo scopo OpenCL definisce il concetto di Work-Group, una matrice di work-item eseguiti in parallelo;<\/p>\n<p>La GPU scheduler\u00e0 intervalli di work-item in gruppi di elementi di elaborazione, i\u00a0<em>work-group<\/em>\u00a0, fino a quando tutti i work-item non sono stati elaborati. I kernel successivi possono essere eseguiti, fino al completamento di tutta l\u2019applicazione.<\/p>\n<h4>Set_index_space_2d()<\/h4>\n<p>La funzione definisce il range delle variazioni delle due coordinate x,y passate a ciascun work-item (uno per ogni pixel destinazione in questo esempio, ma sono possibili criteri diversi..); quindi l&#8217;istruzione:<\/p>\n<pre class=\"lang:default decode:true\" title=\"definizione di NDRange\">if (kernel-&gt;Set_index_space_2d(ret_area-&gt;dx, ret_area-&gt;dy))<\/pre>\n<p>Ha definito che il numero totale di work-item da eseguire sar\u00e0 pari ai valori\u00a0(ret_area-&gt;dx, ret_area-&gt;dy),\u00a0 dimensioni in pixel dell&#8217;immagine destinazione; Non \u00e8 visibile nel codice alcuna istruzione riguardante le dimensioni dei <span style=\"color: #800000;\">work-group<\/span>, il loro numero e le dimensioni in righe e colonne di work-item eseguiti in parallelo (appunto il work-group..): per non complicare la vita del programmatore, la scelta \u00e8 fatta automaticamente dall&#8217;interno della<span style=\"color: #800000;\"> gpu_kernel<\/span>, ad esecuzione della funzione\u00a0\u00a0<em><span style=\"color: #800000;\">Set_index_space_2d()<\/span><\/em>, sulla base delle caratteristiche rilevate del device sul quale viene eseguito il kernel;<\/p>\n<h4>Nota finale<\/h4>\n<p><span style=\"color: #000000;\">Abbiamo presentato, attraverso un frammento di codice, come vengono eseguiti i kernel nel framework di Colibri. Abbiamo mostrato come l&#8217;uso della classe <span style=\"color: #800000;\">gpu_kernel <\/span>e delle classi derivate da <span style=\"color: #800000;\">gpu_mem_object <\/span>semplifichi notevolmente le procedure d&#8217;accesso alle API OpenCL; Con poche righe di codice abbiamo richiamato un programma, compilato lo stesso, passato in forma standardizzata gli argomenti al kernel, eseguito quest&#8217;ultimo ed estratte le informazioni elaborate. Se qualcuno dei lettori volesse cimentarsi con l&#8217;uso diretto, non mediato dalle classi del namespace OPCL di Colibri, potr\u00e0 verificare quanti compiti sono invece svolti in modo completamente nascosto al programmatore dalle classi descritte in questo articolo.<\/span><\/p>\n<p>Da quanto abbiamo scritto forse non si evidenzia abbastanza l&#8217;importanza delle classi derivate da gpu_mem_object; ci limitiamo per ora a ricordare che le classi standardizzano il modo in cui i diversi tipi di informazione sono passati alla gpu; Ad esempio, la classe gpu_matrix_2d per essere passata correttamente al kernel, deve inserire nello stack degli argomenti la mappa dei punti, le dimensioni dx,dy della stessa, le dimensioni in byte dei punti.. la funzione PushKernel() della classe garantisce in un unica istruzione il passaggio di tutte queste informazioni al kernel nell&#8217;ordine atteso:<\/p>\n<p>gli argomenti trasferiti sono in un blocco <span style=\"color: #0000ff;\"><em>type *ptr, size_t nr, size nc, size_t s\u00a0 <\/em><span style=\"color: #000000;\">dove\u00a0<\/span><\/span><br \/>\n<em><span style=\"color: #993300;\">type *ptr<\/span><\/em> sono i punti della matrice ( l&#8217;interpretazione del tipo di puntatore (type *) \u00e8 definita dalla dichiarazione del kernel: ad esempio float *ptr)<br \/>\n<em><span style=\"color: #993300;\">size_t nr<\/span><\/em> numero di righe nella matrice<br \/>\n<span style=\"color: #993300;\"><em>size_t nc<\/em><\/span> numero di colonne<br \/>\n<em><span style=\"color: #993300;\">size_t<\/span><\/em><span style=\"color: #993300;\"> se<\/span> dimensioni in byte di ciascun elemento (di solito \u00e8 inutile, poich\u00e9 l&#8217;interpretazione \u00e8 data dalla dichiarazione del kernel (ad esempio float *ptr corrisponde a un se=4 byte)\u00a0ma viene sempre\u00a0 trasferita per usi particolari;<\/p>\n<p>Vedremo in seguito che questo metodo permette di standardizzare il modo di scrittura degli argomenti in un programma scritto in C-OpenCL, semplificandone la leggibilit\u00e0, anche in presenza di un numero molto elevato di variabili.<\/p>\n<p>Infine, si fa notare che la classe contiene alcune funzioni private, <em>friend<\/em> nei confronti delle classi gpu_program, gpu_buffer e GPU, e pertanto inaccessibili da altri contesti; questo \u00e8 dovuto alla necessit\u00e0 di proteggere la struttura interna della gpu_kernel da usi impropri;\u00a0 La funzione\u00a0<span style=\"color: #993300;\"> Name()<\/span> rimanda il nome assegnato nel codice alla funzione C-Opencl __kernel; le funzioni <span style=\"color: #993300;\">Set_id e Get_id<\/span> hanno l&#8217;evidente significato di assegnare e ottenere l&#8217;id\u00a0cl_kernel di OpenCL; non c&#8217;\u00e8 alcuna necessita d&#8217;uso degli stessi nel codice, pertanto non usarli mai.<\/p>\n<p>Infine, la funzione <span style=\"color: #993300;\">gpu_kernel::Program()<\/span> rimanda l&#8217;istanza <a href=\"https:\/\/www.colorwave.it\/blog\/heterogeneus-computing\/gpu_program\/admin\/2017\/12\/\" target=\"_blank\" rel=\"noopener\">gpu_programm<\/a> che contiene il kernel; anche questa funzione \u00e8 esposta per usi prevalentemente interni alle classi del namespace OPCL, ma pu\u00f2 essere interessante utilizzarla per esplorare le classi incapsulate nella gpu_kernel; ad esempio, questa riga di codice:<\/p>\n<pre class=\"lang:default decode:true\" title=\"ricavare il work-group definito per un kernel\">\tgpu_work_group_str *pdata = &amp;kernel-&gt;(Program()-&gt;Device()-&gt;WorkGroup ( )-&gt;pdata);\r\n<\/pre>\n<p>Permette di risalire alla struttura del work-group\u00a0 usata dal gpu_kernel *kernel;<\/p>\n<p>Nel prossimo articolo mostreremo come scrivere il codice C-OpenCL di un programma compatibile con le classe del namespace OPCL, in particolare con le modalit\u00e0 di definizione degli argomenti di gpu_kernel, mostrando l&#8217;uso delle macro che semplificano la sequenziazione degli stessi.<\/p>\n","protected":false},"excerpt":{"rendered":"<p>Scopo della classe: gpu_kernel \u00e8 stata progettata per facilitare la preparazione ed esecuzione dei kernel OpenCL, nascondendo tutti i dettagli\u00a0 delle operazioni necessarie, la cui scrittura complica notevolmente la stesura del codice: per eseguire una funzione kernel \u00e8 necessario innanzitutto individuare le&#8230;<\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"aside","meta":{"_acf_changed":false,"jetpack_post_was_ever_published":false,"_jetpack_newsletter_access":"","_jetpack_dont_email_post_to_subs":false,"_jetpack_newsletter_tier_id":0,"_jetpack_memberships_contains_paywalled_content":false,"footnotes":"","_jetpack_memberships_contains_paid_content":false,"jetpack_publicize_message":"","jetpack_publicize_feature_enabled":true,"jetpack_social_post_already_shared":true,"jetpack_social_options":{"image_generator_settings":{"template":"highway","enabled":false},"version":2}},"categories":[29,30,107],"tags":[],"jetpack_publicize_connections":[],"acf":[],"jetpack_featured_media_url":"","jetpack_sharing_enabled":true,"jetpack_shortlink":"https:\/\/wp.me\/p9rwsf-Ic","jetpack_likes_enabled":true,"jetpack-related-posts":[],"_links":{"self":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2740"}],"collection":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/comments?post=2740"}],"version-history":[{"count":25,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2740\/revisions"}],"predecessor-version":[{"id":2862,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2740\/revisions\/2862"}],"wp:attachment":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/media?parent=2740"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/categories?post=2740"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/tags?post=2740"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}