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

test

scan-conversion-convex
MMT 6 years ago
parent
65636e447f
commit
5c20289472
5 changed files with 144 additions and 35 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
      header/ScenarioParams.h
  2. 2
      header/Utils.h
  3. 30
      kernels/ScanConversion.cl
  4. 2
      source/FileHelper.cpp
  5. 144
      source/model/processor/strategies/ScanConversion.cpp

1
header/ScenarioParams.h
View File

@ -17,7 +17,6 @@ typedef struct ScenGenOutput_t
field_t<float> depth; field_t<float> depth;
field_t<float> probeRadius; field_t<float> probeRadius;
field_t<float> fieldOfView; field_t<float> fieldOfView;
field_t<float> probeGeometry;
field_t<float> probeFieldOfView; field_t<float> probeFieldOfView;
field_t<float> startDepth; field_t<float> startDepth;
field_t<int> rxLineNo; field_t<int> rxLineNo;

2
header/Utils.h
View File

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

30
kernels/ScanConversion.cl
View File

@ -1,6 +1,6 @@
//#define TEST #define TEST
#define X 0 #define X 500
#define Y 0 #define Y 186
//#define USE_DBL //#define USE_DBL
#ifdef 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, 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* grid_x, global TYPE_FLT* grid_z,
global TYPE_FLT* query_x, global TYPE_FLT* query_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)); int2 gid = (int2)(get_global_id(0), get_global_id(1));
TYPE_FLT dx = grid_x[1] - grid_x[0]; TYPE_FLT dx = grid_x[1] - grid_x[0];
TYPE_FLT dz = grid_z[1] - grid_z[0]; TYPE_FLT dz = grid_z[1] - grid_z[0];
TYPE_FLT x = query_x[gid.y * params.output_width + gid.x]; TYPE_FLT x = query_x[gid.y * params.output_width + gid.x];
TYPE_FLT z = query_z[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_x_index = 0;
int grid_z_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); grid_x_index = floor((x - grid_x[0]) / dx);
if(mabs(z - grid_z[params.grid_z_size - 1]) < EPSILSON) 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; grid_z_index = params.grid_z_size - 2;
}
else 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); grid_z_index = floor((z - grid_z[0]) / dz);
}
TYPE_FLT output_data = 0; 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 && 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) 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++) for(int i = 0; i < sl.length(); i++)
{ {
auto x = sl[i].PARSE; auto x = sl[i].toInt(Q_NULLPTR,16);
arr[index++] = x; arr[index++] = x;
} }
} }

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

