import os
import math
import numpy
import pyopencl as cl
import pyopencl.array
import time
from PIL import Image

TASKS = 1048576
CL_TASKS = int(TASKS / 4)

if __name__ == '__main__':
    print('Cl program betoltese forras allomanybol')
    f = open('blur.cl', 'r', encoding='utf-8')
    kernels = ''.join(f.readlines())
    f.close()

    print('Adatok elokeszitese ... ')
    mask_size = 7
    img = Image.open(os.path.join(os.path.dirname(__file__), 'OpenCV-logo.png'))
    if img.mode != 'RGBA':
        img = img.convert('RGBA')
    img_width = img.size[0]
    img_height = img.size[1]
    img_size = img_width * img_height

    start_time = time.time()
    # Host memória előkészítése
    img_bytes = img.tobytes()

    time_hostdata_loaded = time.time()

    print('Kornyezet letrehozasa ...')
    ctx = cl.create_some_context()
    print('Parancssor letrehozasa ...')
    queue = cl.CommandQueue(ctx, properties=cl.command_queue_properties.PROFILING_ENABLE)
    time_ctx_queue_creation = time.time()

    dev_image_format = cl.ImageFormat(cl.channel_order.RGBA,
                                      cl.channel_type.UNSIGNED_INT8)
    # Eszkoz memoria elokeszitese
    print('Bemeneti es kimeneti eszkoz memoria')
    input_image = cl.Image(ctx,
                           cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR,
                           dev_image_format,
                           img.size,
                           None,
                           img_bytes)
    output_image = cl.Image(ctx,
                            cl.mem_flags.WRITE_ONLY,
                            dev_image_format,
                            img.size)
    time_devicedata_loaded = time.time()

    print('Kernel leforditasa')
    prg = cl.Program(ctx, kernels).build()
    time_kernel_compilation = time.time()

    np_masksize = numpy.int32(mask_size)
    print('Kernel vegrehajtasa')
    evt = prg.to_blur(queue, (img_size, ), (1, ),
                      np_masksize,
                      input_image, output_image)
    print('Kernel vegrehajtasara valo varakozas')
    evt.wait()
    elapsed = 1e-9 * (evt.profile.end - evt.profile.start)


    time_before_readback = time.time()
    buffer = numpy.zeros(img_width * img_height * 4, numpy.uint8)
    origin = (0, 0, 0)
    region = (img_width, img_height, 1)
    cl.enqueue_read_image(queue, output_image,
                          origin, region, buffer).wait()
    time_after_readback = time.time()

    print('Host adatik elokeszitese     : {}'.format(time_hostdata_loaded - start_time))
    print('Kornyezet/Parancssor:        : {}'.format(time_ctx_queue_creation-time_hostdata_loaded))
    print('Adatok feltoltese            : {}'.format(time_devicedata_loaded - time_ctx_queue_creation))
    print('Kernel forditas              : {}'.format(time_kernel_compilation-time_devicedata_loaded))
    print('OpenCL eltelt ido            : {}'.format(elapsed))
    print('Eredmeny letoltese           : {}'.format(time_after_readback - time_before_readback))

    out_filename = os.path.join('out_' + 'OpenCV-logo.png')
    out_im = Image.frombytes("RGBA", img.size, buffer.tobytes())
    if out_im.mode == 'RGBA':
        out_im = out_im.convert('RGB')
    out_im.save(out_filename)

    print('Vege')

Forrás letöltése itt.

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                          CLK_ADDRESS_CLAMP_TO_EDGE |
                          CLK_FILTER_NEAREST;

__kernel void to_blur(int aMaskSize,
                      __read_only image2d_t srcImg,
                      __write_only image2d_t dstImg) {

  int global_id = get_global_id(0);
  int offset = aMaskSize / 2;
  int sum_r = 0;
  int sum_g = 0;
  int sum_b = 0;
  int width = get_image_width(srcImg);
  int height = get_image_height(srcImg); 
  int row_n = global_id / width;
  int col_n = global_id % width;
  uint4 color;
  int2 coord;
  int count = 0;

  for(int y=-offset; y <= offset; y++) {
    for(int x=-offset; x <= offset; x++) {
      coord = (int2)(col_n + x, row_n + y);
      if (coord.x < 0 || coord.x >= width || coord.y < 0 || coord.y >= height) {
        continue;
      }
      color = read_imageui(srcImg, sampler, coord);
      sum_r += color.x;
      sum_g += color.y;
      sum_b += color.z;
      count += 1;
    }
  }
  color.x = sum_r / count;
  color.y = sum_g / count;
  color.z = sum_b / count;
  coord = (int2)(col_n, row_n);
  write_imageui(dstImg, coord, color);
}

Forrás letöltése itt.