mmtalaie
/
kernel-tester
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
Browse Source

Add final changes

test-kernel
sono 5 years ago
parent
d81553a3ba
commit
4914b48062
10 changed files with 524 additions and 45 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 14
      MainWindow.cpp
  2. 3
      header/ScenarioParams.h
  3. 37
      header/model/processor/strategies/Enhance.h
  4. 41
      header/model/processor/strategies/Persist.h
  5. 0
      kernels/DynamicContrast
  6. 81
      kernels/Enhance.cl
  7. 85
      kernels/Persist.cl
  8. 8
      source/FileHelper.cpp
  9. 164
      source/model/processor/strategies/Enhance.cpp
  10. 136
      source/model/processor/strategies/Persist.cpp

14
MainWindow.cpp
View File

@ -12,6 +12,8 @@
#include "model/processor/strategies/Sri.h" #include "model/processor/strategies/Sri.h"
#include "model/processor/strategies/ScanConversion.h" #include "model/processor/strategies/ScanConversion.h"
#include "model/processor/strategies/Rejection.h" #include "model/processor/strategies/Rejection.h"
#include "model/processor/strategies/Enhance.h"
#include "model/processor/strategies/Persist.h"
#include "header/FileHelper.h" #include "header/FileHelper.h"
MainWindow* MainWindow::_instance; MainWindow* MainWindow::_instance;
@ -85,7 +87,8 @@ void MainWindow::registerStrategies()
REGISTER_STRATEGY(TintMap) REGISTER_STRATEGY(TintMap)
REGISTER_STRATEGY(GrayMap) REGISTER_STRATEGY(GrayMap)
REGISTER_STRATEGY(DynCont) REGISTER_STRATEGY(DynCont)
REGISTER_STRATEGY(Enhance)
REGISTER_STRATEGY(Persist)
} }
void MainWindow::pushBackStrategy(const QString strategyName, const QString kernelFolder) void MainWindow::pushBackStrategy(const QString strategyName, const QString kernelFolder)
@ -206,6 +209,15 @@ void MainWindow::readParam(QString path, int mode)
temp = sl[index++].PARSE; temp = sl[index++].PARSE;
update_field<int>(&_scenGenOutput.compressionType, INP2MYFLT(temp)); update_field<int>(&_scenGenOutput.compressionType, INP2MYFLT(temp));
break; break;
case 6:
temp = sl[index++].PARSE;
update_field<int>(&_scenGenOutput.enhance, INP2MYFLT(temp));
update_field<int>(&_scenGenOutput.enhanceAlgorithm, 0);
break;
case 7:
update_field<int>(&_scenGenOutput.persist, 1);
break;
} }
} }

3
header/ScenarioParams.h
View File

@ -39,6 +39,9 @@ typedef struct ScenGenOutput_t
field_t<int> tintMapSelector; field_t<int> tintMapSelector;
field_t<int> sri; field_t<int> sri;
field_t<int> rejectThreshold; field_t<int> rejectThreshold;
field_t<int> enhance;
field_t<int> enhanceAlgorithm;
field_t<int> persist;
}ScenGenOutput_t; }ScenGenOutput_t;
template<typename T> template<typename T>

37
header/model/processor/strategies/Enhance.h
View File

@ -0,0 +1,37 @@
#ifndef ENHANCE_H
#define ENHANCE_H
#include <QObject>
#include <QMetaType>
#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<Image2D, Image2D, Buffer, Enhance_t> _kernelFunctor;
virtual Image* processKernel(Image *frames, Buffer* scratchPad) override;
Enhance_t _kernelParameters;
bool _run;
Buffer* _filter;
};
#endif // ENHANCE_H

41
header/model/processor/strategies/Persist.h
View File

@ -0,0 +1,41 @@
#ifndef PERSIST_H
#define PERSIST_H
#include <QObject>
#include <QMetaType>
#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<Image2D, Image2D, Image2D, Image2D, Persist_t> _kernelFunctor;
virtual Image* processKernel(Image *frames, Buffer* scratchPad) override;
Persist_t _kernelParameters;
QVector<Image*> _queue;
Image* _imageSum;
bool _reCalc;
quint64 _lastWidth;
quint64 _lastHeight;
};
#endif // PERSIST_H

0
kernels/DynamicContrast
View File

81
kernels/Max.cl → kernels/Enhance.cl
View File

