diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..89f64c7 --- /dev/null +++ b/.gitignore @@ -0,0 +1 @@ +*.pro.* diff --git a/MainWindow.cpp b/MainWindow.cpp index becda78..23175a5 100644 --- a/MainWindow.cpp +++ b/MainWindow.cpp @@ -3,6 +3,7 @@ #include #include +#include #include "header/strategies/ScanConversion.h" #include "header/strategies/Rejection.h" @@ -141,61 +142,63 @@ void MainWindow::on_btn_test_clicked() _strategy->finalize(); pushBackStrategy(ui->cb_kernelName->currentText(), workingDir.path() + "/kernels"); - 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; - - _strategy->ReadParams(paramPath, &_scenGenOutput); + 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; + _strategy->ReadParams(paramPath, &_scenGenOutput); - QByteArray arr; - quint64 width; - quint64 height = 0; - if (!FileHelper::ReadInputFile(arr, dataPath, &width, &height)) - { - addToConsole("Some thing wnet wrong in input file"); - abort(); - return; - } + QByteArray arr; + quint64 width; + quint64 height = 0; + if (!FileHelper::ReadInputFile(arr, dataPath, &width, &height)) + { + addToConsole("Some thing wnet wrong in input file"); + abort(); + return; + } - ImageFormat format; - format.image_channel_order = CL_RGBA; + ImageFormat format; + format.image_channel_order = CL_RGBA; #ifdef USE_DBL - format.image_channel_data_type = CL_UNSIGNED_INT16; + format.image_channel_data_type = CL_UNSIGNED_INT16; #else - format.image_channel_data_type = CL_UNSIGNED_INT8; + format.image_channel_data_type = CL_UNSIGNED_INT8; #endif - Image2D* inFrame = static_cast(_openCLHelper - .frame2CLFrame(_CLContext, format, QVector{width, height}, true)); - - 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); - _strategy->cpuProcess(_scenGenOutput); - auto outFrame = _strategy->processKernel(inFrame); - - char *out; - out = (char*)malloc(OUT_WIDTH * OUT_WIDTH * sizeof (char) * sizeof (myflt)); - CLQueue.enqueueReadImage(*outFrame, - CL_TRUE, - array {0, 0, 0}, - array {OUT_WIDTH, OUT_HEIGHT, 1}, - OUT_WIDTH * sizeof (myflt), - 0, - out); - - FileHelper::WriteOutputFile(out, outputDir.path() + "/" + filename, OUT_WIDTH, OUT_HEIGHT); - addToConsole(filename + " has been processed"); - } + Image2D* inFrame = static_cast(_openCLHelper + .frame2CLFrame(_CLContext, format, QVector{width, height}, true)); + + 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); + _strategy->cpuProcess(_scenGenOutput); + auto outFrame = _strategy->processKernel(inFrame); + + char *out; + out = (char*)malloc(OUT_WIDTH * OUT_WIDTH * sizeof (char) * sizeof (myflt)); + CLQueue.enqueueReadImage(*outFrame, + CL_TRUE, + array {0, 0, 0}, + array {OUT_WIDTH, OUT_HEIGHT, 1}, + OUT_WIDTH * sizeof (myflt), + 0, + out); + + FileHelper::WriteOutputFile(out, outputDir.path() + "/" + filename, OUT_WIDTH, OUT_HEIGHT); + addToConsole(filename + " has been processed"); + addToConsole("*************************\r\n"); + } + }); } void MainWindow::on_btn_save_clicked() @@ -225,53 +228,67 @@ void MainWindow::on_btn_compare_clicked() QDir matlabDir(ui->tb_matlabOutp->text()); QDir outputDir(ui->tb_outpDir->text()); - auto totalMax = 0U; - auto outputs = outputDir.entryList(QStringList() << "*.csv", QDir::Files); - foreach(QString filename, outputs) - { - addToConsole("Comaring " + filename); - auto cppout = new myint[OUT_WIDTH * OUT_HEIGHT]; - auto matout = new myint[OUT_WIDTH * OUT_HEIGHT]; - quint64 w1, w2; - quint64 h1 = 0, h2 = 0; + QtConcurrent::run([this, outputDir, matlabDir](){ + auto totalMax = 0UL; + auto outputs = outputDir.entryList(QStringList() << "*.csv", QDir::Files); + foreach(QString filename, outputs) + { + addToConsole("Comaring " + filename); + auto cppout = new myint[OUT_WIDTH * OUT_HEIGHT]; + auto matout = new myint[OUT_WIDTH * OUT_HEIGHT]; + quint64 w1, w2; + quint64 h1 = 0, h2 = 0; - FileHelper::ReadInputFile(cppout, outputDir.path() + "/" + filename, &w1, &h1); + FileHelper::ReadInputFile(cppout, outputDir.path() + "/" + filename, &w1, &h1); - if(!QFile::exists(matlabDir.path() + "/" + filename)) - { - addToConsole("matlab file for " + filename + " does not exist"); - continue; - } - FileHelper::ReadInputFile(matout, matlabDir.path() + "/" + filename, &w2, &h2); + if(!QFile::exists(matlabDir.path() + "/" + filename)) + { + addToConsole("matlab file for " + filename + " does not exist"); + continue; + } + FileHelper::ReadInputFile(matout, matlabDir.path() + "/" + filename, &w2, &h2); - if (w1 != OUT_WIDTH || w2 != OUT_WIDTH) - { - addToConsole("Width mismatch"); - continue; - } - if (h1 != OUT_HEIGHT || h2 != OUT_HEIGHT) - { - addToConsole("Height mismatch"); - continue; - } + if (w1 != OUT_WIDTH || w2 != OUT_WIDTH) + { + addToConsole("Width mismatch"); + continue; + } + if (h1 != OUT_HEIGHT || h2 != OUT_HEIGHT) + { + addToConsole("Height mismatch"); + continue; + } - uint max = 0; - for(int i = 0; i < w1; i++) - { - for(int j = 0; j < h1; j++) + ulong max = 0; + uint imax, jmax; + for(int i = 0; i < w1; i++) { - if(abs(cppout[j * w1 + i] - matout[j * w1 + i]) > max) + for(int j = 0; j < h1; j++) { - max = abs(cppout[j * w1 + i] - matout[j * w1 + i]); + if(abs(cppout[j * w1 + i] - matout[j * w1 + i]) > max) + { + max = abs(cppout[j * w1 + i] - matout[j * w1 + i]); + imax = i; + jmax = j; + } } } + + 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 > totalMax) - totalMax = max; - addToConsole(QString::number(max)); - addToConsole("*************************\r\n"); - } + addToConsole("Compare done, total max: " + QString::number(totalMax)); + }); +} - addToConsole("Compare done, total max: " + QString::number(totalMax)); +void MainWindow::on_btn_clear_clicked() +{ + ui->rtb_console->clear(); } diff --git a/MainWindow.h b/MainWindow.h index 9f8110d..0edd7fb 100644 --- a/MainWindow.h +++ b/MainWindow.h @@ -46,6 +46,8 @@ private slots: void on_btn_compare_clicked(); + void on_btn_clear_clicked(); + private: Ui::MainWindow *ui; QSettings _settings; diff --git a/MainWindow.ui b/MainWindow.ui index faf1a6a..abe5f03 100644 --- a/MainWindow.ui +++ b/MainWindow.ui @@ -7,11 +7,11 @@ 0 0 800 - 404 + 655 - + 0 0 @@ -26,10 +26,13 @@ 10 10 781 - 391 + 641 + + QLayout::SetNoConstraint + @@ -171,6 +174,13 @@ + + + + clear + + + diff --git a/header/Utils.h b/header/Utils.h index 5f2f682..27b62e2 100644 --- a/header/Utils.h +++ b/header/Utils.h @@ -16,7 +16,7 @@ typedef int myint; #define INP2MYFLT(x) *((myflt*)(&x)) #ifdef USE_DBL -#define PARSE toLong(Q_NULLPTR, 16) +#define PARSE toULong(Q_NULLPTR, 16) #else #define PARSE toInt(Q_NULLPTR, 16) #endif diff --git a/kernels/DynamicContrast b/kernels/DynamicContrast new file mode 100644 index 0000000..e69de29 diff --git a/kernels/Rejection.cl b/kernels/Rejection.cl new file mode 100644 index 0000000..ebd3630 --- /dev/null +++ b/kernels/Rejection.cl @@ -0,0 +1,68 @@ +#define TEST +#define X 67 +#define Y 194 + +#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 rejectThr; +}; + +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 Rejection(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 output_data = 0; + TYPE_FLT input = read_data(input_frame, gid.x, gid.y); + output_data = (input - params.rejectThr) * 255 / (255 - params.rejectThr); + 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); + +#ifdef TEST + if(gid.x == X && gid.y == Y) + { + printf("out: %.15f | ", output_data); + } +#endif + + write_imageui(output_frame, gid, pixel); +} + diff --git a/kernels/ScanConversion.cl b/kernels/ScanConversion.cl new file mode 100644 index 0000000..18f25bf --- /dev/null +++ b/kernels/ScanConversion.cl @@ -0,0 +1,127 @@ +#define TEST +#define X 923 +#define Y 436 + +#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 output_width; + int grid_x_size; + int grid_z_size; +}; + +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; +} + + +TYPE_FLT mabs(TYPE_FLT arg) +{ + if (arg < 0) + return -arg; + return arg; +} + +kernel void ScanConversion(read_only image2d_t input_frame, read_write image2d_t output_frame, + global TYPE_FLT* grid_x, global TYPE_FLT* grid_z, + global TYPE_FLT* query_x, global TYPE_FLT* query_z, + struct input params) { + int2 gid = (int2)(get_global_id(0), get_global_id(1)); + TYPE_FLT dx = grid_x[1] - grid_x[0]; + TYPE_FLT dz = grid_z[1] - grid_z[0]; + TYPE_FLT x = query_x[gid.y * params.output_width + gid.x]; + TYPE_FLT z = query_z[gid.y * params.output_width + gid.x]; + + int grid_x_index = 0; + int grid_z_index = 0; + + if(mabs(x - grid_x[params.grid_x_size - 1]) < EPSILSON) + grid_x_index = params.grid_x_size - 2; + else + grid_x_index = floor((x - grid_x[0]) / dx); + + if(mabs(z - grid_z[params.grid_z_size - 1]) < EPSILSON) + grid_z_index = params.grid_z_size - 2; + else + grid_z_index = floor((z - grid_z[0]) / dz); + + TYPE_FLT output_data = 0; + if(grid_x_index >= 0 && grid_x_index < params.grid_x_size - 1 && + grid_z_index >= 0 && grid_z_index < params.grid_z_size - 1) + { + //matrix simplification + TYPE_FLT a11 = grid_z[grid_z_index + 1] - z; + TYPE_FLT a12 = z - grid_z[grid_z_index]; + TYPE_FLT b11 = read_data(input_frame, grid_x_index, grid_z_index); + TYPE_FLT b12 = read_data(input_frame, grid_x_index + 1, grid_z_index); + TYPE_FLT b21 = read_data(input_frame, grid_x_index, grid_z_index + 1); + TYPE_FLT b22 = read_data(input_frame, grid_x_index + 1, grid_z_index + 1); + TYPE_FLT c11 = grid_x[grid_x_index + 1] - x; + TYPE_FLT c21 = x - grid_x[grid_x_index]; + +#ifdef TEST + if(gid.x == X && gid.y == Y) + { + printf("dx: %.9f | ", dx); + printf("dz: %.9f | ", dz); + printf("x: %.9f | ", x); + printf("z: %.9f | ", z); + printf("grid_x_index: %d | ", grid_x_index); + printf("grid_z_index: %d | ", grid_z_index); + printf("a11: %.9f | ", a11); + printf("a12: %.9f | ", a12); + printf("b11: %.9f | ", b11); + printf("b12: %.9f | ", b12); + printf("b21: %.9f | ", b21); + printf("b22: %.9f | ", b22); + printf("c11: %.9f | ", c11); + printf("c21: %.9f | ", c21); + } + #endif + + output_data = 1 / dx / dz * (c11 * (a11 * b11 + a12 * b21) + c21 * (a11 * b12 + a12 * b22)); + } + + 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); + +#ifdef TEST + if(gid.x == X && gid.y == Y) + { + printf("out: %.15f | ", output_data); + } +#endif + + write_imageui(output_frame, gid, pixel); +} + diff --git a/source/strategies/ScanConversion.cpp b/source/strategies/ScanConversion.cpp index 8841feb..4544bdc 100644 --- a/source/strategies/ScanConversion.cpp +++ b/source/strategies/ScanConversion.cpp @@ -73,11 +73,11 @@ void ScanConversion::cpuProcess(ScenGenOutput_t parameters) for(auto i = 0; i < parameters.outputHeight.value; i++) { - auto temp = i * baseZ; + 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] = minZ + temp; + _gridPixelZPos[i * parameters.outputHeight.value + j] = temp; _gridPixelXPos[i * parameters.outputHeight.value + j] = minX + (baseX * j) - tan_result; } }