diff --git a/KernelTester.pro b/KernelTester.pro index 1a6e396..dc002cb 100644 --- a/KernelTester.pro +++ b/KernelTester.pro @@ -20,12 +20,11 @@ SOURCES += $$files(source/*.cpp, true) \ MainWindow.cpp \ LIBS+= -L"$$PWD/header/CL" -lOpenCL -LIBS += -lmatio -LIBS += -L/usr/local/lib + +INCLUDEPATH += "$$PWD/header" HEADERS += $$files(header/*.h, true) \ MainWindow.h \ - header/ScenarioParams.h \ header/CL/cl.h \ header/CL/cl.hpp \ header/CL/cl_d3d9_ext.h \ @@ -42,11 +41,6 @@ HEADERS += $$files(header/*.h, true) \ header/CL/cl2.hpp \ header/CL/opencl.h \ - - -INCLUDEPATH += /home/hasis/Downloads/Compressed/matio-1.5.17/src/ - - FORMS += \ MainWindow.ui diff --git a/MainWindow.cpp b/MainWindow.cpp index e35fde2..b4e3235 100644 --- a/MainWindow.cpp +++ b/MainWindow.cpp @@ -5,13 +5,15 @@ #include #include -#include "header/strategies/ScanConversion.h" -#include "header/strategies/Rejection.h" +#include "model/processor/BIP.h" +#include "model/processor/strategies/DynCont.h" +#include "model/processor/strategies/GrayMap.h" +#include "model/processor/strategies/TintMap.h" +#include "model/processor/strategies/Sri.h" +#include "model/processor/strategies/ScanConversion.h" +#include "model/processor/strategies/Rejection.h" #include "header/FileHelper.h" -#define OUT_WIDTH 1000 -#define OUT_HEIGHT 1000 - MainWindow* MainWindow::_instance; MainWindow::MainWindow(QWidget *parent) @@ -23,17 +25,18 @@ MainWindow::MainWindow(QWidget *parent) ui->setupUi(this); _CLContext = _openCLHelper.getContext(); - CLQueue = _openCLHelper.createCommandQueue(_CLContext, - _openCLHelper - .getDevicesByContext( - _CLContext)[0]); - _scratchPad = new Buffer(_CLContext, CL_MEM_READ_WRITE, 1024); - _scratch = new double[128]; - CLQueue.enqueueWriteBuffer(*_scratchPad, CL_TRUE, 0, 1024, _scratch); + BIP::getInstance()->init(_CLContext, _openCLHelper); + + _scratchPad = new Buffer(_CLContext, CL_MEM_READ_WRITE, SCRATCH_PAD_SIZE * sizeof(float)); +// _scratch = new float[SCRATCH_PAD_SIZE]; +// cl_int error = BIP::getInstance()->CLQueue.enqueueWriteBuffer(*_scratchPad, CL_TRUE, 0, SCRATCH_PAD_SIZE * sizeof(float), _scratch); +// qDebug() << "hi: " << error; registerStrategies(); _strategy = Q_NULLPTR; + + connect(this, &MainWindow::console, this, [this](QString message){ ui->rtb_console->append(">>> " + message);}); } MainWindow::~MainWindow() @@ -59,7 +62,7 @@ void MainWindow::abort() void MainWindow::addToConsole(QString message) { - ui->rtb_console->append(">>> " + message); + emit this->console(message); } bool MainWindow::checkPath(QString path, QString pathName) @@ -78,6 +81,11 @@ void MainWindow::registerStrategies() { REGISTER_STRATEGY(ScanConversion) REGISTER_STRATEGY(Rejection) + REGISTER_STRATEGY(Sri) + REGISTER_STRATEGY(TintMap) + REGISTER_STRATEGY(GrayMap) + REGISTER_STRATEGY(DynCont) + } void MainWindow::pushBackStrategy(const QString strategyName, const QString kernelFolder) @@ -93,12 +101,114 @@ void MainWindow::pushBackStrategy(const QString strategyName, const QString kern static_cast( QMetaType::metaObjectForType(id)->newInstance( Q_ARG(const Context, _CLContext), - Q_ARG(const QString, kernelFolder) + Q_ARG(const QString, kernelFolder), + Q_ARG(const QObject*, this) )); addToConsole("Strategy was created"); } +void MainWindow::readParam(QString path, int mode) +{ + QFile file(path); + file.open(QIODevice::ReadOnly); + + char data[8192]; + //first line is extra + file.readLine(data, 8192); + QString str(data); + str = str.remove(str.length() - 1, 1); + auto sl = str.split(","); + int index = 0; + int temp; + + switch (mode) + { + //sc + case 0: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.linear, INP2MYFLT(temp) == 0); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.depth, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.probeRadius, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.fieldOfView, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.probeFieldOfView, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.virtualOriginalZ, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.minScanAx, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.minScanAz, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.maxScanAx, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.startDepth, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.steering, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.rxLineNo, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.rxFocusPointNo, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.rxLineDaz, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.rxPointDax, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.virtualConvex, INP2MYFLT(temp) != 0); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.vcMaxTheta, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.angle, INP2MYFLT(temp)); + break; + case 1: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.rejectThreshold, INP2MYFLT(temp)); + break; + case 2: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.sri, INP2MYFLT(temp)); + break; + case 3: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.tintMapSelector, INP2MYFLT(temp)); + break; + case 4: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.grayMapSelector, INP2MYFLT(temp)); + break; + case 5: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.dynContSelector, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.dynContGain, INP2MYFLT(temp)); + + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.compressionType, INP2MYFLT(temp)); + break; + } +} + void MainWindow::on_btn_browseInput_clicked() { @@ -154,7 +264,7 @@ void MainWindow::on_btn_test_clicked() auto dataPath = inputDir.path() + "/" + filename; auto paramPath = inputDir.path() + "/params/" + filename; - _strategy->ReadParams(paramPath, &_scenGenOutput); + readParam(paramPath, ui->cb_kernelName->currentIndex()); QByteArray arr; quint64 width; @@ -166,6 +276,9 @@ void MainWindow::on_btn_test_clicked() return; } + OUT_WIDTH = ui->cb_kernelName->currentIndex() ? width : 1000; + OUT_HEIGHT = ui->cb_kernelName->currentIndex() ? height : 1000; + ImageFormat format; format.image_channel_order = CL_RGBA; #ifdef USE_DBL @@ -177,24 +290,24 @@ void MainWindow::on_btn_test_clicked() Image2D* inFrame = static_cast(_openCLHelper .frame2CLFrame(_CLContext, format, QVector{width, height}, true)); - CLQueue.enqueueWriteImage(*inFrame, CL_TRUE, array {0, 0, 0}, + BIP::getInstance()->CLQueue.enqueueWriteImage(*inFrame, CL_TRUE, array {0, 0, 0}, array {width, height, 1}, width * sizeof (myflt), 0, arr.data()); - update_field(&_scenGenOutput.outputWidth, OUT_WIDTH); - update_field(&_scenGenOutput.outputHeight, OUT_HEIGHT); + update_field(&_scenGenOutput.outputWidth, (uint)OUT_WIDTH); + update_field(&_scenGenOutput.outputHeight, (uint)OUT_HEIGHT); _strategy->cpuProcess(_scenGenOutput); auto outFrame = _strategy->processKernel(inFrame, _scratchPad); char *out; - out = (char*)malloc(OUT_WIDTH * OUT_WIDTH * sizeof (char) * sizeof (myflt)); - CLQueue.enqueueReadImage(*outFrame, + out = (char*)malloc(OUT_WIDTH * OUT_HEIGHT * sizeof (myflt)); + BIP::getInstance()->CLQueue.enqueueReadImage(*outFrame, CL_TRUE, array {0, 0, 0}, - array {OUT_WIDTH, OUT_HEIGHT, 1}, + array {(quint64)OUT_WIDTH, (quint64)OUT_HEIGHT, 1}, OUT_WIDTH * sizeof (myflt), 0, out); @@ -238,6 +351,8 @@ void MainWindow::on_btn_compare_clicked() auto outputs = outputDir.entryList(QStringList() << "*.csv", QDir::Files); foreach(QString filename, outputs) { + OUT_HEIGHT = 0; + FileHelper::GetFileSIze(matlabDir.path() + "/" + filename, &OUT_WIDTH, &OUT_HEIGHT); addToConsole("Comaring " + filename); auto cppout = new myint[OUT_WIDTH * OUT_HEIGHT]; auto matout = new myint[OUT_WIDTH * OUT_HEIGHT]; @@ -249,6 +364,8 @@ void MainWindow::on_btn_compare_clicked() if(!QFile::exists(matlabDir.path() + "/" + filename)) { addToConsole("matlab file for " + filename + " does not exist"); + delete [] cppout; + delete [] matout; continue; } FileHelper::ReadInputFile(matout, matlabDir.path() + "/" + filename, &w2, &h2); @@ -256,16 +373,20 @@ void MainWindow::on_btn_compare_clicked() if (w1 != OUT_WIDTH || w2 != OUT_WIDTH) { addToConsole("Width mismatch"); + delete [] cppout; + delete [] matout; continue; } if (h1 != OUT_HEIGHT || h2 != OUT_HEIGHT) { addToConsole("Height mismatch"); + delete [] cppout; + delete [] matout; continue; } ulong max = 0; - uint imax, jmax; + uint imax = 0, jmax = 0; for(int i = 0; i < w1; i++) { for(int j = 0; j < h1; j++) @@ -281,12 +402,17 @@ void MainWindow::on_btn_compare_clicked() if(max > totalMax) totalMax = max; - addToConsole(QString::number(max) - + " ["+ QString::number(jmax) - + ", " + QString::number(imax) + "]"); - addToConsole(QString::number(cppout[jmax * w1 + imax], 16) + " -- " + - QString::number(matout[jmax * w1 + imax], 16)); - addToConsole("*************************\r\n"); + if(max > 0) + { + addToConsole(QString::number(max) + + " ["+ QString::number(jmax) + + ", " + QString::number(imax) + "]"); + addToConsole(QString::number(cppout[jmax * w1 + imax], 16) + " -- " + + QString::number(matout[jmax * w1 + imax], 16)); + } + addToConsole("*************************\r\n"); + delete [] cppout; + delete [] matout; } addToConsole("Compare done, total max: " + QString::number(totalMax)); diff --git a/MainWindow.h b/MainWindow.h index a407011..c301ce8 100644 --- a/MainWindow.h +++ b/MainWindow.h @@ -6,8 +6,9 @@ #include #include -#include "header/OpenCLHelper.h" -#include "header/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" +#include "model/processor/IProcessStrategy.h" +#include "ScenarioParams.h" QT_BEGIN_NAMESPACE @@ -27,7 +28,7 @@ public: static MainWindow* getInstance() { return _instance; } - cl::CommandQueue CLQueue; +// cl::CommandQueue CLQueue; private slots: void on_btn_browseInput_clicked(); @@ -64,8 +65,16 @@ private: IProcessStrategy* _strategy; ScenGenOutput_t _scenGenOutput; + void readParam(QString path, int mode); + Buffer* _scratchPad; - double* _scratch; + float* _scratch; + + quint64 OUT_WIDTH; + quint64 OUT_HEIGHT; + +signals: + void console(QString); }; #endif // MAINWINDOW_H diff --git a/header/FileHelper.h b/header/FileHelper.h index b130b61..b362833 100644 --- a/header/FileHelper.h +++ b/header/FileHelper.h @@ -12,6 +12,7 @@ public: FileHelper(){} static bool ReadInputFile(QByteArray& arr, QString path, quint64* width, quint64* height); static bool ReadInputFile(myint* arr, QString path, quint64* width, quint64* height); + static void GetFileSIze(QString path, quint64* width, quint64* height); static bool WriteOutputFile(char* arr, QString path, quint64 width, quint64 height); static bool WriteOutputFile(myflt* arr, QString path, quint64 width, quint64 height); }; diff --git a/header/ScenarioParams.h b/header/ScenarioParams.h index 1ad13af..7bb8f2c 100644 --- a/header/ScenarioParams.h +++ b/header/ScenarioParams.h @@ -1,6 +1,8 @@ #ifndef SCENARIOPARAMS_H #define SCENARIOPARAMS_H +#include + template struct field_t { @@ -12,20 +14,31 @@ typedef struct ScenGenOutput_t { field_t linear; field_t virtualConvex; - field_t depth; - field_t probeRadius; - field_t fieldOfView; - field_t startDepth; + field_t depth; + field_t probeRadius; + field_t fieldOfView; + field_t probeFieldOfView; + field_t startDepth; field_t rxLineNo; field_t rxFocusPointNo; - field_t rxLineDaz; - field_t rxPointDax; - field_t vcMaxTheta; - field_t angle; - field_t steering; - field_t outputWidth; - field_t outputHeight; - field_t rejectThreshold; + field_t rxLineDaz; + field_t rxPointDax; + field_t vcMaxTheta; + field_t angle; + field_t steering; + field_t minScanAx; + field_t minScanAz; + field_t maxScanAx; + field_t virtualOriginalZ; + field_t outputWidth; + field_t outputHeight; + field_t compressionType; + field_t dynContSelector; + field_t dynContGain; + field_t grayMapSelector; + field_t tintMapSelector; + field_t sri; + field_t rejectThreshold; }ScenGenOutput_t; template diff --git a/header/Utils.h b/header/Utils.h index 27b62e2..4121cd0 100644 --- a/header/Utils.h +++ b/header/Utils.h @@ -2,7 +2,7 @@ #define UTILS_H -#define USE_DBL +//#define USE_DBL #ifdef USE_DBL diff --git a/header/model/processor/BIP.h b/header/model/processor/BIP.h new file mode 100644 index 0000000..dbe52fd --- /dev/null +++ b/header/model/processor/BIP.h @@ -0,0 +1,22 @@ +#ifndef BIP_H +#define BIP_H + +#include "utils/OpenCLHelper.h" + +#define SCRATCH_PAD_SIZE 2000 + +class BIP +{ +public: + BIP(); + static BIP* getInstance(); + + cl::CommandQueue CLQueue; + + void init(cl::Context &context, OpenCLHelper &openCLHelper); + +private: + static BIP* _instance; +}; + +#endif // BIP_H diff --git a/header/IProcessStrategy.h b/header/model/processor/IProcessStrategy.h similarity index 60% rename from header/IProcessStrategy.h rename to header/model/processor/IProcessStrategy.h index 37851ec..ab63e9f 100644 --- a/header/IProcessStrategy.h +++ b/header/model/processor/IProcessStrategy.h @@ -1,12 +1,14 @@ #ifndef IPROCESSSTRATEGY_H #define IPROCESSSTRATEGY_H +//#include "CL/cl2.hpp" #include #include -#include "OpenCLHelper.h" -#include "Utils.h" +//#include "Processor.h" +#include "utils/OpenCLHelper.h" #include "ScenarioParams.h" + #ifdef USE_DBL typedef double myflt; #else @@ -22,28 +24,25 @@ class IProcessStrategy : public QObject public: virtual void cpuProcess(ScenGenOutput_t parameters) = 0; virtual void finalize() = 0; - virtual void ReadParams(QString path, ScenGenOutput_t* params) = 0; - virtual Image* processKernel(Image* inputFrame, Buffer* scratchPad) = 0; + virtual Image* processKernel(Image* inputFrame, Buffer* scrathPad) = 0; +private: +// Context _CLContext; +// Device getDefaultDevice(const Context& context); protected: Kernel _kernel; IProcessStrategy(const Context context, const QString kernelPath, - const QString kernelName); + const QString kernelName, + const QObject* parent); + Context _CLContext; Device getDefaultDevice(const Context& context); - - - + const QObject* _parent; OpenCLHelper _openCLHelper; -signals: - void sgl_outputReady(Image* frames, Buffer* scratchPad); - -public slots: - void slt_process(Image* frames, Buffer* scratchPad); - }; +//Q_DECLARE_METATYPE(IProcessStrategy*); #endif // IPROCESSSTRATEGY_H diff --git a/header/model/processor/strategies/DynCont.h b/header/model/processor/strategies/DynCont.h new file mode 100644 index 0000000..953f3ec --- /dev/null +++ b/header/model/processor/strategies/DynCont.h @@ -0,0 +1,32 @@ +#ifndef DYNCONT_H +#define DYNCONT_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +typedef struct DynCont_t +{ + cl_int state; + cl_int gain; + cl_int dynContSelector; + cl_int compressionType; +}DynCont_t; + + +class DynCont : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE DynCont(const Context context, const QString kernelPath, const QObject *parent); + virtual void cpuProcess(ScenGenOutput_t parameters) override; + virtual void finalize() override; + +private: + KernelFunctor _kernelFunctor; + virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; + DynCont_t _kernelParameters; +}; + +#endif // DYNCONT_H diff --git a/header/model/processor/strategies/GrayMap.h b/header/model/processor/strategies/GrayMap.h new file mode 100644 index 0000000..de394e2 --- /dev/null +++ b/header/model/processor/strategies/GrayMap.h @@ -0,0 +1,29 @@ +#ifndef GRAYMAP_H +#define GRAYMAP_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +typedef struct GrayMap_t +{ + cl_int grayMapSelector; +}GrayMap_t; + + +class GrayMap : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE GrayMap(const Context context, const QString kernelPath, const QObject *parent); + virtual void cpuProcess(ScenGenOutput_t parameters) override; + virtual void finalize() override; + +private: + KernelFunctor _kernelFunctor; + virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; + GrayMap_t _kernelParameters; +}; + +#endif // GRAYMAP_H diff --git a/header/strategies/Rejection.h b/header/model/processor/strategies/Rejection.h similarity index 80% rename from header/strategies/Rejection.h rename to header/model/processor/strategies/Rejection.h index 60d220f..59657cd 100644 --- a/header/strategies/Rejection.h +++ b/header/model/processor/strategies/Rejection.h @@ -3,8 +3,8 @@ #include #include -#include "header/IProcessStrategy.h" -#include "header/OpenCLHelper.h" +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" typedef struct Rejection_t { @@ -18,7 +18,6 @@ public: Q_INVOKABLE Rejection(const Context context, const QString kernelPath); virtual void cpuProcess(ScenGenOutput_t parameters) override; virtual void finalize() override; - virtual void ReadParams(QString path, ScenGenOutput_t *params) override; private: diff --git a/header/strategies/ScanConversion.h b/header/model/processor/strategies/ScanConversion.h similarity index 74% rename from header/strategies/ScanConversion.h rename to header/model/processor/strategies/ScanConversion.h index c46d846..3ef1c13 100644 --- a/header/strategies/ScanConversion.h +++ b/header/model/processor/strategies/ScanConversion.h @@ -3,8 +3,8 @@ #include #include -#include "header/IProcessStrategy.h" -#include "header/OpenCLHelper.h" +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" typedef struct ScanConversion_t { @@ -17,21 +17,23 @@ class ScanConversion : public IProcessStrategy { Q_OBJECT public: - Q_INVOKABLE ScanConversion(const Context context, const QString kernelPath); + Q_INVOKABLE ScanConversion(const Context context, const QString kernelPath, + const QObject* parent); virtual void cpuProcess(ScenGenOutput_t parameters) override; virtual void finalize() override; - virtual void ReadParams(QString path, ScenGenOutput_t *params) override; private: KernelFunctor _kernelFunctor; - virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; + virtual Image* processKernel(Image *frames, Buffer* scrathPad) override; ScanConversion_t _kernelParameters; quint64 _width; quint64 _height; myflt* _gridPixelXPos; myflt* _gridPixelZPos; + myflt* _gridPixelR; + myflt* _gridPixelTheta; myflt* _scanXPos; myflt* _scanZPos; diff --git a/header/model/processor/strategies/Sri.h b/header/model/processor/strategies/Sri.h new file mode 100644 index 0000000..0c06aa6 --- /dev/null +++ b/header/model/processor/strategies/Sri.h @@ -0,0 +1,32 @@ +#ifndef SRI_H +#define SRI_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +typedef struct Sri_t +{ + cl_int state; + cl_int sri; + cl_int width; + cl_int height; +}Sri_t; + + +class Sri : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE Sri(const Context context, const QString kernelPath, const QObject *parent); + virtual void cpuProcess(ScenGenOutput_t parameters) override; + virtual void finalize() override; + +private: + KernelFunctor _kernelFunctor; + virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; + Sri_t _kernelParameters; +}; + +#endif // SRI_H diff --git a/header/model/processor/strategies/TintMap.h b/header/model/processor/strategies/TintMap.h new file mode 100644 index 0000000..f2f999f --- /dev/null +++ b/header/model/processor/strategies/TintMap.h @@ -0,0 +1,28 @@ +#ifndef TINTMAP_H +#define TINTMAP_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +typedef struct TintMap_t +{ + cl_int tintMapSelector; +}TintMap_t; + + +class TintMap : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE TintMap(const Context context, const QString kernelPath, const QObject *parent); + virtual void cpuProcess(ScenGenOutput_t parameters) override; + virtual void finalize() override; + +private: + KernelFunctor _kernelFunctor; + virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; + TintMap_t _kernelParameters; +}; +#endif // TINTMAP_H diff --git a/header/OpenCLHelper.h b/header/utils/OpenCLHelper.h similarity index 100% rename from header/OpenCLHelper.h rename to header/utils/OpenCLHelper.h diff --git a/kernels/DynCont.cl b/kernels/DynCont.cl new file mode 100644 index 0000000..fde4475 --- /dev/null +++ b/kernels/DynCont.cl @@ -0,0 +1,170 @@ +// #define TEST +#define X 100 +#define Y 100 + +//#define USE_DBL +#ifdef USE_DBL +#define TYPE_FLT double +#define TYPE_INT long +#define MASK 0xFFFF +#define SHIFT 16 +#define EPSILSON 4.94065645841247E-324 +#else +#define TYPE_FLT float +#define TYPE_INT int +#define MASK 0x00FF +#define SHIFT 8 +#define EPSILSON 1.401298E-45 +#endif + +constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + + +struct input { + int state; + int gain; + int dyn_cont_selector; + int compression_type; +}; + +TYPE_FLT read_data(image2d_t input_frame, int x, int z) +{ + int2 gid = (int2)(x, z); + uint4 pixel = read_imageui(input_frame, sampler, gid); + TYPE_INT temp = (TYPE_INT)((TYPE_INT)pixel.x & MASK) | + (TYPE_INT)(((TYPE_INT)pixel.y & MASK) << SHIFT) | + (TYPE_INT)(((TYPE_INT)pixel.z & MASK) << (SHIFT * 2)) | + (TYPE_INT)(((TYPE_INT)pixel.w & MASK) << (SHIFT * 3)); + TYPE_FLT raw_data = *((TYPE_FLT*)(&temp)); + return raw_data; +} + + +kernel void DynCont(read_only image2d_t input_frame, read_write image2d_t output_frame, local TYPE_FLT* max, global TYPE_FLT* scratch_pad, struct input params) +{ + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + TYPE_FLT input = read_data(input_frame, gid.x, gid.y); + + if(params.state == 1) + { + uint local_id = get_local_id(0); + uint group_size = get_local_size(0); + max[local_id] = input; + + uint is_odd = group_size % 2; + // Loop for computing localMaxes : divide WorkGroup into 2 parts + for (uint stride = group_size / 2; stride > 0; stride /= 2) + { + // Waiting + barrier(CLK_LOCAL_MEM_FENCE); + + if (local_id < stride) + max[local_id] = max[local_id] > max[local_id + stride] ? max[local_id] : max[local_id + stride]; + + if (local_id == 0) + { + if(is_odd) + max[local_id] = max[local_id] > max[2 * stride] ? max[local_id] : max[2 * stride]; + is_odd = stride % 2; + } + } + + // Write result into scratchPad[nWorkGroups] + if (local_id == 0) + { + scratch_pad[get_group_id(1)] = max[0]; + } + } + else if(params.state == 2) + { + uint local_id = get_local_id(0); + uint group_size = get_local_size(0); + max[local_id] = scratch_pad[gid.x]; + + uint is_odd = group_size % 2; + // Loop for computing localMaxes : divide WorkGroup into 2 parts + for (uint stride = group_size / 2; stride > 0; stride /= 2) + { + // Waiting + barrier(CLK_LOCAL_MEM_FENCE); + + if (local_id < stride) + max[local_id] = max[local_id] > max[local_id + stride] ? max[local_id] : max[local_id + stride]; + + if (local_id == 0) + { + if(is_odd) + max[local_id] = max[local_id] > max[2 * stride] ? max[local_id] : max[2 * stride]; + is_odd = stride % 2; + } + } + + // Write result into scratchPad[nWorkGroups] + if (local_id == 0) + { + scratch_pad[0] = max[0]; + } + } + else if(params.state == 3) + { + int dynamic_range = 60 + params.dyn_cont_selector; + TYPE_FLT max = scratch_pad[0]; + TYPE_FLT output_data = 0; + //log compress + if(params.compression_type == 1) + { + output_data = 20 * log10(input / max); + + if(output_data < -dynamic_range) + output_data = -dynamic_range; + + output_data += params.gain; + + if(output_data > 0) + output_data = 0; + + if(output_data < -dynamic_range) + output_data = -dynamic_range; + + output_data = 255 * (output_data + dynamic_range) / dynamic_range; + } + else + { + output_data = input / max; + + if(output_data < pow(10, (-dynamic_range / 20.0))) + { + output_data = 0; + } + + output_data *= pow(10, (params.gain / 20.0)); + + if(output_data < 0) + { + output_data = 0; + } + + if(output_data > 1) + output_data = 1; + + output_data *= 255; + +#ifdef TEST + if(gid.x == X && gid.y == Y) + printf("%.6f %.6f %.6f ----", output_data, input, max); +#endif + } + + TYPE_INT out = *((TYPE_INT*)(&output_data)); + uint4 pixel; + pixel.x = (TYPE_INT)(out & MASK); + pixel.y = (TYPE_INT)((out >> SHIFT) & MASK); + pixel.z = (TYPE_INT)((out >> (SHIFT *2)) & MASK); + pixel.w = (TYPE_INT)((out >> (SHIFT * 3)) & MASK); + + write_imageui(output_frame, gid, pixel); + } +} + diff --git a/kernels/GrayMap.cl b/kernels/GrayMap.cl new file mode 100644 index 0000000..676a265 --- /dev/null +++ b/kernels/GrayMap.cl @@ -0,0 +1,74 @@ +// #define TEST +#define X 100 +#define Y 100 + +//#define USE_DBL +#ifdef USE_DBL +#define TYPE_FLT double +#define TYPE_INT long +#define MASK 0xFFFF +#define SHIFT 16 +#define EPSILSON 4.94065645841247E-324 +#else +#define TYPE_FLT float +#define TYPE_INT int +#define MASK 0x00FF +#define SHIFT 8 +#define EPSILSON 1.401298E-45 +#endif + +constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + + +struct input { + int gray_map_selector; +}; + +TYPE_FLT read_data(image2d_t input_frame, int x, int z) +{ + int2 gid = (int2)(x, z); + uint4 pixel = read_imageui(input_frame, sampler, gid); + TYPE_INT temp = (TYPE_INT)((TYPE_INT)pixel.x & MASK) | + (TYPE_INT)(((TYPE_INT)pixel.y & MASK) << SHIFT) | + (TYPE_INT)(((TYPE_INT)pixel.z & MASK) << (SHIFT * 2)) | + (TYPE_INT)(((TYPE_INT)pixel.w & MASK) << (SHIFT * 3)); + TYPE_FLT raw_data = *((TYPE_FLT*)(&temp)); + return raw_data; +} + + +kernel void GrayMap(read_only image2d_t input_frame, read_write image2d_t output_frame, struct input params) +{ + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + TYPE_FLT input = read_data(input_frame, gid.x, gid.y); + TYPE_FLT output_data = 0; + + double gama[17] = {0.05, 0.1, 0.2, 0.3, 0.4, 0.5, 0.67, 0.75, 1, 1.33, 1.5, 2, 2.5, 3.33, 5, 10, 20}; + + if(params.gray_map_selector == 17) + { + output_data = 255 * log2(input / 255 + 1); + } + else + { + output_data = 255 * pow((input / 255.0), gama[params.gray_map_selector]); + } + + uchar out = (uchar)output_data; + +#ifdef TEST + if(gid.x == X && gid.y == Y) + printf("%u ---", out); +#endif + + uint4 pixel; + pixel.x = out; + pixel.y = out; + pixel.z = out; + pixel.w = 255; + + write_imageui(output_frame, gid, pixel); +} + diff --git a/kernels/Max.cl b/kernels/Max.cl new file mode 100644 index 0000000..1449df7 --- /dev/null +++ b/kernels/Max.cl @@ -0,0 +1,92 @@ +#define TEST +#define X 0 +#define Y 26 + +#define USE_DBL +#ifdef USE_DBL +#define TYPE_FLT double +#define TYPE_INT long +#define MASK 0xFFFF +#define SHIFT 16 +#define EPSILSON 4.94065645841247E-324 +#else +#define TYPE_FLT float +#define TYPE_INT int +#define MASK 0x00FF +#define SHIFT 8 +#define EPSILSON 1.401298E-45 +#endif + +constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + +TYPE_FLT read_data(image2d_t input_frame, int x, int z) +{ + int2 gid = (int2)(x, z); + uint4 pixel = read_imageui(input_frame, sampler, gid); + TYPE_INT temp = (TYPE_INT)((TYPE_INT)pixel.x & MASK) | + (TYPE_INT)(((TYPE_INT)pixel.y & MASK) << SHIFT) | + (TYPE_INT)(((TYPE_INT)pixel.z & MASK) << (SHIFT * 2)) | + (TYPE_INT)(((TYPE_INT)pixel.w & MASK) << (SHIFT * 3)); + TYPE_FLT raw_data = *((TYPE_FLT*)(&temp)); + return raw_data; +} + + +kernel void Max(read_only image2d_t input_frame, read_write image2d_t output_frame, local double* max, global double* scratch_pad) { + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + + TYPE_FLT output_data = 0; + TYPE_FLT input = read_data(input_frame, gid.x, gid.y); + output_data = input; + + uint local_id = get_local_id(0); + uint group_size = get_local_size(0); + max[local_id] = input; + + if(gid.x == 0 && gid.y == 0) +{ + printf("gid [%d] ", get_group_id(0)); + printf("lid0 [%d] ", get_local_id(0)); + printf("lid1 [%d] ", get_local_id(1)); + printf("lis0 [%d] ", get_local_size(0)); + printf("lis1 [%d] ", get_local_size(1)); +} + uint is_odd = 0; + // Loop for computing localMaxes : divide WorkGroup into 2 parts + for (uint stride = group_size / 2; stride > 0; stride /= 2) + { + // Waiting for each 2x2 max into given workgroup + barrier(CLK_LOCAL_MEM_FENCE); + + // max elements 2 by 2 between local_id and local_id + stride + if (local_id < stride) + max[local_id] = max[local_id] > max[local_id + stride] ? max[local_id] : max[local_id + stride]; + + if (local_id == 0) + { + if(is_odd) + max[local_id] = max[local_id] > max[2 * stride] ? max[local_id] : max[2 * stride]; + is_odd = stride % 2; + } + } + + // Write result into scratchPad[nWorkGroups] + if (local_id == 0) + { + scratch_pad[get_group_id(1)] = max[0]; + printf("out[%d]: %.15f | ", get_group_id(1), scratch_pad[get_group_id(1)]); + } + + + TYPE_INT out = *((TYPE_INT*)(&output_data)); + uint4 pixel; + pixel.x = (TYPE_INT)(out & MASK); + pixel.y = (TYPE_INT)((out >> SHIFT) & MASK); + pixel.z = (TYPE_INT)((out >> (SHIFT *2)) & MASK); + pixel.w = (TYPE_INT)((out >> (SHIFT * 3)) & MASK); + + write_imageui(output_frame, gid, pixel); +} + diff --git a/kernels/ScanConversion.cl b/kernels/ScanConversion.cl index 18f25bf..bc6b78d 100644 --- a/kernels/ScanConversion.cl +++ b/kernels/ScanConversion.cl @@ -1,8 +1,8 @@ -#define TEST -#define X 923 -#define Y 436 +//#define TEST +#define X 0 +#define Y 0 -#define USE_DBL +//#define USE_DBL #ifdef USE_DBL #define TYPE_FLT double #define TYPE_INT long @@ -123,5 +123,4 @@ kernel void ScanConversion(read_only image2d_t input_frame, read_write image2d_t #endif write_imageui(output_frame, gid, pixel); -} - +} \ No newline at end of file diff --git a/kernels/Sri.cl b/kernels/Sri.cl new file mode 100644 index 0000000..fbf7302 --- /dev/null +++ b/kernels/Sri.cl @@ -0,0 +1,199 @@ +// #define TEST +#define X 95 +#define Y 22 + +//#define USE_DBL +#ifdef USE_DBL +#define TYPE_FLT double +#define TYPE_INT long +#define MASK 0xFFFF +#define SHIFT 16 +#define EPSILSON 4.94065645841247E-324 +#else +#define TYPE_FLT float +#define TYPE_INT int +#define MASK 0x00FF +#define SHIFT 8 +#define EPSILSON 1.401298E-45 +#endif + +constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + + +struct input { + int state; + int sri; + int width; + int height; +}; + +TYPE_FLT read_data(image2d_t input_frame, int x, int z) +{ + int2 gid = (int2)(x, z); + uint4 pixel = read_imageui(input_frame, sampler, gid); + TYPE_INT temp = (TYPE_INT)((TYPE_INT)pixel.x & MASK) | + (TYPE_INT)(((TYPE_INT)pixel.y & MASK) << SHIFT) | + (TYPE_INT)(((TYPE_INT)pixel.z & MASK) << (SHIFT * 2)) | + (TYPE_INT)(((TYPE_INT)pixel.w & MASK) << (SHIFT * 3)); + TYPE_FLT raw_data = *((TYPE_FLT*)(&temp)); + return raw_data; +} + + +kernel void Sri(read_only image2d_t input_frame, read_write image2d_t output_frame, local TYPE_FLT* sum, global TYPE_FLT* scratch_pad, struct input params) +{ + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + TYPE_FLT input = read_data(input_frame, gid.x, gid.y); + TYPE_FLT output_data = 0; + + if(params.sri == 0) + { + output_data = input; + } + else if(params.state == 1 || params.state == 4) + { + uint local_id = get_local_id(0); + uint group_size = get_local_size(0); + sum[local_id] = input; + + uint is_odd = group_size % 2; + // Loop for computing localMaxes : divide WorkGroup into 2 parts + for (uint stride = group_size / 2; stride > 0; stride /= 2) + { + // Waiting + barrier(CLK_LOCAL_MEM_FENCE); + + if (local_id < stride) + sum[local_id] += sum[local_id + stride]; + + if (local_id == 0) + { + if(is_odd) + sum[local_id] += sum[2 * stride]; + is_odd = stride % 2; + } + } + + // Write result into scratchPad[nWorkGroups] + if (local_id == 0) + { + scratch_pad[get_group_id(1)] = sum[0]; + } + + return; + } + else if(params.state == 2 || params.state == 5) + { + uint local_id = get_local_id(0); + uint group_size = get_local_size(0); + sum[local_id] = scratch_pad[gid.x]; + + uint is_odd = group_size % 2; + // Loop for computing localMaxes : divide WorkGroup into 2 parts + for (uint stride = group_size / 2; stride > 0; stride /= 2) + { + // Waiting + barrier(CLK_LOCAL_MEM_FENCE); + + if (local_id < stride) + sum[local_id] += sum[local_id + stride]; + + if (local_id == 0) + { + if(is_odd) + sum[local_id] += sum[2 * stride]; + is_odd = stride % 2; + } + } + + // Write result into scratchPad[nWorkGroups] + if (local_id == 0) + { + scratch_pad[0] = sum[0] / (params.width * params.height); + } + + return; + } + else if(params.state == 3) + { + TYPE_FLT mean = scratch_pad[0]; + output_data = pow((input - mean), 2); + } + else if(params.state == 6) + { + int temp[5] = {3, 3, 5, 7, 7}; + int p = temp[params.sri - 1]; + int q = params.width > 1 ? p : 1; + + float alpha = 1 - 0.25 * (params.sri - 1); + TYPE_FLT rho2 = scratch_pad[0]; + + int lag1 = (p - 1) / 2; + int lag2 = (q - 1) / 2; + + float window[49]; + + int window_size = p * q; + for(int i = 0; i < window_size; i++) + window[i] = 0; + if(gid.x >= lag2 && gid.x < params.width - lag2 && gid.y >= lag1 && gid.y < params.height - lag1) + { + for(int i = -lag1; i <= lag1 ; i++) + { + for(int j = -lag2; j <= lag2 ; j++) + { + window[(i + lag1) * q + (j + lag2)] = read_data(input_frame, gid.x + j, gid.y + i); + } + } + +#ifdef TEST + if(gid.x == X && gid.y == Y) + { + for(int i = 0; i < window_size; i++) + printf("[%d] > %f; ", i, window[i]); + } +#endif + + + float mean = 0; + for(int i = 0; i < window_size; i++) + { + mean += window[i]; + } + mean /= window_size; + + for(int i = 0; i < window_size; i++) + { + window[i] = pow((window[i] - mean), 2); + } + + float sigma2 = 0; + for(int i = 0; i < window_size; i++) + { + sigma2 += window[i]; + } + sigma2 /= window_size; + + output_data = mean + alpha * (sigma2 / (sigma2 + rho2)) * (input - mean); + } + else + { + output_data = input; + } + + if(output_data < 0) + output_data = 0; + } + + TYPE_INT out = *((TYPE_INT*)(&output_data)); + uint4 pixel; + pixel.x = (TYPE_INT)(out & MASK); + pixel.y = (TYPE_INT)((out >> SHIFT) & MASK); + pixel.z = (TYPE_INT)((out >> (SHIFT *2)) & MASK); + pixel.w = (TYPE_INT)((out >> (SHIFT * 3)) & MASK); + + write_imageui(output_frame, gid, pixel); +} + diff --git a/kernels/TintMap.cl b/kernels/TintMap.cl new file mode 100644 index 0000000..4454c49 --- /dev/null +++ b/kernels/TintMap.cl @@ -0,0 +1,100 @@ +// #define TEST +#define X 100 +#define Y 100 + +//#define USE_DBL +#ifdef USE_DBL +#define TYPE_FLT double +#define TYPE_INT long +#define MASK 0xFFFF +#define SHIFT 16 +#define EPSILSON 4.94065645841247E-324 +#else +#define TYPE_FLT float +#define TYPE_INT int +#define MASK 0x00FF +#define SHIFT 8 +#define EPSILSON 1.401298E-45 +#endif + +constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + + +struct input { + int tint_map_selector; +}; + + +kernel void TintMap(read_only image2d_t input_frame, read_write image2d_t output_frame, struct input params) +{ + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + uint4 pixel = read_imageui(input_frame, sampler, gid); + + private float r_step, g_step, b_step; + private int r_base, g_base, b_base; + + switch(params.tint_map_selector) + { + //gray + case 0: + r_base = g_base = b_base = 0; + r_step = g_step = b_step = 1.0; + break; + //blue + case 1: + r_base = g_base = b_base = 0; + r_step = g_step = 0 ; + b_step = 1.0f; + break; + //cool blue + case 2: + r_base = 0; + g_base = b_base = 255; + r_step = 1.0f; + g_step = -1.0f; + b_step = 0; + break; + //copper + case 3: + r_base = g_base = b_base = 0; + r_step = 1.0f; + g_step = 0.7812f; + b_step = 0.4975f; + break; + //otherwise + case 4: + r_base = g_base = b_base = 0; + r_step = g_step = b_step = 1.0f; + break; + } + + uchar data = pixel.x; + if(data > 0) + { + float r = r_base + data * r_step; + if(r < 0) r = 0; + if(r > 255) r = 255; + + float g = g_base + data * g_step; + if(g < 0) g = 0; + if(g > 255) g = 255; + + float b = b_base + data * b_step; + if(b < 0) b = 0; + if(b > 255) b = 255; + + #ifdef TEST + if(gid.x == X && gid.y == Y) + printf("%u ---", out); + #endif + + pixel.x = (uchar)b; + pixel.y = (uchar)g; + pixel.z = (uchar)r; + } + + write_imageui(output_frame, gid, pixel); +} + diff --git a/source/FileHelper.cpp b/source/FileHelper.cpp index 61a9d0f..86cac45 100644 --- a/source/FileHelper.cpp +++ b/source/FileHelper.cpp @@ -54,11 +54,11 @@ bool FileHelper::ReadInputFile(myint *arr, QString path, quint64 *width, quint64 QFile file(path); file.open(QIODevice::ReadOnly); - char data[200000]; + char data[8192]; bool first = true; int index = 0; - while (file.readLine(data, 200000)) + while (file.readLine(data, 8192)) { QString str(data); if(str.length() == 0) @@ -85,6 +85,26 @@ bool FileHelper::ReadInputFile(myint *arr, QString path, quint64 *width, quint64 return true; } +void FileHelper::GetFileSIze(QString path, quint64 *width, quint64 *height) +{ + QFile file(path); + file.open(QIODevice::ReadOnly); + + char data[20000]; + + bool first = true; + while (file.readLine(data, 20000)) + { + QString str(data); + if(str.length() == 0) + break; + (*height)++; + str = str.remove(str.length() - 1, 1); + auto sl = str.split(","); + *width = sl.length(); + } +} + bool FileHelper::WriteOutputFile(char* arr, QString path, quint64 width, quint64 height) { QFile file(path); @@ -109,7 +129,7 @@ bool FileHelper::WriteOutputFile(char* arr, QString path, quint64 width, quint64 | (((quint64)(arr[index++] & 0xFF)) << 16) | (((quint64)(arr[index++] & 0xFF)) << 24); auto d = *((float*)(&temp)); #endif - str += QString("%1").arg(temp, 16, 16, QChar('0')); + str += QString("%1").arg(temp, 8, 16, QChar('0')); if(j != width - 1) str += ","; } diff --git a/source/model/processor/BIP.cpp b/source/model/processor/BIP.cpp new file mode 100644 index 0000000..bfd0b2c --- /dev/null +++ b/source/model/processor/BIP.cpp @@ -0,0 +1,24 @@ +#include "model/processor/BIP.h" + +BIP* BIP::_instance; + +BIP::BIP() +{ + +} + +BIP *BIP::getInstance() +{ + if(_instance == Q_NULLPTR) + _instance = new BIP(); + + return _instance; +} + +void BIP::init(cl::Context &context, OpenCLHelper &openCLHelper) +{ + CLQueue = openCLHelper.createCommandQueue(context, + openCLHelper + .getDevicesByContext( + context)[0]); +} diff --git a/source/IProcessStrategy.cpp b/source/model/processor/IProcessStrategy.cpp similarity index 71% rename from source/IProcessStrategy.cpp rename to source/model/processor/IProcessStrategy.cpp index a660c4e..11f3166 100644 --- a/source/IProcessStrategy.cpp +++ b/source/model/processor/IProcessStrategy.cpp @@ -1,15 +1,11 @@ -#include "header/IProcessStrategy.h" +#include "model/processor/IProcessStrategy.h" #include #include -void IProcessStrategy::slt_process(Image* inputFrame, Buffer* scratchPad) -{ - -} - IProcessStrategy::IProcessStrategy(const Context context, const QString kernelPath, - const QString kernelName) + const QString kernelName, + const QObject* parent) { try { diff --git a/source/model/processor/strategies/DynCont.cpp b/source/model/processor/strategies/DynCont.cpp new file mode 100644 index 0000000..5642d1a --- /dev/null +++ b/source/model/processor/strategies/DynCont.cpp @@ -0,0 +1,88 @@ +#include "model/processor/strategies/DynCont.h" +#include "model/processor/BIP.h" +#include +#include + +DynCont::DynCont(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "DynCont", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (DynCont_t)); +} + +void DynCont::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.dynContGain.isUpdated) + { + _kernelParameters.gain = parameters.dynContGain.value; + } + + if(parameters.dynContSelector.isUpdated) + { + _kernelParameters.dynContSelector = parameters.dynContSelector.value; + } + + + if(parameters.compressionType.isUpdated) + { + _kernelParameters.compressionType = parameters.compressionType.value; + } +} + +void DynCont::finalize() +{ +} + + +Image* DynCont::processKernel(Image *frames, Buffer* scratchPad) +{ + auto format = frames->getImageInfo(); + auto width = frames->getImageInfo(); + auto height = frames->getImageInfo(); + + auto imageOutput = new Image2D(_CLContext, + CL_MEM_READ_WRITE, + ImageFormat(format.image_channel_order, format.image_channel_data_type), + width, + height); + //width should be lass than CL_DEVICE_MAX_WORK_GROUP_SIZE(1024 for current machine) or this will break + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height), cl::NDRange(width, 1)); + + _kernelParameters.state = 1; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + Local(width * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + + cl::EnqueueArgs eargs2(BIP::getInstance()->CLQueue, cl::NDRange(height, 1), cl::NDRange(height, 1)); + + _kernelParameters.state = 2; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs2, + *static_cast(frames), + *imageOutput, + Local(height * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + + _kernelParameters.state = 3; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + Local(1), + *scratchPad, + _kernelParameters); + + delete frames; + + return imageOutput; +} diff --git a/source/model/processor/strategies/GrayMap.cpp b/source/model/processor/strategies/GrayMap.cpp new file mode 100644 index 0000000..9de4f31 --- /dev/null +++ b/source/model/processor/strategies/GrayMap.cpp @@ -0,0 +1,57 @@ +#include "model/processor/strategies/GrayMap.h" +#include "model/processor/BIP.h" +#include +#include + +GrayMap::GrayMap(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "GrayMap", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (GrayMap_t)); +} + +void GrayMap::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.grayMapSelector.isUpdated) + { + _kernelParameters.grayMapSelector = parameters.grayMapSelector.value; + } +} + +void GrayMap::finalize() +{ +} + + +Image* GrayMap::processKernel(Image *frames, Buffer* scratchPad) +{ + auto width = frames->getImageInfo(); + auto height = frames->getImageInfo(); + + ImageFormat format; + format.image_channel_order = CL_RGBA; + format.image_channel_data_type = CL_UNSIGNED_INT8; + + auto imageOutput = new Image2D(_CLContext, + CL_MEM_READ_WRITE, + format, + width, + height); + //width should be lass than CL_DEVICE_MAX_WORK_GROUP_SIZE(1024 for current machine) or this will break + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); + + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + _kernelParameters); + + delete frames; + + return imageOutput; +} + + diff --git a/source/strategies/Rejection.cpp b/source/model/processor/strategies/Rejection.cpp similarity index 57% rename from source/strategies/Rejection.cpp rename to source/model/processor/strategies/Rejection.cpp index ceebd0e..5a4be4f 100644 --- a/source/strategies/Rejection.cpp +++ b/source/model/processor/strategies/Rejection.cpp @@ -1,11 +1,11 @@ -#include "header/strategies/Rejection.h" -#include "MainWindow.h" +#include "model/processor/strategies/Rejection.h" +#include "model/processor/BIP.h" #include #include Rejection::Rejection(const Context context, const QString kernelPath) : - IProcessStrategy(context, kernelPath, "Rejection"), + IProcessStrategy(context, kernelPath, "Rejection", Q_NULLPTR), _kernelFunctor(KernelFunctor(_kernel)) { @@ -14,35 +14,13 @@ Rejection::Rejection(const Context context, void Rejection::cpuProcess(ScenGenOutput_t parameters) { - _rejectionThr = parameters.rejectThreshold.value; + //_rejectionThr = parameters.rejectThreshold.value; } void Rejection::finalize() { } -void Rejection::ReadParams(QString path, ScenGenOutput_t *params) -{ - QFile file(path); - file.open(QIODevice::ReadOnly); - - char data[8192]; - //first line is extra - file.readLine(data, 8192); - QString str(data); - str = str.remove(str.length() - 1, 1); - auto sl = str.split(","); - - //decode here - /***scenario specific code***/ - int index = 0; - - auto temp = sl[index++].PARSE; - update_field(¶ms->rejectThreshold, INP2MYFLT(temp)); - /***End of scenario specific code***/ -} - - Image* Rejection::processKernel(Image *frames, Buffer* scratchPad) { Context context = _openCLHelper.getContext(); @@ -54,7 +32,7 @@ Image* Rejection::processKernel(Image *frames, Buffer* scratchPad) ImageFormat(format.image_channel_order, format.image_channel_data_type), width, height); - cl::EnqueueArgs eargs(MainWindow::getInstance()->CLQueue, cl::NDRange(width, height)); + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); _openCLHelper.runKernelFunctor(_kernelFunctor, eargs, diff --git a/source/model/processor/strategies/ScanConversion.cpp b/source/model/processor/strategies/ScanConversion.cpp new file mode 100644 index 0000000..c241695 --- /dev/null +++ b/source/model/processor/strategies/ScanConversion.cpp @@ -0,0 +1,424 @@ +#include "model/processor/strategies/ScanConversion.h" +#include +#include +#include "model/processor/BIP.h" + +ScanConversion::ScanConversion(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "ScanConversion", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + _gridPixelXPos = Q_NULLPTR; + _gridPixelZPos = Q_NULLPTR; + _scanXPos = Q_NULLPTR; + _scanZPos = Q_NULLPTR; + + _gridX = Q_NULLPTR; + _gridZ = Q_NULLPTR; + _queryX = Q_NULLPTR; + _queryZ = Q_NULLPTR; + + _gridPixelR = Q_NULLPTR; + _gridPixelTheta = Q_NULLPTR; +} + +void ScanConversion::cpuProcess(ScenGenOutput_t params) +{ + auto context = _openCLHelper.getContext(); + auto queue = BIP::getInstance()->CLQueue; + //virtual convex + if(params.virtualConvex.value) + { + if(params.outputWidth.isUpdated|| params.outputHeight.isUpdated) + { + if(_gridPixelXPos) + delete[] _gridPixelXPos; + + if(_gridPixelZPos) + delete[] _gridPixelZPos; + + _gridPixelXPos = new myflt[params.outputWidth.value * + params.outputHeight.value]; + _gridPixelZPos = new myflt[params.outputWidth.value * + params.outputHeight.value]; + + _width = params.outputWidth.value; + _height = params.outputHeight.value; + + _kernelParameters.width = _width; + } + + if(params.depth.isUpdated || params.startDepth.isUpdated || + params.fieldOfView.isUpdated || params.vcMaxTheta.isUpdated || + params.outputWidth.isUpdated || params.outputHeight.isUpdated) + { + auto maxZ = params.depth.value; + auto minZ = params.startDepth.value * cosf(params.vcMaxTheta.value); + auto maxX = params.fieldOfView.value / 2 + maxZ * sinf(params.vcMaxTheta.value); + auto minX = -1 * maxX; + + auto baseZ = (maxZ - minZ) / (params.outputHeight.value - 1); + auto baseX = (maxX - minX) / (params.outputWidth.value - 1); + + for(auto i = 0; i < params.outputHeight.value; i++) + { + auto temp = i * baseZ; + for(auto j = 0; j < params.outputWidth.value; j++) + { + _gridPixelZPos[i * params.outputHeight.value + j] = minZ + temp; + _gridPixelXPos[i * params.outputHeight.value + j] = minX + (baseX * j); + } + } + + auto virtualOriginZ = params.fieldOfView.value / 2 / tanf(params.vcMaxTheta.value); + + if(_gridPixelR) + delete[] _gridPixelR; + + if(_gridPixelTheta) + delete[] _gridPixelTheta; + + _gridPixelR = new myflt[params.outputWidth.value * + params.outputHeight.value]; + _gridPixelTheta = new myflt[params.outputWidth.value * + params.outputHeight.value]; + + auto strTan = tanf(params.steering.value); + + for(auto i = 0; i < params.outputWidth.value * + params.outputHeight.value; i++) + { + auto x = _gridPixelXPos[i]; + auto z = _gridPixelZPos[i]; + + auto pixelTheta = atanf(x / (z + virtualOriginZ)); + + if(pixelTheta >= -params.vcMaxTheta.value - abs(params.steering.value) && + pixelTheta <= params.vcMaxTheta.value + abs(params.steering.value)) + { + if(params.steering.value == 0) + { + auto gridPixelAx = sqrtf(powf(x - virtualOriginZ * tanf(pixelTheta), 2) + powf(z, 2)); + if(gridPixelAx >= params.startDepth.value && + gridPixelAx <= params.depth.value) + { + _gridPixelR[i] = gridPixelAx; + _gridPixelTheta[i] = pixelTheta; + } + else + { + _gridPixelR[i] = 0; + _gridPixelTheta[i] = 0; + } + } + else + { + auto a = virtualOriginZ * strTan; + auto b = x * strTan + virtualOriginZ + z; + auto c = x - z * strTan; + auto interceptTheta = atanf((b + sqrtf(powf(b, 2) - 4 * a * c)) / (2 * a)); + if(interceptTheta > params.vcMaxTheta.value || + interceptTheta < -params.vcMaxTheta.value) + { + interceptTheta = atanf((b - sqrtf(powf(b, 2) - 4 * a * c)) / (2 * a)); + if(interceptTheta > params.vcMaxTheta.value || + interceptTheta < -params.vcMaxTheta.value) + { + _gridPixelR[i] = 0; + _gridPixelTheta[i] = 0; + } + else + { + auto gridPixelAx = sqrtf(powf(x - virtualOriginZ * tanf(interceptTheta), 2) + powf(z, 2)); + if(gridPixelAx >= params.startDepth.value && + gridPixelAx <= params.depth.value) + { + _gridPixelR[i] = gridPixelAx; + _gridPixelTheta[i] = interceptTheta; + } + else + { + _gridPixelR[i] = 0; + _gridPixelTheta[i] = 0; + } + } + } + else + { + auto gridPixelAx = sqrtf(powf(x - virtualOriginZ * tanf(interceptTheta), 2) + powf(z, 2)); + if(gridPixelAx >= params.startDepth.value && + gridPixelAx <= params.depth.value) + { + _gridPixelR[i] = gridPixelAx; + _gridPixelTheta[i] = interceptTheta; + } + else + { + _gridPixelR[i] = 0; + _gridPixelTheta[i] = 0; + } + } + } + } + else + { + _gridPixelR[i] = 0; + _gridPixelTheta[i] = 0; + } + } + + if(_queryX) + delete _queryX; + + if(_queryZ) + delete _queryZ; + _queryX = _openCLHelper.array2CLBuffer(context, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt)); + _queryZ = _openCLHelper.array2CLBuffer(context, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt)); + + BIP::getInstance()->CLQueue.enqueueWriteBuffer(*_queryZ, CL_TRUE, 0, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt), + _gridPixelR); + queue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt), + _gridPixelTheta); + } + + if(params.minScanAz.isUpdated || params.rxLineNo.isUpdated || + params.rxLineDaz.isUpdated) + { + if(_scanXPos) + delete [] _scanXPos; + + if(_gridX) + delete _gridX; + + _scanXPos = new myflt[params.rxLineNo.value]; + for(auto i = 0; i < params.rxLineNo.value; i++) + { + _scanXPos[i] = params.minScanAz.value + i * params.rxLineDaz.value; + } + + _gridX = _openCLHelper.array2CLBuffer(context, + static_cast(params.rxLineNo.value) + * sizeof (myflt)); + + queue.enqueueWriteBuffer(*_gridX, CL_TRUE, 0, + static_cast(params.rxLineNo.value) * sizeof (myflt), + _scanXPos); + + _kernelParameters.gridXSize = params.rxLineNo.value; + } + + if(params.rxFocusPointNo.isUpdated || params.minScanAx.isUpdated || + params.rxPointDax.isUpdated) + { + if(_scanZPos) + delete[] _scanZPos; + + if(_gridZ) + delete _gridZ; + + _scanZPos = new myflt [params.rxFocusPointNo.value]; + for(auto i = 0; i < params.rxFocusPointNo.value; i++) + { + _scanZPos[i] = params.minScanAx.value + i * params.rxPointDax.value; + } + + _gridZ = _openCLHelper.array2CLBuffer(context, + static_cast(params.rxFocusPointNo.value) + * sizeof (myflt)); + + queue.enqueueWriteBuffer(*_gridZ, CL_TRUE, 0, + static_cast(params.rxFocusPointNo.value) * sizeof (myflt), + _scanZPos); + + _kernelParameters.gridZSize = params.rxFocusPointNo.value; + } + } + //linear non convex + else + { + if(params.outputWidth.isUpdated || params.outputHeight.isUpdated) + { + if(_gridPixelXPos) + delete[] _gridPixelXPos; + + if(_gridPixelZPos) + delete[] _gridPixelZPos; + + _gridPixelXPos = new myflt[params.outputWidth.value * + params.outputHeight.value]; + _gridPixelZPos = new myflt[params.outputWidth.value * + params.outputHeight.value]; + + if(_queryX) + delete _queryX; + + if(_queryZ) + delete _queryZ; + + _queryX = _openCLHelper.array2CLBuffer(context, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt)); + _queryZ = _openCLHelper.array2CLBuffer(context, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt)); + + _width = params.outputWidth.value; + _height = params.outputHeight.value; + + _kernelParameters.width = _width; + } + + if(params.depth.isUpdated || params.startDepth.isUpdated || + params.fieldOfView.isUpdated || params.steering.isUpdated || + params.outputWidth.isUpdated || params.outputHeight.isUpdated) + { + auto maxZ = params.depth.value; + auto minZ = params.startDepth.value; + auto maxX = params.fieldOfView.value / 2; + auto minX = -1 * maxX; + + auto baseZ = (maxZ - minZ) / (params.outputHeight.value - 1); + auto baseX = (maxX - minX) / (params.outputWidth.value - 1); + + for(auto i = 0; i < params.outputHeight.value; i++) + { + auto temp = i * baseZ; + auto tan_result = temp * tanf(params.steering.value); + for(auto j = 0; j < params.outputWidth.value; j++) + { + _gridPixelZPos[i * params.outputHeight.value + j] = minZ + temp; + _gridPixelXPos[i * params.outputHeight.value + j] = minX + (baseX * j) - tan_result; + } + } + + BIP::getInstance()->CLQueue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt), + _gridPixelXPos); + queue.enqueueWriteBuffer(*_queryZ, CL_TRUE, 0, + params.outputWidth.value * + params.outputHeight.value * + sizeof (myflt), + _gridPixelZPos); + } + + if(params.fieldOfView.isUpdated || params.rxLineNo.isUpdated || + params.rxLineDaz.isUpdated) + { + if(_scanXPos) + delete [] _scanXPos; + + if(_gridX) + delete _gridX; + + _scanXPos = new myflt[params.rxLineNo.value]; + for(auto i = 0; i < params.rxLineNo.value; i++) + { + _scanXPos[i] = -params.fieldOfView.value / 2 + i * params.rxLineDaz.value; + } + + _gridX = _openCLHelper.array2CLBuffer(context, + static_cast(params.rxLineNo.value) + * sizeof (myflt)); + + queue.enqueueWriteBuffer(*_gridX, CL_TRUE, 0, + static_cast(params.rxLineNo.value) * sizeof (myflt), + _scanXPos); + + _kernelParameters.gridXSize = params.rxLineNo.value; + } + + if(params.rxFocusPointNo.isUpdated || params.rxPointDax.isUpdated) + { + if(_scanZPos) + delete[] _scanZPos; + + if(_gridZ) + delete _gridZ; + + _scanZPos = new myflt [params.rxFocusPointNo.value]; + for(auto i = 0; i < params.rxFocusPointNo.value; i++) + { + _scanZPos[i] = params.startDepth.value + i * params.rxPointDax.value; + } + + _gridZ = _openCLHelper.array2CLBuffer(context, + static_cast(params.rxFocusPointNo.value) + * sizeof (myflt)); + + queue.enqueueWriteBuffer(*_gridZ, CL_TRUE, 0, + static_cast(params.rxFocusPointNo.value) * sizeof (myflt), + _scanZPos); + + _kernelParameters.gridZSize = params.rxFocusPointNo.value; + } + } +} + +void ScanConversion::finalize() +{ + delete _gridPixelXPos; + delete _gridPixelZPos; + delete _scanXPos; + delete _scanZPos; + + delete _gridX; + delete _gridZ; + delete _queryX; + delete _queryZ; + + _gridPixelXPos = Q_NULLPTR; + _gridPixelZPos = Q_NULLPTR; + _scanXPos = Q_NULLPTR; + _scanZPos = Q_NULLPTR; + + _gridX = Q_NULLPTR; + _gridZ = Q_NULLPTR; + _queryX = Q_NULLPTR; + _queryZ = Q_NULLPTR; +} + +Image* ScanConversion::processKernel(Image *frames, Buffer* scrathPad) +{ + Context context = _openCLHelper.getContext(); + auto format = frames->getImageInfo(); + auto imageOutput = new Image2D(context, + CL_MEM_READ_WRITE, + ImageFormat(format.image_channel_order, format.image_channel_data_type), + _width, + _height); + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(_width, _height)); + + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *(Image2D*)(frames), + *imageOutput, + *_gridX, + *_gridZ, + *_queryX, + *_queryZ, + _kernelParameters); + delete frames; + + return imageOutput; +} + + diff --git a/source/model/processor/strategies/Sri.cpp b/source/model/processor/strategies/Sri.cpp new file mode 100644 index 0000000..08a87d5 --- /dev/null +++ b/source/model/processor/strategies/Sri.cpp @@ -0,0 +1,126 @@ +#include "model/processor/strategies/Sri.h" +#include "model/processor/BIP.h" +#include +#include + +Sri::Sri(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "Sri", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (Sri_t)); +} + +void Sri::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.sri.isUpdated) + { + _kernelParameters.sri = parameters.sri.value; + } +} + +void Sri::finalize() +{ +} + + +Image* Sri::processKernel(Image *frames, Buffer* scratchPad) +{ + auto format = frames->getImageInfo(); + auto width = frames->getImageInfo(); + auto height = frames->getImageInfo(); + + _kernelParameters.width = width; + _kernelParameters.height = height; + + auto imageOutput = new Image2D(_CLContext, + CL_MEM_READ_WRITE, + ImageFormat(format.image_channel_order, format.image_channel_data_type), + width, + height); + if(_kernelParameters.sri != 0) + { + //width should be lass than CL_DEVICE_MAX_WORK_GROUP_SIZE(1024 for current machine) or this will break + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height), cl::NDRange(width, 1)); + + _kernelParameters.state = 1; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + Local(width * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + cl::EnqueueArgs eargs2(BIP::getInstance()->CLQueue, cl::NDRange(height, 1), cl::NDRange(height, 1)); + + _kernelParameters.state = 2; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs2, + *static_cast(frames), + *imageOutput, + Local(height * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + + _kernelParameters.state = 3; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + Local(1), + *scratchPad, + _kernelParameters); + + + _kernelParameters.state = 4; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(imageOutput), + *imageOutput, + Local(width * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + + _kernelParameters.state = 5; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs2, + *static_cast(frames), + *imageOutput, + Local(height * sizeof (myflt)), + *scratchPad, + _kernelParameters); + + cl::EnqueueArgs eargs3(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); + _kernelParameters.state = 6; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs3, + *static_cast(frames), + *imageOutput, + Local(1), + *scratchPad, + _kernelParameters); + + } + else + { + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + Local(49 * sizeof (myflt)), + *scratchPad, + _kernelParameters); + } + + + delete frames; + return imageOutput; +} + + diff --git a/source/model/processor/strategies/TintMap.cpp b/source/model/processor/strategies/TintMap.cpp new file mode 100644 index 0000000..8492676 --- /dev/null +++ b/source/model/processor/strategies/TintMap.cpp @@ -0,0 +1,57 @@ +#include "model/processor/strategies/TintMap.h" +#include "model/processor/BIP.h" +#include +#include + +TintMap::TintMap(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "TintMap", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (TintMap_t)); +} + +void TintMap::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.tintMapSelector.isUpdated) + { + _kernelParameters.tintMapSelector = parameters.tintMapSelector.value; + } +} + +void TintMap::finalize() +{ +} + + +Image* TintMap::processKernel(Image *frames, Buffer* scratchPad) +{ + auto width = frames->getImageInfo(); + auto height = frames->getImageInfo(); + + ImageFormat format; + format.image_channel_order = CL_RGBA; + format.image_channel_data_type = CL_UNSIGNED_INT8; + + auto imageOutput = new Image2D(_CLContext, + CL_MEM_READ_WRITE, + format, + width, + height); + //width should be lass than CL_DEVICE_MAX_WORK_GROUP_SIZE(1024 for current machine) or this will break + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); + + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + _kernelParameters); + + delete frames; + + return imageOutput; +} + + diff --git a/source/strategies/ScanConversion.cpp b/source/strategies/ScanConversion.cpp deleted file mode 100644 index 3bd169d..0000000 --- a/source/strategies/ScanConversion.cpp +++ /dev/null @@ -1,255 +0,0 @@ -#include "header/strategies/ScanConversion.h" -#include "MainWindow.h" -#include -#include - -ScanConversion::ScanConversion(const Context context, - const QString kernelPath) : - IProcessStrategy(context, kernelPath, "ScanConversion"), - _kernelFunctor(KernelFunctor(_kernel)) -{ - _gridPixelXPos = Q_NULLPTR; - _gridPixelZPos = Q_NULLPTR; - _scanXPos = Q_NULLPTR; - _scanZPos = Q_NULLPTR; - - _gridX = Q_NULLPTR; - _gridZ = Q_NULLPTR; - _queryX = Q_NULLPTR; - _queryZ = Q_NULLPTR; -} - -void ScanConversion::cpuProcess(ScenGenOutput_t parameters) -{ - auto context = _CLContext; - auto queue = MainWindow::getInstance()->CLQueue; - - if(parameters.outputWidth.isUpdated || parameters.outputHeight.isUpdated) - { - if(_gridPixelXPos) - delete[] _gridPixelXPos; - - if(_gridPixelZPos) - delete[] _gridPixelZPos; - - _gridPixelXPos = new myflt[parameters.outputWidth.value * - parameters.outputHeight.value]; - _gridPixelZPos = new myflt[parameters.outputWidth.value * - parameters.outputHeight.value]; - - if(_queryX) - delete _queryX; - - if(_queryZ) - delete _queryZ; - - _queryX = _openCLHelper.array2CLBuffer(context, - parameters.outputWidth.value * - parameters.outputHeight.value * - sizeof (myflt)); - _queryZ = _openCLHelper.array2CLBuffer(context, - parameters.outputWidth.value * - parameters.outputHeight.value * - sizeof (myflt)); - - _width = parameters.outputWidth.value; - _height = parameters.outputHeight.value; - - _kernelParameters.width = _width; - } - - if(parameters.depth.isUpdated || parameters.startDepth.isUpdated || - parameters.fieldOfView.isUpdated || parameters.steering.isUpdated || - parameters.outputWidth.isUpdated || parameters.outputHeight.isUpdated) - { - auto maxZ = parameters.depth.value; - auto minZ = parameters.startDepth.value; - auto maxX = parameters.fieldOfView.value / 2; - auto minX = -1 * maxX; - - auto baseZ = (maxZ - minZ) / (parameters.outputHeight.value - 1); - auto baseX = (maxX - minX) / (parameters.outputWidth.value - 1); - - for(auto i = 0; i < parameters.outputHeight.value; i++) - { - auto temp = minZ + i * baseZ; - auto tan_result = temp * tan(parameters.steering.value); - for(auto j = 0; j < parameters.outputWidth.value; j++) - { - _gridPixelZPos[i * parameters.outputHeight.value + j] = temp; - _gridPixelXPos[i * parameters.outputHeight.value + j] = minX + (baseX * j) - tan_result; - } - } - - queue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0, - parameters.outputWidth.value * - parameters.outputHeight.value * - sizeof (myflt), - _gridPixelXPos); - queue.enqueueWriteBuffer(*_queryZ, CL_TRUE, 0, - parameters.outputWidth.value * - parameters.outputHeight.value * - sizeof (myflt), - _gridPixelZPos); - } - - if(parameters.fieldOfView.isUpdated || parameters.rxLineNo.isUpdated || - parameters.rxLineDaz.isUpdated) - { - if(_scanXPos) - delete [] _scanXPos; - - if(_gridX) - delete _gridX; - - _scanXPos = new myflt[parameters.rxLineNo.value]; - for(auto i = 0; i < parameters.rxLineNo.value; i++) - { - _scanXPos[i] = -parameters.fieldOfView.value / 2 + i * parameters.rxLineDaz.value; - } - - _gridX = _openCLHelper.array2CLBuffer(context, - static_cast(parameters.rxLineNo.value) - * sizeof (myflt)); - - queue.enqueueWriteBuffer(*_gridX, CL_TRUE, 0, - static_cast(parameters.rxLineNo.value) * sizeof (myflt), - _scanXPos); - - _kernelParameters.gridXSize = parameters.rxLineNo.value; - } - - if(parameters.rxFocusPointNo.isUpdated || parameters.rxPointDax.isUpdated) - { - if(_scanZPos) - delete[] _scanZPos; - - if(_gridZ) - delete _gridZ; - - _scanZPos = new myflt [parameters.rxFocusPointNo.value]; - for(auto i = 0; i < parameters.rxFocusPointNo.value; i++) - { - _scanZPos[i] = parameters.startDepth.value + i * parameters.rxPointDax.value; - } - - _gridZ = _openCLHelper.array2CLBuffer(context, - static_cast(parameters.rxFocusPointNo.value) - * sizeof (myflt)); - - queue.enqueueWriteBuffer(*_gridZ, CL_TRUE, 0, - static_cast(parameters.rxFocusPointNo.value) * sizeof (myflt), - _scanZPos); - - _kernelParameters.gridZSize = parameters.rxFocusPointNo.value; - } -} - -void ScanConversion::finalize() -{ - delete _gridPixelXPos; - delete _gridPixelZPos; - delete _scanXPos; - delete _scanZPos; - - delete _gridX; - delete _gridZ; - delete _queryX; - delete _queryZ; - - _gridPixelXPos = Q_NULLPTR; - _gridPixelZPos = Q_NULLPTR; - _scanXPos = Q_NULLPTR; - _scanZPos = Q_NULLPTR; - - _gridX = Q_NULLPTR; - _gridZ = Q_NULLPTR; - _queryX = Q_NULLPTR; - _queryZ = Q_NULLPTR; -} - -void ScanConversion::ReadParams(QString path, ScenGenOutput_t *params) -{ - QFile file(path); - file.open(QIODevice::ReadOnly); - - char data[8192]; - //first line is extra - file.readLine(data, 8192); - QString str(data); - str = str.remove(str.length() - 1, 1); - auto sl = str.split(","); - - //decode here - /***scenario specific code***/ - int index = 0; - auto temp = sl[index++].PARSE; - update_field(¶ms->linear, INP2MYFLT(temp) == 0); - - temp = sl[index++].PARSE; - update_field(¶ms->depth, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->probeRadius, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->fieldOfView, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->startDepth, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->steering, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->rxLineNo, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->rxFocusPointNo, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->rxLineDaz, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->rxPointDax, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->virtualConvex, INP2MYFLT(temp) != 0); - - temp = sl[index++].PARSE; - update_field(¶ms->vcMaxTheta, INP2MYFLT(temp)); - - temp = sl[index++].PARSE; - update_field(¶ms->angle, INP2MYFLT(temp)); - /***End of scenario specific code***/ -} - - -Image* ScanConversion::processKernel(Image *frames, Buffer* scratchPad) -{ - Context context = _openCLHelper.getContext(); - auto format = frames->getImageInfo(); - auto imageOutput = new Image2D(context, - CL_MEM_READ_WRITE, - ImageFormat(format.image_channel_order, format.image_channel_data_type), - _width, - _height); - cl::EnqueueArgs eargs(MainWindow::getInstance()->CLQueue, cl::NDRange(_width, _height)); - - _openCLHelper.runKernelFunctor(_kernelFunctor, - eargs, - *static_cast(frames), - *imageOutput, - *_gridX, - *_gridZ, - *_queryX, - *_queryZ, - _kernelParameters); - delete frames; - - return imageOutput; -} - -