Un esempio d'uso di Push_kernel..

Analizziamo il seguente segmento di codice:

STRING prog_name = L"image";
STRING kernel_name = L"test_kernel";

OPCL::gpu_kernel *kernel = OPCL::GPU::OpenKernel(prog_name, kernel_name);
if (kernel)
{
	#define vs 1024
	float v[vs];

	for (int i = 0; i < vs; i++)
		v[i] = i;

	gpu_buf *src = new gpu_buf;
	buf->Init(kernel, vs*sizeof(float), CL_MEM_READ_ONLY, src);

	gpu_buf *dest = new gpu_buf;
	buf->Init(kernel, vs*sizeof(float), CL_MEM_WRITE_ONLY, 0);

	src->Push_kernel();
	src->Push_kernel();
	if (kernel->Set_index_space_1d(vs)
	{
		if (kernel->Exec())
		{
			//  __trace(L" kernel eseguito correttamente")
			bool res = kernel->PopArg(dest);
			................................................
				...................................
		}
		else
			OPCL::GPU::Display_error_info();
	}
	kernel->FlushMemObjects();

Il codice chiama la funzione statica OPCL::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ò notare, con una sola riga di programma sino state svolte tutte le funzioni di inizializzazione che richiedono normalmente molto più codice!

Quindi vengono allocati due gpu_buffer, il primo dei quali è inizializzato ai suoi indici, esi trasferiscono i buffer alla coda parametri del kernel (funzione Push_kernel); il kernel è stato scritto in modo da ricevere questi parametri

nel modo atteso:

__kernel test_kernel(__global float *src, int n_src, __global float *dest, int n_dest)
{
	int idx = get_global_id(0);
	if (idx < n_src)
		dest[idx] = src[idx] * src[idx]
}

Il kernel banalmente inserisce nel vettore dest i quadrati degli elementi di src ( chiamato dalla funzione kernel->Exec()) .

Il vettore aggiornato può quindi essere richiamato dalla GPU mediante la kernel->PopArg(dest) per gli usi successivi. L’esempio è volutamente banale, e serve ad evidenziare che, per ciascuna chiamata  a gpu_buffer::Push_kernel(), vengono trasferiti sequenzialmente un buffer di memoria e le delle sue dimensioni in byte: _global float *src, int n_src,    __global float *dest, int n_dest;

L’ultima istruzione del segmento di codice riportato è kernel->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.