OpenCL Fractals
OpenCL Κλασματομορφές (Φράκταλς)
IntroductionΕισαγωγή
In this article, we present a Python application that computes Fractals (for the time being, only the Mandelbrot set) using OpenCL, in order to parallelize and distribute the work across different devices.
OpenCL (Open Computing Language) is an open standard for writing code that is executed in parallel across heterogeneous platforms, such as:
- Central Processing Units (CPUs)
- Graphics Processing Units (GPUs)
- Field-Programmable Gate Arrays (FPGAs)
- Digital Signal Processors (DSPs)
Our application divides each fractal frame into blocks that are processed separately (each one by a different OpenCL device). The results from each device are combined to produce the final complete image. We also make use of coarse-grained SVM (Shared Virtual Memory) buffers to avoid copying large amount of data, back and forth.
Σε αυτό το άρθρο, παρουσιάζουμε μια εφαρμογή Python που υπολογίζει κλασματομορφές (προς το παρόν, μόνο το σύνολο Μάντελμπροτ) χρησιμοποιώντας το OpenCL, ώστε να παραλληλοποιήσει και να κατανήμει τη δουλειά σε διαφορετικές συσκευές.
Το OpenCL (Open Computing Language) είναι ένα ανοιχτό πρότυπο για τη συγγραφή κώδικα που εκτελείται με παραλληλία σε ετερογενείς πλατφόρμες, όπως λ.χ.:
- Κεντρικές Μονάδες Επεξεργασίας (CPUs)
- Γραφικές Μονάδες Επεξεργασίας (GPUs)
- Επιτόπια-Προγραμματιζόμενες Συστοιχίες Πυλών (FPGAs)
- Επεξεργαστές Ψηφιακού Σήματος (DSPs)
Η εφαρμογή μας διαιρεί την κάθε εικόνα κλασματομορφής σε επιμέρους τμήματα που υπολογίζονται χωριστά (καθένα από μια διαφορετική συσκευή OpenCL). Τα αποτελέσματα από την κάθε συσκευή συνδυάζονται για να παράγουν την τελική συνολική εικόνα. Χρησιμοποιούμε επίσης άδρες περιοχές Kοινής Eικονικής Mνήμης (coarse-grained Shared Virtual Memory) για να αποφύγουμε την αντιγραφή, πέρα δώθε, μεγάλων ποσοτήτων δεδομένων.
RequirementsΠροαπαιτούμενα
First of all, let's see the hardware / software requirements and the python packages we need:
Πρώτα από όλα, ας δούμε τις απαιτήσεις υλικού / λογισμικού και τα πακέτα Python που χρειαζόμαστε:
- At least one platform / device (and its OpenCL driver) that support OpenCL >= 2.0 and coarse-grained SVM.
- Kivy: An open source Python framework that facilitates the development of cross-platform GUI apps.
- NumPy: A Python library that is used for working with multi-dimensional arrays.
- PyOpenCL: Let us access the OpenCL computation API from Python. (Note that we need a version compiled with OpenCL >= 2.0 support).
- Τουλάχιστον μια πλατφόρμα / συσκευή (και τον OpenCL οδηγό της) που να υποστηρίζει OpenCL >= 2.0 και αδρές περιοχές Κοινής Εικονικής Μνήμης (coarse-grained SVM).
- Kivy: Ένα ανοιχτού κώδικα πακέτο λογισμικού Python που διευκολύνει την ανάπτυξη γραφικών εφαρμογών σε πολλές διαφορετικές πλατφόρμες.
- NumPy: Μια βιβλιοθήκη Python που χρησιμοποιείται για την εργασία με πολυδιάστατους πίνακες.
- PyOpenCL: Μας επιτρέπει να προσπελάσουμε την προγραμματιστική διεπαφή του OpenCL από την Python. (Σημειώστε ότι χρειαζόμαστε μια έκδοση που να έχει μεταγλωττιστεί να υποστηρίζει OpenCL >= 2.0).
CodeΚώδικας
OpenCL kernelΠυρήνας OpenCL
This is our OpenCL code («kernel») for the Mandelbrot set:
Αυτός είναι ο OpenCL κώδικας μας («kernel») για το σύνολο Μάντελμπροτ:
// Enable double-precision decimals (64-bit floats)
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
__kernel void mandelbrot (const double x0, const double y0,
const double step,
const unsigned int width,
const unsigned int max_iters,
const unsigned char red_factor,
const unsigned char green_factor,
const unsigned char blue_factor,
__global unsigned int *restrict framebuf)
{
double x = 0.0;
double y = 0.0;
double x2 = 0.0;
double y2 = 0.0;
unsigned int iters = 0;
const size_t xpos = get_global_id(0);
const size_t ypos = get_global_id(1);
const double xstep = x0 + (xpos * step);
const double ystep = y0 - (ypos * step);
#pragma unroll
while ((x2 + y2 < 4) && (iters < max_iters))
{
y = 2 * x * y + ystep;
x = x2 - y2 + xstep;
x2 = x * x;
y2 = y * y;
iters++;
}
if (iters >= max_iters)
{
framebuf[width * ypos + xpos] = 0xFF000000; // (fully opaque black)
}
else
{
framebuf[width * ypos + xpos] = (0xFF000000 |
(iters * blue_factor) << 16 |
(iters * green_factor) << 8 |
(iters * red_factor));
}
}
Device enumerationΑπαρίθμηση συσκευών
Here, we enumerate the devices in the platforms and we look for those that support coarse-grained SVM (Shared Virtual Memory). We get «Command Queues» for them:
Εδώ, απαριθμούμε τις συσκευές στις πλατφόρμες και αναζητούμε αυτές που υποστηρίζουν αδρές περιοχές Κοινής Εικονικής Μνήμης (coarse-grained SVM). Λαμβάνουμε «Ουρές Εντολών» («Command Queues») για αυτές:
def get_csvm_queues(self):
''' Find devices in platforms that support coarse-grained SVM
(Shared Virtual Memory) and get Command Queues for them. '''
csvm_queues = []
for platform in cl.get_platforms():
Logger.info("APP: Platform: %s (%s)",
platform.name, platform.version.strip())
context = cl.Context(dev_type=cl.device_type.ALL,
properties=[(cl.context_properties.PLATFORM,
platform)])
queue = cl.CommandQueue(context)
context = queue.context # A fix for some buggy implementations
device = queue.device
Logger.info("APP: -> Device: %s", device.name)
# Is coarse-grained buffer SVM (Shared Virtual Memory) supported?
csvm = characterize.has_coarse_grain_buffer_svm(device)
Logger.info("APP: * Coarse-grained buffer SVM: %s", csvm)
if csvm: # If yes, we save the queue
csvm_queues.append(queue)
num_queues = len(csvm_queues)
Logger.info("APP: %d workers will be created.", num_queues)
if num_queues == 0:
Logger.error("APP: No device found to support coarse-grained SVM.")
sys.exit(-1)
return csvm_queues
Image segmentationΚατάτμηση της εικόνας
As we mentioned before, we divide each fractal frame into blocks that are processed separately (each one by a different OpenCL device). We implement this using a Kivy vertical box layout, in which we add a number of image blocks equal to the number of OpenCL devices / workers:
Όπως αναφέραμε προηγουμένως, χωρίζουμε την κάθε εικόνα κλασματομορφής σε τμήματα που υπολογίζονται χωριστά, (καθένα από μια διαφορετική συσκευή OpenCL). Αυτό το υλοποιούμε χρησιμοποιώντας μια δίαταξη κάθετης στοίχισης του Kivy (vertical box layout), στην οποία προσθέτουμε αριθμό τμημάτων με εικόνες ίσο με τον αριθμό των συσκευών / εργατών OpenCL:
class MyBoxLayout(BoxLayout):
""" Class representing a vertical box layout that contains
the image blocks that compose a fractal frame. """
def __init__(self, n_blocks, **kwargs):
super().__init__(**kwargs)
self.rectx = 0
self.recty = 0
self.images = []
# We distribute the calculation of fractal image in distinct
# image blocks (number of blocks = number of OpenCL devices / workers)
for _ in range(n_blocks):
self.images.append(Image())
self.add_widget(self.images[-1])
# Area selection rectangle (for zoom in, initially hidden)
with self.canvas.after:
Color(1, 1, 1)
self.rect = Line(width=1, rectangle=(0, 0, 0, 0))
![Image divided in blocks in a vertical layout Image divided in blocks in a vertical layout](/images/image_blocks.jpg)
Worker initializationΑρχικοποιήση των εργατών
Here, we initialize the workers:
Εδώ, αρχικοποιούμε τους εργάτες:
def init_workers(self):
''' Initialization of workers '''
for idx, worker in enumerate(self.workers):
# Build the OpenCL program for each worker
program = cl.Program(worker["queue"].context,
MANDELBROT_OPENCL_CODE).build()
# Allocate coarse-grained SVM buffer for each worker
# Size is multiplied by 4 because we have 4 bytes per pixel (RGBA)
buffer_size = 4 * self.width * self.block_height
svm_buffer = cl.SVM(cl.csvm_empty(worker["queue"].context,
buffer_size, np.ubyte))
assert isinstance(svm_buffer.mem, np.ndarray)
# Create a texture for each worker's image block
texture = Texture.create(size=(self.width, self.block_height),
colorfmt="rgba")
self.screen.images[idx].texture = texture
# Calculate y0 of fractal
block_heights = (len(self.workers) - idx - 1) * self.block_height
y_dim = (self.height / 2) - block_heights
y_0 = np.double(args.y + self.step_size * y_dim)
# Save worker parameters
worker["program"] = program
worker["svm_buffer"] = svm_buffer
worker["svm_buffer_shape"] = (self.width, self.block_height)
worker["texture"] = self.screen.images[idx].texture
worker["y_0"] = y_0
Fractal renderingAπόδοση της κλασματομορφής
Now, we are ready to calculate the Mandelbrot set by executing the OpenCL kernel for each worker. Then, we collect the results and we render the complete fractal frame by blitting each individual result in the texture of each individual image block:
Τώρα, είμαστε έτοιμοι να υπολογίσουμε το σύνολο Μάντελμπροτ, εκτελώντας τον πυρήνα OpenCL για κάθε εργάτη. Έπειτα, συλλέγουμε τα αποτελέσματα και αποδίδουμε τη συνολική εικόνα της κλασματομορφής, μεταφέροντας το κάθε επιμέρους αποτέλεσμα στην υφή (texture) της κάθε επιμέρους εικόνας:
def update_fractal(self):
''' Draw a fractal frame. '''
for worker in self.workers:
# Execute the OpenCL kernel for each worker
worker["program"].mandelbrot(worker["queue"],
worker["svm_buffer_shape"], None,
self.x_0, worker["y_0"],
self.step_size, self.width,
self.max_iters,
PALETTES[self.palette_index]["RF"],
PALETTES[self.palette_index]["GF"],
PALETTES[self.palette_index]["BF"],
worker["svm_buffer"])
for worker in self.workers:
# Collect the results
worker["queue"].finish()
# We map the Coarse-grained buffer SVM (Shared Virtual Memory)
with worker["svm_buffer"].map_rw(worker["queue"]) as svm_buffer:
# We blit each worker result to the texture of its image block
worker["texture"].blit_buffer(svm_buffer,
colorfmt="rgba",
bufferfmt='ubyte')
self.screen.canvas.ask_update()
User interactionΑλληλεπίδραση χρήστη
MouseΠοντίκι
- Left click: Select area for zoom in
- Middle click: Use next color palette
- Right click: Return to the previous zoom (i.e. zoom out)
- Αριστερό κλικ: Επιλογή περιοχής για μεγένθυση
- Μεσαίο κλικ: Χρήση της επόμενης παλέττας χρωμάτων
- Δεξιό κλικ: Επιστροφή στην προηγούμενη μεγένθυση
KeyboardΠληκτρολόγιο
- WASD or arrows keys: Move the fractal
- Plus (+) / Minus (-) keys: Zoom in/out
- PageUp or c: Use next color palette
- PageDn: Use previous color palette
- Backspace: Return to the previous zoom
- F5: Increase maximum iterations by 1000 (may improve current frame rendering)
- Esc or q: Exit
- WASD ή βελάκια: Μετακίνηση στoν χώρο της κλασματομορφής
- Συν (+) / Πλην (-): Μεγένθυση/Σμίκρυνση
- PageUp or c: Χρήση της επόμενης παλέττας χρωμάτων
- PageDn: Χρήση της προηγούμενης παλέττας χρωμάτων
- Backspace: Επιστροφή στην προηγούμενη μεγένθυση
- F5: Αύξηση του μέγιστου αριθμού διαδοχικών υπολογισμών κατά 1000 (ίσως βελτιώσει την απόδοση της τρέχουσας εικόνας)
- Esc or q: Exit
ScreenshotsΣτιγμιότυπα
Full source codeΠλήρης πηγαίος κώδικας
You can download the full source code from here.
Μπορείτε να κατεβάσετε τον πλήρη πηγαίο κώδικα από εδώ.