@ -1,8 +1,8 @@
#define TEST // #define TEST
#define X 0 #define X 95
#define Y 26 #define Y 22
#define USE_DBL //#define USE_DBL
#ifdef USE_DBL #ifdef USE_DBL
#define TYPE_FLT double #define TYPE_FLT double
#define TYPE_INT long #define TYPE_INT long
@ -21,6 +21,14 @@ constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE | CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_LINEAR; 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) TYPE_FLT read_data(image2d_t input_frame, int x, int z)
{ {
int2 gid = (int2)(x, 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)); 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); 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); int p = params.filter_height;
uint group_size = get_local_size(0); int q = params.filter_width;
max[local_id] = input;
if(gid.x == 0 && gid.y == 0) int lag1 = (p - 1) / 2;
{ int lag2 = (q - 1) / 2;
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(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) for(int j = -lag2; j <= lag2 ; j++)
max[local_id] = max[local_id] > max[2 * stride] ? max[local_id] : max[2 * stride]; {
is_odd = stride % 2; if(gid.x == 947 && gid.y == 1)
} {
} printf("input: %f; ", output_data);
}
// Write result into scratchPad[nWorkGroups] output_data += read_data(input_frame, gid.x + j, gid.y + i) * filter[(i + lag1) * q + (j + lag2)];
if (local_id == 0) }
}
}
else
{ {
scratch_pad[get_group_id(1)] = max[0]; output_data = input;
printf("out[%d]: %.15f | ", get_group_id(1), scratch_pad[get_group_id(1)]);
} }
if(output_data < 0)
output_data = 0;
if(output_data > 255)
output_data = 255;
TYPE_INT out = *((TYPE_INT*)(&output_data)); TYPE_INT out = *((TYPE_INT*)(&output_data));
uint4 pixel; uint4 pixel;
pixel.x = (TYPE_INT)(out & MASK); pixel.x = (TYPE_INT)(out & MASK);

85
kernels/Persist.cl
View File

@ -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);
}

8
source/FileHelper.cpp
View File

@ -9,10 +9,10 @@ bool FileHelper::ReadInputFile(QByteArray& arr, QString path, quint64 *width, qu
QFile file(path); QFile file(path);
file.open(QIODevice::ReadOnly); file.open(QIODevice::ReadOnly);
char data[20000]; char data[32000];
bool first = true; bool first = true;
while (file.readLine(data, 20000)) while (file.readLine(data, 32000))
{ {
QString str(data); QString str(data);
if(str.length() == 0) if(str.length() == 0)
@ -54,11 +54,11 @@ bool FileHelper::ReadInputFile(myint *arr, QString path, quint64 *width, quint64
QFile file(path); QFile file(path);
file.open(QIODevice::ReadOnly); file.open(QIODevice::ReadOnly);
char data[8192]; char data[32000];
bool first = true; bool first = true;
int index = 0; int index = 0;
while (file.readLine(data, 8192)) while (file.readLine(data, 32000))
{ {
QString str(data); QString str(data);
if(str.length() == 0) if(str.length() == 0)

164
source/model/processor/strategies/Enhance.cpp
View File

@ -0,0 +1,164 @@
#include "model/processor/strategies/Enhance.h"
#include "model/processor/BIP.h"
#include <QPixmap>
#include <QImage>
Enhance::Enhance(const Context context,
const QString kernelPath,
const QObject *parent = Q_NULLPTR) :
IProcessStrategy(context, kernelPath, "Enhance", parent),
_kernelFunctor(KernelFunctor<Image2D, Image2D,
Buffer, Enhance_t>(_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<CL_IMAGE_FORMAT>();
auto width = frames->getImageInfo<CL_IMAGE_WIDTH>();
auto height = frames->getImageInfo<CL_IMAGE_HEIGHT>();
_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<Image2D, Image2D, Buffer, Enhance_t>(_kernelFunctor,
eargs,
*static_cast<Image2D*>(frames),
*imageOutput,
*_filter,
_kernelParameters);
delete frames;
return imageOutput;
}
else
{
return frames;
}
}

136
source/model/processor/strategies/Persist.cpp
View File

@ -0,0 +1,136 @@
#include "model/processor/strategies/Persist.h"
#include "model/processor/BIP.h"
#include <QPixmap>
#include <QImage>
Persist::Persist(const Context context,
const QString kernelPath,
const QObject *parent = Q_NULLPTR) :
IProcessStrategy(context, kernelPath, "Persist", parent),
_kernelFunctor(KernelFunctor<Image2D, Image2D,
Image2D, Image2D, Persist_t>(_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<CL_IMAGE_FORMAT>();
auto width = frames->getImageInfo<CL_IMAGE_WIDTH>();
auto height = frames->getImageInfo<CL_IMAGE_HEIGHT>();
//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<Image2D, Image2D, Image2D, Image2D, Persist_t>(_kernelFunctor,
eargs,
*static_cast<Image2D*>(_queue[i]),
*outputImage,
*static_cast<Image2D*>(_imageSum),
*static_cast<Image2D*>(_queue[_kernelParameters.persist - 1]),
_kernelParameters);
}
}
_kernelParameters.mode = 2;
_openCLHelper.runKernelFunctor<Image2D, Image2D, Image2D, Image2D, Persist_t>(_kernelFunctor,
eargs,
*static_cast<Image2D*>(frames),
*outputImage,
*static_cast<Image2D*>(_imageSum),
*static_cast<Image2D*>(_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;
}
Write Preview
Loading…
Cancel
Save