Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • gpgpu/parallel_worlds
  • plustig/parallel_worlds
  • peter.faber/parallel-worlds-cuda
  • gfx/parallel_worlds
  • ai01031/parallel_worlds
5 results
Show changes
Hide whitespace changes
Inline Side-by-side
parallel_worlds_3/src/SobelFilter.hpp
View file @ bf365ef0
......@@ -84,7 +84,6 @@ public:
}
};
SobelFilter sobelFilter(defaultDevice,"src/greyImageFilters.cl", "sobel", 1, 1);
#endif /* SOBELFILTER_HPP_ */
parallel_worlds_3/src/cl2.hpp 0 → 100644
View file @ bf365ef0
Source diff could not be displayed: it is too large. Options to address this: view the blob.
parallel_worlds_3/src/greyImageFilters.cl
View file @ bf365ef0
/* start the kernel with one work-item per pixel
** first work-dimension (0) is image width (x)
// rot13
/* fgneg gur xreary jvgu bar jbex-vgrz cre cvkry
** svefg jbex-qvzrafvba (0) vf vzntr jvqgu (k)
*/
__kernel void grey(__global unsigned char *inImg,
__global unsigned char *outImg,
__private unsigned int w,__private unsigned int h) {
__private unsigned int x;__private unsigned int y;
x = get_global_id(0);
y = get_global_id(1);
if(y<h) {
if(x<w) {
// greyscale conversion (c.f. http://en.wikipedia.org/wiki/Grayscale)
// Y = 0.2126R + 0.7152G + 0.0722B
outImg[x+w*y] = 0.0722 * inImg[3*(x+w*y)] /* blue */
+ 0.7152 * inImg[3*(x+w*y)+1] /* green */
+ 0.2126 * inImg[3*(x+w*y)+2]; /* red */
__xreary ibvq terl(__tybony hafvtarq pune *vaVzt,
__tybony hafvtarq pune *bhgVzt,
__cevingr hafvtarq vag j,__cevingr hafvtarq vag u) {
__cevingr hafvtarq vag k;__cevingr hafvtarq vag l;
k = trg_tybony_vq(0);
l = trg_tybony_vq(1);
vs(l<u) {
vs(k<j) {
// terlfpnyr pbairefvba (p.s. uggc://ra.jvxvcrqvn.bet/jvxv/Tenlfpnyr)
// L = 0.2126E + 0.7152T + 0.0722O
bhgVzt[k+j*l] = 0.0722 * vaVzt[3*(k+j*l)] /* oyhr */
+ 0.7152 * vaVzt[3*(k+j*l)+1] /* terra */
+ 0.2126 * vaVzt[3*(k+j*l)+2]; /* erq */
}
}
}
#define ARR(A,x,y,maxX) (A[(x)+(y)*(maxX)])
/// sobel filter (cf. http://en.wikipedia.org/wiki/Sobel_operator):
// detect edges by computing the convolution
// with matrix {{-1,0,1},{-2,0,2},{-1,0,1}} in x- and y- direction;
// the result is computed as c*sqrt(G_x^2 + G_y^2) (where G_x/G_y
// is the convolution with the above matrix);
// this computation is only done for interior pixels -- the edges
// of the image are blacked out;
// @param inImg pointer to the input grey image in device memory
// @param outImg pointer to the output grey image in device memory
// @param w width of image
// @param h height of image
// @param c coefficient by which to multiply the actual convolution
// @param img image portion for computation -- to be shared between
// work-items of a work-group (each work-item writes exactly
// 1 pixel of img)
// Note: img has to be passed via Kernel::setArg(), because its size
// depends on the size of the work-group (otherwise it could have been
// defined inside the kernel)
__kernel void sobel(__global unsigned char *inImg,
__global unsigned char *outImg,
unsigned int w,unsigned int h,
float c,
__local unsigned char *img){
// coordinates of input pixel in cache array img
unsigned int xCache;unsigned int yCache;
// coordinates of pixel in input/output image
unsigned int x;unsigned int y;
// number of output pixels per work-group in x/y direction
// will evaluate to 8, since the kernel will be started on a
// 10 * 10 work-group
unsigned int numOutX; unsigned int numOutY;
numOutX = get_local_size(0) - 2; numOutY = get_local_size(1) - 2;
x = get_group_id(0) * numOutX + get_local_id(0);
y = get_group_id(1) * numOutY + get_local_id(1);
xCache = get_local_id(0); yCache = get_local_id(1);
if(x<w && y<h){
// read pixels from original image into cache
ARR(img,xCache,yCache,get_local_size(0)) = ARR(inImg,x,y,w);
// border pixels are all black
if(0==x||0==y||w-1==x||h-1==y){
ARR(outImg,x,y,w) = 0;
#qrsvar NEE(N,k,l,znkK) (N[(k)+(l)*(znkK)])
/// fbory svygre (ps. uggc://ra.jvxvcrqvn.bet/jvxv/Fbory_bcrengbe):
// qrgrpg rqtrf ol pbzchgvat gur pbaibyhgvba
// jvgu zngevk {{-1,0,1},{-2,0,2},{-1,0,1}} va k- naq l- qverpgvba;
// gur erfhyg vf pbzchgrq nf p*fdeg(T_k^2 + T_l^2) (jurer T_k/T_l
// vf gur pbaibyhgvba jvgu gur nobir zngevk);
// guvf pbzchgngvba vf bayl qbar sbe vagrevbe cvkryf -- gur rqtrf
// bs gur vzntr ner oynpxrq bhg;
// @cnenz vaVzt cbvagre gb gur vachg terl vzntr va qrivpr zrzbel
// @cnenz bhgVzt cbvagre gb gur bhgchg terl vzntr va qrivpr zrzbel
// @cnenz j jvqgu bs vzntr
// @cnenz u urvtug bs vzntr
// @cnenz p pbrssvpvrag ol juvpu gb zhygvcyl gur npghny pbaibyhgvba
// @cnenz vzt vzntr cbegvba sbe pbzchgngvba -- gb or funerq orgjrra
// jbex-vgrzf bs n jbex-tebhc (rnpu jbex-vgrz jevgrf rknpgyl
// 1 cvkry bs vzt)
// Abgr: vzt unf gb or cnffrq ivn Xreary::frgNet(), orpnhfr vgf fvmr
// qrcraqf ba gur fvmr bs gur jbex-tebhc (bgurejvfr vg pbhyq unir orra
// qrsvarq vafvqr gur xreary)
__xreary ibvq fbory(__tybony hafvtarq pune *vaVzt,
__tybony hafvtarq pune *bhgVzt,
hafvtarq vag j,hafvtarq vag u,
sybng p,
__ybpny hafvtarq pune *vzt){
// pbbeqvangrf bs vachg cvkry va pnpur neenl vzt
hafvtarq vag kPnpur;hafvtarq vag lPnpur;
// pbbeqvangrf bs cvkry va vachg/bhgchg vzntr
hafvtarq vag k;hafvtarq vag l;
// ahzore bs bhgchg cvkryf cre jbex-tebhc va k/l qverpgvba
// jvyy rinyhngr gb 8, fvapr gur xreary jvyy or fgnegrq ba n
// 10 * 10 jbex-tebhc
hafvtarq vag ahzBhgK; hafvtarq vag ahzBhgL;
ahzBhgK = trg_ybpny_fvmr(0) - 2; ahzBhgL = trg_ybpny_fvmr(1) - 2;
k = trg_tebhc_vq(0) * ahzBhgK + trg_ybpny_vq(0);
l = trg_tebhc_vq(1) * ahzBhgL + trg_ybpny_vq(1);
kPnpur = trg_ybpny_vq(0); lPnpur = trg_ybpny_vq(1);
vs(k<j && l<u){
// ernq cvkryf sebz bevtvany vzntr vagb pnpur
NEE(vzt,kPnpur,lPnpur,trg_ybpny_fvmr(0)) = NEE(vaVzt,k,l,j);
// obeqre cvkryf ner nyy oynpx
vs(0==k||0==l||j-1==k||u-1==l){
NEE(bhgVzt,k,l,j) = 0;
}
}
// wait for all work-items to finish copying
barrier(CLK_LOCAL_MEM_FENCE);
if(x<w-1 && y<h-1){
// compute result value and write it back to device memory
// (but only for interior pixels, i.e. 1<=id<=max-1)
if(xCache > 0 && xCache < get_local_size(0) - 1){
if(yCache > 0 && yCache < get_local_size(1) - 1){
__private float G_x =
-ARR(img,xCache-1,yCache-1,get_local_size(0))
-2*ARR(img,xCache-1,yCache,get_local_size(0))
-ARR(img,xCache-1,yCache+1,get_local_size(0))
+ARR(img,xCache+1,yCache-1,get_local_size(0))
+2*ARR(img,xCache+1,yCache,get_local_size(0))
+ARR(img,xCache+1,yCache+1,get_local_size(0));
__private float G_y =
-ARR(img,xCache-1,yCache-1,get_local_size(0))
-2*ARR(img,xCache,yCache-1,get_local_size(0))
-ARR(img,xCache+1,yCache-1,get_local_size(0))
+ARR(img,xCache-1,yCache+1,get_local_size(0))
+2*ARR(img,xCache,yCache+1,get_local_size(0))
+ARR(img,xCache+1,yCache+1,get_local_size(0));
// sqrt is a predefined OpenCL function!
ARR(outImg,x,y,w) = (unsigned char) (c * sqrt(G_x*G_x + G_y*G_y));
// jnvg sbe nyy jbex-vgrzf gb svavfu pbclvat
oneevre(PYX_YBPNY_ZRZ_SRAPR);
vs(k<j-1 && l<u-1){
// pbzchgr erfhyg inyhr naq jevgr vg onpx gb qrivpr zrzbel
// (ohg bayl sbe vagrevbe cvkryf, v.r. 1<=vq<=znk-1)
vs(kPnpur > 0 && kPnpur < trg_ybpny_fvmr(0) - 1){
vs(lPnpur > 0 && lPnpur < trg_ybpny_fvmr(1) - 1){
__cevingr sybng T_k =
-NEE(vzt,kPnpur-1,lPnpur-1,trg_ybpny_fvmr(0))
-2*NEE(vzt,kPnpur-1,lPnpur,trg_ybpny_fvmr(0))
-NEE(vzt,kPnpur-1,lPnpur+1,trg_ybpny_fvmr(0))
+NEE(vzt,kPnpur+1,lPnpur-1,trg_ybpny_fvmr(0))
+2*NEE(vzt,kPnpur+1,lPnpur,trg_ybpny_fvmr(0))
+NEE(vzt,kPnpur+1,lPnpur+1,trg_ybpny_fvmr(0));
__cevingr sybng T_l =
-NEE(vzt,kPnpur-1,lPnpur-1,trg_ybpny_fvmr(0))
-2*NEE(vzt,kPnpur,lPnpur-1,trg_ybpny_fvmr(0))
-NEE(vzt,kPnpur+1,lPnpur-1,trg_ybpny_fvmr(0))
+NEE(vzt,kPnpur-1,lPnpur+1,trg_ybpny_fvmr(0))
+2*NEE(vzt,kPnpur,lPnpur+1,trg_ybpny_fvmr(0))
+NEE(vzt,kPnpur+1,lPnpur+1,trg_ybpny_fvmr(0));
// fdeg vf n cerqrsvarq BcraPY shapgvba!
NEE(bhgVzt,k,l,j) = (hafvtarq pune) (p * fdeg(T_k*T_k + T_l*T_l));
}
}
}
}
parallel_worlds_3/src/main.cpp
View file @ bf365ef0
......@@ -4,6 +4,11 @@
#include "OpenCLInterface.hpp"
#include "EffectFilter.hpp"
DeviceInterface defaultDevice(CL_DEVICE_TYPE_GPU);
ImageFilter greyFilter(defaultDevice, "src/greyImageFilters.cl", "grey", 3, 1,true);
SobelFilter sobelFilter(defaultDevice,"src/greyImageFilters.cl", "sobel", 1, 1);
EffectFilter effectFilter(defaultDevice,"src/effectFilter.cl", "effectFilter", 3, 1, 3, true);
int main(int argc, const char** argv) {
cv::VideoCapture capture(0); //0=default, -1=any camera, 1..99=your camera
cv::Mat frame;
......