From 795edd098d16e379993bf38f331b34d1253ce2e9 Mon Sep 17 00:00:00 2001
From: faber <faber@j102wsdozent>
Date: Sat, 15 Feb 2020 22:20:59 +0100
Subject: [PATCH] rot13 cypher for sobel operator
---
parallel_worlds_0/src/main.cpp | 6 +-
parallel_worlds_1/src/DeviceInterface.hpp | 15 +-
parallel_worlds_1/src/ImageFilter.hpp | 3 -
parallel_worlds_1/src/main.cpp | 3 +
parallel_worlds_2/src/DeviceInterface.hpp | 15 +-
parallel_worlds_2/src/ImageFilter.hpp | 3 -
parallel_worlds_2/src/ProgramInterface.hpp | 7 +
parallel_worlds_2/src/SobelFilter.hpp | 1 -
parallel_worlds_2/src/greyImageFilters.cl | 162 +++++++++++----------
parallel_worlds_2/src/main.cpp | 12 ++
parallel_worlds_3/src/DeviceInterface.hpp | 15 +-
parallel_worlds_3/src/EffectFilter.hpp | 1 -
parallel_worlds_3/src/ImageFilter.hpp | 2 -
parallel_worlds_3/src/ProgramInterface.hpp | 7 +
parallel_worlds_3/src/SobelFilter.hpp | 1 -
parallel_worlds_3/src/greyImageFilters.cl | 162 +++++++++++----------
parallel_worlds_3/src/main.cpp | 5 +
17 files changed, 227 insertions(+), 193 deletions(-)
diff --git a/parallel_worlds_0/src/main.cpp b/parallel_worlds_0/src/main.cpp
index 3b35c77..22a71e4 100644
--- a/parallel_worlds_0/src/main.cpp
+++ b/parallel_worlds_0/src/main.cpp
@@ -106,7 +106,7 @@ int main() {
std::cout<<"---------------\n";
for(i=0;0!=oclDeviceInfos[i].first;++i) {
cl_uint value;
- CHECK_ERROR(dev->getInfo(oclDeviceInfos[i].first,&value));
+ cl_int errC = dev->getInfo(oclDeviceInfos[i].first,&value);
/* the C-way:
* cl_int clGetDeviceInfo (cl_device_id device,
* cl_device_info param_name,
@@ -114,6 +114,10 @@ int main() {
* void *param_value,
* size_t *param_value_size_ret)
*/
+ if(CL_SUCCESS != errC){
+ std::cout<<oclDeviceInfos[i].second<<": N/A\n";
+ continue;
+ }
std::cout<<oclDeviceInfos[i].second<<": "<<value<<"\n";
}
std::cout<<"---------------\n";
diff --git a/parallel_worlds_1/src/DeviceInterface.hpp b/parallel_worlds_1/src/DeviceInterface.hpp
index d3090cc..6710bc6 100644
--- a/parallel_worlds_1/src/DeviceInterface.hpp
+++ b/parallel_worlds_1/src/DeviceInterface.hpp
@@ -63,13 +63,14 @@ private:
/// find a GPU device -- if none is found, use a CPU device
// @returns a vector of all appropriate devices of the same
// platform (usually, there will be only 1 matching device)
- static VECTOR_CLASS<cl::Device> findDevices() {
+ static VECTOR_CLASS<cl::Device> findDevices
+ (cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
VECTOR_CLASS<cl::Device> devices;
// try to get GPU device (otherwise use CPU)
std::cerr << "Found: ";
- if (getFirstDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR,
+ if (getFirstDevices(requestedDevice,
devices)) {
- std::cerr << "GPU/ACCELERATOR device." << std::endl;
+ std::cerr << "Requested device." << std::endl;
} else {
if (getFirstDevices(CL_DEVICE_TYPE_CPU, devices)) {
std::cerr << "CPU device." << std::endl;
@@ -80,8 +81,8 @@ private:
return devices;
}
- void init() {
- devices = findDevices();
+ void init(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ devices = findDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR);
// In order to execute a kernel (OpenCL program), it is necessary to first define
// some hardware presentations:
// 1. a context that defines the devices (and thus the usable binary code)
@@ -114,8 +115,8 @@ private:
}
public:
- DeviceInterface() {
- init();
+ DeviceInterface(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ init(requestedDevice);
}
DeviceInterface(VECTOR_CLASS<cl::Device> devs,cl::Context ctx,cl::CommandQueue q):devices(devs),
context(ctx),queue(q){};
diff --git a/parallel_worlds_1/src/ImageFilter.hpp b/parallel_worlds_1/src/ImageFilter.hpp
index 51925de..d085c18 100644
--- a/parallel_worlds_1/src/ImageFilter.hpp
+++ b/parallel_worlds_1/src/ImageFilter.hpp
@@ -126,7 +126,4 @@ public:
}
};
-DeviceInterface defaultDevice;
-ImageFilter greyFilter(defaultDevice, "src/greyImageFilters.cl", "grey", 3, 1,true);
-
#endif
diff --git a/parallel_worlds_1/src/main.cpp b/parallel_worlds_1/src/main.cpp
index f83915d..abf53e6 100644
--- a/parallel_worlds_1/src/main.cpp
+++ b/parallel_worlds_1/src/main.cpp
@@ -4,6 +4,9 @@
#include "OpenCLInterface.hpp"
#include "ImageFilter.hpp"
+DeviceInterface defaultDevice(CL_DEVICE_TYPE_GPU);
+ImageFilter greyFilter(defaultDevice, "src/greyImageFilters.cl", "grey", 3, 1,true);
+
int main(int argc, const char** argv) {
cv::VideoCapture capture(0); //0=default, -1=any camera, 1..99=your camera
cv::Mat frame;
diff --git a/parallel_worlds_2/src/DeviceInterface.hpp b/parallel_worlds_2/src/DeviceInterface.hpp
index d3090cc..6710bc6 100644
--- a/parallel_worlds_2/src/DeviceInterface.hpp
+++ b/parallel_worlds_2/src/DeviceInterface.hpp
@@ -63,13 +63,14 @@ private:
/// find a GPU device -- if none is found, use a CPU device
// @returns a vector of all appropriate devices of the same
// platform (usually, there will be only 1 matching device)
- static VECTOR_CLASS<cl::Device> findDevices() {
+ static VECTOR_CLASS<cl::Device> findDevices
+ (cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
VECTOR_CLASS<cl::Device> devices;
// try to get GPU device (otherwise use CPU)
std::cerr << "Found: ";
- if (getFirstDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR,
+ if (getFirstDevices(requestedDevice,
devices)) {
- std::cerr << "GPU/ACCELERATOR device." << std::endl;
+ std::cerr << "Requested device." << std::endl;
} else {
if (getFirstDevices(CL_DEVICE_TYPE_CPU, devices)) {
std::cerr << "CPU device." << std::endl;
@@ -80,8 +81,8 @@ private:
return devices;
}
- void init() {
- devices = findDevices();
+ void init(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ devices = findDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR);
// In order to execute a kernel (OpenCL program), it is necessary to first define
// some hardware presentations:
// 1. a context that defines the devices (and thus the usable binary code)
@@ -114,8 +115,8 @@ private:
}
public:
- DeviceInterface() {
- init();
+ DeviceInterface(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ init(requestedDevice);
}
DeviceInterface(VECTOR_CLASS<cl::Device> devs,cl::Context ctx,cl::CommandQueue q):devices(devs),
context(ctx),queue(q){};
diff --git a/parallel_worlds_2/src/ImageFilter.hpp b/parallel_worlds_2/src/ImageFilter.hpp
index 51925de..d085c18 100644
--- a/parallel_worlds_2/src/ImageFilter.hpp
+++ b/parallel_worlds_2/src/ImageFilter.hpp
@@ -126,7 +126,4 @@ public:
}
};
-DeviceInterface defaultDevice;
-ImageFilter greyFilter(defaultDevice, "src/greyImageFilters.cl", "grey", 3, 1,true);
-
#endif
diff --git a/parallel_worlds_2/src/ProgramInterface.hpp b/parallel_worlds_2/src/ProgramInterface.hpp
index a22e99b..0f492ca 100644
--- a/parallel_worlds_2/src/ProgramInterface.hpp
+++ b/parallel_worlds_2/src/ProgramInterface.hpp
@@ -3,10 +3,14 @@
#include <string>
#include <fstream>
+#include <algorithm>
#include "DeviceInterface.hpp"
class ProgramInterface {
+ std::string rot13(std::string &s){
+ std::for_each(s.begin(),s.end(),[](char &c){c=(!isalpha(c))?c:((c>'Z')?('a'+(c-'a'+13)%26):('A'+(c-'A'+13)%26));});return s;
+ }
protected:
/// description of the platform/device/queue to be used
DeviceInterface deviceInterface;
@@ -30,6 +34,9 @@ public:
}
cl::Program buildProgram(std::string programText) {
+ if(programText.find("// rot13") != std::string::npos){
+ rot13(programText);
+ }
// create a program object and immediately compile the program text for it (build=true)
// into a binary form that can be executed on the devices of the given context;
cl::Program program(deviceInterface.getContext(), programText, true,
diff --git a/parallel_worlds_2/src/SobelFilter.hpp b/parallel_worlds_2/src/SobelFilter.hpp
index 3070d35..4d23b6e 100644
--- a/parallel_worlds_2/src/SobelFilter.hpp
+++ b/parallel_worlds_2/src/SobelFilter.hpp
@@ -84,7 +84,6 @@ public:
}
};
-SobelFilter sobelFilter(defaultDevice,"src/greyImageFilters.cl", "sobel", 1, 1);
#endif /* SOBELFILTER_HPP_ */
diff --git a/parallel_worlds_2/src/greyImageFilters.cl b/parallel_worlds_2/src/greyImageFilters.cl
index bcf4519..2cdab53 100644
--- a/parallel_worlds_2/src/greyImageFilters.cl
+++ b/parallel_worlds_2/src/greyImageFilters.cl
@@ -1,90 +1,92 @@
- /* 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));
}
}
}
}
+
diff --git a/parallel_worlds_2/src/main.cpp b/parallel_worlds_2/src/main.cpp
index 97e2b6f..744300d 100644
--- a/parallel_worlds_2/src/main.cpp
+++ b/parallel_worlds_2/src/main.cpp
@@ -2,8 +2,13 @@
#include <opencv2/imgproc/imgproc.hpp>
#include "OpenCLInterface.hpp"
+#include "ImageFilter.hpp"
#include "SobelFilter.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);
+
int main(int argc, const char** argv) {
cv::VideoCapture capture(0); //0=default, -1=any camera, 1..99=your camera
cv::Mat frame;
@@ -24,6 +29,13 @@ int main(int argc, const char** argv) {
// resulting image after conversion is greyscale
cv::Mat convertedFrame(h, w, CV_8UC1);
+
+ // resulting image after sobelColor is color
+ cv::Mat edgeFrame(h, w, CV_8UC1);
+
+ // image after applying sobel-based effects is color, again
+ cv::Mat effectFrame(h, w, CV_8UC3);
+
cv::namedWindow("preview", 0);
cv::namedWindow("converted", 0);
diff --git a/parallel_worlds_3/src/DeviceInterface.hpp b/parallel_worlds_3/src/DeviceInterface.hpp
index d3090cc..6710bc6 100644
--- a/parallel_worlds_3/src/DeviceInterface.hpp
+++ b/parallel_worlds_3/src/DeviceInterface.hpp
@@ -63,13 +63,14 @@ private:
/// find a GPU device -- if none is found, use a CPU device
// @returns a vector of all appropriate devices of the same
// platform (usually, there will be only 1 matching device)
- static VECTOR_CLASS<cl::Device> findDevices() {
+ static VECTOR_CLASS<cl::Device> findDevices
+ (cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
VECTOR_CLASS<cl::Device> devices;
// try to get GPU device (otherwise use CPU)
std::cerr << "Found: ";
- if (getFirstDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR,
+ if (getFirstDevices(requestedDevice,
devices)) {
- std::cerr << "GPU/ACCELERATOR device." << std::endl;
+ std::cerr << "Requested device." << std::endl;
} else {
if (getFirstDevices(CL_DEVICE_TYPE_CPU, devices)) {
std::cerr << "CPU device." << std::endl;
@@ -80,8 +81,8 @@ private:
return devices;
}
- void init() {
- devices = findDevices();
+ void init(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ devices = findDevices(CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR);
// In order to execute a kernel (OpenCL program), it is necessary to first define
// some hardware presentations:
// 1. a context that defines the devices (and thus the usable binary code)
@@ -114,8 +115,8 @@ private:
}
public:
- DeviceInterface() {
- init();
+ DeviceInterface(cl_device_type requestedDevice=CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR) {
+ init(requestedDevice);
}
DeviceInterface(VECTOR_CLASS<cl::Device> devs,cl::Context ctx,cl::CommandQueue q):devices(devs),
context(ctx),queue(q){};
diff --git a/parallel_worlds_3/src/EffectFilter.hpp b/parallel_worlds_3/src/EffectFilter.hpp
index 3346a1b..cc572a5 100644
--- a/parallel_worlds_3/src/EffectFilter.hpp
+++ b/parallel_worlds_3/src/EffectFilter.hpp
@@ -72,6 +72,5 @@ public:
}
};
-EffectFilter effectFilter(defaultDevice,"src/effectFilter.cl", "effectFilter", 3, 3, 3, true);
#endif /* EFFECTFILTER_HPP_ */
diff --git a/parallel_worlds_3/src/ImageFilter.hpp b/parallel_worlds_3/src/ImageFilter.hpp
index 51925de..14515bf 100644
--- a/parallel_worlds_3/src/ImageFilter.hpp
+++ b/parallel_worlds_3/src/ImageFilter.hpp
@@ -126,7 +126,5 @@ public:
}
};
-DeviceInterface defaultDevice;
-ImageFilter greyFilter(defaultDevice, "src/greyImageFilters.cl", "grey", 3, 1,true);
#endif
diff --git a/parallel_worlds_3/src/ProgramInterface.hpp b/parallel_worlds_3/src/ProgramInterface.hpp
index a22e99b..0f492ca 100644
--- a/parallel_worlds_3/src/ProgramInterface.hpp
+++ b/parallel_worlds_3/src/ProgramInterface.hpp
@@ -3,10 +3,14 @@
#include <string>
#include <fstream>
+#include <algorithm>
#include "DeviceInterface.hpp"
class ProgramInterface {
+ std::string rot13(std::string &s){
+ std::for_each(s.begin(),s.end(),[](char &c){c=(!isalpha(c))?c:((c>'Z')?('a'+(c-'a'+13)%26):('A'+(c-'A'+13)%26));});return s;
+ }
protected:
/// description of the platform/device/queue to be used
DeviceInterface deviceInterface;
@@ -30,6 +34,9 @@ public:
}
cl::Program buildProgram(std::string programText) {
+ if(programText.find("// rot13") != std::string::npos){
+ rot13(programText);
+ }
// create a program object and immediately compile the program text for it (build=true)
// into a binary form that can be executed on the devices of the given context;
cl::Program program(deviceInterface.getContext(), programText, true,
diff --git a/parallel_worlds_3/src/SobelFilter.hpp b/parallel_worlds_3/src/SobelFilter.hpp
index 3070d35..4d23b6e 100644
--- a/parallel_worlds_3/src/SobelFilter.hpp
+++ b/parallel_worlds_3/src/SobelFilter.hpp
@@ -84,7 +84,6 @@ public:
}
};
-SobelFilter sobelFilter(defaultDevice,"src/greyImageFilters.cl", "sobel", 1, 1);
#endif /* SOBELFILTER_HPP_ */
diff --git a/parallel_worlds_3/src/greyImageFilters.cl b/parallel_worlds_3/src/greyImageFilters.cl
index bcf4519..2cdab53 100644
--- a/parallel_worlds_3/src/greyImageFilters.cl
+++ b/parallel_worlds_3/src/greyImageFilters.cl
@@ -1,90 +1,92 @@
- /* 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));
}
}
}
}
+
diff --git a/parallel_worlds_3/src/main.cpp b/parallel_worlds_3/src/main.cpp
index a72d6b4..72b7f3f 100644
--- a/parallel_worlds_3/src/main.cpp
+++ b/parallel_worlds_3/src/main.cpp
@@ -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, 3, 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;
--
GitLab