|
| 1 | +#!/usr/bin/env python3 |
| 2 | +# this example shows how to blur, gray scale or brightness adjust the picture |
| 3 | +# using pyopencl and image2d_t buffer objects |
| 4 | +# this python script will ask you for: |
| 5 | +# input image path |
| 6 | +# number of which filter you want to apply to |
| 7 | +# if gaussian blur is chosen it will ask you for parameters of gaussian kernel |
| 8 | +# if brightness adjustment is chosen it will ask you for a scaler that will |
| 9 | +# determines how much to adjust it |
| 10 | +# this script was my homework that was given on course in Parallel Algorithms |
| 11 | +# for more info contact me at bbozic13023rn@raf.rs |
| 12 | +__copyright__ = "Copyright (C) 2025 Bogdan Bozic" |
| 13 | + |
| 14 | +__license__ = """ |
| 15 | +Permission is hereby granted, free of charge, to any person |
| 16 | +obtaining a copy of this software and associated documentation |
| 17 | +files (the "Software"), to deal in the Software without |
| 18 | +restriction, including without limitation the rights to use, |
| 19 | +copy, modify, merge, publish, distribute, sublicense, and/or sell |
| 20 | +copies of the Software, and to permit persons to whom the |
| 21 | +Software is furnished to do so, subject to the following |
| 22 | +conditions: |
| 23 | +
|
| 24 | +The above copyright notice and this permission notice shall be |
| 25 | +included in all copies or substantial portions of the Software. |
| 26 | +
|
| 27 | +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, |
| 28 | +EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES |
| 29 | +OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND |
| 30 | +NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT |
| 31 | +HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, |
| 32 | +WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING |
| 33 | +FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR |
| 34 | +OTHER DEALINGS IN THE SOFTWARE. |
| 35 | +""" |
| 36 | +import os |
| 37 | +import numpy as np |
| 38 | +from PIL import Image |
| 39 | +import pyopencl as cl |
| 40 | + |
| 41 | + |
| 42 | +platform = cl.get_platforms()[0] |
| 43 | +device = platform.get_devices()[0] |
| 44 | +context = cl.Context([device]) |
| 45 | +queue = cl.CommandQueue(context) |
| 46 | +n_threads = device.max_work_group_size |
| 47 | + |
| 48 | + |
| 49 | +def calc_kernel(size: int, sigma: float) -> np.ndarray: |
| 50 | + x = np.linspace(-(size // 2), size // 2, size) |
| 51 | + x /= np.sqrt(2) * sigma |
| 52 | + x2 = x**2 |
| 53 | + kernel = np.exp(-x2[:, None] - x2[None, :]) |
| 54 | + return kernel / kernel.sum() |
| 55 | + |
| 56 | + |
| 57 | +def sum_arr(arr: np.array) -> int: |
| 58 | + # N is first bigger number that is multiple of n_threads |
| 59 | + n = (np.ceil(len(arr) / n_threads) * n_threads).astype(np.int32) |
| 60 | + # this fills the array with trailing zeros so that the kernel |
| 61 | + # doesnt access random memory |
| 62 | + arr = np.concatenate((arr, np.zeros(n - len(arr)))) |
| 63 | + mf = cl.mem_flags |
| 64 | + # creates a openCL buffer obj that has the copy of an array |
| 65 | + arr_buf = cl.Buffer( |
| 66 | + context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=arr.astype(np.int32) |
| 67 | + ) |
| 68 | + # in reduce var will be sum(arr[0:n_threads]) for 0th index and so on |
| 69 | + reduce = np.zeros(n // n_threads).astype(np.int32) |
| 70 | + # same as prev openCL buf this one will be filled by kernel |
| 71 | + reduce_buf = cl.Buffer(context, mf.READ_WRITE, size=reduce.nbytes) |
| 72 | + # in local buf, reduction is applied on local memory |
| 73 | + local_buf = cl.LocalMemory(4 * n_threads) |
| 74 | + program.reduce_sum(queue, (n,), (n_threads,), arr_buf, reduce_buf, local_buf).wait() |
| 75 | + cl.enqueue_copy(queue, reduce, reduce_buf) |
| 76 | + # this process could be applied more than once on an array but for |
| 77 | + # images applying reduction_sum once is more than enough |
| 78 | + return np.sum(reduce) |
| 79 | + |
| 80 | + |
| 81 | +def sum_matrix(matrix: np.array) -> np.array: |
| 82 | + # this function calculates avg values for every pixel value |
| 83 | + # when you flatten a RGBA image every 4th index is red value |
| 84 | + # starting at 0. For blue value it is every 4th index but |
| 85 | + # starting from 1 and so on |
| 86 | + return np.array( |
| 87 | + [ |
| 88 | + sum_arr(np.array(matrix[::4])), |
| 89 | + sum_arr(np.array(matrix[1::4])), |
| 90 | + sum_arr(np.array(matrix[2::4])), |
| 91 | + ] |
| 92 | + ) |
| 93 | + |
| 94 | + |
| 95 | +def save(filter: str): |
| 96 | + dest = np.empty_like(image) |
| 97 | + cl.enqueue_copy(queue, dest, dest_buf, origin=(0, 0), region=shape) |
| 98 | + file_name, file_extension = os.path.splitext(image_path) |
| 99 | + Image.fromarray(dest).save(f"{file_name}_{filter}{file_extension}", "PNG") |
| 100 | + |
| 101 | + |
| 102 | +program = cl.Program( |
| 103 | + context, |
| 104 | + """ |
| 105 | +__kernel void gray_scale(read_only const image2d_t src, |
| 106 | + write_only image2d_t dest, |
| 107 | + const sampler_t sampler){ |
| 108 | + int2 pos = (int2)(get_global_id(0),get_global_id(1)); |
| 109 | + uint4 pixel = read_imageui(src,sampler,pos); |
| 110 | + uint g=pixel.x*299/1000+pixel.y*578/1000+pixel.z*114/1000; |
| 111 | + write_imageui(dest,pos,(uint4)(g,g,g,255)); |
| 112 | +} |
| 113 | +__kernel void brightness_adj(read_only const image2d_t src, |
| 114 | + write_only image2d_t dest, |
| 115 | + const sampler_t sampler, |
| 116 | + const float scalar, |
| 117 | + const float4 mean_intensity){ |
| 118 | + int2 pos = (int2)(get_global_id(0),get_global_id(1)); |
| 119 | + uint4 pixel = read_imageui(src,sampler,pos); |
| 120 | + uint4 new_pixel = convert_uint4((float4)scalar*(convert_float4(pixel)\ |
| 121 | + -mean_intensity)+mean_intensity); |
| 122 | + uint4 overflow = convert_uint4(new_pixel>(uint4)255); |
| 123 | + write_imageui(dest,pos,select(new_pixel,(uint4)255,overflow)); |
| 124 | +} |
| 125 | +__kernel void reduce_sum(__global int *in, |
| 126 | + __global int *reduce, |
| 127 | + __local int *buffer) |
| 128 | +{ |
| 129 | + uint gid = get_global_id(0); |
| 130 | + uint wid = get_group_id(0); |
| 131 | + uint lid = get_local_id(0); |
| 132 | + uint gs = get_local_size(0); |
| 133 | + buffer[lid] = in[gid]; |
| 134 | +
|
| 135 | + for(uint s = gs/2; s > 0; s >>= 1) { |
| 136 | + barrier(CLK_LOCAL_MEM_FENCE); |
| 137 | + if(lid < s) { |
| 138 | + buffer[lid] += buffer[lid+s]; |
| 139 | + } |
| 140 | + } |
| 141 | + if(lid == 0) reduce[wid] = buffer[lid]; |
| 142 | +} |
| 143 | +__kernel void gaussian_blur(read_only const image2d_t src, |
| 144 | + write_only image2d_t dest, |
| 145 | + const sampler_t sampler, |
| 146 | + read_only const image2d_t gauss_kernel, |
| 147 | + const short dim){ |
| 148 | + int2 pos = (int2)(get_global_id(0),get_global_id(1)); |
| 149 | + uint4 pixel=(uint4)0; |
| 150 | + float4 src_pixel,gauss_pixel; |
| 151 | + for(int i = -dim/2;i<dim/2;i++){ |
| 152 | + for(int j = -dim/2;j<dim/2;j++){ |
| 153 | + src_pixel=convert_float4(read_imageui(src,sampler,(int2)(i,j)+pos)); |
| 154 | + gauss_pixel=convert_float4(read_imageui(gauss_kernel,sampler,(int2)(i+dim/2,j+dim/2)))/(float4)255; |
| 155 | + pixel=pixel+convert_uint4(src_pixel*gauss_pixel); |
| 156 | + } |
| 157 | + } |
| 158 | + uint4 overflow = convert_uint4(pixel>(uint4)255); |
| 159 | + write_imageui(dest,pos,select(pixel,(uint4)255,overflow)); |
| 160 | +} |
| 161 | +""", |
| 162 | +).build() |
| 163 | + |
| 164 | +image_path = "./noisyImage.jpg" |
| 165 | +# intel-compile-runtime supports RGBA so im using that format |
| 166 | +image = np.array(Image.open(image_path).convert("RGBA")) |
| 167 | + |
| 168 | +src_buf = cl.image_from_array(context, image, 4) |
| 169 | + |
| 170 | +image_format = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8) |
| 171 | + |
| 172 | +shape = (image.shape[1], image.shape[0]) |
| 173 | + |
| 174 | +dest_buf = cl.create_image(context, cl.mem_flags.WRITE_ONLY, image_format, shape) |
| 175 | + |
| 176 | +sampler = cl.Sampler( |
| 177 | + context, False, cl.addressing_mode.CLAMP_TO_EDGE, cl.filter_mode.NEAREST |
| 178 | +) |
| 179 | + |
| 180 | +program.gray_scale(queue, shape, None, src_buf, dest_buf, sampler) |
| 181 | +save("gray_scale") |
| 182 | + |
| 183 | +ker = calc_kernel(7, 1) |
| 184 | +# line below creates a gauss kernel but stacks it in Z axis |
| 185 | +# so that it can be interpreted as an Image |
| 186 | +ker = np.array( |
| 187 | + np.stack((ker, ker, ker, np.ones((len(ker), len(ker)))), axis=2) * 255 |
| 188 | +).astype(np.uint8) |
| 189 | + |
| 190 | +gauss_buf = cl.image_from_array(context, ker, 4) |
| 191 | + |
| 192 | +program.gaussian_blur( |
| 193 | + queue, |
| 194 | + shape, |
| 195 | + None, |
| 196 | + src_buf, |
| 197 | + dest_buf, |
| 198 | + sampler, |
| 199 | + gauss_buf, |
| 200 | + np.int32(len(ker)), |
| 201 | +) |
| 202 | +save("gauss") |
| 203 | + |
| 204 | +# turn the image into 1D array so that it can be sum reduced in opencl |
| 205 | +# it could have been 2D or even 3D but this was simpler |
| 206 | +temp_image = image.flatten().astype(np.int32) |
| 207 | +# concatenates [avg(red_pixel_val),avg(blue_pixel_val),avg(green_pixel_val)] |
| 208 | +# with 255 for the alpha value |
| 209 | +pixel = np.concatenate( |
| 210 | + ( |
| 211 | + (sum_matrix(temp_image) / (len(temp_image) // 4)).astype(np.float32), |
| 212 | + np.array([255.0]).astype(np.float32), |
| 213 | + ) |
| 214 | +) |
| 215 | + |
| 216 | +# scale = input("Enter the scale of brightness adjustment:") |
| 217 | +scale = 1.5 |
| 218 | +program.brightness_adj( |
| 219 | + queue, shape, None, src_buf, dest_buf, sampler, np.float32(scale), pixel |
| 220 | +) |
| 221 | +save("brightness") |
0 commit comments