@ -20,7 +20,7 @@ ScanConversion::ScanConversion(const Context context,
_queryX = Q_NULLPTR; _queryX = Q_NULLPTR;
_queryZ = Q_NULLPTR; _queryZ = Q_NULLPTR;
_gridPixelR = Q_NULLPTR; _gridPixelR = Q_NULLPTR;
_gridPixelTheta = Q_NULLPTR; _gridPixelTheta = Q_NULLPTR;
} }
@ -29,14 +29,14 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto context = _openCLHelper.getContext(); auto context = _openCLHelper.getContext();
auto queue = BIP::getInstance()->CLQueue; auto queue = BIP::getInstance()->CLQueue;
//convex //convex
if (params.probeGeometry.value) if (!params.linear.value)
{ {
if(params.outputWidth.isUpdated || params.outputHeight.isUpdated) if(params.outputWidth.isUpdated || params.outputHeight.isUpdated)
{ {
if(_gridPixelXPos) if(_gridPixelXPos)
delete[] _gridPixelXPos; delete[] _gridPixelXPos;
if(_gridPixelXPos) if(_gridPixelZPos)
delete[] _gridPixelXPos; delete[] _gridPixelZPos;
_gridPixelXPos = new myflt[params.outputWidth.value * params.outputHeight.value]; _gridPixelXPos = new myflt[params.outputWidth.value * params.outputHeight.value];
_gridPixelZPos = new myflt[params.outputWidth.value * params.outputHeight.value]; _gridPixelZPos = new myflt[params.outputWidth.value * params.outputHeight.value];
@ -58,6 +58,21 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto maxZ = params.depth.value + params.probeRadius.value; 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 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 baseZ = (maxZ - minZ) / (params.outputHeight.value - 1);
auto baseX = (maxX - minX) / (params.outputWidth.value - 1); auto baseX = (maxX - minX) / (params.outputWidth.value - 1);
@ -67,15 +82,23 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto baseIndex = i * params.outputHeight.value; auto baseIndex = i * params.outputHeight.value;
for (auto j = 0; j < params.outputWidth.value; ++j) for (auto j = 0; j < params.outputWidth.value; ++j)
{ {
_gridPixelXPos[baseIndex+j] = minX + baseX * j;
_gridPixelZPos[baseIndex+j] = minZ + temp; 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) if(_gridPixelR)
delete[] _gridPixelR; delete[] _gridPixelR;
_gridPixelR = new myflt[params.outputWidth.value * params.outputHeight.value];
if(_gridPixelTheta) if(_gridPixelTheta)
delete[] _gridPixelTheta; delete[] _gridPixelTheta;
_gridPixelTheta = new myflt[params.outputWidth.value * params.outputHeight.value];
auto virtualOriginalZ = params.probeRadius.value * (cosf(params.fieldOfView.value/2) - auto virtualOriginalZ = params.probeRadius.value * (cosf(params.fieldOfView.value/2) -
sinf(params.fieldOfView.value/2) / sinf(params.fieldOfView.value/2) /
@ -90,7 +113,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto strSin = 2 * sinf(params.steering.value); auto strSin = 2 * sinf(params.steering.value);
auto maxInterceptTheta = params.fieldOfView.value / 2; auto maxInterceptTheta = params.fieldOfView.value / 2;
auto radius2 = powf(params.probeRadius.value, 2.0); auto radius2 = powf(params.probeRadius.value, 2);
if(params.steering.value == 0.0f) if(params.steering.value == 0.0f)
{ {
@ -115,14 +138,14 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
{ {
interceptX = (alpha * beta + sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1); interceptX = (alpha * beta + sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX; interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptTheta); interceptTheta = atan2f(interceptX, interceptZ);
} }
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta)) if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{ {
interceptX = (alpha * beta + sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1); interceptX = (alpha * beta - sqrtf(-1 * powf(alpha, 2) + (powf(beta, 2) + 1)* radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX; interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptTheta); interceptTheta = atan2f(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta)) if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{ {
@ -136,7 +159,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptZ; _gridPixelTheta[i] = interceptTheta;
} }
else else
{ {
@ -152,7 +175,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
_gridPixelTheta[i] = interceptZ; _gridPixelTheta[i] = interceptTheta;
} }
else else
{ {
@ -177,7 +200,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
auto pixelTheta = atan2f(x, z - virtualOriginalZ); auto pixelTheta = atan2f(x, z - virtualOriginalZ);
auto pixelR = sqrtf(powf(x, 2) + powf((z - virtualOriginalZ), 2)); auto pixelR = sqrtf(powf(x, 2) + powf((z - virtualOriginalZ), 2));
if(pixelR >= minR && pixelR <= maxR && pixelTheta >= maxTheta) if(pixelR >= minR && pixelR <= maxR && pixelTheta >= minTheta && pixelTheta <=maxTheta)
{ {
auto ro = pixelR / strSin; auto ro = pixelR / strSin;
auto xo = ro * cosf(params.steering.value - pixelTheta); auto xo = ro * cosf(params.steering.value - pixelTheta);
@ -198,8 +221,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
} }
else else
{ {
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value) * sinf(interceptTheta), 2) + auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf((z - params.probeRadius.value) * cosf(interceptTheta), 2)); powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
@ -214,8 +237,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
} }
else else
{ {
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value) * sinf(interceptTheta), 2) + auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf((z - params.probeRadius.value) * cosf(interceptTheta), 2)); powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
@ -240,7 +263,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta)) if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{ {
interceptX = (alpha * beta - sqrtf(-1 * powf(alpha, 2) + (powf(beta,2) + 1) * radius2)) / (powf(beta, 2) + 1); interceptX = (alpha * beta - sqrtf(-1 * powf(alpha, 2) + (powf(beta,2) + 1) * radius2)) / (powf(beta, 2) + 1);
interceptZ = alpha - beta * interceptX; // TODO:: can delete this line interceptZ = alpha - beta * interceptX;
interceptTheta = atan2f(interceptX, interceptZ); interceptTheta = atan2f(interceptX, interceptZ);
if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta)) if(interceptTheta > maxInterceptTheta || interceptTheta < (-1 * maxInterceptTheta))
{ {
@ -249,8 +272,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
} }
else else
{ {
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value) * sinf(interceptTheta), 2) + auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf((z - params.probeRadius.value) * cosf(interceptTheta), 2)); powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
@ -267,8 +290,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
} }
else else
{ {
auto gridPixelAx = sqrtf(powf((x - params.probeRadius.value) * sinf(interceptTheta), 2) + auto gridPixelAx = sqrtf(powf(x - params.probeRadius.value * sinf(interceptTheta), 2) +
powf((z - params.probeRadius.value) * cosf(interceptTheta), 2)); powf(z - params.probeRadius.value * cosf(interceptTheta), 2));
if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value) if(gridPixelAx >= params.startDepth.value && gridPixelAx <= params.depth.value)
{ {
_gridPixelR[i] = params.probeRadius.value + gridPixelAx; _gridPixelR[i] = params.probeRadius.value + gridPixelAx;
@ -289,13 +312,79 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
} }
} }
} }
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 //TODO:: write after convex
// if(params.rxLineNo.isUpdated || params.rxFocusPointNo.isUpdated || //scanTheta
// params.minScanAx.isUpdated || params.rxPointDax) 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
//virtual convex
else if(params.virtualConvex.value) else if(params.virtualConvex.value)
{ {
if(params.outputWidth.isUpdated|| params.outputHeight.isUpdated) if(params.outputWidth.isUpdated|| params.outputHeight.isUpdated)
@ -481,7 +570,7 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
static_cast<quint64>(params.rxLineNo.value) static_cast<quint64>(params.rxLineNo.value)
* sizeof (myflt)); * sizeof (myflt));
queue.enqueueWriteBuffer(*_gridX, CL_TRUE, 0, queue.enqueueWriteBuffer(*_gridX, CL_TRUE, 0,
static_cast<unsigned long>(params.rxLineNo.value) * sizeof (myflt), static_cast<unsigned long>(params.rxLineNo.value) * sizeof (myflt),
_scanXPos); _scanXPos);
@ -514,7 +603,8 @@ void ScanConversion::cpuProcess(ScenGenOutput_t params)
_kernelParameters.gridZSize = params.rxFocusPointNo.value; _kernelParameters.gridZSize = params.rxFocusPointNo.value;
} }
} }
//linear non convex
//linear non convex
else else
{ {
if(params.outputWidth.isUpdated || params.outputHeight.isUpdated) if(params.outputWidth.isUpdated || params.outputHeight.isUpdated)

Write Preview
Loading…
Cancel
Save