From 2711067741435d0be09e6d6fc89d8f15fe93057a Mon Sep 17 00:00:00 2001 From: mmtalaie Date: Wed, 26 Aug 2020 18:58:00 +0430 Subject: [PATCH] CRi init --- MainWindow.cpp | 29 ++++-- header/model/processor/strategies/Cri.h | 4 +- kernels/Cri.cl | 109 +++++++++++++++++++++- source/model/processor/strategies/Cri.cpp | 6 +- 4 files changed, 134 insertions(+), 14 deletions(-) diff --git a/MainWindow.cpp b/MainWindow.cpp index b9c2265..b27a1ff 100644 --- a/MainWindow.cpp +++ b/MainWindow.cpp @@ -215,11 +215,19 @@ void MainWindow::readParam(QString path, int mode) temp = sl[index++].PARSE; update_field(&_scenGenOutput.enhance, INP2MYFLT(temp)); - update_field(&_scenGenOutput.enhanceAlgorithm, 0); + update_field(&_scenGenOutput.enhanceAlgorithm, 0); break; case 7: update_field(&_scenGenOutput.persist, 1); - break; + break; + case 8: + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.criFilterMode, INP2MYFLT(temp)); + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.frameCntr, INP2MYFLT(temp)); + temp = sl[index++].PARSE; + update_field(&_scenGenOutput.scenarioFrameNo, INP2MYFLT(temp)); + break; } } @@ -255,7 +263,7 @@ void MainWindow::on_btn_test_clicked() if (!checkPath(ui->tb_workingDir->text(), "Working")) return; - QDir inputDir(ui->tb_inpDir->text()); + QDir inputDir(ui->tb_inpDir->text()); QDir outputDir(ui->tb_outpDir->text()); QDir workingDir(ui->tb_workingDir->text()); @@ -270,13 +278,22 @@ void MainWindow::on_btn_test_clicked() _strategy->finalize(); pushBackStrategy(ui->cb_kernelName->currentText(), workingDir.path() + "/kernels"); - QtConcurrent::run([this, inputDir, outputDir](){ + QtConcurrent::run([this, inputDir, outputDir]() + { auto inputs = inputDir.entryList(QStringList() << "*.csv", QDir::Files); foreach(QString filename, inputs) { addToConsole("processing " + filename); auto dataPath = inputDir.path() + "/" + filename; auto paramPath = inputDir.path() + "/params/" + filename; + qDebug()< 0; fp--) + { + paramPath.replace(QStringLiteral("_%1").arg(fp),""); + } + qDebug()<cb_kernelName->currentIndex()); @@ -315,7 +332,7 @@ void MainWindow::on_btn_test_clicked() auto bufferPath = inputDir.path() + "/buffer"; for (int i = 1; i <= 16; ++i) { - auto bfname = QString("%1\%2\") + // auto bfname = QString("%1\%2\"); } @@ -340,7 +357,7 @@ void MainWindow::on_btn_test_clicked() addToConsole(filename + " has been processed"); addToConsole("*************************\r\n"); } - }); + }); } void MainWindow::on_btn_save_clicked() diff --git a/header/model/processor/strategies/Cri.h b/header/model/processor/strategies/Cri.h index ea8a5ae..58db251 100644 --- a/header/model/processor/strategies/Cri.h +++ b/header/model/processor/strategies/Cri.h @@ -7,7 +7,7 @@ #include "model/processor/IProcessStrategy.h" #include "utils/OpenCLHelper.h" -#define CRI_MAX_BUFFER_SIZE 16 +#define CRI_MAX_BUFFER_SIZE 11 typedef struct Cri_t @@ -31,5 +31,7 @@ private: virtual Image* processKernel(Image *frames, Buffer* scratchPad) override; Cri_t _kernelParameters; Image2D* _criBuffer[CRI_MAX_BUFFER_SIZE]; + // + int counter = 0; }; #endif // CRI_H diff --git a/kernels/Cri.cl b/kernels/Cri.cl index 2fcc43e..87a8d77 100644 --- a/kernels/Cri.cl +++ b/kernels/Cri.cl @@ -1,16 +1,117 @@ +// #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 criFilterMode; int frameCntr; int scenariFrameNo; + float zeroSteeringFrameWeight; + float nonZeroSteeringFrameWeight; }; -kernel void Cri(read_only image2d_array_t input_frame, read_write image2d_t output_frame, struct input params) + +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 Cri(read_only image2d_array_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[0], gid.x, gid.y); +// int4 gid = ((int4)(get_global_size(0), get_global_size(1),get_global_size(2) , 0)); +// // printf(" :hi! I am the cri kernel 1"); +// // const int rows = get_image_height(input_frame); +// // const int cols = get_image_width(input_frame); +// printf("gid : %d\n",gid.z); +// // printf(input_frame); +// // int4 f = read_imagei(input_frame, (rows,cols,1,1)); +// // printf("out: %a | ", f.x); +// // output_frame = input_frame; +// printf(" :hi! I am the cri kernel"); +// } +int isNonzeropixel(int x, int y, image2d_t arrayInput) { + if(read_data(x,y,input_frame) == 0.0F) + return 0; + else + return 1; +} + +kernel void Cri(read_only image2d_t input_frame1, read_only image2d_t input_frame2, read_only image2d_t input_frame3, + read_only image2d_t input_frame4, read_only image2d_t input_frame5, read_only image2d_t input_frame6, + read_only image2d_t input_frame7, read_only image2d_t input_frame8, read_only image2d_t input_frame9, + read_only image2d_t input_frame10, read_only image2d_t input_frame11, read_write image2d_t output_frame, struct input params) +{ + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + image2d_t imgArray[11] = {input_frame1, input_frame2, input_frame3, + input_frame4, input_frame5, input_frame6, + input_frame7, input_frame8, input_frame9, + input_frame10, input_frame11}; + const int rows = get_image_height(input_frame); const int cols = get_image_width(input_frame); - int4 f = read_imagei(input_frame, (rows,cols,1,1)); - printf("out: %a | ", f.x); - // printf(" :hi! I am the cri kernel"); + + int nonZeroFrameNo = 0; + int i = 1; + for(i= 1; i < 11;i++) + { + nonZeroFrameNo = nonZeroFrameNo + isNonzeropixel(gid.x,gid.y,imgArray[i]); + } + + float additionalPoint = (1 - params.zeroSteeringFrameWeight - nonZeroFrameNo * params.nonZeroSteeringFrameWeight); + float pixelWeight[12] = {0.0F}; + float sum = 0; + for(i = 1; i < 11; i++) + { + pixelWeight[i] = params.nonZeroSteeringFrameWeight + additionalPoint * params.nonZeroSteeringFrameWeight/(params.nonZeroSteeringFrameWeight * nonZeroFrameNo + params.zeroSteeringFrameWeight); + sum = pixelWeight[i]; + } + pixelWeight[0] = 1 - sum; + + TYPE_FLT output_data = 0; + for(i = 0; i < 11; i++) + { + output_data = output_data + (pixelWeight[i] * read_data(gid.x,gid.y,imgArray[i]) + } + + 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); } \ No newline at end of file diff --git a/source/model/processor/strategies/Cri.cpp b/source/model/processor/strategies/Cri.cpp index fabaf50..fe1e9e8 100644 --- a/source/model/processor/strategies/Cri.cpp +++ b/source/model/processor/strategies/Cri.cpp @@ -32,7 +32,7 @@ Image* Cri::processKernel(Image *frames, Buffer* scratchPad) auto width = frames->getImageInfo(); auto height = frames->getImageInfo(); - _criBuffer[_kernelParameters.scenariFrameNo] = static_cast(frames); + _criBuffer[counter++] = static_cast(frames); int err = 100; // Image2DArray bufferframes = Image2DArray(_CLContext, CL_MEM_READ_WRITE, ImageFormat(format.image_channel_order, format.image_channel_data_type), @@ -46,10 +46,10 @@ Image* Cri::processKernel(Image *frames, Buffer* scratchPad) // 0, // _criBuffer); auto imgs = new Image2DArray(_CLContext, CL_MEM_READ_WRITE, ImageFormat(format.image_channel_order, format.image_channel_data_type), - 16, + 11, width, height, - width *4* sizeof (myflt), + 0, 0, _criBuffer, &err);