From 4914b48062e8df9cc4cdf74a22a407a341c5594b Mon Sep 17 00:00:00 2001 From: sono Date: Sun, 12 Apr 2020 11:38:36 +0430 Subject: [PATCH] Add final changes --- MainWindow.cpp | 14 +- header/ScenarioParams.h | 3 + header/model/processor/strategies/Enhance.h | 37 ++++ header/model/processor/strategies/Persist.h | 41 +++++ kernels/DynamicContrast | 0 kernels/{Max.cl => Enhance.cl} | 81 ++++----- kernels/Persist.cl | 85 +++++++++ source/FileHelper.cpp | 8 +- source/model/processor/strategies/Enhance.cpp | 164 ++++++++++++++++++ source/model/processor/strategies/Persist.cpp | 136 +++++++++++++++ 10 files changed, 524 insertions(+), 45 deletions(-) create mode 100644 header/model/processor/strategies/Enhance.h create mode 100644 header/model/processor/strategies/Persist.h delete mode 100644 kernels/DynamicContrast rename kernels/{Max.cl => Enhance.cl} (51%) create mode 100644 kernels/Persist.cl create mode 100644 source/model/processor/strategies/Enhance.cpp create mode 100644 source/model/processor/strategies/Persist.cpp diff --git a/MainWindow.cpp b/MainWindow.cpp index b4e3235..0e79d3e 100644 --- a/MainWindow.cpp +++ b/MainWindow.cpp @@ -12,6 +12,8 @@ #include "model/processor/strategies/Sri.h" #include "model/processor/strategies/ScanConversion.h" #include "model/processor/strategies/Rejection.h" +#include "model/processor/strategies/Enhance.h" +#include "model/processor/strategies/Persist.h" #include "header/FileHelper.h" MainWindow* MainWindow::_instance; @@ -85,7 +87,8 @@ void MainWindow::registerStrategies() REGISTER_STRATEGY(TintMap) REGISTER_STRATEGY(GrayMap) REGISTER_STRATEGY(DynCont) - + REGISTER_STRATEGY(Enhance) + REGISTER_STRATEGY(Persist) } void MainWindow::pushBackStrategy(const QString strategyName, const QString kernelFolder) @@ -206,6 +209,15 @@ void MainWindow::readParam(QString path, int mode) temp = sl[index++].PARSE; update_field(&_scenGenOutput.compressionType, INP2MYFLT(temp)); break; + case 6: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.enhance, INP2MYFLT(temp)); + + update_field(&_scenGenOutput.enhanceAlgorithm, 0); + break; + case 7: + update_field(&_scenGenOutput.persist, 1); + break; } } diff --git a/header/ScenarioParams.h b/header/ScenarioParams.h index 7bb8f2c..a951f15 100644 --- a/header/ScenarioParams.h +++ b/header/ScenarioParams.h @@ -39,6 +39,9 @@ typedef struct ScenGenOutput_t field_t tintMapSelector; field_t sri; field_t rejectThreshold; + field_t enhance; + field_t enhanceAlgorithm; + field_t persist; }ScenGenOutput_t; template diff --git a/header/model/processor/strategies/Enhance.h b/header/model/processor/strategies/Enhance.h new file mode 100644 index 0000000..4a29ad1 --- /dev/null +++ b/header/model/processor/strategies/Enhance.h @@ -0,0 +1,37 @@ +#ifndef ENHANCE_H +#define ENHANCE_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +typedef struct Enhance_t +{ + cl_int filterWidth; + cl_int filterHeight; + cl_int width; + cl_int height; +}Enhance_t; + + +class Enhance : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE Enhance(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; + + Enhance_t _kernelParameters; + + bool _run; + + Buffer* _filter; +}; + +#endif // ENHANCE_H diff --git a/header/model/processor/strategies/Persist.h b/header/model/processor/strategies/Persist.h new file mode 100644 index 0000000..2febf62 --- /dev/null +++ b/header/model/processor/strategies/Persist.h @@ -0,0 +1,41 @@ +#ifndef PERSIST_H +#define PERSIST_H + +#include +#include +#include "model/processor/IProcessStrategy.h" +#include "utils/OpenCLHelper.h" + +#define PERSIST_MAX_SIZE 8 + +typedef struct Persist_t +{ + cl_int persist; + cl_int mode; +}Persist_t; + + +class Persist : public IProcessStrategy +{ + Q_OBJECT +public: + Q_INVOKABLE Persist(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; + + Persist_t _kernelParameters; + + QVector _queue; + Image* _imageSum; + + bool _reCalc; + quint64 _lastWidth; + quint64 _lastHeight; +}; + + +#endif // PERSIST_H diff --git a/kernels/DynamicContrast b/kernels/DynamicContrast deleted file mode 100644 index e69de29..0000000 diff --git a/kernels/Max.cl b/kernels/Enhance.cl similarity index 51% rename from kernels/Max.cl rename to kernels/Enhance.cl index 1449df7..e0e16fa 100644 --- a/kernels/Max.cl +++ b/kernels/Enhance.cl @@ -1,8 +1,8 @@ -#define TEST -#define X 0 -#define Y 26 +// #define TEST +#define X 95 +#define Y 22 -#define USE_DBL +//#define USE_DBL #ifdef USE_DBL #define TYPE_FLT double #define TYPE_INT long @@ -21,6 +21,14 @@ constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; + +struct input { + int filter_width; + int filter_height; + int width; + int height; +}; + TYPE_FLT read_data(image2d_t input_frame, int x, int z) { int2 gid = (int2)(x, z); @@ -34,52 +42,45 @@ TYPE_FLT read_data(image2d_t input_frame, int x, int z) } -kernel void Max(read_only image2d_t input_frame, read_write image2d_t output_frame, local double* max, global double* scratch_pad) { +kernel void Enhance(read_only image2d_t input_frame, read_write image2d_t output_frame, global TYPE_FLT* filter, struct input params) +{ 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; + TYPE_FLT output_data = 0; - uint local_id = get_local_id(0); - uint group_size = get_local_size(0); - max[local_id] = input; + int p = params.filter_height; + int q = params.filter_width; - 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]; + int lag1 = (p - 1) / 2; + int lag2 = (q - 1) / 2; - if (local_id == 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++) { - if(is_odd) - max[local_id] = max[local_id] > max[2 * stride] ? max[local_id] : max[2 * stride]; - is_odd = stride % 2; - } - } + for(int j = -lag2; j <= lag2 ; j++) + { + if(gid.x == 947 && gid.y == 1) + { + printf("input: %f; ", output_data); + } - // Write result into scratchPad[nWorkGroups] - if (local_id == 0) + output_data += read_data(input_frame, gid.x + j, gid.y + i) * filter[(i + lag1) * q + (j + lag2)]; + } + } + } + else { - scratch_pad[get_group_id(1)] = max[0]; - printf("out[%d]: %.15f | ", get_group_id(1), scratch_pad[get_group_id(1)]); + output_data = input; } - + if(output_data < 0) + output_data = 0; + + if(output_data > 255) + output_data = 255; + + TYPE_INT out = *((TYPE_INT*)(&output_data)); uint4 pixel; pixel.x = (TYPE_INT)(out & MASK); diff --git a/kernels/Persist.cl b/kernels/Persist.cl new file mode 100644 index 0000000..8e8fd0d --- /dev/null +++ b/kernels/Persist.cl @@ -0,0 +1,85 @@ +// #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 persist; + int mode; +}; + +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 Persist(read_only image2d_t input_frame, read_write image2d_t output_frame, read_write image2d_t sum_frame, read_write image2d_t removed_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; + TYPE_FLT sum = 0; + + if(params.mode == 0) + { + sum = input; + } + else if(params.mode == 1) + { + sum = read_data(sum_frame, gid.x, gid.y); + sum += input; + } + else if(params.mode == 2) + { + sum = read_data(sum_frame, gid.x, gid.y); + sub = read_data(removed_frame, gid.x, gid.y); + sum = sum + input - sub; + } + + output_data = sum / params.persist; + + 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); + + out = *((TYPE_INT*)(&sum)); + 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(sum_frame, gid, pixel); +} + diff --git a/source/FileHelper.cpp b/source/FileHelper.cpp index 86cac45..eb14064 100644 --- a/source/FileHelper.cpp +++ b/source/FileHelper.cpp @@ -9,10 +9,10 @@ bool FileHelper::ReadInputFile(QByteArray& arr, QString path, quint64 *width, qu QFile file(path); file.open(QIODevice::ReadOnly); - char data[20000]; + char data[32000]; bool first = true; - while (file.readLine(data, 20000)) + while (file.readLine(data, 32000)) { QString str(data); if(str.length() == 0) @@ -54,11 +54,11 @@ bool FileHelper::ReadInputFile(myint *arr, QString path, quint64 *width, quint64 QFile file(path); file.open(QIODevice::ReadOnly); - char data[8192]; + char data[32000]; bool first = true; int index = 0; - while (file.readLine(data, 8192)) + while (file.readLine(data, 32000)) { QString str(data); if(str.length() == 0) diff --git a/source/model/processor/strategies/Enhance.cpp b/source/model/processor/strategies/Enhance.cpp new file mode 100644 index 0000000..5a876cf --- /dev/null +++ b/source/model/processor/strategies/Enhance.cpp @@ -0,0 +1,164 @@ +#include "model/processor/strategies/Enhance.h" +#include "model/processor/BIP.h" +#include +#include + +Enhance::Enhance(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "Enhance", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (Enhance_t)); + _filter = Q_NULLPTR; +} + +void Enhance::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.enhance.isUpdated || parameters.enhanceAlgorithm.isUpdated) + { + _run = false; + if(parameters.enhance.value == 0) + { + return; + } + + float *hpf; + int *backgroundRecovery; + + //sharpen + if(parameters.enhanceAlgorithm.value == 0) + { + auto alpha = 0.25f * (parameters.enhance.value - 1); + + _kernelParameters.filterWidth = 3; + _kernelParameters.filterHeight = 3; + + hpf = new float[9]; + backgroundRecovery = new int[9]; + + auto factor = 4 / (alpha + 1); + int index = 0; + + hpf[index++] = factor * (alpha / 4); + hpf[index++] = factor * ((1 - alpha) / 4); + hpf[index++] = factor * (alpha / 4); + + hpf[index++] = factor * ((1 - alpha) / 4); + hpf[index++] = factor * -1; + hpf[index++] = factor * ((1 - alpha) / 4); + + hpf[index++] = factor * (alpha / 4); + hpf[index++] = factor * ((1 - alpha) / 4); + hpf[index++] = factor * (alpha / 4); + + index = 0; + + memset(backgroundRecovery, 0, 9 * sizeof (int)); + backgroundRecovery[4] = 1; + + _run = true; + } + //log + else if(parameters.enhanceAlgorithm.value == 1) + { + _kernelParameters.filterWidth = 5; + _kernelParameters.filterHeight = 5; + + hpf = new float[25]; + backgroundRecovery = new int[25]; + + auto sigma = 0.5f + 0.25f * (parameters.enhance.value - 1); + float sumHg = 0; + float h[25]; + + for(int i = -2; i <= 2; i++) + { + for(int j = -2; j <= 2; j++) + { + float temp = expf(-(powf(i, 2) + powf(j, 2)) + / 2 / powf(sigma, 2)); + h[(i + 2) * 5 + (j + 2)] = (powf(i, 2) + powf(j, 2) - 2 * powf(sigma, 2)) * + temp; + + sumHg += temp; + } + } + + for(int i = 0; i < 25 ; i++) + { + hpf[i] = h[i] / powf(sigma, 4) / sumHg; + } + + memset(backgroundRecovery, 0, 25 * sizeof (int)); + backgroundRecovery[12] = 1; + + _run = true; + } + + if(_run) + { + auto filterSize = _kernelParameters.filterWidth * _kernelParameters.filterHeight; + float filter[filterSize]; + for(int i = 0; i < filterSize ; i++) + { + filter[i] = backgroundRecovery[i] - hpf[i]; + } + + delete [] hpf; + delete [] backgroundRecovery; + + if(_filter) + delete _filter; + + _filter = _openCLHelper.array2CLBuffer(_CLContext, filterSize * sizeof (float)); + BIP::getInstance()->CLQueue.enqueueWriteBuffer(*_filter, CL_TRUE, + 0, filterSize * sizeof (float), + filter); + } + } +} + +void Enhance::finalize() +{ + if(_filter) + delete _filter; +} + + +Image* Enhance::processKernel(Image *frames, Buffer* scratchPad) +{ + if(_run) + { + 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); + cl::EnqueueArgs eargs(BIP::getInstance()->CLQueue, cl::NDRange(width, height)); + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *imageOutput, + *_filter, + _kernelParameters); + + + delete frames; + return imageOutput; + } + else + { + return frames; + } +} + + diff --git a/source/model/processor/strategies/Persist.cpp b/source/model/processor/strategies/Persist.cpp new file mode 100644 index 0000000..2c97f30 --- /dev/null +++ b/source/model/processor/strategies/Persist.cpp @@ -0,0 +1,136 @@ +#include "model/processor/strategies/Persist.h" +#include "model/processor/BIP.h" +#include +#include + +Persist::Persist(const Context context, + const QString kernelPath, + const QObject *parent = Q_NULLPTR) : + IProcessStrategy(context, kernelPath, "Persist", parent), + _kernelFunctor(KernelFunctor(_kernel)) +{ + memset(&_kernelParameters, 0, sizeof (Persist_t)); + _imageSum = Q_NULLPTR; + _reCalc = true; +} + +void Persist::cpuProcess(ScenGenOutput_t parameters) +{ + if(parameters.persist.isUpdated) + { + _kernelParameters.persist = parameters.persist.value; + _reCalc = true; + } +} + +void Persist::finalize() +{ + foreach (auto frame, _queue) + { + delete frame; + } + _queue.clear(); + + if(_imageSum) + delete _imageSum; +} + + +Image* Persist::processKernel(Image *frames, Buffer* scratchPad) +{ + Image* temp; + + if(_queue.length() < PERSIST_MAX_SIZE) + { + _queue.push_front(frames); + return frames; + } + + if(_kernelParameters.persist != 1) + { + auto format = frames->getImageInfo(); + auto width = frames->getImageInfo(); + auto height = frames->getImageInfo(); + + //if width or height changes we should erase everything + if(_lastWidth != width || _lastHeight != height) + { + _reCalc = false; + + if(_imageSum) + delete _imageSum; + + _imageSum = new Image2D(_CLContext, + CL_MEM_READ_WRITE, + ImageFormat(format.image_channel_order, format.image_channel_data_type), + width, + height); + + foreach (auto frame, _queue) + { + delete frame; + } + _queue.clear(); + + _lastWidth = width; + _lastHeight = height; + } + + + auto outputImage = new Image2D(_CLContext, + 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)); + + if(_reCalc) + { + _reCalc = false; + for(auto i = 0; i < _kernelParameters.persist; i++) + { + if(i == 0) + _kernelParameters.mode = 0; + else if(i < _kernelParameters.persist) + _kernelParameters.mode = 1; + + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(_queue[i]), + *outputImage, + *static_cast(_imageSum), + *static_cast(_queue[_kernelParameters.persist - 1]), + _kernelParameters); + } + } + + _kernelParameters.mode = 2; + _openCLHelper.runKernelFunctor(_kernelFunctor, + eargs, + *static_cast(frames), + *outputImage, + *static_cast(_imageSum), + *static_cast(_queue[_kernelParameters.persist - 1]), + _kernelParameters); + + temp = outputImage; + } + else + { + temp = frames; + } + + delete _queue[PERSIST_MAX_SIZE - 1]; + for(auto i = 0; i < PERSIST_MAX_SIZE - 1; i++) + { + _queue[i + 1] = _queue[i]; + } + + _queue[0] = frames; + + + return temp; +}