Skip to content

I
imagej-elphel
Commit 9b711009 authored by Andrey Filippov's avatar Andrey Filippov
Browse files
Options

more debugging

parent 24112d36
Hide whitespace changes
Inline Side-by-side
Showing with 158 additions and 22 deletions
src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java
View file @ 9b711009
......@@ -116,7 +116,7 @@ public class GPUTileProcessor {
static int THREADSX = DTT_SIZE;
public static int NUM_CAMS = 4;
public static int NUM_PAIRS = 6; // top hor, bottom hor, left vert, right vert, main diagonal, other diagonal
static int NUM_COLORS = 3;
public static int NUM_COLORS = 3;
public static int IMG_WIDTH = 2592;
public static int IMG_HEIGHT = 1936;
static int KERNELS_HOR = 164;
......
src/main/java/com/elphel/imagej/tileprocessor/GeometryCorrection.java
View file @ 9b711009
......@@ -3376,13 +3376,15 @@ matrix([[-0.125, -0.125, 0.125, 0.125, -0.125, 0.125, -0. , -0. , -0.
double ri_scale = 0.001 * this.pixelSize / this.distortionRadius;
System.out.println("fl_pix="+fl_pix+", ri_scale="+ri_scale);
double [] xyz = (disparity > 0) ? getWorldCoordinates( // USED in lwir
double [] xyz = ((disparity > 0) ? getWorldCoordinates( // USED in lwir
px, // double px,
py, // double py,
disparity, // double disparity,
true) : null; // boolean correctDistortions)
true) : null); // boolean correctDistortions)
System.out.println("xyz[0]="+xyz[0]+", xyz[1]="+xyz[1]+", xyz[2]="+xyz[2]);
if (xyz != null) {
System.out.println("xyz[0]="+xyz[0]+", xyz[1]="+xyz[1]+", xyz[2]="+xyz[2]);
}
for (int i = 0; i < numSensors; i++){
// non-distorted XY of the shifted location of the individual sensor
......
src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java
View file @ 9b711009
......@@ -4071,6 +4071,32 @@ public class ImageDtt {
}
}
}
if (debug_gpu) {
System.out.println("--- color_avg ---");
for (int ncol = 0; ncol < numcol; ncol++) if (iclt_tile[0][ncol] != null) {
System.out.println("\n --- color_avg ["+ncol+"] ---");
for (int ii = 0; ii < 2 * transform_size; ii++) {
for (int jj = 0; jj < 2 * transform_size; jj++) {
System.out.print(String.format("%10.4f ", color_avg[ncol][2* transform_size * ii + jj]));
}
System.out.println();
}
for (int ccam = 0; ccam < port_weights.length; ccam ++) {
System.out.println("--- max_diff ["+ccam+"] ---");
System.out.println("\n --- imclt ["+ccam+"]["+ncol+"] ---");
for (int ii = 0; ii < 2 * transform_size; ii++) {
for (int jj = 0; jj < 2 * transform_size; jj++) {
System.out.print(String.format("%10.4f ", iclt_tile[ccam][ncol][2* transform_size * ii + jj]));
}
System.out.println();
}
}
}
System.out.println("--- max_diff: "+max_diff[0]+","+max_diff[1]+","+max_diff[2]+","+max_diff[3]);
System.out.println("--- R: "+ports_rgb[ 0]+", "+ports_rgb[ 1]+", "+ports_rgb[ 2]+", "+ports_rgb[ 3]);
System.out.println("--- B: "+ports_rgb[ 4]+", "+ports_rgb[ 5]+", "+ports_rgb[ 6]+", "+ports_rgb[ 7]);
System.out.println("--- G: "+ports_rgb[ 8]+", "+ports_rgb[ 9]+", "+ports_rgb[10]+", "+ports_rgb[11]);
}
return rgba;
}
......@@ -8973,6 +8999,8 @@ public class ImageDtt {
lpf_rgb[i] = filter;
}
generateTextureTiles(// not used in lwir
null, // final double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
null, // final double [] max_diff, // maximal (weighted) deviation of each channel from the average
clt_parameters,
extra_disparity,
quad, // number of subcameras
......@@ -8993,6 +9021,8 @@ public class ImageDtt {
}
public void generateTextureTiles(// not used in lwir
final double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
final double [] max_diff, // maximal (weighted) deviation of each channel from the average
final CLTParameters clt_parameters,
final double extra_disparity,
final int quad, // number of subcameras
......@@ -9187,8 +9217,8 @@ public class ImageDtt {
texture_tiles[tileY][tileX] = tile_combine_rgba(
tiles_debayered, // iclt_tile, // [port][numcol][256]
null, // double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
null, // max_diff, // maximal (weighted) deviation of each channel from the average
ports_rgb, // null, // double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
max_diff, // null, // max_diff, // maximal (weighted) deviation of each channel from the average
lt_window2, // [256]
port_offsets, // [port]{x_off, y_off}
img_mask, // which port to use, 0xf - all 4 (will modify as local variable)
......@@ -9454,6 +9484,9 @@ public class ImageDtt {
final double [][] disparity_bimap, // [23][tilesY][tilesX], only [6][] is needed on input or null - do not calculate
// last 2 - contrast, avg/ "geometric average)
final double [][] ml_data, // data for ML - 18 layers - 4 center areas (3x3, 5x5,..) per camera-per direction, 1 - composite, and 1 with just 1 data (target disparity)
// added in debug version
final double [][] disparity_map, // [8][tilesY][tilesX], only [6][] is needed on input or null - do not calculate
final double [][][][] texture_tiles_main, // [tilesY][tilesX]["RGBA".length()][]; null - will skip images combining
final double [][][][] texture_tiles_aux, // [tilesY][tilesX]["RGBA".length()][]; null - will skip images combining
final int width, // may be not multiple of 8, same for the height
......@@ -9476,6 +9509,7 @@ public class ImageDtt {
final int debug_tileY = clt_parameters.tileY;
final int quad_main = image_data_main.length; // number of subcameras
final int quad_aux = image_data_aux.length; // number of subcameras
final int quad = 4; // number of subcameras
final int numcol = 3; // number of colors
final int nChn = image_data_main[0].length;
final int height=image_data_main[0][0].length/width;
......@@ -9541,6 +9575,27 @@ public class ImageDtt {
final double [] filter = doubleGetCltLpfFd(clt_parameters.getCorrSigma(isMonochrome()));
dbg_filter_corr = filter;
// add optional initialization of debug layers here
if (disparity_map != null){
for (int i = 0; i<disparity_map.length;i++){
if (i < OVEREXPOSED) {
disparity_map[i] = new double [tilesY*tilesX];
} else if (i == OVEREXPOSED) {
// if (saturation_imp!= null) {
// disparity_map[i] = new double [tilesY*tilesX];
// }
} else if (i >= IMG_TONE_RGB) {
// if (texture_tiles != null) { // for now - enable 12 tone layers only together with texture tiles
disparity_map[i] = new double [tilesY*tilesX];
// }
}
}
}
final double [][] lpf_rgb = new double[isMonochrome()?1:3][];
if (isMonochrome()) {
lpf_rgb[0] = doubleGetCltLpfFd(clt_parameters.gpu_sigma_m);
......@@ -10091,8 +10146,23 @@ public class ImageDtt {
}
}
double [] max_diff = null;
if ((disparity_map != null) && (disparity_map.length >= (IMG_DIFF0_INDEX + quad))){
max_diff = new double[quad];
}
double [] ports_rgb = null;
int ports_rgb_len = quad*numcol; // 12
if ((disparity_map != null) && (disparity_map.length >= (IMG_TONE_RGB + ports_rgb_len))) {
ports_rgb = new double[ports_rgb_len];
}
if (texture_tiles_main !=null) {
generateTextureTiles (
ports_rgb, //final double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
max_diff, // final double [] max_diff, // maximal (weighted) deviation of each channel from the average
clt_parameters, // final EyesisCorrectionParameters.CLTParameters clt_parameters,
extra_disparity_main, // final double extra_disparity,
quad_main, // final int quad, // number of subcameras
......@@ -10113,6 +10183,8 @@ public class ImageDtt {
}
if (texture_tiles_aux !=null) {
generateTextureTiles (
null, // final double [] ports_rgb, // average values of R,G,B for each camera (R0,R1,...,B2,B3)
null, // final double [] max_diff, // maximal (weighted) deviation of each channel from the average
clt_parameters, // final EyesisCorrectionParameters.CLTParameters clt_parameters,
extra_disparity_aux, // final double extra_disparity,
quad_aux, // final int quad, // number of subcameras
......@@ -10131,6 +10203,8 @@ public class ImageDtt {
tile_lma_debug_level, // final int debugLevel);
false); // debug_gpu); //boolean debug_gpu
}
if (debug_gpu) {
double [][] texture_tile = texture_tiles_main[clt_parameters.tileY][clt_parameters.tileX];
int tile = +clt_parameters.tileY * tilesX + +clt_parameters.tileX;
......@@ -10148,8 +10222,16 @@ public class ImageDtt {
}
}
if ((disparity_map != null) && (disparity_map.length >= (IMG_DIFF0_INDEX + quad))){
for (int i = 0; i < max_diff.length; i++){
disparity_map[IMG_DIFF0_INDEX + i][tIndex] = max_diff[i];
}
}
if (ports_rgb != null) {
for (int i = 0; i < ports_rgb.length; i++){
disparity_map[IMG_TONE_RGB + i][tIndex] = ports_rgb[i];
}
}
// Save channel tiles to result
if (clt_bidata != null) {
......
src/main/java/com/elphel/imagej/tileprocessor/TwoQuadCLT.java
View file @ 9b711009
......@@ -1390,6 +1390,7 @@ public class TwoQuadCLT {
double [][][] port_xy_main_dbg = new double [tilesX*tilesY][][];
double [][][] port_xy_aux_dbg = new double [tilesX*tilesY][][];
// double [][][] corr2ddata = new double [1][][];
double [][] disparity_map = new double [ImageDtt.DISPARITY_TITLES.length][];
final double [][][][][][][] clt_bidata = // new double[2][quad][nChn][tilesY][tilesX][][]; // first index - main/aux
image_dtt.clt_bi_quad_dbg (
......@@ -1411,6 +1412,7 @@ public class TwoQuadCLT {
// types: 0 - selected correlation (product+offset), 1 - sum
disparity_bimap, // final double [][] disparity_bimap, // [23][tilesY][tilesX]
ml_data, // final double [][] ml_data, // data for ML - 10 layers - 4 center areas (3x3, 5x5,..) per camera-per direction, 1 - composite, and 1 with just 1 data (target disparity)
disparity_map, // final double [][] disparity_map, // [8][tilesY][tilesX], only [6][] is needed on input or null - do not calculate
texture_tiles_main, // final double [][][][] texture_tiles_main, // [tilesY][tilesX]["RGBA".length()][]; null - will skip images combining
texture_tiles_aux, // final double [][][][] texture_tiles_aux, // [tilesY][tilesX]["RGBA".length()][]; null - will skip images combining
imp_quad_main[0].getWidth(), // final int width,
......@@ -1426,7 +1428,29 @@ public class TwoQuadCLT {
port_xy_main_dbg, // final double [][][] port_xy_main_dbg, // for each tile/port save x,y pixel coordinates (gpu code development)
port_xy_aux_dbg); // final double [][][] port_xy_aux_dbg) // for each tile/port save x,y pixel coordinates (gpu code development)
///// double [][] disparity_map = new double [ImageDtt.DISPARITY_TITLES.length][]; //[0] -residual disparity, [1] - orthogonal (just for debugging)
String [] sub_titles = new String [GPUTileProcessor.NUM_CAMS * (GPUTileProcessor.NUM_COLORS+1)];
double [][] sub_disparity_map = new double [sub_titles.length][];
for (int ncam = 0; ncam < GPUTileProcessor.NUM_CAMS; ncam++) {
sub_disparity_map[ncam] = disparity_map[ncam + ImageDtt.IMG_DIFF0_INDEX];
sub_titles[ncam] = ImageDtt.DISPARITY_TITLES[ncam + ImageDtt.IMG_DIFF0_INDEX];
for (int ncol = 0; ncol < GPUTileProcessor.NUM_COLORS; ncol++) {
sub_disparity_map[ncam + (ncol + 1)* GPUTileProcessor.NUM_CAMS] =
disparity_map[ncam +ncol* GPUTileProcessor.NUM_CAMS+ ImageDtt.IMG_TONE_RGB];
sub_titles[ncam + (ncol + 1)* GPUTileProcessor.NUM_CAMS] =
ImageDtt.DISPARITY_TITLES[ncam +ncol* GPUTileProcessor.NUM_CAMS+ ImageDtt.IMG_TONE_RGB];
}
}
// String [] sub_titles = {ImageDtt.DISPARITY_TITLES[ImageDtt.IMG_DIFF0_INDEX]
(new ShowDoubleFloatArrays()).showArrays(
sub_disparity_map,
tilesX,
tilesY,
true,
name + "-CPU-EXTRA-D"+clt_parameters.disparity,
sub_titles);
// Create list of all correlation pairs
double [][][][][][] clt_data = clt_bidata[0];
int numTiles = tilesX * tilesY;
......@@ -1463,10 +1487,9 @@ public class TwoQuadCLT {
wh[0],
wh[1],
true,
"CORR2D_CPU",
name + "-CPU-CORR2D-D"+clt_parameters.disparity,
GPUTileProcessor.getCorrTitles());
if ((save_prefix != null) && (save_prefix != "")) {
if (debugLevel < -1000) {
......
src/main/resources/kernels/TileProcessor.cuh
View file @ 9b711009
......@@ -2149,7 +2149,8 @@ __global__ void textures_accumulate( // (8,4,1) (N,1,1)
(float*) shr.mclt_debayer, // float * mclt_tile, // debayer // has gaps to align with union !
(float*) mclt_tiles, // float * rbg_tile, // if not null - original (not-debayered) rbg tile to use for the output
(float *) shr1.rgbaw, // float * rgba,
// result
// if calc_extra, rbg_tile will be ignored and output generated with blurred (debayered) data. Done so as debayered data is needed
// to calculate max_diff_shared
calc_extra, // int calc_extra, // 1 - calcualate ports_rgb, max_diff
ports_rgb_shared, // float ports_rgb_shared [NUM_COLORS][NUM_CAMS], // return to system memory (optionally pass null to skip calculation)
......@@ -3512,7 +3513,7 @@ __device__ void tile_combine_rgba(
float * mclt_tile, // debayer // has gaps to align with union !
float * rbg_tile, // if not null - original (not-debayered) rbg tile to use for the output
float * rgba, // result
int calc_extra, // 1 - calcualate ports_rgb, max_diff
int calc_extra, // 1 - calculate ports_rgb, max_diff (if not null - will ignore rbg_tile !)
float ports_rgb_shared [NUM_COLORS][NUM_CAMS], // return to system memory (optionally pass null to skip calculation)
float max_diff_shared [NUM_CAMS], // return to system memory (optionally pass null to skip calculation)
float max_diff_tmp [NUM_CAMS][TEXTURE_THREADS_PER_TILE],
......@@ -3846,7 +3847,7 @@ __device__ void tile_combine_rgba(
#endif // #ifdef DEBUG9
///
if (rbg_tile) {
if (rbg_tile && (calc_extra == 0)) { // will keep debayered if (calc_extra == 0)
float k = 0.0;
int rbga_offset = colors * (DTT_SIZE2*DTT_SIZE21); // padded in union !
#pragma unroll
......@@ -3923,12 +3924,12 @@ __device__ void tile_combine_rgba(
int row = (pass >> 1);
int col = ((pass & 1) << 3) + threadIdx.x;
int i = row * DTT_SIZE21 + col;
int row_sym = row ^ ((row & 8)? 0xf : 0);
int col_sym = col ^ ((col & 8)? 0xf : 0);
/// int row_sym = row ^ ((row & 8)? 0xf : 0);
/// int col_sym = col ^ ((col & 8)? 0xf : 0);
// Was it a bug?
// float wnd2 = HWINDOW_SQ[row_sym] * HWINDOW_SQi[col_sym];
float wnd2 = HWINDOW_SQ[row_sym] * HWINDOW_SQ[col_sym];
/// float wnd2 = HWINDOW_SQ[row_sym] * HWINDOW_SQ[col_sym];
float * mclt_cam_i = mclt_tile + colors_offset * cam + i;
// float * mclt_cam_i = rbg_tile + colors_offset * cam + i;
//
......@@ -3939,7 +3940,7 @@ __device__ void tile_combine_rgba(
float dc = *(mclt_cam_i + (DTT_SIZE2*(DTT_SIZE21 + 1)) * ncol) - *(rgba + (DTT_SIZE2*DTT_SIZE21) * ncol + i);
d2 += *(chn_weights + ncol) * dc * dc;
}
// d2 *= wnd2;
/// d2 *= wnd2;
max_diff_tmp[cam][threadIdx.x] = fmaxf(max_diff_tmp[cam][threadIdx.x], d2);
}
__syncthreads();
......@@ -3959,8 +3960,29 @@ __device__ void tile_combine_rgba(
for (int i = 0; i < TEXTURE_THREADS_PER_TILE; i++){
printf("tmp[%d] %f %f %f %f\n",i, max_diff_tmp[0][i],max_diff_tmp[1][i],max_diff_tmp[2][i],max_diff_tmp[3][i]);
}
for (int ncol = 0; ncol < colors; ncol++){
printf("\n average for color %d\n",ncol);
debug_print_mclt(
rgba + (DTT_SIZE2*DTT_SIZE21) * ncol,
-1);
for (int ncam = 0; ncam < NUM_CAMS;ncam ++){
printf("\n mclt for color %d, camera %d\n",ncol,ncam);
debug_print_mclt(
mclt_tile + (DTT_SIZE2*(DTT_SIZE21 + 1)) * ncol + colors_offset * ncam,
-1);
#if 0
printf("\n rgb_tile for color %d, camera %d\n",ncol,ncam);
if (rgb_tile) {
debug_print_mclt(
rbg_tile + (DTT_SIZE2*(DTT_SIZE21 + 1)) * ncol + colors_offset * ncam,
-1);
}
#endif
}
}
}
__syncthreads();// __syncwarp();
#endif // #ifdef DEBUG22
}
if (calc_extra) {
......
src/main/resources/kernels/geometry_correction.cu
View file @ 9b711009
......@@ -389,6 +389,13 @@ extern "C" __global__ void get_tiles_offsets(
float disparity = gpu_tasks[task_num].target_disparity;
int tileX = (cxy & 0xffff);
int tileY = (cxy >> 16);
#ifdef DEBUG23
if ((ncam == 0) && (tileX == DBG_TILE_X) && (tileY == DBG_TILE_Y)){
printf ("\n get_tiles_offsets(): Debugging tileX=%d, tileY=%d, ncam = %d\n", tileX,tileY,ncam);
printf("\n");
__syncthreads();
}
#endif //#ifdef DEBUG23
float px = tileX * DTT_SIZE + DTT_SIZE/2; // - shiftX;
float py = tileY * DTT_SIZE + DTT_SIZE/2; // - shiftY;
......@@ -524,7 +531,7 @@ extern "C" __global__ void get_tiles_offsets(
float dpXci_droll = drvi_drl[0] * norm_z - pXci * drvi_drl[2] / rvi[2];
float dpYci_droll = drvi_drl[1] * norm_z - pYci * drvi_drl[2] / rvi[2];
#ifdef DEBUG21
#ifdef DEBUG210
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("drvi_daz[0] = %f, drvi_daz[1] = %f, drvi_daz[2] = %f\n", drvi_daz[0], drvi_daz[1], drvi_daz[2]);
printf("drvi_dtl[0] = %f, drvi_dtl[1] = %f, drvi_dtl[2] = %f\n", drvi_dtl[0], drvi_dtl[1], drvi_dtl[2]);
......@@ -554,7 +561,7 @@ extern "C" __global__ void get_tiles_offsets(
dd1[1][0] = (-rot_deriv.rots[ncam][1][0]*rXY[0] -rot_deriv.rots[ncam][1][1]*rXY[1])*norm_z;
dd1[1][1] = ( rot_deriv.rots[ncam][1][0]*rXY[1] -rot_deriv.rots[ncam][1][1]*rXY[0])*norm_z;
#ifdef DEBUG21
#ifdef DEBUG210
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("dd1[0][0] = %f, dd1[0][1] = %f\n",dd1[0][0],dd1[0][1]);
printf("dd1[1][0] = %f, dd1[1][1] = %f\n",dd1[1][0],dd1[1][1]);
......@@ -622,7 +629,7 @@ extern "C" __global__ void get_tiles_offsets(
disp_dist[2] = s_dist * scale_distortXrot2Xdd1[0][0] + c_dist * scale_distortXrot2Xdd1[1][0];
disp_dist[3] = s_dist * scale_distortXrot2Xdd1[0][1] + c_dist * scale_distortXrot2Xdd1[1][1];
#ifdef DEBUG21
#ifdef DEBUG210
if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
printf("scale_distortXrot2Xdd1[0][0] = %f, scale_distortXrot2Xdd1[0][1] = %f\n",scale_distortXrot2Xdd1[0][0],scale_distortXrot2Xdd1[0][1]);
printf("scale_distortXrot2Xdd1[1][0] = %f, scale_distortXrot2Xdd1[1][1] = %f\n",scale_distortXrot2Xdd1[1][0],scale_distortXrot2Xdd1[1][1]);
......@@ -654,7 +661,7 @@ 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 DEBUG21
#ifdef DEBUG210
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);
}
......@@ -677,7 +684,7 @@ extern "C" __global__ void get_tiles_offsets(
pXY[0] += ers_Xci * rD2rND; // added correction to pixel X
pXY[1] += ers_Yci * rD2rND; // added correction to pixel Y
#ifdef DEBUG21
#ifdef DEBUG210
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]);
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment