OpenCL Fractals

OpenCL Κλασματομορφές (Φράκταλς)

07 Oct 2023 07 Οκτ 2023 · Coding Προγραμματισμός · graphics γραφικά algorithm αλγόριθμος opencl opencl parallelism παραλληλισμός python python

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

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