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

Merge branch 'scan-conversion-convex' of http://194.5.205.38:3000/mmtalaie/kernel-tester

master
sono 5 years ago
parent
66ebeb24fb 5c20289472
commit
ca8edc11ff
5 changed files with 390 additions and 13 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      header/ScenarioParams.h
  2. 2
      header/Utils.h
  3. 30
      kernels/ScanConversion.cl
  4. 2
      source/FileHelper.cpp
  5. 367
      source/model/processor/strategies/ScanConversion.cpp

2
header/ScenarioParams.h
View File

@ -16,7 +16,7 @@ typedef struct ScenGenOutput_t
field_t<bool> virtualConvex;
field_t<float> depth;
field_t<float> probeRadius;
field_t<float> fieldOfView;
field_t<float> fieldOfView;
field_t<float> probeFieldOfView;
field_t<float> startDepth;
field_t<int> rxLineNo;

2
header/Utils.h
View File

@ -18,7 +18,7 @@ typedef int myint;
#ifdef USE_DBL
#define PARSE toULong(Q_NULLPTR, 16)
#else
#define PARSE toInt(Q_NULLPTR, 16)
#define PARSE toUInt(Q_NULLPTR, 16)
#endif

30
kernels/ScanConversion.cl
View File

@ -1,6 +1,6 @@
//#define TEST
#define X 0
#define Y 0
#define TEST
#define X 500
#define Y 186
//#define USE_DBL
#ifdef USE_DBL
@ -51,13 +51,15 @@ TYPE_FLT mabs(TYPE_FLT 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) {
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;
@ -67,11 +69,29 @@ kernel void ScanConversion(read_only image2d_t input_frame, read_write image2d_t
grid_x_index = floor((x - grid_x[0]) / dx);
if(mabs(z - grid_z[params.grid_z_size - 1]) < EPSILSON)
{
if(gid.x == X && gid.y == Y)
printf("here 1 ");
grid_z_index = params.grid_z_size - 2;
}
else
{
if(gid.x == X && gid.y == Y)
printf("here 2 %0.9f",grid_z[0]);
grid_z_index = floor((z - grid_z[0]) / dz);
}
TYPE_FLT output_data = 0;
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);
}
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)
{

2
source/FileHelper.cpp
View File

@ -78,7 +78,7 @@ bool FileHelper::ReadInputFile(myint *arr, QString path, quint64 *width, quint64
}
for(int i = 0; i < sl.length(); i++)
{
auto x = sl[i].PARSE;
auto x = sl[i].toInt(Q_NULLPTR,16);
arr[index++] = x;
}
}

367
source/model/processor/strategies/ScanConversion.cpp
View File

