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

Fixed bug in input file parsing

test-kernel
sono 5 years ago
parent
e0e42a310b
commit
2b85cdb078
9 changed files with 317 additions and 92 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. 1
      .gitignore
  2. 25
      MainWindow.cpp
  3. 2
      MainWindow.h
  4. 16
      MainWindow.ui
  5. 2
      header/Utils.h
  6. 0
      kernels/DynamicContrast
  7. 68
      kernels/Rejection.cl
  8. 127
      kernels/ScanConversion.cl
  9. 4
      source/strategies/ScanConversion.cpp

1
.gitignore
View File

@ -0,0 +1 @@
*.pro.*

25
MainWindow.cpp
View File

@ -3,6 +3,7 @@
#include <QFileDialog> #include <QFileDialog>
#include <QFile> #include <QFile>
#include <QtConcurrent/QtConcurrent>
#include "header/strategies/ScanConversion.h" #include "header/strategies/ScanConversion.h"
#include "header/strategies/Rejection.h" #include "header/strategies/Rejection.h"
@ -141,6 +142,7 @@ void MainWindow::on_btn_test_clicked()
_strategy->finalize(); _strategy->finalize();
pushBackStrategy(ui->cb_kernelName->currentText(), workingDir.path() + "/kernels"); pushBackStrategy(ui->cb_kernelName->currentText(), workingDir.path() + "/kernels");
QtConcurrent::run([this, inputDir, outputDir](){
auto inputs = inputDir.entryList(QStringList() << "*.csv", QDir::Files); auto inputs = inputDir.entryList(QStringList() << "*.csv", QDir::Files);
foreach(QString filename, inputs) foreach(QString filename, inputs)
{ {
@ -150,7 +152,6 @@ void MainWindow::on_btn_test_clicked()
_strategy->ReadParams(paramPath, &_scenGenOutput); _strategy->ReadParams(paramPath, &_scenGenOutput);
QByteArray arr; QByteArray arr;
quint64 width; quint64 width;
quint64 height = 0; quint64 height = 0;
@ -195,7 +196,9 @@ void MainWindow::on_btn_test_clicked()
FileHelper::WriteOutputFile(out, outputDir.path() + "/" + filename, OUT_WIDTH, OUT_HEIGHT); FileHelper::WriteOutputFile(out, outputDir.path() + "/" + filename, OUT_WIDTH, OUT_HEIGHT);
addToConsole(filename + " has been processed"); addToConsole(filename + " has been processed");
addToConsole("*************************\r\n");
} }
});
} }
void MainWindow::on_btn_save_clicked() void MainWindow::on_btn_save_clicked()
@ -225,7 +228,8 @@ void MainWindow::on_btn_compare_clicked()
QDir matlabDir(ui->tb_matlabOutp->text()); QDir matlabDir(ui->tb_matlabOutp->text());
QDir outputDir(ui->tb_outpDir->text()); QDir outputDir(ui->tb_outpDir->text());
auto totalMax = 0U; QtConcurrent::run([this, outputDir, matlabDir](){
auto totalMax = 0UL;
auto outputs = outputDir.entryList(QStringList() << "*.csv", QDir::Files); auto outputs = outputDir.entryList(QStringList() << "*.csv", QDir::Files);
foreach(QString filename, outputs) foreach(QString filename, outputs)
{ {
@ -255,7 +259,8 @@ void MainWindow::on_btn_compare_clicked()
continue; continue;
} }
uint max = 0; ulong max = 0;
uint imax, jmax;
for(int i = 0; i < w1; i++) for(int i = 0; i < w1; i++)
{ {
for(int j = 0; j < h1; j++) for(int j = 0; j < h1; j++)
@ -263,15 +268,27 @@ void MainWindow::on_btn_compare_clicked()
if(abs(cppout[j * w1 + i] - matout[j * w1 + i]) > max) if(abs(cppout[j * w1 + i] - matout[j * w1 + i]) > max)
{ {
max = abs(cppout[j * w1 + i] - matout[j * w1 + i]); max = abs(cppout[j * w1 + i] - matout[j * w1 + i]);
imax = i;
jmax = j;
} }
} }
} }
if(max > totalMax) if(max > totalMax)
totalMax = max; totalMax = max;
addToConsole(QString::number(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"); 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();
} }

2
MainWindow.h
View File

@ -46,6 +46,8 @@ private slots:
void on_btn_compare_clicked(); void on_btn_compare_clicked();
void on_btn_clear_clicked();
private: private:
Ui::MainWindow *ui; Ui::MainWindow *ui;
QSettings _settings; QSettings _settings;

16
MainWindow.ui
View File

@ -7,11 +7,11 @@
<x>0</x> <x>0</x>
<y>0</y> <y>0</y>
<width>800</width> <width>800</width>
<height>404</height> <height>655</height>
</rect> </rect>
</property> </property>
<property name="sizePolicy"> <property name="sizePolicy">
<sizepolicy hsizetype="Fixed" vsizetype="Fixed"> <sizepolicy hsizetype="Expanding" vsizetype="Expanding">
<horstretch>0</horstretch> <horstretch>0</horstretch>
<verstretch>0</verstretch> <verstretch>0</verstretch>
</sizepolicy> </sizepolicy>
@ -26,10 +26,13 @@
<x>10</x> <x>10</x>
<y>10</y> <y>10</y>
<width>781</width> <width>781</width>
<height>391</height> <height>641</height>
</rect> </rect>
</property> </property>
<layout class="QVBoxLayout" name="verticalLayout"> <layout class="QVBoxLayout" name="verticalLayout">
<property name="sizeConstraint">
<enum>QLayout::SetNoConstraint</enum>
</property>
<item> <item>
<layout class="QHBoxLayout" name="horizontalLayout"> <layout class="QHBoxLayout" name="horizontalLayout">
<item> <item>
@ -171,6 +174,13 @@
</property> </property>
</widget> </widget>
</item> </item>
<item>
<widget class="QPushButton" name="btn_clear">
<property name="text">
<string>clear</string>
</property>
</widget>
</item>
<item> <item>
<spacer name="horizontalSpacer"> <spacer name="horizontalSpacer">
<property name="orientation"> <property name="orientation">

2
header/Utils.h
View File

@ -16,7 +16,7 @@ typedef int myint;
#define INP2MYFLT(x) *((myflt*)(&x)) #define INP2MYFLT(x) *((myflt*)(&x))
#ifdef USE_DBL #ifdef USE_DBL
#define PARSE toLong(Q_NULLPTR, 16) #define PARSE toULong(Q_NULLPTR, 16)
#else #else
#define PARSE toInt(Q_NULLPTR, 16) #define PARSE toInt(Q_NULLPTR, 16)
#endif #endif

0
kernels/DynamicContrast
View File

68
kernels/Rejection.cl
View File

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

127
kernels/ScanConversion.cl
View File

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

4
source/strategies/ScanConversion.cpp
View File

@ -73,11 +73,11 @@ void ScanConversion::cpuProcess(ScenGenOutput_t parameters)
for(auto i = 0; i < parameters.outputHeight.value; i++) 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); auto tan_result = temp * tan(parameters.steering.value);
for(auto j = 0; j < parameters.outputWidth.value; j++) 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; _gridPixelXPos[i * parameters.outputHeight.value + j] = minX + (baseX * j) - tan_result;
} }
} }

Write Preview
Loading…
Cancel
Save