Első példa
A példaprogam C++ nyelven keresztül használja ki az OpenCL szolgáltatásait. Az OpenCL kód elhanyagolható a gazdaprogram kódjához képest. Egyszerűség kedvéért karaktersorozatként beépítésre került a kódba. A kernel két char típusú tömb összegzését végzi. Az olvashatóság végett kiemelve:
__kernel void oclSample(__global char* a, __global char* b, __global char* c)
{
int i = get_global_id(0) ;
c[i] = a[i] + b[i] ;
}
És a teljes példaprogram:
//
// OpenCL példaprogram
//
// Készítette: Hapák József, ELTE-IK MSc hallgató
// Év: 2011
// Környezet: Visual Studio 2010 Ultimate
//
#include <iostream>
#include <CL/opencl.h>
#define BUFFER_SIZE 4096
#define SOURCE_LINE_COUNT 5
const char* oclSource[] = {
"__kernel void oclSample(__global char* a, __global char* b, __global char* c)",
"{",
" int i = get_global_id(0) ;",
" c[i] = a[i] + b[i] ;",
"}"
} ;
char a[] = {39,60,54,45,90,0,36,100,56,87,14,62,-21,74,58,79,44,76,1,42} ;
char b[] = {33,41,54,63,21,32,43,12,45,23,53,14,53,13,53,35,64,24,32,-42} ;
char c[20] ;
using namespace std ;
void PrintPlatformInfo(cl_platform_id platform)
{
char buffer[BUFFER_SIZE] ;
clGetPlatformInfo(platform,CL_PLATFORM_VERSION,BUFFER_SIZE,buffer,NULL) ;
cout<<"OpenCL Plaform version: "<<buffer<<endl ;
clGetPlatformInfo(platform,CL_PLATFORM_NAME,BUFFER_SIZE,buffer,NULL) ;
cout<<"OpenCL Plaform name: "<<buffer<<endl ;
clGetPlatformInfo(platform,CL_PLATFORM_VENDOR ,BUFFER_SIZE,buffer,NULL) ;
cout<<"OpenCL Plaform vendor: "<<buffer<<"\n"<<endl ;
}
void PrintProgramBuildInfo(cl_program program, cl_device_id device)
{
size_t bufferLength ;
clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,NULL,NULL,&bufferLength) ;
char* buffer = new char[bufferLength] ;
clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,bufferLength,buffer,NULL) ;
cout<<"BuildLog: "<<buffer<<"\n"<<endl ;
delete[] buffer ;
}
int main()
{
// Platformok lekérése
cl_uint platformCount ;
clGetPlatformIDs(0,NULL,&platformCount) ;
cout<<"OpenCL platform count: "<<platformCount<<"\n"<<endl ;
cl_platform_id* platforms = new cl_platform_id[platformCount] ;
clGetPlatformIDs(platformCount,platforms,NULL) ;
// Ha nem találunk platformot kilépünk
if (platformCount == 0) exit(-1) ;
// Információk kiirása a platformokról
for (unsigned int i = 0 ; i < platformCount; ++i)
{
cout<<i+1<<". platform"<<endl ;
PrintPlatformInfo(platforms[i]) ;
}
// Használandó platform bekérérése
unsigned int platformID = 1 ;
do {
cin>>platformID ;
} while ( (platformID < 1) || (platformID > platformCount) ) ;
platformID-- ;
// Eszközök lekérése
cl_uint deviceCount ;
cl_device_id* devices ;
clGetDeviceIDs(platforms[platformID],CL_DEVICE_TYPE_ALL,0,NULL,&deviceCount) ;
devices = new cl_device_id[deviceCount] ;
clGetDeviceIDs(platforms[platformID],CL_DEVICE_TYPE_ALL,deviceCount,devices,NULL) ;
// Kontextus és parancslista létrehozása
cl_context clContext = clCreateContext(NULL,deviceCount,devices,NULL,NULL,NULL) ;
cl_command_queue clCommandQueue = clCreateCommandQueue(clContext,devices[0],NULL,NULL) ;
// Bufferek létrehozása
cl_mem a_buf = clCreateBuffer(clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(a),a,NULL) ;
cl_mem b_buf = clCreateBuffer(clContext,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(b),b,NULL) ;
cl_mem c_buf = clCreateBuffer(clContext,CL_MEM_READ_WRITE,sizeof(c),NULL,NULL) ;
// Program létrehozása forráskódból, fordítás, fordítás eredményének kiirása
cl_program clProgram = clCreateProgramWithSource(clContext,
SOURCE_LINE_COUNT,oclSource,NULL,NULL) ;
clBuildProgram(clProgram,deviceCount,devices,NULL,NULL,NULL) ;
PrintProgramBuildInfo(clProgram,devices[0]) ;
// kernel objektum létrehozása, paraméterek beállítása
cl_kernel clKernel = clCreateKernel(clProgram,"oclSample",NULL) ;
clSetKernelArg(clKernel,0,sizeof(cl_mem),&a_buf) ;
clSetKernelArg(clKernel,1,sizeof(cl_mem),&b_buf) ;
clSetKernelArg(clKernel,2,sizeof(cl_mem),&c_buf) ;
// kernel futtatása
const size_t kernelRange[] = { sizeof(a) } ;
clEnqueueNDRangeKernel(clCommandQueue,clKernel,1,NULL,kernelRange,NULL,NULL,NULL,NULL) ;
// adatok visszaolvasása, parancslista készrefuttatása
clEnqueueReadBuffer(clCommandQueue,c_buf,CL_TRUE,0,sizeof(c),c,NULL,NULL,NULL) ;
clFinish(clCommandQueue) ;
cout<<"Result:\n"<<c<<endl ;
delete[] platforms ;
delete[] devices ;
return 0 ;
}
A program megfelően működött a következő platformokon:
- AMD StreamSDK 2.3 x86
- Intel OpenCL alpha SDK x86
- nVidia 9600mGT
Mandelbrot fraktált generáló program
A program elsődleges célja az OpenCL és az OpenGL együttes használatának bemutatása.
A program végrehajtása során GLUT-tal OpenGL ablakot és kontextust hozunk létra az alapértelmezett eszközön, majd ebből a kontextusból OpenCL kontextust készítünk a Mandel osztály egy objektumában, amely osztály a továbbiakban elfedi előlünk az OpenCL erőforrások kezelését.
A megjelenítést az ablakot teljesen kitöltő négyzetlap textúrázásával oldjuk meg, ami a lineáris interpoláció miatt átméretezéskor is elfogadható minőséget eredményez.
A textúrát egy pixel pufferből töltjük fel, ami a Mandel osztály belső pufferének OpenGL oldali párja, így lényegében egyazon puffert megosztva használjuk a két programozási platform között. GPU-s implementáció esetén ez a PCI-Express adatmozgatás elhagyását jelentheti, így mindenképpen effektív technológia.
/*
Fraktál generálása.
*/
void Mandel::Process()
{
// Függő OpenGL műveletek befejezése.
glFinish();
// Megosztott OpenCL/OpenGL objektum lefoglalása az OpenCL számára.
clEnqueueAcquireGLObjects(queue, 1, &buffer, 0, nullptr, nullptr);
// Kernel futtatása a képterület egészére, width*height rácson.
clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, dimensions, nullptr, 0, nullptr, nullptr);
// Megosztott objektum elengedése.
clEnqueueReleaseGLObjects(queue, 1, &buffer, 0, nullptr, nullptr);
// Függő OpenCL műveletek befejezése.
clFinish(queue);
}
A fraktál generálása és a textúra feltöltése képkockánként újra megtörténik, így a Fraps-el kapott FPS érték releváns összehasonlítást ad a sebesség tekintetében.
/*
A GLUT-nak átadott, ciklikusan meghívott renderfüggvény.
*/
void Display()
{
// A fraktál kiszámítása.
mandel->Process();
// A textúre frissítése a pixel buffer object tartalmával.
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, GL_RGB, GL_UNSIGNED_BYTE, 0);
// Háttérszín hardveres gyorstörlése (elhagyható).
glClear(GL_COLOR_BUFFER_BIT);
// Teljes képernyőre való feszítés (alapértelmezett vetítés mellett).
glBegin(GL_TRIANGLE_STRIP);
glTexCoord2f(0.0f, 0.0f); glVertex2f(-1.0f, -1.0f);
glTexCoord2f(1.0f, 0.0f); glVertex2f( 1.0f, -1.0f);
glTexCoord2f(0.0f, 1.0f); glVertex2f(-1.0f, 1.0f);
glTexCoord2f(1.0f, 1.0f); glVertex2f( 1.0f, 1.0f);
glEnd();
// Dupla pufferelésnél puffercsere.
glutSwapBuffers();
// Új kirajzolási esemény elhelyezése.
glutPostRedisplay();
}
Az OpenGL kontextus és eszköz lekérdezése Win32 platformfüggő módon történik meg.
/*
OpenCL kontextus inicializálása.
*/
void Mandel::InitializeContext()
{
// OpenGL kontextus és eszköz lekérése WGL API hívásokkal.
HGLRC contextGL = wglGetCurrentContext();
HDC deviceGL = wglGetCurrentDC();
// Az OpenCL kontextus létrehozása az OpenGL alapján.
cl_context_properties properties[5] = { CL_GL_CONTEXT_KHR,
reinterpret_cast(contextGL),
CL_WGL_HDC_KHR, reinterpret_cast(deviceGL),
0
};
clGetGLContextInfoKHR(properties, CL_DEVICES_FOR_GL_CONTEXT_KHR, 0, nullptr, &deviceNum);
devices = new cl_device_id[deviceNum];
clGetGLContextInfoKHR(properties, CL_DEVICES_FOR_GL_CONTEXT_KHR, deviceNum, devices, nullptr);
context = clCreateContext(properties, 1, devices, nullptr, nullptr, nullptr);
}
Magának a "névadó" Mandelbrot fraktálnak a generálása a következőképpen fest OpenCL C-ben:
__kernel void GenerateMandelbrot(__global uchar *rgb)
{
const uint MaxIterations = 60,
x = get_global_id(0),
y = get_global_id(1),
_width = get_global_size(0),
_height = get_global_size(1),
i = y*_width + x;
const float MinRe = -2.0f,
MaxRe = 1.0f,
MinIm = -1.0f,
MaxIm = MinIm+(MaxRe-MinRe)*convert_float(_height)/convert_float(_width),
Re_factor = (MaxRe-MinRe)/convert_float(_width-1),
Im_factor = (MaxIm-MinIm)/convert_float(_height-1);
const float2 c = (float2)(MinRe + convert_float(x)*Re_factor, MaxIm - convert_float(y)*Im_factor);
float2 Z = c;
uint n;
for(n = 0; n < MaxIterations; ++n)
{
float2 Z2 = Z*Z;
if(Z2.s0 + Z2.s1 > 4.0f)
{
break;
}
Z = (float2)(Z2.s0 - Z2.s1, 2*Z.s0*Z.s1) + c;
}
uchar color = (n == MaxIterations) ? 0x00
: convert_uchar(convert_float(0xFF) * convert_float(n) / convert_float(MaxIterations));
rgb[3*i + 0] = color;
rgb[3*i + 1] = color;
rgb[3*i + 2] = color;
}
A program futtatásához meglévő OpenCL implementáció szükséges. Egyebek híján, az AMD APP SDK 2.4 CPU implementációja bármely SSE2-t támogató processzoron alkalmas erre a célra.
WebCL példák
Néhány WebCL-ben írt példaprogram. Segítségükkel interaktívan ismerkedhetünk az OpenCL-lel bármilyen SDK, illetve fejlesztő eszköz telepítése nélkül. Futtatásához azonban Firefox 4-re illetve a WebCL oldalán taláható pluginra lesz szükségünk.
FIGYELEM! Mivel ez egy előzetes implementációja a szabványnak lehet, hogy nem minden környezetben működőképes, illetve használata a rendszer működését instabillá teheti!
- 1. Példa - WebCL platform meglétének ellenőrzése a böngészőben.
- 2. Példa - Információk lekérése a rendszerben található platformokról és eszközökről.
- 3. Példa - Kernel fordítás, futtatás, memória műveletek.
- 4. Példa - Képek kezelése. Mandelbrot halmaz rajzolása Keszthelyi Balázs példájából kiindulva.
Julia halmaz ábrázolása
Júlia halmaz WebCL alatt Keszthelyi Balázs példájából kiindulva.
Készítette: Ujj László - 2012