@ -20,7 +20,7 @@ ScanConversion::ScanConversion(const Context context,
_queryX = Q_NULLPTR;
_queryZ = Q_NULLPTR;
_gridPixelR = Q_NULLPTR;
_gridPixelR = Q_NULLPTR;
_gridPixelTheta = Q_NULLPTR;
}
@ -30,8 +30,364 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto context = _openCLHelper.getContext();
auto queue = BIP::getInstance()->CLQueue;
//virtual convex
if(params.virtualConvex.value)
//convex
if (!params.linear.value)
{
if(params.outputWidth.isUpdated || params.outputHeight.isUpdated)
{
if(_gridPixelXPos)
delete[] _gridPixelXPos;
if(_gridPixelZPos)
delete[] _gridPixelZPos;
_gridPixelXPos = new myflt[params.outputWidth.value * params.outputHeight.value];
_gridPixelZPos = new myflt[params.outputWidth.value * params.outputHeight.value];
_width = params.outputWidth.value;
_height = params.outputHeight.value;
_kernelParameters.width = _width;
}
if(params.depth.isUpdated || params.probeRadius.isUpdated ||
params.fieldOfView.isUpdated || params.angle.isUpdated ||
params.startDepth.isUpdated || params.startDepth.isUpdated ||
params.outputWidth.isUpdated || params.outputHeight.isUpdated||
params.steering.isUpdated)
{
auto maxX = params.probeRadius.value*sinf(params.fieldOfView.value/2)+params.depth.value*sinf(params.angle.value/2);
auto minX = -1 * maxX;
auto maxZ = params.depth.value + params.probeRadius.value;
auto minZ = params.probeRadius.value*cosf(params.fieldOfView.value/2)+params.startDepth.value*cosf(params.angle.value/2);
// auto baseZ = (maxZ - minZ) / (params.outputHeight.value - 1);
// auto baseX = (maxX - minX) / (params.outputWidth.value - 1);
// for (auto i = 0; i < params.outputHeight.value; ++i)
// {
// auto temp = i * baseZ;
// auto baseIndex = i * params.outputHeight.value;
// for (auto j = 0; j < params.outputWidth.value; ++j)
// {
// auto b = minX + baseX * j;
// _gridPixelXPos[baseIndex+j] = b;
// auto a = minZ + temp;
// _gridPixelZPos[baseIndex+j] = a;
// }
// }
auto baseZ = (maxZ - minZ) / (params.outputHeight.value - 1);
auto baseX = (maxX - minX) / (params.outputWidth.value - 1);
for (auto i = 0; i < params.outputHeight.value; ++i)
{
auto temp = i * baseZ;
auto baseIndex = i * params.outputHeight.value;
for (auto j = 0; j < params.outputWidth.value; ++j)
{
auto a = minZ + i * (maxZ - minZ) / (params.outputHeight.value - 1);
_gridPixelZPos[baseIndex+j] = a;
auto b = minX + baseX * j;
_gridPixelXPos[baseIndex+j] = b;
}
}
if(_gridPixelR)
delete[] _gridPixelR;
_gridPixelR = new myflt[params.outputWidth.value * params.outputHeight.value];
if(_gridPixelTheta)
delete[] _gridPixelTheta;
_gridPixelTheta = new myflt[params.outputWidth.value * params.outputHeight.value];
auto virtualOriginalZ = params.probeRadius.value * (cosf(params.fieldOfView.value/2) -
sinf(params.fieldOfView.value/2) /
tanf(params.angle.value/2));
auto virtualOriginalZ2 = powf(virtualOriginalZ, 2);
auto minR = params.probeRadius.value - virtualOriginalZ + params.startDepth.value;
auto maxR = params.probeRadius.value + params.depth.value;
auto minTheta = (-1 * params.angle.value) / 2 - abs(params.steering.value);
auto maxTheta = params.angle.value / 2 + abs(params.steering.value);
auto strSin = 2 * sinf(params.steering.value);
auto maxInterceptTheta = params.fieldOfView.value / 2;
auto radius2 = powf(params.probeRadius.value, 2);
if(params.steering.value == 0.0f)
{
for (auto i = 0; i < params.outputWidth.value * params.outputHeight.value; ++i)
{
auto x = _gridPixelXPos[i];
auto z = _gridPixelZPos[i];
auto pixelTheta = atan2f(x, z - virtualOriginalZ);
auto pixelR = sqrtf(powf(x, 2) + powf((z - virtualOriginalZ), 2));
if(pixelR >= minR && pixelR <= maxR && pixelTheta >= minTheta && pixelTheta <= maxTheta)
{
auto interceptTheta = 0.0f;
auto interceptX = 0.0f;
auto interceptZ = 0.0f;
auto alpha = virtualOriginalZ;
auto beta = (virtualOriginalZ - z) / x;
if (x == 0.0f)
interceptTheta = 0;
else
{
interceptX = (alpha * beta + sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptZ);
}
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
interceptX = (alpha * beta - sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
else
{
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value * sinf(interceptTheta)), 2) +
powf((z - params.probeRadius.value * cosf(interceptTheta)), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value * sinf(interceptTheta)), 2) +
powf((z - params.probeRadius.value * cosf(interceptTheta)), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
for (auto i = 0; i < params.outputWidth.value * params.outputHeight.value; ++i)
{
auto x = _gridPixelXPos[i];
auto z = _gridPixelZPos[i];
auto pixelTheta = atan2f(x, z - virtualOriginalZ);
auto pixelR = sqrtf(powf(x, 2) + powf((z - virtualOriginalZ), 2));
if(pixelR >= minR && pixelR <= maxR && pixelTheta >= minTheta && pixelTheta <=maxTheta)
{
auto ro = pixelR / strSin;
auto xo = ro * cosf(params.steering.value - pixelTheta);
auto zo = ro * sinf(params.steering.value - pixelTheta) + virtualOriginalZ;
if(zo == 0.0f)
{
auto interceptX = (radius2 - virtualOriginalZ2) / 2 / xo;
auto interceptZ = sqrtf(radius2 - powf(interceptX , 2));
auto interceptTheta = atan2f(interceptX, interceptZ);
if (interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
interceptZ = -1 * sqrtf(radius2 - powf(interceptX , 2));
interceptTheta = atan2(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
else
{
auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
auto alpha = (radius2 - virtualOriginalZ2 + 2 * zo * virtualOriginalZ) / 2 / zo;
auto beta = xo / zo;
auto interceptX = (alpha * beta + sqrtf(-1 * powf(alpha, 2) + (powf(beta,2) + 1) * radius2)) / (powf(beta, 2) + 1);
auto interceptZ = alpha - beta * interceptX;
auto interceptTheta = atan2f(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
interceptX = (alpha * beta - sqrtf(-1 * powf(alpha, 2) + (powf(beta,2) + 1) * radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
else
{
auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
else
{
auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{
_gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptTheta;
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
}
else
{
_gridPixelR[i] = 0;
_gridPixelTheta[i] = 0;
}
}
}
if(_queryX)
delete _queryX;
if(_queryZ)
delete _queryZ;
_queryX = _openCLHelper.array2CLBuffer(context,
params.outputWidth.value *
params.outputHeight.value *
sizeof (myflt));
_queryZ = _openCLHelper.array2CLBuffer(context,
params.outputWidth.value *
params.outputHeight.value *
sizeof (myflt));
BIP::getInstance()->CLQueue.enqueueWriteBuffer(*_queryZ, CL_TRUE, 0,
params.outputWidth.value *
params.outputHeight.value *
sizeof (myflt),
_gridPixelR);
queue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0,
params.outputWidth.value *
params.outputHeight.value *
sizeof (myflt),
_gridPixelTheta);
}
//TODO:: write after convex
//scanTheta
if(params.rxLineNo.isUpdated || params.minScanAz.isUpdated ||
params.rxLineDaz.isUpdated)
{
if(_scanXPos)
delete [] _scanXPos;
if(_gridX)
delete _gridX;
_scanXPos = new myflt[params.rxLineNo.value];
for (auto i = 0; i < params.rxLineNo.value ; ++i)
{
_scanXPos[i] = params.minScanAz.value + i * params.rxLineDaz.value;
}
_gridX = _openCLHelper.array2CLBuffer(context, static_cast<qint64>(params.rxLineNo.value * sizeof (myflt)));
auto er = queue.enqueueWriteBuffer(*_gridX,CL_TRUE, 0, static_cast<unsigned long>(params.rxLineNo.value * sizeof (myflt)), _scanXPos);
_kernelParameters.gridXSize = params.rxLineNo.value;
}
//scanR
if(params.rxFocusPointNo.isUpdated || params.minScanAx.isUpdated ||
params.rxPointDax.isUpdated)
{
if(_scanZPos)
delete [] _scanZPos;
if(_gridZ)
delete _gridZ;
_scanZPos = new myflt[params.rxFocusPointNo.value];
for (auto i = 0; i < params.rxFocusPointNo.value; ++i)
{
_scanZPos[i] = (params.minScanAx.value + i * params.rxPointDax.value ) +params.probeRadius.value;
}
_gridZ = _openCLHelper.array2CLBuffer(context, static_cast<qint64>(params.rxFocusPointNo.value * sizeof (myflt)));
auto er = queue.enqueueWriteBuffer(*_gridZ, CL_TRUE, 0, static_cast<unsigned long>(params.rxFocusPointNo.value * sizeof (myflt)),_scanZPos);
_kernelParameters.gridZSize = params.rxFocusPointNo.value;
}
}
//virtual convex
else if(params.virtualConvex.value)
{
if(params.outputWidth.isUpdated|| params.outputHeight.isUpdated)
{
@ -190,7 +546,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
params.outputHeight.value *
sizeof (myflt),
_gridPixelR);
queue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0,
queue.enqueueWriteBuffer(*_queryX, CL_TRUE, 0,
params.outputWidth.value *
params.outputHeight.value *
sizeof (myflt),
@ -249,7 +605,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
_kernelParameters.gridZSize = params.rxFocusPointNo.value;
}
}
//linear non convex
//linear non convex
else
{
if(params.outputWidth.isUpdated || params.outputHeight.isUpdated)

Write Preview
Loading…
Cancel
Save