{"id":2571,"date":"2017-12-20T15:45:43","date_gmt":"2017-12-20T15:45:43","guid":{"rendered":"http:\/\/www.colorwave.it\/blog\/?p=2571"},"modified":"2018-01-29T06:19:57","modified_gmt":"2018-01-29T06:19:57","slug":"gpu_buffer","status":"publish","type":"post","link":"https:\/\/www.colorwave.it\/blog\/framework\/gpu_buffer\/admin\/2017\/12\/","title":{"rendered":"gpu_buffer"},"content":{"rendered":"

gpu_buffer \u00e8, come evidenziato nel nome, il gestore di un buffer di memoria per i\/o su una GPU;\u00a0 appartiene al gruppo delle classi derivate da\u00a0 gpu_mem_obj<\/a>, e gestisce un blocco di memoria assegnata ad un context OpenCL. Come tutte le altre classi del namespace OPCL\u00a0 \u00e8 stata progettata per semplificare l’uso delle API OpenCL.\u00a0 In sintesi, gpu_buffer<\/em> \u00e8 la classe che gestisce la corrispondenza fra un blocco di memoria su CPU ed il suo i\/o sulle GPU gestite dalle API OpenCL. nascondendo al programmatore tutte le complessit\u00e0 legate a queste operazioni.<\/p>\n

\tclass AFX_EXT_CLASS gpu_buffer : public gpu_mem_object\r\n\t{\r\n\tpublic:\r\n\t\tgpu_buffer ( );\r\n\t\t~gpu_buffer ( );\r\n\r\n\t\tbool Init ( class gpu_kernel *kernel, cl_uint memsize, cl_mem_flags mode = CL_MEM_READ_ONLY, void *host_buf=0 );\r\n\t\tvoid *Ptr ( );\r\n\t\tsize_t Bytes ( );\r\n\r\n\t\tbool Push_kernel(gpu_kernel *kernel = 0);\r\n\r\n\t\t__use_private_struct\r\n\t};\r\n\r\n<\/pre>\n
void *Ptr() e size_t Bytes()<\/h4>\n
le due funzioni Ptr() e Bytes() rimandano il puntatore della memoria host\u00a0 e le dimensioni della memoria allocata. La memoria host \u00e8 allocata sulla CPU e pu\u00f2 essere definita esternamente alla classe GPU, assegnando il puntatore attraverso la funzione gpu_buffer::Init, o allocata automaticamente nella gpu_buffer, ed accedendovi attraverso la funzione Ptr().\u00a0 di seguito, in questo articolo, vedremo come ( e perch\u00e8) utilizzare questi due diversi metodi di allocazione.<\/p>\n
La classe contiene un puntatore alla memoria allocata su CPU, e nasconde un corrispondente handle OpenCL cl_mem definito attraverso la funzione OpenCL clCreateBuffer:<\/a><\/em>\u00a0 l’handle cl_mem \u00e8 accessibile mediante la funzione\u00a0Get_id<\/span>\u00a0() ereditata da gpu_mem_object, ma non sar\u00e0 mai necessario utilizzare informazioni e funzioni OpenCL direttamente, a meno che non si debba costruire nuove classe del namespace OPCL,<\/p>\n
bool Init()<\/h4>\n
La prima cosa da notare \u00e8 il primo parametro della funzione , gpu_kernel *kernel; ci\u00f2 pu\u00f2 apparire strano a\u00a0 chi ha dimestichezza con\u00a0 le API OpenCL, poich\u00e9 la memoria \u00e8 sempre\u00a0 gestita su un context,<\/a>\u00a0ed in seguito eventualmente passata ad un kernel eseguito su GPU: poich\u00e8 un kernel<\/a> \u00e8 sempre associato ad un programma<\/a>\u00a0, a sua volta associato ad un context, \u00e8 evidente\u00a0 che l’associazione fra gpu_buffer e context \u00e8 eseguito nella Init in modo completamente nascosto. Vedremo che questa scelta non comporta alcuna limitazione alla programmazione, anzi semplifica l’interfaccia di programmazione.<\/p>\n
Il secondo parametri della funzione \u00e8\u00a0memsize, cio\u00e8 la dimensione del buffer host da gestire; in un prossimo articolo sulla creazione di nuove classi derivate da gpu_buffer vedremo che questo parametro, per un uso particolare \u00e8 passato a o, e ne analizzeremo il significato; per l’uso standard invece memsize deve essere >0;<\/p>\n
il terzo parametro\u00a0cl_mem_flags mode \u00e8 uno dei tipi numerali di OpenCL, e definisce il tipo di trattamento che del buffer dovr\u00e0 fare l’API OpenCL; i valori possibili sono<\/p>\n
CL_MEM_READ_WRITE \r\nCL_MEM_WRITE_ONLY \r\nCL_MEM_READ_ONLY \r\nCL_MEM_USE_HOST_PTR \r\nCL_MEM_ALLOC_HOST_PTR \r\nCL_MEM_COPY_HOST_PTR\r\n<\/pre>\n
Per un’approfondimento del significato di questi valori leggere qui<\/a>.<\/p>\n
Il quarto ed ultimo parametro \u00e8 host_buf, un puntatore alla memoria che definisce la memoria allocata sulla CPU da associare alla host_buf\u00a0 lato GPU;<\/p>\n
\"\"<\/div><\/div>Una particolarit\u00e0 di questo parametro \u00e8 che pu\u00f2 essere passato con valore 0; In questo caso la funzione Init procede all’allocazione di un buffer di\u00a0memsize byte; L’accesso al buffer sar\u00e0 successivamente effettuato utilizzando il puntatore restituito dalla funzione Ptr() per inizializzare il suo contenuto; In questo caso la distruzione del buffer sar\u00e0 effettuata dal distruttore del gpu_buffer quando chiamata la funzione gpu_kernel::FlushMemObjects<\/em><\/p>\n
bool Push_kernel()<\/h4>\n
La funzione \u00e8 virtuale ed in genere ridefinita per le classi derivate da gpu_buffer; attraverso questa funzione viene trasferita l’informazione organizzata nella gpu_buffer<\/em> (o classi derivate) nella queue dei parametri del kernel. In genere non \u00e8 necessario ridefinire il kernel, in quanto gi\u00e0 assegnato nella funzione Init(); tuttavia \u00e8 possibile passare gli stessi dati ad un kernel aggiuntivo da eseguire in parallelo o in sequenza con quello di default;<\/p>\n
Compiti svolti da Push_kernel<\/h5>\n
Per capire come la funzione trasferisca le informazioni nella coda dei parametri del kernel, viene di seguito riportato il codice della funzione di gpu_buffer;<\/p>\n
\tbool gpu_buffer::Push_kernel(gpu_kernel *kernel)\r\n\t{\r\n\t\t__getstr\r\n\t\t_test_alt_kernel\t\t\r\n\t\treturn\r\n\t\t(\r\n\t\t\tkernel->PushArg(this) &&\r\n\t\t\tkernel->PushArg(_$str->bytes)\t\t\/\/ numero di byte inseriti\r\n\t\t);\r\n\t}\r\n<\/pre>\n
la macro _test_alt_str<\/em> permette di assegnare un kernel alternativo, eventualmente passato alla funzione, e la sua discussione non serve a questa esposizione; la riga kernel->Push(arg) trasferisce il buffer allocato in gpu_buffer (quello ottenibile dalla funzione gpu_buffer::Ptr() sulla coda dei parametri del kernel; con metodo simile le dimensioni del buffer sono aggiunte alla coda come secondo parametro;
\n<\/p>\n
Analizziamo il seguente segmento di codice:<\/p>\n
STRING prog_name = L\"image\";\r\nSTRING kernel_name = L\"test_kernel\";\r\n\r\nOPCL::gpu_kernel *kernel = OPCL::GPU::OpenKernel(prog_name, kernel_name);\r\nif (kernel)\r\n{\r\n\t#define vs 1024\r\n\tfloat v[vs];\r\n\r\n\tfor (int i = 0; i < vs; i++)\r\n\t\tv[i] = i;\r\n\r\n\tgpu_buf *src = new gpu_buf;\r\n\tbuf->Init(kernel, vs*sizeof(float), CL_MEM_READ_ONLY, src);\r\n\r\n\tgpu_buf *dest = new gpu_buf;\r\n\tbuf->Init(kernel, vs*sizeof(float), CL_MEM_WRITE_ONLY, 0);\r\n\r\n\tsrc->Push_kernel();\r\n\tsrc->Push_kernel();\r\n\tif (kernel->Set_index_space_1d(vs)\r\n\t{\r\n\t\tif (kernel->Exec())\r\n\t\t{\r\n\t\t\t\/\/  __trace(L\" kernel eseguito correttamente\")\r\n\t\t\tbool res = kernel->PopArg(dest);\r\n\t\t\t................................................\r\n\t\t\t\t...................................\r\n\t\t}\r\n\t\telse\r\n\t\t\tOPCL::GPU::Display_error_info();\r\n\t}\r\n\tkernel->FlushMemObjects();<\/pre>\n
Il codice chiama la funzione statica\u00a0OPCL::GPU::OpenKernel la quale, individuata una piattaforma ed un device, crea un context, richiama un programma contenente un kernel e genera l’oggetto gpu_kernel, pronto per essere usato. Come si pu\u00f2 notare, con una sola riga di programma sino state svolte tutte le funzioni di inizializzazione che richiedono normalmente molto pi\u00f9 codice!<\/p>\n
Quindi vengono allocati due gpu_buffer, il primo dei quali \u00e8 inizializzato ai suoi indici, esi trasferiscono i buffer alla coda parametri del kernel (funzione Push_kernel); il kernel \u00e8 stato scritto in modo da ricevere questi parametri<\/p>\n
nel modo atteso:<\/p>\n
__kernel test_kernel(__global float *src, int n_src, __global float *dest, int n_dest)\r\n{\r\n\tint idx = get_global_id(0);\r\n\tif (idx < n_src)\r\n\t\tdest[idx] = src[idx] * src[idx]\r\n}<\/pre>\n
Il kernel banalmente inserisce nel vettore dest i quadrati degli elementi di src ( chiamato dalla funzione kernel->Exec()) .<\/p>\n
Il vettore aggiornato pu\u00f2 quindi essere richiamato dalla GPU mediante la\u00a0kernel->PopArg(dest) per gli usi successivi. L’esempio \u00e8 volutamente banale, e serve ad evidenziare che, per ciascuna chiamata\u00a0 a gpu_buffer::Push_kernel(), vengono trasferiti sequenzialmente un buffer di memoria e le delle sue dimensioni in byte:\u00a0_global float *src, int n_src,\u00a0 \u00a0 __global float *dest, int n_dest;<\/em><\/p>\n
L’ultima istruzione del segmento di codice riportato \u00e8\u00a0kernel->FlushMemObjects(); questa funzione distrugge tutti gli oggetti di memoria associati al kernel (gpu_buf *dest e gpu_buf *src): le istanze dest e buf risulteranno automaticamente distrutti dopo questa chiamata di funzione.<\/p>\n
 <\/p>\n
 <\/p>\n
 <\/p>\n","protected":false},"excerpt":{"rendered":"
gpu_buffer \u00e8, come evidenziato nel nome, il gestore di un buffer di memoria per i\/o su una GPU;\u00a0 appartiene al gruppo delle classi derivate da\u00a0 gpu_mem_obj, e gestisce un blocco di memoria assegnata ad un context OpenCL. Come tutte le altre classi…<\/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":[19,29,108,107],"tags":[],"jetpack_publicize_connections":[],"acf":[],"jetpack_featured_media_url":"","jetpack_sharing_enabled":true,"jetpack_shortlink":"https:\/\/wp.me\/p9rwsf-Ft","jetpack_likes_enabled":true,"jetpack-related-posts":[],"_links":{"self":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2571"}],"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=2571"}],"version-history":[{"count":13,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2571\/revisions"}],"predecessor-version":[{"id":2864,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/posts\/2571\/revisions\/2864"}],"wp:attachment":[{"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/media?parent=2571"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/categories?post=2571"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.colorwave.it\/blog\/wp-json\/wp\/v2\/tags?post=2571"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}