Last active
March 27, 2021 17:28
-
-
Save OluwoleOyetoke/b97a08670c673efc2e02cbd9e53f5b0b to your computer and use it in GitHub Desktop.
OpenCL Kernel for Circle Detection
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
//OpenCL kernel for circle detection using Circular Hough Transform | |
//pragmas to enable floating point operations | |
#pragma OPENCL EXTENSION cl_khr_fp64 : enable | |
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable | |
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable | |
//Already created Sine and Cosine Look-up table for improved speed | |
__constant double cosineTableGlobal[360] = { 1, 0.540302, -0.416147, -0.989992, -0.653644, 0.283662, 0.96017, 0.753902, -0.1455, -0.91113, -0.839072, 0.0044257, 0.843854, 0.907447, 0.136737, -0.759688, -0.957659, -0.275163, 0.660317, 0.988705, 0.408082, -0.547729, -0.999961, -0.532833, 0.424179, 0.991203, 0.646919, -0.292139, -0.962606, -0.748058, 0.154251, 0.914742, 0.834223, -0.0132767, -0.84857, -0.903692, -0.127964, 0.765414, 0.955074, 0.266643, -0.666938, -0.987339, -0.399985, 0.555113, 0.999843, 0.525322, -0.432178, -0.992335, -0.640144, 0.300593, 0.964966, 0.742154, -0.162991, -0.918283, -0.82931, 0.0221268, 0.85322, 0.899867, 0.11918, -0.77108, -0.952413, -0.258102, 0.673507, 0.985897, 0.391857, -0.562454, -0.999647, -0.51777, 0.440143, 0.99339, 0.633319, -0.309023, -0.967251, -0.736193, 0.171717, 0.921751, 0.824331, -0.030975, -0.857803, -0.895971, -0.110387, 0.776686, 0.949678, 0.24954, -0.680023, -0.984377, -0.383698, 0.56975, 0.999373, 0.510177, -0.448074, -0.994367, -0.626444, 0.317429, | |
0.969459, 0.730174, -0.18043, -0.925148, -0.819288, 0.0398209, 0.862319, 0.892005, 0.101586, -0.782231, -0.946868, -0.240959, 0.686487, 0.98278, 0.37551, -0.577002, -0.999021, -0.502544, 0.455969, 0.995267, 0.619521, -0.32581, -0.971592, -0.724097, 0.189129, 0.928471, 0.814181, -0.0486636, -0.866767, -0.887969, -0.0927762, 0.787715, 0.943984, 0.232359, -0.692896, -0.981106, -0.367291, 0.584209, 0.99859, 0.494872, -0.463829, -0.996088, -0.612548, 0.334165, 0.973649, 0.717964, -0.197814, -0.931722, -0.80901, 0.0575025, 0.871147, 0.883863, 0.0839594, -0.793136, -0.941026, -0.223741, 0.699251, 0.979355, 0.359044, -0.59137, -0.998081, -0.487161, 0.471652, 0.996831, 0.605528, -0.342495, -0.975629, -0.711775, 0.206482, 0.9349, 0.803775, -0.0663369, -0.875459, -0.879689, -0.0751361, 0.798496, 0.937995, 0.215105, -0.705551, -0.977527, -0.350769, 0.598484, 0.997494, 0.479412, -0.479439, -0.997496, -0.59846, 0.350797, 0.977533, 0.70553, -0.215135, -0.938005, -0.798478, | |
0.0751662, 0.879703, 0.875445, 0.0663069, -0.803793, -0.93489, -0.206453, 0.711796, 0.975623, 0.342466, -0.605552, -0.996829, -0.471626, 0.487188, 0.998083, 0.591345, -0.359072, -0.979361, -0.699229, 0.22377, 0.941037, 0.793118, -0.0839895, -0.883877, -0.871133, -0.0574724, 0.809028, 0.931711, 0.197784, -0.717985, -0.973642, -0.334137, 0.612572, 0.996085, 0.463802, -0.494898, -0.998592, -0.584184, 0.367319, 0.981111, 0.692874, -0.232388, -0.943994, -0.787696, 0.0928062, 0.887983, 0.866752, 0.0486335, -0.814198, -0.92846, -0.1891, 0.724118, 0.971585, 0.325781, -0.619544, -0.995264, -0.455942, 0.50257, 0.999022, 0.576978, -0.375538, -0.982785, -0.686465, 0.240988, 0.946878, 0.782212, -0.101616, -0.892018, -0.862304, -0.0397908, 0.819306, 0.925136, 0.180401, -0.730194, -0.969452, -0.3174, 0.626468, 0.994364, 0.448047, -0.510203, -0.999374, -0.569726, 0.383726, 0.984382, 0.680001, -0.249569, -0.949687, -0.776667, 0.110417, 0.895984, 0.857788, 0.0309449, -0.824348, | |
-0.92174, -0.171688, 0.736213, 0.967243, 0.308994, -0.633343, -0.993387, -0.440116, 0.517796, 0.999648, 0.562429, -0.391885, -0.985902, -0.673485, 0.258131, 0.952422, 0.771061, -0.11921, -0.89988, -0.853204, -0.0220966, 0.829327, 0.918271, 0.162961, -0.742174, -0.964958, -0.300564, 0.640167, 0.992332, 0.432151, -0.525348, -0.999844, -0.555088, 0.400013, 0.987344, 0.666916, -0.266672, -0.955083, -0.765395, 0.127994, 0.903705, 0.848554, 0.0132466, -0.83424, -0.91473, -0.154222, 0.748078, 0.962598, 0.29211, -0.646942, -0.991199, -0.424152, 0.532859, 0.999961, 0.547704, -0.40811, -0.988709, -0.660294, 0.275192, 0.957668, 0.759668, -0.136767, -0.907459, -0.843838, -0.00439555, 0.839088, 0.911118, 0.14547, -0.753922, -0.960162, -0.283633, 0.653666, 0.989988, 0.416119, -0.540328, -1, -0.540277, 0.416174, 0.989997, 0.653621 }; | |
__constant double sineTableGlobal[360] = { 0, 0.841471, 0.909297, 0.14112, -0.756802, -0.958924, -0.279415, 0.656987, 0.989358, 0.412118, -0.544021, -0.99999, -0.536573, 0.420167, 0.990607, 0.650288, -0.287903, -0.961397, -0.750987, 0.149877, 0.912945, 0.836656, -0.00885131, -0.84622, -0.905578, -0.132352, 0.762558, 0.956376, 0.270906, -0.663634, -0.988032, -0.404038, 0.551427, 0.999912, 0.529083, -0.428183, -0.991779, -0.643538, 0.296369, 0.963795, 0.745113, -0.158623, -0.916522, -0.831775, 0.0177019, 0.850904, 0.901788, 0.123573, -0.768255, -0.953753, -0.262375, 0.670229, 0.986628, 0.395925, -0.558789, -0.999755, -0.521551, 0.436165, 0.992873, 0.636738, -0.304811, -0.966118, -0.739181, 0.167356, 0.920026, 0.826829, -0.0265512, -0.85552, -0.897928, -0.114785, 0.773891, 0.951055, 0.253823, -0.676772, -0.985146, -0.387782, 0.566108, 0.99952, 0.513978, -0.444113, -0.993889, -0.629888, 0.313229, 0.968364, 0.73319, -0.176076, -0.923458, -0.821818, 0.0353983, 0.860069, 0.893997, 0.105988, -0.779466, -0.948282, | |
-0.245252, 0.683262, 0.983588, 0.379608, -0.573382, -0.999207, -0.506366, 0.452026, 0.994827, 0.622989, -0.321622, -0.970535, -0.727143, 0.184782, 0.926819, 0.816743, -0.0442427, -0.864551, -0.889996, -0.0971819, 0.78498, 0.945435, 0.236661, -0.689698, -0.981952, -0.371404, 0.580611, 0.998815, 0.498713, -0.459903, -0.995687, -0.61604, 0.329991, 0.97263, 0.721038, -0.193473, -0.930106, -0.811603, 0.0530836, 0.868966, 0.885925, 0.0883687, -0.790433, -0.942514, -0.228052, 0.69608, 0.98024, 0.363171, -0.587795, -0.998345, -0.491022, 0.467745, 0.996469, 0.609044, -0.338333, -0.974649, -0.714876, 0.20215, 0.933321, 0.806401, -0.0619203, -0.873312, -0.881785, -0.0795485, 0.795824, 0.93952, 0.219425, -0.702408, -0.97845, -0.35491, 0.594933, 0.997797, 0.483292, -0.47555, -0.997173, -0.602, 0.346649, 0.976591, 0.708659, -0.210811, -0.936462, -0.801135, 0.0707522, 0.87759, 0.877575, 0.0707222, -0.801153, -0.936451, -0.210781, 0.70868, 0.976584, 0.346621, -0.602024, -0.997171, | |
-0.475524, 0.483318, 0.997799, 0.594909, -0.354938, -0.978457, -0.702386, 0.219455, 0.93953, 0.795806, -0.0795786, -0.881799, -0.873297, -0.0618903, 0.806418, 0.93331, 0.20212, -0.714898, -0.974642, -0.338305, 0.609068, 0.996467, 0.467719, -0.491048, -0.998347, -0.587771, 0.363199, 0.980246, 0.696058, -0.228082, -0.942525, -0.790415, 0.0883987, 0.885939, 0.868951, 0.0530535, -0.811621, -0.930095, -0.193444, 0.721059, 0.972623, 0.329962, -0.616064, -0.995684, -0.459877, 0.498739, 0.998817, 0.580587, -0.371432, -0.981958, -0.689676, 0.236691, 0.945445, 0.784962, -0.0972119, -0.890009, -0.864536, -0.0442126, 0.81676, 0.926807, 0.184752, -0.727163, -0.970528, -0.321594, 0.623012, 0.994824, 0.451999, -0.506392, -0.999208, -0.573357, 0.379636, 0.983593, 0.68324, -0.245281, -0.948292, -0.779447, 0.106017, 0.89401, 0.860054, 0.0353682, -0.821835, -0.923447, -0.176046, 0.733211, 0.968357, 0.3132, -0.629911, -0.993885, -0.444086, 0.514004, 0.999521, 0.566083, -0.387809, -0.985151, | |
-0.67675, 0.253853, 0.951064, 0.773872, -0.114815, -0.897941, -0.855504, -0.026521, 0.826846, 0.920014, 0.167326, -0.739201, -0.96611, -0.304782, 0.636761, 0.992869, 0.436138, -0.521577, -0.999756, -0.558764, 0.395953, 0.986633, 0.670207, -0.262404, -0.953762, -0.768235, 0.123603, 0.901801, 0.850888, 0.0176718, -0.831791, -0.916509, -0.158593, 0.745133, 0.963787, 0.29634, -0.643561, -0.991775, -0.428155, 0.529108, 0.999912, 0.551402, -0.404065, -0.988036, -0.663611, 0.270935, 0.956385, 0.762539, -0.132382, -0.905591, -0.846204, -0.00882117, 0.836672, 0.912933, 0.149847, -0.751007, -0.961389, -0.287874, 0.650311, 0.990603, 0.42014, -0.536598, -0.99999, -0.543996, 0.412146, 0.989363, 0.656964, -0.279444, -0.958933, -0.756783, 0.14115, 0.90931, 0.841455, -3.01444e-05, -0.841487, -0.909285, -0.14109, 0.756822 }; | |
__kernel void circularHough( | |
__global float *img, | |
__global int *accumulator, | |
__global int *xDimension, | |
__global int *yDimension, | |
__global int *range, | |
__global float *maxPixVal) | |
{ | |
//Define Variables | |
int id = get_global_id(0); //Get Current Global ID | |
int minRadius = (int) ceil(((double)((*yDimension)*0.1))); | |
int radius=0; | |
int a=0; | |
int b=0; | |
int pos = 0; | |
//Get x,y,z index | |
int idx = id; | |
int zed = floor(((double)id) / (((double)(*xDimension)) * ((double)(*yDimension)))); | |
idx = id-(zed * (*xDimension) * (*yDimension)); | |
int wy = floor(((double)idx) / ((double)(*xDimension))); | |
int ex = idx % (*xDimension); | |
int val1 = (int) (wy + ((*xDimension)*ex)); | |
float val2= (float) (0.7*((float)(*maxPixVal))); | |
if ((img[val1] > 0) && (img[val1] >= val2)) { | |
for (int rad = 0; rad < (*range); rad++) { //Loop through radius range | |
radius = minRadius + rad - 1; | |
for (int tetha = 0; tetha < 360; tetha++) { //Every possible angle | |
a = (int) ( ceil( ((double)ex) - (((double)radius) * ((double)cosineTableGlobal[tetha])) )); | |
b = (int) ( ceil( ((double)wy) - (((double)radius) * ((double)sineTableGlobal[tetha])) )); | |
if (a > 0 && b > 0 && a <= (*xDimension) && b <= (*yDimension)) { | |
pos = (int) ((rad * (*xDimension) * (*yDimension)) + (b * (*xDimension)) + a); | |
atom_inc(&accumulator[pos]); | |
} | |
}//end of theta loop | |
}//end of radius loop | |
}//end of if | |
} |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
May I ask why you took the voting approach (1-read, many-writes) versus the accumulation approach (many-reads, 1-write).
In your case the many writes require an
atom_inc
, doesn't that hurt performances?