...
 
Commits (2)
......@@ -107,21 +107,17 @@
</dependency>
<!-- https://mvnrepository.com/artifact/net.sf.ehcache/ehcache-core -->
<!--
<dependency>
<groupId>net.sf.ehcache</groupId>
<artifactId>ehcache-core</artifactId>
<version>2.6.2</version>
</dependency>\
-->
</dependency>
<!-- https://mvnrepository.com/artifact/org.slf4j/jcl-over-slf4j -->
<!--
<dependency>
<groupId>org.slf4j</groupId>
<artifactId>jcl-over-slf4j</artifactId>
<version>1.7.5</version>
</dependency>
-->
</dependencies>
......
......@@ -135,7 +135,7 @@ public class CLTParameters {
public double fcorr_disp_diff = 1.5; // consider only tiles with absolute residual disparity lower than
public boolean fcorr_quadratic = true; // Use quadratic polynomial for fine correction (false - only linear)
public boolean fcorr_ignore = false; // Ignore currently calculated fine correction
public double fcorr_inf_strength = 0.20 ; // Minimal correlation strength to use for infinity correction
public double fcorr_inf_strength = 0.15 ; // Minimal correlation strength to use for infinity correction
public double fcorr_inf_diff = 0.2; // Disparity half-range for infinity
public boolean fcorr_inf_quad = true; // Use quadratic polynomial for infinity correction (false - only linear)
public boolean fcorr_inf_vert = false; // Correct infinity in vertical direction (false - only horizontal)
......
......@@ -651,6 +651,7 @@ public class GPUTileProcessor {
private boolean geometry_correction_set = false;
private boolean geometry_correction_vector_set = false;
public int gpu_debug_level = 1;
public GpuQuad(
final QuadCLT quadCLT,
// final int img_width,
......@@ -878,9 +879,16 @@ public class GPUTileProcessor {
public void resetGeometryCorrection() {
geometry_correction_set = false;
geometry_correction_vector_set = false;
if (gpu_debug_level > -1) {
System.out.println("======resetGeometryCorrection()");
}
}
public void resetGeometryCorrectionVector() {
geometry_correction_vector_set = false;
if (gpu_debug_level > -1) {
System.out.println("======resetGeometryCorrectionVector()");
}
}
public int getImageWidth() {return this.img_width;}
public int getImageHeight() {return this.img_height;}
......@@ -897,6 +905,9 @@ public class GPUTileProcessor {
public void setGeometryCorrectionVector() { // will reset geometry_correction_vector_set when running GPU kernel
setExtrinsicsVector(
quadCLT.getGeometryCorrection().getCorrVector());
if (gpu_debug_level > -1) {
System.out.println("======setGeometryCorrectionVector()");
}
}
public void setGeometryCorrection(GeometryCorrection gc,
......@@ -912,6 +923,9 @@ public class GPUTileProcessor {
}
cuMemcpyHtoD(gpu_geometry_correction, Pointer.to(fgc), fgc.length * Sizeof.FLOAT);
cuMemAlloc (gpu_rot_deriv, 5 * num_cams *3 *3 * Sizeof.FLOAT); // NCAM of 3x3 rotation matrices, plus 4 derivative matrices for each camera
if (gpu_debug_level > -1) {
System.out.println("======setGeometryCorrection()");
}
}
/**
* Copy extrinsic correction vector to the GPU memory
......@@ -924,6 +938,9 @@ public class GPUTileProcessor {
fcv[i] = (float) dcv[i];
}
cuMemcpyHtoD(gpu_correction_vector, Pointer.to(fcv), fcv.length * Sizeof.FLOAT);
if (gpu_debug_level > -1) {
System.out.println("======setExtrinsicsVector()");
}
}
/**
......@@ -945,6 +962,10 @@ public class GPUTileProcessor {
tile_tasks[i].asFloatArray(ftasks, i* TPTASK_SIZE, use_aux);
}
cuMemcpyHtoD(gpu_tasks, Pointer.to(ftasks), TPTASK_SIZE * num_task_tiles * Sizeof.FLOAT);
if (gpu_debug_level > -1) {
System.out.println("======setTasks()");
}
}
void checkTasks (TpTask [] tile_tasks) {
......@@ -1103,6 +1124,9 @@ public class GPUTileProcessor {
ncam); // int ncam)
}
kernels_set = true;
if (gpu_debug_level > -1) {
System.out.println("======setConvolutionKernels()");
}
}
public void setBayerImage(
......@@ -1160,6 +1184,9 @@ public class GPUTileProcessor {
ncam); // int ncam)
}
bayer_set = true;
if (gpu_debug_level > -1) {
System.out.println("======setBayerImages()");
}
}
// prepare tasks for full frame, same dispaity.
......@@ -1536,6 +1563,9 @@ public class GPUTileProcessor {
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize(); // remove later
geometry_correction_vector_set = true;
if (gpu_debug_level > -1) {
System.out.println("======execRotDerivs()");
}
}
/**
......@@ -1565,6 +1595,9 @@ public class GPUTileProcessor {
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize(); // remove later
geometry_correction_set = true;
if (gpu_debug_level > -1) {
System.out.println("======execCalcReverseDistortions()");
}
}
/**
......@@ -1595,6 +1628,9 @@ public class GPUTileProcessor {
0, null, // Shared memory size and stream (shared - only dynamic, static is in code)
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize(); // remove later
if (gpu_debug_level > -1) {
System.out.println("======execSetTilesOffsets()");
}
}
/**
......@@ -1639,6 +1675,9 @@ public class GPUTileProcessor {
0, null, // Shared memory size and stream (shared - only dynamic, static is in code)
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize(); // remove later
if (gpu_debug_level > -1) {
System.out.println("======execConvertDirect()");
}
}
......
......@@ -516,7 +516,8 @@ public class AlignmentCorrection {
for (int nTile = 0; nTile < numTiles; nTile++) if (center_mask[nTile]){
if (disp_strength[str_index][nTile] > min_strength) {
//clt_parameters.fcorr_inf_diff
if (Math.abs(disp_strength[disp_index][nTile]) <= max_diff) {
// if (Math.abs(disp_strength[disp_index][nTile]) <= max_diff) {
if ((disp_strength[disp_index][nTile] <= max_diff) && (disp_strength[disp_index][nTile] >= 5* max_diff)) { // asymmetry
double weight= disp_strength[str_index][nTile];
sdw += weight * disp_strength[disp_index][nTile];
sw += weight;
......@@ -2176,7 +2177,7 @@ B = |+dy0 -dy1 -2*dy3 |
magic_coeff, // magic_coeff, // still not understood coefficient that reduces reported disparity value. Seems to be around 8.5
debugLevel);
if (debugLevel > -1) {
if (debugLevel > -3) { // -1
double inf_weight = 0.0;
for (Sample s: inf_samples_list) {
inf_weight += s.weight;
......
......@@ -2222,7 +2222,7 @@ public class Correlation2d {
return mosaic;
}
public Corr2dLMA corrLMA2( // multi-tile
public Corr2dLMA corrLMA2Multi( // multi-tile
ImageDttParameters imgdtt_params,
int clust_width,
double [][] corr_wnd, // correlation window to save on re-calculation of the window
......@@ -2453,7 +2453,7 @@ public class Correlation2d {
imgdtt_params.lma_str_scale, // convert lma-generated strength to match previous ones - scale
imgdtt_params.lma_str_offset // convert lma-generated strength to match previous ones - add to result
);
if (debug_level > -1) lma.printStats(ds,clust_width);
if (debug_level > 0) lma.printStats(ds,clust_width);
if (debug_level > 1) {
......@@ -2551,7 +2551,7 @@ public class Correlation2d {
imgdtt_params.lma_str_scale, // convert lma-generated strength to match previous ones - scale
imgdtt_params.lma_str_offset // convert lma-generated strength to match previous ones - add to result
);
if (debug_level > -2) {
if (debug_level > 0) { // -2) {
lma.printStats(ds,clust_width);
}
}
......@@ -2619,8 +2619,9 @@ public class Correlation2d {
return lmaSuccess? lma: null;
}
public Corr2dLMA corrLMA2( // single tile
public Corr2dLMA corrLMA2Single( // single tile
ImageDttParameters imgdtt_params,
boolean adjust_ly, // adjust Lazy Eye
double [][] corr_wnd, // correlation window to save on re-calculation of the window
double [] corr_wnd_inv_limited, // correlation window, limited not to be smaller than threshold - used for finding max/convex areas (or null)
double [][] corrs,
......@@ -2758,12 +2759,12 @@ public class Correlation2d {
imgdtt_params.lmas_adjust_wm, // boolean adjust_width, // adjust width of the maximum - lma_adjust_wm
imgdtt_params.lmas_adjust_ag, // boolean adjust_scales, // adjust 2D correlation scales - lma_adjust_ag
imgdtt_params.lmas_adjust_wy, // boolean adjust_ellipse, // allow non-circular correlation maximums lma_adjust_wy
false, //imgdtt_params.lma_adjust_wxy, // boolean adjust_lazyeye_par, // adjust disparity corrections parallel to disparities lma_adjust_wxy
false, // imgdtt_params.lma_adjust_ly1, // boolean adjust_lazyeye_ortho, // adjust disparity corrections orthogonal to disparities lma_adjust_ly1
(adjust_ly ? imgdtt_params.lma_adjust_wxy : false), //imgdtt_params.lma_adjust_wxy, // boolean adjust_lazyeye_par, // adjust disparity corrections parallel to disparities lma_adjust_wxy
(adjust_ly ? imgdtt_params.lma_adjust_ly1: false), // imgdtt_params.lma_adjust_ly1, // boolean adjust_lazyeye_ortho, // adjust disparity corrections orthogonal to disparities lma_adjust_ly1
disp_str2, // xcenter, // double disp0, // initial value of disparity
imgdtt_params.lma_half_width, // double half_width, // A=1/(half_widh)^2 lma_half_width
0.0, // imgdtt_params.lma_cost_wy, // double cost_lazyeye_par, // cost for each of the non-zero disparity corrections lma_cost_wy
0.0 //imgdtt_params.lma_cost_wxy // double cost_lazyeye_odtho // cost for each of the non-zero ortho disparity corrections lma_cost_wxy
(adjust_ly ? imgdtt_params.lma_cost_wy : 0.0), // imgdtt_params.lma_cost_wy, // double cost_lazyeye_par, // cost for each of the non-zero disparity corrections lma_cost_wy
(adjust_ly ? imgdtt_params.lma_cost_wxy : 0.0) //imgdtt_params.lma_cost_wxy // double cost_lazyeye_odtho // cost for each of the non-zero ortho disparity corrections lma_cost_wxy
);
lma.setMatrices(disp_dist);
lma.initMatrices(); // should be called after initVector and after setMatrices
......
......@@ -1678,11 +1678,13 @@ public class ImageDtt extends ImageDttCPU {
final float [][][][] fcorr_td_centers = new float [tilesY][tilesX][][]; // sparse, only in cluster centers
final int [][] num_in_cluster = new int [clustersY][clustersX]; // only in cluster centers
final double [][][][] disp_dist = new double [clustersY][clustersX][][];
final double [][][][] pxpy = new double [clustersY][clustersX][][];
final double [][][] clust_pY = new double [clustersY][clustersX][];
final double [][][] pxpy = new double [clustersY][clustersX][];
final double [][] disparity_array_center = new double [clustersY][clustersX];
final Thread[] threads = newThreadArray(threadsMax);
final AtomicInteger ai = new AtomicInteger(0);
final double shiftX = 0.0;
final double shiftY = 0.0;
// TODO: Maybe calculate full energy in each TD tile for normalization
for (int ithread = 0; ithread < threads.length; ithread++) {
threads[ithread] = new Thread() {
......@@ -1701,7 +1703,7 @@ public class ImageDtt extends ImageDttCPU {
boolean debugCluster = (clustX == debug_clustX) && (clustY == debug_clustY);
/// boolean debugCluster1 = (Math.abs(clustX - debug_clustX) < 10) && (Math.abs(clustY - debug_clustY) < 10);
if (debugCluster) {
System.out.println("debugCluster");
System.out.println("debugCluster1");
}
/// int clust_lma_debug_level = debugCluster? imgdtt_params.lma_debug_level : -5;
// filter only tiles with similar disparity to enable lazy eye for the ERS.
......@@ -1738,13 +1740,13 @@ public class ImageDtt extends ImageDttCPU {
}
}
}
avg /= num_good_tiles;
if (num_good_tiles ==0) {
break;
}
if ((mx-mn) <= imgdtt_params.lma_disp_range ) {
break;
}
avg /= num_good_tiles;
if ((mx-avg) > (avg-mn)) {
fcorr_td[mxTy][mxTx] = null;
} else {
......@@ -1758,9 +1760,13 @@ public class ImageDtt extends ImageDttCPU {
float [][] ftd = new float [num_pairs][4* transform_size * transform_size];
if (num_good_tiles > 0) {
double dscale = 1.0/num_good_tiles;
// average fdisp_dist and fpxpy over remaining tiles
//fpxpy
double [] avg_py = new double [quad];
double [][] avg_disp_dist = new double [quad][4];
double [][] avg_pxpy = new double [quad][2];
// double [][] avg_pxpy = new double [quad][2];
double [] avg_pxpy = new double [2];
for (int cTileY = 0; cTileY < tileStep; cTileY++) {
int tileY = clustY * tileStep + cTileY ;
if (tileY < tilesY) {
......@@ -1771,10 +1777,12 @@ public class ImageDtt extends ImageDttCPU {
for (int i = 0; i < 4; i++) {
avg_disp_dist[nc][i] += fdisp_dist[tileY][tileX][nc][i];
}
for (int i = 0; i < 2; i++) {
avg_pxpy[nc][i] += fpxpy[tileY][tileX][nc][i];
}
avg_py[nc] += fpxpy[tileY][tileX][nc][1]; // averaging some tiles, but it will be the same for all channels
}
double centerX = tileX * transform_size + transform_size/2 - shiftX;
double centerY = tileY * transform_size + transform_size/2 - shiftY;
avg_pxpy[0] += centerX;
avg_pxpy[1] += centerY;
}
}
}
......@@ -1796,12 +1804,15 @@ public class ImageDtt extends ImageDttCPU {
for (int i = 0; i < 4; i++) {
avg_disp_dist[nc][i] *= dscale;
}
for (int i = 0; i < 2; i++) {
avg_pxpy[nc][i] *= dscale;
}
avg_py[nc] *= dscale;
}
for (int i = 0; i < 2; i++) {
avg_pxpy[i] *= dscale;
}
disp_dist[clustY][clustX] = avg_disp_dist;
pxpy [clustY][clustX] = avg_pxpy;
clust_pY [clustY][clustX] = avg_py; // to use for ERS
// accumulate TD tiles
// If needed - save average
for (int cTileY = 0; cTileY < tileStep; cTileY++) {
......@@ -1881,6 +1892,7 @@ public class ImageDtt extends ImageDttCPU {
int tileY = nt / tilesX;
int clustX = tileX/tileStep;
int clustY = tileY/tileStep;
/// boolean debugCluster1 = (Math.abs(clustX - debug_clustX) < 10) && (Math.abs(clustY - debug_clustY) < 10);
double [] disp_str = null;
for (int indx_pair = 0; indx_pair < num_pairs; indx_pair++) {
......@@ -1905,14 +1917,18 @@ public class ImageDtt extends ImageDttCPU {
// debug new LMA correlations
boolean debugCluster = (clustX == debug_clustX) && (clustY == debug_clustY);
if (debugCluster) {
System.out.println("debugCluster2");
}
int tile_lma_debug_level = ((tileX == debug_tileX) && (tileY == debug_tileY))? imgdtt_params.lma_debug_level : -1;
int tdl = debugCluster ? tile_lma_debug_level : -3;
if (debugCluster && (globalDebugLevel > -1)) { // -2)) {
System.out.println("Will run new LMA for tileX="+tileX+", tileY="+tileY);
}
double [] poly_disp = {Double.NaN, 0.0};
Corr2dLMA lma2 = corr2d.corrLMA2( // num_tiles == 1
Corr2dLMA lma2 = corr2d.corrLMA2Single( // multitile num_tiles == 1
imgdtt_params, // ImageDttParameters imgdtt_params,
true, // boolean adjust_ly, // adjust Lazy Eye
corr_wnd, // double [][] corr_wnd, // correlation window to save on re-calculation of the window
corr_wnd_inv_limited, // corr_wnd_limited, // correlation window, limited not to be smaller than threshold - used for finding max/convex areas (or null)
corrs, // double [][] corrs,
......@@ -1926,6 +1942,28 @@ public class ImageDtt extends ImageDttCPU {
tileX, // int tileX, // just for debug output
tileY); // int tileY
/*
double [][] poly_disp2 = {{Double.NaN, 0.0}};
double [][][] corrs2 = {corrs};
double [][][] tile_disp_dist2 = {tile_disp_dist};
// TODO: maybe use corrLMA2Single again, but take care of initVector!
Corr2dLMA lma2 = corr2d.corrLMA2Multi( // multitile num_tiles == 1
imgdtt_params, // ImageDttParameters imgdtt_params,
1, // int clust_width,
corr_wnd, // double [][] corr_wnd, // correlation window to save on re-calculation of the window
corr_wnd_inv_limited, // corr_wnd_limited, // correlation window, limited not to be smaller than threshold - used for finding max/convex areas (or null)
corrs2, // corrs, // double [][] corrs,
tile_disp_dist2,
rXY, // double [][] rXY, // non-distorted X,Y offset per nominal pixel of disparity
imgdtt_params.dbg_pair_mask, // int pair_mask, // which pairs to process
// null, // disp_str[cTile], //corr_stat[0], // double xcenter, // preliminary center x in pixels for largest baseline
poly_disp2, // double[] poly_ds, // null or pair of disparity/strength
imgdtt_params.ortho_vasw_pwr, // double vasw_pwr, // value as weight to this power,
tdl, // tile_lma_debug_level, //+2, // int debug_level,
tileX, // int tileX, // just for debug output
tileY); // int tileY
*/
if (lma2 != null) {
// was for single tile
disp_str = lma2.lmaDisparityStrength(
......@@ -1966,15 +2004,19 @@ public class ImageDtt extends ImageDttCPU {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DISP] += (lma_ds[0][0] + disparity_array_center[clustY][clustX] + disparity_corr) * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_TARGET] += (disparity_array_center[clustY][clustX] + disparity_corr) * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DIFF] += lma_ds[0][0] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 0] += pxpy[clustY][clustX][0][0] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 1] += pxpy[clustY][clustX][0][1] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 0] += pxpy[clustY][clustX][0] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 1] += pxpy[clustY][clustX][1] * w;
for (int cam = 0; cam < quad; cam++) {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DYDDISP0 + cam] += tile_disp_dist[cam][2] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PYDIST + cam] += clust_pY [clustY][clustX][cam] * w;
}
sum_w += w;
}
if (sum_w > 0.0) { // may be simplified as there is only one tile now
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_STRENGTH] = stats[0];
double strengh_k = 0.2*Math.sqrt(num_in_cluster[clustY][clustX]); // approximately matching old/multitile
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_STRENGTH] =
stats[0] * num_in_cluster[clustY][clustX] * strengh_k;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DISP] /= sum_w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_TARGET] /= sum_w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DIFF] /= sum_w;
......@@ -1982,6 +2024,7 @@ public class ImageDtt extends ImageDttCPU {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 1] /= sum_w;
for (int cam = 0; cam < quad; cam++) {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DYDDISP0 + cam] /= sum_w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PYDIST + cam] /= sum_w;
}
for (int cam = 0; cam < ddnd.length; cam++) {
......
......@@ -1743,6 +1743,7 @@ public class ImageDttCPU {
double [][] disp_str = new double [clustSize][];
/// double [][] dY_dD = new double [clustSize][quad];
double [][] pxpy = new double [clustSize][2];
// double [][] pYdist= new double [clustSize][quad]; // average pY for ERS calculation
boolean debugCluster = (clustX == debug_clustX) && (clustY == debug_clustY);
/// boolean debugCluster1 = (Math.abs(clustX - debug_clustX) < 10) && (Math.abs(clustY - debug_clustY) < 10);
......@@ -2125,8 +2126,9 @@ public class ImageDttCPU {
System.out.println("Will run new LMA for tileX="+tileX+", tileY="+tileY);
}
double [] poly_disp = {Double.NaN, 0.0};
Corr2dLMA lma2 = corr2d.corrLMA2(
Corr2dLMA lma2 = corr2d.corrLMA2Single(
imgdtt_params, // ImageDttParameters imgdtt_params,
false, // boolean adjust_ly, // adjust Lazy Eye
corr_wnd, // double [][] corr_wnd, // correlation window to save on re-calculation of the window
corr_wnd_inv_limited, // corr_wnd_limited, // correlation window, limited not to be smaller than threshold - used for finding max/convex areas (or null)
corrs[cTile], // double [][] corrs,
......@@ -2179,7 +2181,7 @@ public class ImageDttCPU {
System.out.println("Will run new LMA for clustX="+clustX+", clustY="+clustY);
}
Corr2dLMA lma2 = corr2d.corrLMA2(
Corr2dLMA lma2 = corr2d.corrLMA2Multi(
imgdtt_params, // ImageDttParameters imgdtt_params,
tileStep, // int clust_width,
corr_wnd, // double [][] corr_wnd, // correlation window to save on re-calculation of the window
......@@ -2225,6 +2227,7 @@ public class ImageDttCPU {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 1] += pxpy[cTile][1] * w;
for (int cam = 0; cam < quad; cam++) {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DYDDISP0 + cam] += disp_dist[cTile][cam][2] * w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PYDIST + cam] += centersXY[cTile][cam][1] * w;
}
sum_w += w;
}
......@@ -2241,6 +2244,7 @@ public class ImageDttCPU {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PX + 1] /= sum_w;
for (int cam = 0; cam < quad; cam++) {
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_DYDDISP0 + cam] /= sum_w;
lazy_eye_data[nCluster][ExtrinsicAdjustment.INDX_PYDIST + cam] /= sum_w;
}
for (int cam = 0; cam < ddnd.length; cam++) {
......@@ -3318,8 +3322,9 @@ public class ImageDttCPU {
System.out.println("Will run new LMA for tileX="+tileX+", tileY="+tileY);
double [] poly_disp = {Double.NaN, 0.0};
Corr2dLMA lma2 = corr2d.corrLMA2(
Corr2dLMA lma2 = corr2d.corrLMA2Single(
imgdtt_params, // ImageDttParameters imgdtt_params,
false, // boolean adjust_ly, // adjust Lazy Eye
corr_wnd, // double [][] corr_wnd, // correlation window to save on re-calculation of the window
corr_wnd_inv_limited, // corr_wnd_limited, // correlation window, limited not to be smaller than threshold - used for finding max/convex areas (or null)
corrs, // double [][] corrs,
......
......@@ -348,8 +348,8 @@ public class QuadCLT extends QuadCLTCPU {
}
if (resetEV) {
if (quadCLT_main.getGPU() != null) quadCLT_main.getGPU().resetGeometryCorrectionVector();
if (quadCLT_aux.getGPU() != null) quadCLT_aux.getGPU().resetGeometryCorrectionVector();
if (quadCLT_main.getGPU() != null) quadCLT_main.gpuResetCorrVector(); // getGPU().resetGeometryCorrectionVector();
if (quadCLT_aux.getGPU() != null) quadCLT_aux.gpuResetCorrVector(); // .getGPU().resetGeometryCorrectionVector();
}
// get fat_zero (absolute) and color scales
......@@ -2328,7 +2328,7 @@ public class QuadCLT extends QuadCLTCPU {
null, // final float [][][][] corr_combo_td, // [4][tilesY][tilesX][pair][4*64] TD of combo corrs: qud, cross, hor,vert
// each of the top elements may be null to skip particular combo type
fdisp_dist, // final float [][][][] fdisp_dist, // [tilesY][tilesX][cams][4], // disparity derivatives vectors or null
fpxpy, // final float [][][][] fdisp_dist, // [tilesY][tilesX][cams][2], tile {pX,pY}
fpxpy, // final float [][][][] fpxpy, // [tilesY][tilesX][cams][2], tile {pX,pY}
//// Uses quadCLT from gpuQuad
// correlation results - final and partial
null, // [type][tilesY][tilesX][(2*transform_size-1)*(2*transform_size-1)] // if null - will not calculate
......@@ -2380,7 +2380,7 @@ public class QuadCLT extends QuadCLTCPU {
disparity_array, // final double [][] disparity_array, // [tilesY][tilesX] - individual per-tile expected disparity
fcorr_td, // final float [][][][] fcorr_td, // [tilesY][tilesX][pair][4*64] transform domain representation of 6 corr
fdisp_dist, // final float [][][][] fdisp_dist, // [tilesY][tilesX][cams][4], // disparity derivatives vectors or null
fpxpy, // final float [][][][] fpxpy, // [tilesY][tilesX][cams][2], tile {pX,pY}
fpxpy, // final float [][][][] fpxpy, // [tilesY][tilesX][cams][2], tile {pX,pY}
clt_parameters.gpu_corr_scale, // gpu_corr_scale, // final double gpu_corr_scale, // 0.75; // reduce GPU-generated correlation values
gpu_fat_zero, // final double gpu_fat_zero, // clt_parameters.getGpuFatZero(is_mono);absolute == 30.0\
geometryCorrection, // final GeometryCorrection geometryCorrection,
......@@ -2399,4 +2399,16 @@ public class QuadCLT extends QuadCLTCPU {
scan.resetProcessed();
return scan;
}
public void gpuResetCorrVector() {
if (getGPU() != null) {
getGPU().resetGeometryCorrectionVector();
getGPU().resetGeometryCorrection(); // testing, just in case, should be removed
}
}
public void gpuResetGeometryCorrection() {
if (getGPU() != null) {
getGPU().resetGeometryCorrection();
}
}
}
......@@ -7758,6 +7758,8 @@ if (debugLevel > -100) return true; // temporarily !
boolean ok = quadCLT_main.extrinsicsCLT(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
false, // adjust_poly,
-1.0, // double inf_min,
1.0, // double inf_max,
threadsMax, //final int threadsMax, // maximal number of threads to launch
updateStatus,// final boolean updateStatus,
debugLevelInner); // final int debugLevel)
......@@ -7791,9 +7793,18 @@ if (debugLevel > -100) return true; // temporarily !
System.out.println("Adjusting aux camera image set for "+quadCLT_main.image_name+
" (w/o rig), pass "+(num_adjust_aux+1)+" of "+quadCLT_main.correctionsParameters.rig_batch_adjust_aux);
}
double inf_min = -1.0;
double inf_max = 1.0;
if (num_adjust_aux >= (quadCLT_main.correctionsParameters.rig_batch_adjust_aux/2)) {
inf_min = -0.2;
inf_max = 0.05;
}
boolean ok = quadCLT_aux.extrinsicsCLT(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
false, // adjust_poly,
inf_min, // double inf_min,
inf_max, // double inf_max,
threadsMax, //final int threadsMax, // maximal number of threads to launch
updateStatus, // final boolean updateStatus,
debugLevelInner); // final int debugLevel)
......@@ -8214,11 +8225,11 @@ if (debugLevel > -100) return true; // temporarily !
{
if ((quadCLT_main != null) && (quadCLT_main.getGPU() != null)) {
quadCLT_main.getGPU().resetGeometryCorrection();
quadCLT_main.getGPU().resetGeometryCorrectionVector();
quadCLT_main.gpuResetCorrVector(); // .getGPU().resetGeometryCorrectionVector();
}
if ((quadCLT_aux != null) && (quadCLT_aux.getGPU() != null)) {
quadCLT_aux.getGPU().resetGeometryCorrection();
quadCLT_aux.getGPU().resetGeometryCorrectionVector();
quadCLT_aux.gpuResetCorrVector(); // .getGPU().resetGeometryCorrectionVector();
}
// final boolean batch_mode = clt_parameters.batch_run;
......@@ -8303,7 +8314,7 @@ if (debugLevel > -100) return true; // temporarily !
System.out.println("Adjusting main camera image set for "+quadCLT_main.image_name+
", pass "+(num_adjust_main+1)+" of "+adjust_main);
}
if (debugLevel > -5){
if (debugLevel > -1){
int scan_index = quadCLT_main.tp.clt_3d_passes.size() -1;
quadCLT_main.tp.showScan(
quadCLT_main.tp.clt_3d_passes.get(scan_index), // CLTPass3d scan,
......@@ -8314,10 +8325,17 @@ if (debugLevel > -100) return true; // temporarily !
"pre-adjust-extrinsic-scan-"+s); //String title)
}
}
double inf_min = -1.0;
double inf_max = 1.0;
if (num_adjust_main >= (adjust_main/2)) {
inf_min = -0.2;
inf_max = 0.05;
}
boolean ok = quadCLT_main.extrinsicsCLT(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
false, // adjust_poly,
inf_min, // double inf_min,
inf_max, // double inf_max,
threadsMax, //final int threadsMax, // maximal number of threads to launch
updateStatus,// final boolean updateStatus,
debugLevelInner); // final int debugLevel)
......@@ -8567,10 +8585,18 @@ if (debugLevel > -100) return true; // temporarily !
}
if (quadCLT_aux.ds_from_main == null) {
System.out.println("BUG: quadCLT_aux.ds_from_main should be not null here!");
double inf_min = -1.0;
double inf_max = 1.0;
if (num_adjust_aux >= (adjust_aux/2)) {
inf_min = -0.2;
inf_max = 0.05;
}
// adjust w/o main camera - maybe will be used in the future
boolean ok = quadCLT_aux.extrinsicsCLT(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
false, // adjust_poly,
inf_min, // double inf_min,
inf_max, // double inf_max,
threadsMax, //final int threadsMax, // maximal number of threads to launch
updateStatus,// final boolean updateStatus,
debugLevelInner); // final int debugLevel)
......
......@@ -48,7 +48,7 @@ extern "C" __global__ void convert_direct( // called with a single block, single
float ** gpu_kernels, // [NUM_CAMS],
float ** gpu_images, // [NUM_CAMS],
struct tp_task * gpu_tasks,
float ** gpu_clt, // [NUM_CAMS][TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS][TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
size_t dstride, // in floats (pixels)
int num_tiles, // number of tiles in task
int lpf_mask, // apply lpf to colors : bit 0 - red, bit 1 - blue, bit2 - green. Now - always 0 !
......@@ -57,10 +57,12 @@ extern "C" __global__ void convert_direct( // called with a single block, single
int kernels_hor,
int kernels_vert,
int * gpu_active_tiles, // pointer to the calculated number of non-zero tiles
int * pnum_active_tiles); // indices to gpu_tasks
int * pnum_active_tiles, // indices to gpu_tasks
int tilesx);
extern "C" __global__ void correlate2D(
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS] ->[TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
int colors, // number of colors (3/1)
float scale0, // scale for R
float scale1, // scale for B
......@@ -68,41 +70,58 @@ extern "C" __global__ void correlate2D(
float fat_zero, // here - absolute
struct tp_task * gpu_tasks, // array of per-tile tasks (now bits 4..9 - correlation pairs)
int num_tiles, // number of tiles in task
int tilesx, // number of tile rows
int * gpu_corr_indices, // packed tile+pair
int * pnum_corr_tiles, // pointer to a number of correlation tiles to process
const size_t corr_stride, // in floats
int corr_radius, // radius of the output correlation (7 for 15x15)
float * gpu_corrs); // correlation output data
extern "C" __global__ void corr2D_normalize(
int num_corr_tiles, // number of correlation tiles to process
const size_t corr_stride_td, // in floats
float * gpu_corrs_td, // correlation tiles in transform domain
const size_t corr_stride, // in floats
float * gpu_corrs, // correlation output data (either pixel domain or transform domain
float fat_zero, // here - absolute
int corr_radius); // radius of the output correlation (7 for 15x15)
extern "C" __global__ void corr2D_combine(
int num_tiles, // number of tiles to process (each with num_pairs)
int num_pairs, // num pairs per tile (should be the same)
int init_output, // !=0 - reset output tiles to zero before accumulating
int pairs_mask, // selected pairs (0x3 - horizontal, 0xc - vertical, 0xf - quad, 0x30 - cross)
int * gpu_corr_indices, // packed tile+pair
int * gpu_combo_indices, // output if noty null: packed tile+pairs_mask (will point to the first used pair
const size_t corr_stride, // (in floats) stride for the input TD correlations
float * gpu_corrs, // input correlation tiles
const size_t corr_stride_combo, // (in floats) stride for the output TD correlations (same as input)
float * gpu_corrs_combo); // combined correlation output (one per tile)
extern "C" __global__ void textures_nonoverlap(
struct tp_task * gpu_tasks,
int num_tiles, // number of tiles in task list
// int num_tilesx, // number of tiles in a row
// declare arrays in device code?
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int * pnum_texture_tiles, // returns total number of elements in gpu_texture_indices array
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS] ->[TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
// TODO: use geometry_correction rXY !
struct gc * gpu_geometry_correction,
int colors, // number of colors (3/1)
int is_lwir, // do not perform shot correction
float params[5],
// float min_shot, // 10.0
// float scale_shot, // 3.0
// float diff_sigma, // pixel value/pixel change
// float diff_threshold, // pixel value/pixel change
// float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
// int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
// combining both non-overlap and overlap (each calculated if pointer is not null )
size_t texture_stride, // in floats (now 256*4 = 1024) // may be 0 if not needed
float * gpu_texture_tiles, // (number of colors +1 + ?)*16*16 rgba texture tiles // may be 0 if not needed
float * gpu_diff_rgb_combo); // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS] // may be 0 if not needed
float * gpu_diff_rgb_combo, //); // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS] // may be 0 if not needed
int num_tilesx);
extern "C"
__global__ void imclt_rbg_all(
float ** gpu_clt, // [NUM_CAMS][TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS][TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_corr_images, // [NUM_CAMS][WIDTH, 3 * HEIGHT]
int apply_lpf,
int colors,
......@@ -111,7 +130,7 @@ __global__ void imclt_rbg_all(
const size_t dstride); // in floats (pixels)
extern "C" __global__ void imclt_rbg(
float * gpu_clt, // [TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float * gpu_clt, // [TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float * gpu_rbg, // WIDTH, 3 * HEIGHT
int apply_lpf,
int mono, // defines lpf filter
......@@ -130,22 +149,15 @@ extern "C" __global__ void generate_RBGA(
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int * num_texture_tiles, // number of texture tiles to process (8 separate elements for accumulation)
int * woi, // x,y,width,height of the woi
int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
int height, // <= TILESY, use for faster processing of LWIR images
int width, // <= TILES-X, use for faster processing of LWIR images (should be actual + 1)
int height, // <= TILES-Y, use for faster processing of LWIR images
// Parameters for the texture generation
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS] ->[TILES-Y][TILES-X][NUM_COLORS][DTT_SIZE*DTT_SIZE]
// TODO: use geometry_correction rXY !
struct gc * gpu_geometry_correction,
int colors, // number of colors (3/1)
int is_lwir, // do not perform shot correction
float params[5], // mitigating CUDA_ERROR_INVALID_PTX
/*
float min_shot, // 10.0
float scale_shot, // 3.0
float diff_sigma, // pixel value/pixel change
float diff_threshold, // pixel value/pixel change
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
*/
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // return channel weights after A in RGBA (was removed)
......
......@@ -294,7 +294,6 @@ extern "C" __global__ void get_tiles_offsets(
float * gpu_rByRDist, // length should match RBYRDIST_LEN
trot_deriv * gpu_rot_deriv)
{
// int task_num = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx.x * TILES_PER_BLOCK_GEOM + threadIdx.x
int task_num = blockIdx.x * blockDim.y + threadIdx.y; // blockIdx.x * TILES_PER_BLOCK_GEOM + threadIdx.y
if (task_num >= num_tiles){
return;
......@@ -306,6 +305,7 @@ extern "C" __global__ void get_tiles_offsets(
__shared__ float rByRDist [RBYRDIST_LEN];
__shared__ struct corr_vector extrinsic_corr;
__shared__ trot_deriv rot_deriv;
__shared__ float pY_offsets[TILES_PER_BLOCK_GEOM][NUM_CAMS];
float pXY[2]; // result to be copied to task
// copy data common to all threads
{
......@@ -362,8 +362,7 @@ extern "C" __global__ void get_tiles_offsets(
(extrinsic_corr.imu_move[0] != 0.0) ||
(extrinsic_corr.imu_move[1] != 0.0) ||
(extrinsic_corr.imu_move[2] != 0.0);
// Temporary
imu_exists = 0;
#ifdef DEBUG21
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("\nTile = %d, camera= %d\n", task_num, ncam);
......@@ -373,6 +372,9 @@ extern "C" __global__ void get_tiles_offsets(
}
__syncthreads();// __syncwarp();
#endif // DEBUG21
// String dbg_s = corr_vector.toString();
/* Starting with required tile center X, Y and nominal distortion, for each sensor port:
* 1) unapply common distortion (maybe for different - master camera)
......@@ -401,15 +403,10 @@ extern "C" __global__ void get_tiles_offsets(
float pXcd = px - 0.5 * geometry_correction.pixelCorrectionWidth;
float pYcd = py - 0.5 * geometry_correction.pixelCorrectionHeight;
// float rXY [NUM_CAMS][2];
float rXY [2];
// for (int i = 0; i < NUM_CAMS;i++){
// rXY[ncam][0] = geometry_correction.rXY[ncam][0];
// rXY[ncam][1] = geometry_correction.rXY[ncam][1];
rXY[0] = geometry_correction.rXY[ncam][0];
rXY[1] = geometry_correction.rXY[ncam][1];
// }
float rD = sqrtf(pXcd*pXcd + pYcd*pYcd)*0.001*geometry_correction.pixelSize; // distorted radius in a virtual center camera
float rND2R=getRByRDist(rD/geometry_correction.distortionRadius, rByRDist);
......@@ -489,9 +486,17 @@ extern "C" __global__ void get_tiles_offsets(
float pYid = pYci * rD2rND;
pXY[0] = pXid + geometry_correction.pXY0[ncam][0];
pXY[1] = pYid + geometry_correction.pXY0[ncam][1];
// new for ERS
pY_offsets[threadIdx.y][ncam] = pXY[1] - geometry_correction.woi_tops[ncam];
__syncthreads();
// Each thread re-calculate same sum
float lines_avg = 0;
for (int i = 0; i < NUM_CAMS; i ++){
lines_avg += pY_offsets[threadIdx.y][i];
}
lines_avg *= (1.0/NUM_CAMS);
// used when calculating derivatives, TODO: combine calculations !
float pY_offset = pY_offsets[threadIdx.y][ncam] - lines_avg;
#ifdef DEBUG21
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("pXci0 = %f, pYci0 = %f\n", pXci0, pYci0);
......@@ -501,6 +506,7 @@ extern "C" __global__ void get_tiles_offsets(
printf("rD2rND = %f\n", rD2rND);
printf("pXid = %f, pYid = %f\n", pXid, pYid);
printf("pXY[0] = %f, pXY[1] = %f\n", pXY[0], pXY[1]); // OK
printf("lines_avg = %f, pY_offset = %f\n", lines_avg, pY_offset);
}
__syncthreads();// __syncwarp();
#endif // DEBUG21
......@@ -514,14 +520,10 @@ extern "C" __global__ void get_tiles_offsets(
#pragma unroll
for (int j = 0; j< 3; j++){
// drvi_daz[j] = rot_deriv.d_daz[ncam][j][0] * rvi[0] + rot_deriv.d_daz[ncam][j][1] * rvi[1] + rot_deriv.d_daz[ncam][j][2] * rvi[2];
// drvi_dtl[j] = rot_deriv.d_tilt[ncam][j][0] * rvi[0] + rot_deriv.d_tilt[ncam][j][1] * rvi[1] + rot_deriv.d_tilt[ncam][j][2] * rvi[2];
// drvi_drl[j] = rot_deriv.d_roll[ncam][j][0] * rvi[0] + rot_deriv.d_roll[ncam][j][1] * rvi[1] + rot_deriv.d_roll[ncam][j][2] * rvi[2];
drvi_daz[j] = rot_deriv.d_daz[ncam][j][0] * pXci0 + rot_deriv.d_daz[ncam][j][1] * pYci0 + rot_deriv.d_daz[ncam][j][2] * fl_pix;
drvi_dtl[j] = rot_deriv.d_tilt[ncam][j][0] * pXci0 + rot_deriv.d_tilt[ncam][j][1] * pYci0 + rot_deriv.d_tilt[ncam][j][2] * fl_pix;
drvi_drl[j] = rot_deriv.d_roll[ncam][j][0] * pXci0 + rot_deriv.d_roll[ncam][j][1] * pYci0 + rot_deriv.d_roll[ncam][j][2] * fl_pix;
}
// double [][] avi = {{pXci0}, {pYci0},{fl_pix}};
float dpXci_dazimuth = drvi_daz[0] * norm_z - pXci * drvi_daz[2] / rvi[2];
float dpYci_dazimuth = drvi_daz[1] * norm_z - pYci * drvi_daz[2] / rvi[2];
......@@ -573,25 +575,6 @@ extern "C" __global__ void get_tiles_offsets(
// unity vector in the direction of radius
float c_dist = pXci/rNDi;
float s_dist = pYci/rNDi;
/*
double [][] arot2= {
{c_dist, s_dist},
{-s_dist, c_dist}};
Matrix rot2 = new Matrix(arot2); // convert from non-distorted X,Y to parallel and perpendicular (CCW) to the radius
double [][] ascale_distort = {
{rD2rND + ri* drD2rND_dri, 0 },
{0, rD2rND}};
Matrix scale_distort = new Matrix(ascale_distort); // scale component parallel to radius as distortion derivative, perpendicular - as distortion
Matrix dd2 = rot2.transpose().times(scale_distort).times(rot2).times(dd1);
disp_dist[i][0] = dd2.get(0, 0);
disp_dist[i][1] = dd2.get(0, 1);
disp_dist[i][2] = dd2.get(1, 0); // d_py/d_disp
disp_dist[i][3] = dd2.get(1, 1);
*/
//#undef NVRTC_BUG
float drD2rND_dri = 0.0;
{
......@@ -612,11 +595,6 @@ extern "C" __global__ void get_tiles_offsets(
}
float scale_distort00 = rD2rND + ri* drD2rND_dri;
float scale_distort11 = rD2rND;
// float rot2Xdd1[2][2];
// rot2Xdd1[0][0] = c_dist * dd1[0][0] + s_dist * dd1[1][0];
// rot2Xdd1[0][1] = c_dist * dd1[0][1] + s_dist * dd1[1][1];
// rot2Xdd1[1][0] = -s_dist * dd1[0][0] + c_dist * dd1[1][0];
// rot2Xdd1[1][1] = -s_dist * dd1[0][1] + c_dist * dd1[1][1];
float scale_distortXrot2Xdd1[2][2];
scale_distortXrot2Xdd1[0][0] = ( c_dist * dd1[0][0] + s_dist * dd1[1][0]) * scale_distort00;
scale_distortXrot2Xdd1[0][1] = ( c_dist * dd1[0][1] + s_dist * dd1[1][1]) * scale_distort00;
......@@ -651,6 +629,7 @@ extern "C" __global__ void get_tiles_offsets(
// float imu_move[3]; // dx/dt, dy/dt, dz/dt 16..19 geometry_correction.imu_move
// ERS linear does not yet use per-port rotations, probably not needed
if (imu_exists){
/*
float delta_t = disp_dist[2] * disparity * geometry_correction.line_time; // positive for top cameras, negative - for bottom //disp_dist[2]=dd2.get(1, 0)
float ers_Xci = delta_t * (
dpXci_dtilt * extrinsic_corr.imu_rot[0] +
......@@ -660,9 +639,22 @@ extern "C" __global__ void get_tiles_offsets(
dpYci_dtilt * extrinsic_corr.imu_rot[0] +
dpYci_dazimuth * extrinsic_corr.imu_rot[1] +
dpYci_droll * extrinsic_corr.imu_rot[2]);
#ifdef DEBUG210
*/
float ers_x =
dpXci_dtilt * extrinsic_corr.imu_rot[0] +
dpXci_dazimuth * extrinsic_corr.imu_rot[1] +
dpXci_droll * extrinsic_corr.imu_rot[2];
float ers_y =
dpYci_dtilt * extrinsic_corr.imu_rot[0] +
dpYci_dazimuth * extrinsic_corr.imu_rot[1] +
dpYci_droll * extrinsic_corr.imu_rot[2];
#ifdef DEBUG21
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("delta_t = %f, ers_Xci = %f, ers_Yci = %f\n", delta_t, ers_Xci, ers_Yci);
// printf("delta_t = %f, ers_Xci = %f, ers_Yci = %f\n", delta_t, ers_Xci, ers_Yci);
printf("ers_x = %f, ers_y = %f\n", ers_x, ers_y);
}
__syncthreads();// __syncwarp();
#endif // DEBUG21
......@@ -674,22 +666,30 @@ extern "C" __global__ void get_tiles_offsets(
dpXci_pYci_imu_lin[1][1] = wdisparity / k; // dpy/ dworld_Y
dpXci_pYci_imu_lin[0][2] = (xyz[0] / k) * dwdisp_dz; // dpx/ dworld_Z
dpXci_pYci_imu_lin[1][2] = (xyz[1] / k) * dwdisp_dz; // dpy/ dworld_Z
/*
ers_Xci += delta_t* (
dpXci_pYci_imu_lin[0][0] * extrinsic_corr.imu_move[0] +
dpXci_pYci_imu_lin[0][2] * extrinsic_corr.imu_move[2]);
ers_Yci += delta_t* (
dpXci_pYci_imu_lin[1][1] * extrinsic_corr.imu_move[1] +
dpXci_pYci_imu_lin[1][2] * extrinsic_corr.imu_move[2]);
pXY[0] += ers_Xci * rD2rND; // added correction to pixel X
pXY[1] += ers_Yci * rD2rND; // added correction to pixel Y
*/
ers_x += dpXci_pYci_imu_lin[0][0] * extrinsic_corr.imu_move[0] +
dpXci_pYci_imu_lin[0][2] * extrinsic_corr.imu_move[2];
ers_y += dpXci_pYci_imu_lin[1][1] * extrinsic_corr.imu_move[1] +
dpXci_pYci_imu_lin[1][2] * extrinsic_corr.imu_move[2];
float delta_t = (pY_offset/ (1.0 - geometry_correction.line_time * ers_y)) * geometry_correction.line_time; // positive for top cameras, negative - for bottom //disp_dist[2]=dd2.get(1, 0)
#ifdef DEBUG210
pXY[0] += delta_t * ers_x * rD2rND; // added correction to pixel X
pXY[1] += delta_t * ers_y * rD2rND; // added correction to pixel Y
#ifdef DEBUG21
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("k = %f, wdisparity = %f, dwdisp_dz = %f\n", k, wdisparity, dwdisp_dz);
printf("dpXci_pYci_imu_lin[0][0] = %f, dpXci_pYci_imu_lin[0][2] = %f\n", dpXci_pYci_imu_lin[0][0],dpXci_pYci_imu_lin[0][2]);
printf("dpXci_pYci_imu_lin[1][1] = %f, dpXci_pYci_imu_lin[1][2] = %f\n", dpXci_pYci_imu_lin[1][1],dpXci_pYci_imu_lin[1][2]);
printf("delta_t = %f, ers_Xci = %f, ers_Yci = %f\n", delta_t, ers_Xci, ers_Yci);
printf("delta_t = %f, ers_x = %f, ers_y = %f\n", delta_t, ers_x, ers_y);
printf("pXY[0] = %f, pXY[1] = %f\n", pXY[0], pXY[1]); // OK
}
__syncthreads();// __syncwarp();
......@@ -703,6 +703,7 @@ extern "C" __global__ void get_tiles_offsets(
}
extern "C" __global__ void calcReverseDistortionTable(
struct gc * geometry_correction,
float * rByRDist)
......@@ -841,6 +842,7 @@ __device__ void printGeometryCorrection(struct gc * g){
printf("%22s: %f\n","cameraRadius", g->cameraRadius);
printf("%22s: %f\n","disparityRadius", g->disparityRadius);
printf("%22s: %f, %f, %f, %f \n","woi_tops", g->woi_tops[0], g->woi_tops[1], g->woi_tops[2], g->woi_tops[3]);
#endif //ifndef JCUDA
}
......
......@@ -138,6 +138,7 @@ struct gc {
// only used for the multi-quad systems
float cameraRadius; // =0; // average distance from the "mass center" of the sensors to the sensors
float disparityRadius; // =150.0; // distance between cameras to normalize disparity units to. sqrt(2)*disparityRadius for quad
float woi_tops [NUM_CAMS]; // used to calculate scanline timing
};
#define RAD_COEFF_LEN 7
extern "C" __global__ void get_tiles_offsets(
......