Progression du développement du module 'OpenCL d'Abstrasy 2.0 beta-3
Je continue de travailler au développement du module 'OpenCL d'Abstrasy 2.0 beta-3, et ce dernier commence à prendre forme.
Je vous propose aujourd'hui un petit aperçu qui nécessite une bonne puissance de calculs : Explorer l'ensemble de Mandelbrot.
Pour cela, nous écrirons ce petit script qui comprend également le code source du kernel OpenCL:
(import 'G2d)
(import 'OpenCL)
(display-clr)
(define 'displ (G2d:Display "Mandelbrot"))
(define 'WIDTH 512)
(define 'HEIGHT 512)
(define 'im (G2d:Image WIDTH HEIGHT '!ARGB))
(displ:set-size WIDTH HEIGHT)
(displ:center)
(displ:show)
(define 'source `
// A very simple OpenCL kernel for computing the mandelbrot set
//
// output : A buffer with sizeX*sizeY elements, storing
// the colors as RGB ints
// sizeX, sizeX : The width and height of the buffer
// x0,y0,x1,y1 : The rectangle in which the mandelbrot
// set will be computed
// maxIterations : The maximum number of iterations
// colorMask : Mask of the colors elements,
// containing the pixel colors
__kernel void computeMandelbrot(
__global uint *output,
int sizeX, int sizeY,
float x0, float y0,
float x1, float y1,
int maxIterations,
int colorMask
)
{
unsigned int ix = get_global_id(0);
unsigned int iy = get_global_id(1);
float r = x0 + ix * (x1 - x0) / sizeX;
float i = y0 + iy * (y1 - y0) / sizeY;
float x = 0;
float y = 0;
float magnitudeSquared = 0;
int iteration = 0;
while ( iteration<maxIterations && magnitudeSquared<4 ) {
float xx = x * x;
float yy = y * y;
y = 2 * x * y + i;
x = xx - yy + r;
magnitudeSquared = xx + yy;
iteration++;
}
if (iteration == maxIterations) {
output[ iy * sizeX + ix ] = 0;
}
else {
float alpha = (float) iteration / maxIterations;
int colorIndex = (int) (alpha * colorMask);
output[ iy * sizeX + ix ] = 0xff000000 | colorIndex;
}
}
`)
(define 'device ((OpenCL:Platforms:get-all-gpu-devices) 0))
(define 'context (OpenCL:Context device))
(define 'command-queue (context:command-queue device))
(define 'program (context:program source))
(define 'mandelbrot_kernel (program:kernel "computeMandelbrot"))
(define 'im2d_array (array-of-ints (im:get-ARGB-ints)))
(define 'im2d_buffer (context:buffer im2d_array))
(mandelbrot_kernel:set-arg 0 im2d_buffer)
(mandelbrot_kernel:set-arg 1 WIDTH)
(mandelbrot_kernel:set-arg 2 HEIGHT)
(mandelbrot_kernel:set-arg 7 8192)
(mandelbrot_kernel:set-arg 8 0xffffff)
(define 'x0 -1.63248f)
(define 'y0 0.0f)
(define 'w (float WIDTH))
(define '$ws (/ 1.0f w))
(define '$s -1.0f)
(define '$wk w)
(define 'ndrange (OpenCL:NDRange (long-ints WIDTH HEIGHT)))
(while{displ}
{
(set! $wk (+ $wk (* $wk $ws $s)))
(if{<=? $wk 0.00009f}
{ if{negative? $s} {set! $s (- $s)} }
elif{>=? $wk w}
{ if{positive? $s} {set! $s (- $s)} }
)
(mandelbrot_kernel:set-arg 3 (- x0 $wk))
(mandelbrot_kernel:set-arg 4 (- y0 $wk))
(mandelbrot_kernel:set-arg 5 (+ x0 $wk))
(mandelbrot_kernel:set-arg 6 (+ y0 $wk))
(command-queue:enqueue-NDRange-kernel mandelbrot_kernel ndrange)
(command-queue:enqueue-read-buffer im2d_buffer)
(im:set-ARGB-ints im2d_array)
(displ:update im)
}
)
(displ:dispose)
En bref, ce script utilise le module 'G2d pour créer une image 2D et l'afficher à l'écran, et le module 'OpenCL pour bénéficier de l'accélération matérielle.
Ainsi la plus grande partie des calculs sont réalisés directement par le premier GPU trouvé (dans mon cas, le GK104s de la NVIDIA Quadro K5000). Voici le code du programme qui sera exécuté par le matériel:
__kernel void computeMandelbrot( __global uint *output, int sizeX, int sizeY, float x0, float y0, float x1, float y1, int maxIterations, int colorMask ) { unsigned int ix = get_global_id(0); unsigned int iy = get_global_id(1); float r = x0 + ix * (x1 - x0) / sizeX; float i = y0 + iy * (y1 - y0) / sizeY; float x = 0; float y = 0; float magnitudeSquared = 0; int iteration = 0; while ( iteration<maxIterations && magnitudeSquared<4 ) { float xx = x * x; float yy = y * y; y = 2 * x * y + i; x = xx - yy + r; magnitudeSquared = xx + yy; iteration++; } if (iteration == maxIterations) { output[ iy * sizeX + ix ] = 0; } else { float alpha = (float) iteration / maxIterations; int colorIndex = (int) (alpha * colorMask); output[ iy * sizeX + ix ] = 0xff000000 | colorIndex; } }
Le kernel reçoit tout d'abord un pointeur vers un buffer de sortie, puis les paramètres qui délimitent la zone à explorer ainsi que le mode de coloration. En fonction du nombre d'itérations, le kernel affecte une couleur (du type aRGB) dans le buffer de sortie.
Le script Abstrasy se charge de la préparation du traitement et de la mise en place de l'interopérabilité avec l'environnement OpenCL.
La boucle finale consiste tout simplement à lancer l'exécution du kernel pour chaque image qui constituent le résultat de l'exploration de la zone dont les limites sont ajustées au fur et à mesure. Le résultat se trouve dans un array-of-ints que l'on utilise comme raster pour actualiser l'image du rendu final.
La boucle s'arrête lorsqu'on ferme le fenêtre.
Voyons à présent le fichier log obtenu lorsque nous lançons l'exécution de notre script:
### NEW INTERPRETER INSTANCE STARTING...
#
# SUPERVISOR-THREAD -> register Thread[__Supervisor__,6,main]
#
# PARSING:
# . Parsing source code : 2951 characters...
# . Parsing time : 23ms.
# . Generate AST :
. {(import '@G2d) (import '@OpenCL) (display-clr) (define '@displ (@G2d:Display "Mandelbrot")) ... })
#
# RUNNING :
# . Searching library G2d :
# . file:/home/luc/workspace/8_Abstrasy_2_0/Z1_TESTS/G2d.jar
# . file:/home/luc/workspace/8_Abstrasy_2_0/bin/abstrasy/libs/G2d.jar
# . abstrasy.libs.g2d
# . Add built-in package : abstrasy.libs.g2d as G2d:*
# . Add 'G2d:Font -> (external "abstrasy.libs.g2d.AExt_Font" #|0x742D52DB|#)
# . Add 'G2d:TexturePaint -> (external "abstrasy.libs.g2d.AExt_TexturePaint" #|0x377CDA0B|#)
# . Add 'G2d:CubicCurve -> (external "abstrasy.libs.g2d.AExt_CubicCurve" #|0x571DFFE9|#)
# . Add 'G2d:Ellipse -> (external "abstrasy.libs.g2d.AExt_Ellipse" #|0x18C5E3E5|#)
# . Add 'G2d:Point -> (external "abstrasy.libs.g2d.AExt_Point" #|0x42F215F9|#)
# . Add 'G2d:Rectangle -> (external "abstrasy.libs.g2d.AExt_Rectangle" #|0x7AA858F9|#)
# . Add 'G2d:Polygon -> (external "abstrasy.libs.g2d.AExt_Polygon" #|0x75BDD1EB|#)
# . Add 'G2d:GradientPaint -> (external "abstrasy.libs.g2d.AExt_GradientPaint" #|0x42D0DE72|#)
# . Add 'G2d:Arc -> (external "abstrasy.libs.g2d.AExt_Arc" #|0x10DF560C|#)
# . Add 'G2d:Line -> (external "abstrasy.libs.g2d.AExt_Line" #|0x2888AB7E|#)
# . Add 'G2d:Text -> (external "abstrasy.libs.g2d.AExt_Text" #|0x33A1150D|#)
# . Add 'G2d:RoundRectangle -> (external "abstrasy.libs.g2d.AExt_RoundRectangle" #|0x77467E45|#)
# . Add 'G2d:QuadCurve -> (external "abstrasy.libs.g2d.AExt_QuadCurve" #|0x1DC33997|#)
# . Add 'G2d:Canvas -> (external "abstrasy.libs.g2d.AExt_Canvas" #|0x707DB20A|#)
# . Add 'G2d:Stroke -> (external "abstrasy.libs.g2d.AExt_Stroke" #|0x4B4FC875|#)
# . Add 'G2d:Chord -> (external "abstrasy.libs.g2d.AExt_Chord" #|0x0F3F230F|#)
# . Add 'G2d:Display -> (external "abstrasy.libs.g2d.AExt_Display" #|0x0B34BCFE|#)
# . Add 'G2d:Image -> (external "abstrasy.libs.g2d.AExt_Image" #|0x237DCCAD|#)
# . Add 'G2d:Kernel -> (external "abstrasy.libs.g2d.AExt_Kernel" #|0x552CB496|#)
# . Add 'G2d:Tk -> (G2d:Tk #|0x0560EEB6|#)
# . Add 'G2d:Pie -> (external "abstrasy.libs.g2d.AExt_Pie" #|0x35D87288|#)
# . Add 'G2d:Path -> (external "abstrasy.libs.g2d.AExt_Path" #|0x47F7909B|#)
# . Searching library OpenCL :
# . file:/home/luc/workspace/8_Abstrasy_2_0/Z1_TESTS/OpenCL.jar
# . file:/home/luc/workspace/8_Abstrasy_2_0/bin/abstrasy/libs/OpenCL.jar
# . abstrasy.libs.opencl
# . Create new resources cache :
# . Resources cache /tmp/rc-0522572c-1793-4cd9-83c6-e03a0842665e.tmp created.
# . Cache resources /JOCL-0.1.9.jar from file:/home/luc/workspace/8_Abstrasy_2_0/bin/abstrasy/libs/opencl/extras/JOCL-0.1.9.jar
# . Link external resources from /tmp/rc-0522572c-1793-4cd9-83c6-e03a0842665e.tmp/JOCL-0.1.9.jar
# . Cache resources /JOCL-Utils.jar from file:/home/luc/workspace/8_Abstrasy_2_0/bin/abstrasy/libs/opencl/extras/JOCL-Utils.jar
# . Link external resources from /tmp/rc-0522572c-1793-4cd9-83c6-e03a0842665e.tmp/JOCL-Utils.jar
# . Add built-in package : abstrasy.libs.opencl as OpenCL:*
# . Add 'OpenCL:Kernel-Trait -> (external "abstrasy.libs.opencl.AExt_Kernel_Trait" #|0x42530DDC|#)
# . Add 'OpenCL:Image2D -> (external "abstrasy.libs.opencl.AExt_Image2D" #|0x669B050A|#)
# . Add 'OpenCL:Program -> (external "abstrasy.libs.opencl.AExt_Program" #|0x105A14E3|#)
# . Add 'OpenCL:Platform -> (external "abstrasy.libs.opencl.AExt_Platform" #|0x0B309363|#)
# . Add 'OpenCL:Context -> (external "abstrasy.libs.opencl.AExt_Context" #|0x7DE88C1A|#)
# . Add 'OpenCL:Command-Queue -> (external "abstrasy.libs.opencl.AExt_Command_Queue" #|0x72CCB34F|#)
# . Add 'OpenCL:Platforms -> (OpenCL:Platforms #|0x50D3C840|#)
# . Add 'OpenCL:CL -> (OpenCL:CL #|0x5D09C785|#)
# . Add 'OpenCL:Buffer -> (external "abstrasy.libs.opencl.AExt_Buffer" #|0x13C4B1D5|#)
# . Add 'OpenCL:Kernel -> (external "abstrasy.libs.opencl.AExt_Kernel" #|0x62573DE1|#)
# . Add 'OpenCL:Device -> (external "abstrasy.libs.opencl.AExt_Device" #|0x571AFF26|#)
# . Image2d : Start refreshing View2D implementation.
# . Image2d : Stop refreshing View2D implementation.
# . Image2d : Start refreshing View2D implementation.
# . Finalize abstrasy.libs.g2d.AExt_Display<0x0AEAB445>
# FINALIZE RESIDUAL RESOURCES :
# . Finalize abstrasy.libs.opencl.AExt_Buffer<0x4B4E1B60>
# . Finalize abstrasy.libs.opencl.AExt_Kernel<0x6B894814>
# . Finalize abstrasy.libs.opencl.AExt_Program<0x2C237B7D>
# . Finalize abstrasy.libs.opencl.AExt_Command_Queue<0x7334AA20>
# . Finalize abstrasy.libs.opencl.AExt_Context<0x481A2A3D>
# . Image2d : Stop refreshing View2D implementation.
# . Running time : 3858ms.
#
#
# SUPERVISOR-THREAD -> emptying messages.
# SUPERVISOR-THREAD -> emptying signals.
# SUPERVISOR-THREAD -> emptying uncaught exceptions.
# SUPERVISOR-THREAD -> unregister Thread[__Supervisor__,6,main]
# SUPERVISOR-THREAD -> END.
Comme vous pouvez le constater, l'interpréteur extrait et charge automatiquement les ressources nécessaires pour l'exécution de notre script. Cela inclut les fichiers «JOCL-0.1.9.jar» et «JOCL-Utils.jar». Bien sûr, comme le test est réalisé sur une machine équipée de Linux, la bibliothèque «libJOCL_0_1_9-linux-x86_64.so» est aussi automatiquement extraite lorsque le paquetage «JOCL-0.1.9.jar» est chargé dans le classpath.
L'interpréteur explore ensuite le contenu du module et inscrit ses composants dans le registre.
On peut donc les utiliser à partir de l'espace de noms OpenCL:.
Le programme est lancé et s'exécute jusqu'à ce qu'on ferme la fenêtre G2d:Display.
On constate que l'interpréteur s'occupe de finaliser et de restituer les ressources qui n'ont pas été fermée explicitement dans le script. L'exécution s'arrête ensuite.
Mais, pour le plaisir voici le rendu obtenu à l'écran :
Résultat entre 90 et 200 fps avec une carte NVIDIA Quadro K5000, soit 52,8 Mpixels/s.
Qui a dit qu'un langage interprété devait être lent?… 
