Skip to content

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

refactoring, combining with the GPU

parent 4a49fd62
Hide whitespace changes
Inline Side-by-side
Showing with 12521 additions and 11814 deletions
src/main/java/com/elphel/imagej/cameras/CLTParameters.java
View file @ ac185fe7
......@@ -779,11 +779,22 @@ public class CLTParameters {
public double gpu_fatz = 30.0;
public double gpu_fatz_m = 30.0;
public int gpu_woi_tx = 0;
public int gpu_woi_ty = 0;
public int gpu_woi_twidth = 324;
public int gpu_woi_theight = 242;
public boolean gpu_woi_round = false;
public boolean gpu_woi = false; // if true - use gpu_woi_tx, ...
public int gpu_woi_tx = 0;
public int gpu_woi_ty = 0;
public int gpu_woi_twidth = 324;
public int gpu_woi_theight = 242;
public boolean gpu_woi_round = false;
public boolean gpu_save_ports_xy = false; // debug feature - save calculated ports X,Y to compare with Java-generated
public boolean gpu_show_jtextures = true; // debug feature - show Java-generated textures from non-overlapping in GPU (will not generate if false)
public boolean gpu_show_extra = true; // show low-res data for macro
public boolean gpu_use_main = false; // accelerate tile processor for the main quad camera
public boolean gpu_use_main_macro = false; // accelerate tile processor for the main quad camera in macro mode
public boolean gpu_use_main_adjust = false; // accelerate tile processor for the main quad camera for field calibration
public boolean gpu_use_aux = false; // accelerate tile processor for the aux (lwir) quad camera
public boolean gpu_use_aux_macro = false; // accelerate tile processor for the aux (lwir) quad camera in macro mode
public boolean gpu_use_aux_adjust = false; // accelerate tile processor for the aux (lwir) quad camera for field calibration
public boolean replaceWeakOutliers = true; // false;
......@@ -830,6 +841,17 @@ public class CLTParameters {
public boolean batch_run = false; // turned on only while running in batch mode
public boolean useGPU() {
return useGPU(false) || useGPU(true);
}
public boolean useGPU(boolean aux) {
return aux? (gpu_use_aux || gpu_use_aux_macro || gpu_use_aux_adjust)
: (gpu_use_main || gpu_use_main_macro || gpu_use_main_adjust);
}
public double getCorrSigma(boolean monochrome) {
return monochrome ? corr_sigma_mono : corr_sigma;
}
......@@ -1558,14 +1580,25 @@ public class CLTParameters {
properties.setProperty(prefix+"gpu_fatz", this.gpu_fatz +"");
properties.setProperty(prefix+"gpu_fatz_m", this.gpu_fatz_m +"");
properties.setProperty(prefix+"gpu_woi", this.gpu_woi +"");
properties.setProperty(prefix+"gpu_woi_tx", this.gpu_woi_tx +"");
properties.setProperty(prefix+"gpu_woi_ty", this.gpu_woi_ty +"");
properties.setProperty(prefix+"gpu_woi_twidth", this.gpu_woi_twidth +"");
properties.setProperty(prefix+"gpu_woi_theight", this.gpu_woi_theight +"");
properties.setProperty(prefix+"gpu_woi_round", this.gpu_woi_round +"");
properties.setProperty(prefix+"debug_initial_discriminate", this.debug_initial_discriminate+"");
properties.setProperty(prefix+"dbg_migrate", this.dbg_migrate+"");
properties.setProperty(prefix+"gpu_save_ports_xy", this.gpu_save_ports_xy +"");
properties.setProperty(prefix+"gpu_show_jtextures", this.gpu_show_jtextures +"");
properties.setProperty(prefix+"gpu_show_extra", this.gpu_show_extra +"");
properties.setProperty(prefix+"gpu_use_main", this.gpu_use_main +"");
properties.setProperty(prefix+"gpu_use_main_macro", this.gpu_use_main_macro +"");
properties.setProperty(prefix+"gpu_use_main_adjust", this.gpu_use_main_adjust +"");
properties.setProperty(prefix+"gpu_use_aux", this.gpu_use_aux +"");
properties.setProperty(prefix+"gpu_use_aux_macro", this.gpu_use_aux_macro +"");
properties.setProperty(prefix+"gpu_use_aux_adjust", this.gpu_use_aux_adjust +"");
properties.setProperty(prefix+"debug_initial_discriminate", this.debug_initial_discriminate+"");
properties.setProperty(prefix+"dbg_migrate", this.dbg_migrate+"");
properties.setProperty(prefix+"dbg_early_exit", this.dbg_early_exit+"");
properties.setProperty(prefix+"show_first_bg", this.show_first_bg+"");
......@@ -2332,14 +2365,25 @@ public class CLTParameters {
if (properties.getProperty(prefix+"gpu_fatz")!=null) this.gpu_fatz=Double.parseDouble(properties.getProperty(prefix+"gpu_fatz"));
if (properties.getProperty(prefix+"gpu_fatz_m")!=null) this.gpu_fatz_m=Double.parseDouble(properties.getProperty(prefix+"gpu_fatz_m"));
if (properties.getProperty(prefix+"gpu_woi")!=null) this.gpu_woi=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_woi"));
if (properties.getProperty(prefix+"gpu_woi_tx")!=null) this.gpu_woi_tx=Integer.parseInt(properties.getProperty(prefix+"gpu_woi_tx"));
if (properties.getProperty(prefix+"gpu_woi_ty")!=null) this.gpu_woi_ty=Integer.parseInt(properties.getProperty(prefix+"gpu_woi_ty"));
if (properties.getProperty(prefix+"gpu_woi_twidth")!=null) this.gpu_woi_twidth=Integer.parseInt(properties.getProperty(prefix+"gpu_woi_twidth"));
if (properties.getProperty(prefix+"gpu_woi_theight")!=null) this.gpu_woi_theight=Integer.parseInt(properties.getProperty(prefix+"gpu_woi_theight"));
if (properties.getProperty(prefix+"gpu_woi_round")!=null) this.gpu_woi_round=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_woi_round"));
if (properties.getProperty(prefix+"debug_initial_discriminate")!=null) this.debug_initial_discriminate=Boolean.parseBoolean(properties.getProperty(prefix+"debug_initial_discriminate"));
if (properties.getProperty(prefix+"dbg_migrate")!=null) this.dbg_migrate=Boolean.parseBoolean(properties.getProperty(prefix+"dbg_migrate"));
if (properties.getProperty(prefix+"gpu_save_ports_xy")!=null) this.gpu_save_ports_xy=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_save_ports_xy"));
if (properties.getProperty(prefix+"gpu_show_jtextures")!=null) this.gpu_show_jtextures=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_show_jtextures"));
if (properties.getProperty(prefix+"gpu_show_extra")!=null) this.gpu_show_extra=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_show_extra"));
if (properties.getProperty(prefix+"gpu_use_main")!=null) this.gpu_use_main=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_main"));
if (properties.getProperty(prefix+"gpu_use_main_macro")!=null) this.gpu_use_main_macro=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_main_macro"));
if (properties.getProperty(prefix+"gpu_use_main_adjust")!=null) this.gpu_use_main_adjust=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_main_adjust"));
if (properties.getProperty(prefix+"gpu_use_aux")!=null) this.gpu_use_aux=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_aux"));
if (properties.getProperty(prefix+"gpu_use_aux_macro")!=null) this.gpu_use_aux_macro=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_aux_macro"));
if (properties.getProperty(prefix+"gpu_use_aux_adjust")!=null) this.gpu_use_aux_adjust=Boolean.parseBoolean(properties.getProperty(prefix+"gpu_use_aux_adjust"));
if (properties.getProperty(prefix+"debug_initial_discriminate")!=null) this.debug_initial_discriminate=Boolean.parseBoolean(properties.getProperty(prefix+"debug_initial_discriminate"));
if (properties.getProperty(prefix+"dbg_migrate")!=null) this.dbg_migrate=Boolean.parseBoolean(properties.getProperty(prefix+"dbg_migrate"));
if (properties.getProperty(prefix+"dbg_early_exit")!=null) this.dbg_early_exit=Integer.parseInt(properties.getProperty(prefix+"dbg_early_exit"));
if (properties.getProperty(prefix+"show_first_bg")!=null) this.show_first_bg=Boolean.parseBoolean(properties.getProperty(prefix+"show_first_bg"));
......@@ -3263,6 +3307,8 @@ public class CLTParameters {
"Add squared fat zero to the sum of squared amplitudes, monochrome images");
gd.addMessage ("--- GPU WOI selection ---");
gd.addCheckbox ("Use following WOI for GPU processing (unchecked - ignore WOI dimensions)", this.gpu_woi);
gd.addNumericField("WOI left", this.gpu_woi_tx, 0, 6,"tiles",
"Left WOI margin, in tiles (0..323");
gd.addNumericField("WOI top", this.gpu_woi_ty, 0, 6,"tiles",
......@@ -3271,8 +3317,18 @@ public class CLTParameters {
"WOI width, in tiles (1..324");
gd.addNumericField("WOI height", this.gpu_woi_theight, 0, 6,"tiles",
"WOI height, in tiles (1..242");
gd.addCheckbox ("Select circle/ellipse within the rectanghular WOI ", this.gpu_woi_round);
gd.addCheckbox ("Select circle/ellipse within the rectanghular WOI", this.gpu_woi_round);
gd.addCheckbox ("Debug feature - save calculated ports X,Y to compare with Java-generated", this.gpu_save_ports_xy);
gd.addCheckbox ("Show Java-generated textures from non-overlapping in GPU (will not generate if false)", this.gpu_show_jtextures);
gd.addCheckbox ("Show low-res data for macro (will not generate if false)", this.gpu_show_extra);
gd.addCheckbox ("Accelerate tile processor for the main quad camera", this.gpu_use_main);
gd.addCheckbox ("Accelerate tile processor for the main quad camera in macro mode", this.gpu_use_main_macro);
gd.addCheckbox ("Accelerate tile processor for the main quad camera for field calibration", this.gpu_use_main_adjust);
gd.addCheckbox ("Accelerate tile processor for the aux (lwir) quad camera", this.gpu_use_aux);
gd.addCheckbox ("Accelerate tile processor for the aux (lwir) quad camera in macro mode", this.gpu_use_aux_macro);
gd.addCheckbox ("Accelerate tile processor for the aux (lwir) quad camera for field calibration", this.gpu_use_aux_adjust);
gd.addTab ("LWIR", "parameters for LWIR/EO 8-camera rig");
this.lwir.dialogQuestions(gd);
gd.addTab ("Debug", "Other debug images");
......@@ -4014,12 +4070,23 @@ public class CLTParameters {
this.gpu_fatz = gd.getNextNumber();
this.gpu_fatz_m = gd.getNextNumber();
this.gpu_woi= gd.getNextBoolean();
this.gpu_woi_tx = (int) gd.getNextNumber();
this.gpu_woi_ty = (int) gd.getNextNumber();
this.gpu_woi_twidth = (int) gd.getNextNumber();
this.gpu_woi_theight =(int) gd.getNextNumber();
this.gpu_woi_round= gd.getNextBoolean();
this.gpu_save_ports_xy= gd.getNextBoolean();
this.gpu_show_jtextures= gd.getNextBoolean();
this.gpu_show_extra= gd.getNextBoolean();
this.gpu_use_main= gd.getNextBoolean();
this.gpu_use_main_macro= gd.getNextBoolean();
this.gpu_use_main_adjust= gd.getNextBoolean();
this.gpu_use_aux= gd.getNextBoolean();
this.gpu_use_aux_macro= gd.getNextBoolean();
this.gpu_use_aux_adjust= gd.getNextBoolean();
this.lwir.dialogAnswers(gd);
this.debug_initial_discriminate= gd.getNextBoolean();
......
src/main/java/com/elphel/imagej/correction/Eyesis_Correction.java
View file @ ac185fe7
......@@ -4906,7 +4906,7 @@ private Panel panel1,
importAux();
return;
/* ======================================================================== */
} else if (label.equals("CLT 2*4 images")) {
} else if (label.equals("CLT 2*4 images")) { // never used
DEBUG_LEVEL=MASTER_DEBUG_LEVEL;
EYESIS_CORRECTIONS.setDebug(DEBUG_LEVEL);
getPairImages();
......@@ -5416,7 +5416,7 @@ private Panel panel1,
return true;
}
public boolean getPairImages() {
public boolean getPairImages() { // Never Used
if (QUAD_CLT == null){
QUAD_CLT = new QuadCLT (
QuadCLT.PREFIX,
......@@ -5810,10 +5810,7 @@ private Panel panel1,
if (GPU_QUAD == null) {
try {
GPU_QUAD = GPU_TILE_PROCESSOR. new GpuQuad(
2592,
1936,
164, // final int kernels_hor,
123, // final int kernels_vert,
QUAD_CLT,
4,
3);
} catch (Exception e) {
......@@ -5822,6 +5819,7 @@ private Panel panel1,
e.printStackTrace();
return false;
} //final int debugLevel);
QUAD_CLT.setGPU(GPU_QUAD);
}
// For now keep GPU_QUAD_AUX==null
......@@ -6266,6 +6264,8 @@ private Panel panel1,
public boolean batchLwir() {
long startTime=System.nanoTime();
// load needed sensor and kernels files
if (!prepareRigImages()) return false;
String configPath=getSaveCongigPath();
......@@ -6278,8 +6278,52 @@ private Panel panel1,
/// if (COLOR_PROC_PARAMETERS_AUX == null) {
/// COLOR_PROC_PARAMETERS_AUX = COLOR_PROC_PARAMETERS.clone();
/// }
if (CLT_PARAMETERS.useGPU()) { // only init GPU instances if it is used
if (GPU_TILE_PROCESSOR == null) {
try {
GPU_TILE_PROCESSOR = new GPUTileProcessor(CORRECTION_PARAMETERS.tile_processor_gpu);
} catch (Exception e) {
System.out.println("Failed to initialize GPU class");
// TODO Auto-generated catch block
e.printStackTrace();
return false;
} //final int debugLevel);
}
if (CLT_PARAMETERS.useGPU(false) && (QUAD_CLT != null) && (GPU_QUAD == null)) { // if GPU main is needed
try {
GPU_QUAD = GPU_TILE_PROCESSOR.new GpuQuad(
QUAD_CLT,
4,
3);
} catch (Exception e) {
System.out.println("Failed to initialize GpuQuad class");
// TODO Auto-generated catch block
e.printStackTrace();
return false;
} //final int debugLevel);
QUAD_CLT.setGPU(GPU_QUAD);
}
if (CLT_PARAMETERS.useGPU(true) && (QUAD_CLT_AUX != null) && (GPU_QUAD_AUX == null)) { // if GPU AUX is needed
try {
GPU_QUAD_AUX = GPU_TILE_PROCESSOR. new GpuQuad(//
QUAD_CLT_AUX,
4,
3);
} catch (Exception e) {
System.out.println("Failed to initialize GpuQuad class");
// TODO Auto-generated catch block
e.printStackTrace();
return false;
} //final int debugLevel);
QUAD_CLT_AUX.setGPU(GPU_QUAD_AUX);
}
}
try {
TWO_QUAD_CLT.batchLwirRig(
// GPU_QUAD, // GPUTileProcessor.GpuQuad gpuQuad_main, // may be null if GPU for MAIN is not use3d
// GPU_QUAD_AUX, // GPUTileProcessor.GpuQuad gpuQuad_aux, // may be null if GPU for AUX is not use3d
QUAD_CLT, // QuadCLT quadCLT_main,
QUAD_CLT_AUX, // QuadCLT quadCLT_aux,
CLT_PARAMETERS, // EyesisCorrectionParameters.DCTParameters dct_parameters,
......
src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java
View file @ ac185fe7
......@@ -69,6 +69,7 @@ import java.util.concurrent.atomic.AtomicInteger;
import com.elphel.imagej.tileprocessor.DttRad2;
import com.elphel.imagej.tileprocessor.GeometryCorrection;
import com.elphel.imagej.tileprocessor.ImageDtt;
import com.elphel.imagej.tileprocessor.QuadCLT;
import Jama.Matrix;
import ij.IJ;
......@@ -162,7 +163,6 @@ public class GPUTileProcessor {
private CUfunction GPU_CALC_REVERSE_DISTORTION_kernel = null;
CUmodule module; // to access constants memory
public class TpTask {
public int task; // [0](+1) - generate 4 images, [4..9]+16..+512 - correlation pairs, 2 - generate texture tiles
public float target_disparity;
......@@ -202,78 +202,19 @@ public class GPUTileProcessor {
}
}
flt[indx++] = this.target_disparity;
/*
for (int i = 0; i < NUM_CAMS; i++) { // actually disp_dist will be initialized by the GPU
indx+= 4;
/*
flt[indx++] = disp_dist[i][0];
flt[indx++] = disp_dist[i][1];
flt[indx++] = disp_dist[i][2];
flt[indx++] = disp_dist[i][3];
*/
}
*/
return flt;
}
}
/*
public class CltExtra{ // never used?
public float data_x; // kernel data is relative to this displacement X (0.5 pixel increments)
public float data_y; // kernel data is relative to this displacement Y (0.5 pixel increments)
public float center_x; // actual center X (use to find derivatives)
public float center_y; // actual center X (use to find derivatives)
public float dxc_dx; // add this to data_x per each pixel X-shift relative to the kernel center location
public float dxc_dy; // same per each Y-shift pixel
public float dyc_dx;
public float dyc_dy;
public CltExtra() {}
public CltExtra(
float data_x, // kernel data is relative to this displacement X (0.5 pixel increments)
float data_y, // kernel data is relative to this displacement Y (0.5 pixel increments)
float center_x, // actual center X (use to find derivatives)
float center_y, // actual center X (use to find derivatives)
float dxc_dx, // add this to data_x per each pixel X-shift relative to the kernel center location
float dxc_dy, // same per each Y-shift pixel
float dyc_dx,
float dyc_dy)
{
this.data_x = data_x;
this.data_y = data_y;
this.center_x = center_x;
this.center_y = center_y;
this.dxc_dx = dxc_dx;
this.dxc_dy = dxc_dy;
this.dyc_dx = dyc_dx;
this.dyc_dy = dyc_dy;
}
public CltExtra(float [] data, int indx)
{
this.data_x = data[indx++];
this.data_y = data[indx++];
this.center_x = data[indx++];
this.center_y = data[indx++];
this.dxc_dx = data[indx++];
this.dxc_dy = data[indx++];
this.dyc_dx = data[indx++];
this.dyc_dy = data[indx++];
}
public float [] asFloatArray() {
float [] flt = new float [8];
return asFloatArray(flt, 0);
}
public float [] asFloatArray(float [] flt, int indx) {
flt[indx++] = this.data_x;
flt[indx++] = this.data_y;
flt[indx++] = this.center_x;
flt[indx++] = this.center_y;
flt[indx++] = this.dxc_dx;
flt[indx++] = this.dxc_dy;
flt[indx++] = this.dyc_dx;
flt[indx++] = this.dyc_dy;
return flt;
}
}
*/
private static long getPointerAddress(CUdeviceptr p)
{
......@@ -299,10 +240,6 @@ public class GPUTileProcessor {
"#define NUM_CAMS " + NUM_CAMS+"\n"+
"#define NUM_PAIRS " + NUM_PAIRS+"\n"+
"#define NUM_COLORS " + NUM_COLORS+"\n"+
// "#define IMG_WIDTH " + IMG_WIDTH+"\n"+
// "#define IMG_HEIGHT " + IMG_HEIGHT+"\n"+
// "#define KERNELS_HOR " + KERNELS_HOR+"\n"+
// "#define KERNELS_VERT " + KERNELS_VERT+"\n"+
"#define KERNELS_LSTEP " + KERNELS_LSTEP+"\n"+
"#define THREADS_PER_TILE " + THREADS_PER_TILE+"\n"+
"#define TILES_PER_BLOCK " + TILES_PER_BLOCK+"\n"+
......@@ -331,7 +268,8 @@ public class GPUTileProcessor {
}
public GPUTileProcessor(String cuda_project_directory) throws IOException
public GPUTileProcessor(
String cuda_project_directory) throws IOException
{
// From code by Marco Hutter - http://www.jcuda.org
......@@ -568,6 +506,7 @@ public class GPUTileProcessor {
}
public class GpuQuad{ // quad camera description
public final QuadCLT quadCLT;
public final int img_width;
public final int img_height;
public final int kernels_hor; // int kernels_hor,
......@@ -620,21 +559,25 @@ public class GPUTileProcessor {
private int num_task_tiles;
private int num_corr_tiles;
private int num_texture_tiles;
private boolean geometry_correction_set = false;
private boolean geometry_correction_vector_set = false;
public GpuQuad(
final int img_width,
final int img_height,
final int kernels_hor,
final int kernels_vert,
final QuadCLT quadCLT,
// final int img_width,
// final int img_height,
// final int kernels_hor,
// final int kernels_vert,
final int num_cams,
final int num_colors
) {
this.img_width = img_width;
this.img_height = img_height;
this.quadCLT = quadCLT;
int [] wh = quadCLT.getGeometryCorrection().getSensorWH();
this.img_width = wh[0]; // img_width;
this.img_height = wh[1]; // img_height;
this.num_cams = num_cams;
this.num_colors = num_colors; // maybe should always be 3?
this.kernels_hor = kernels_hor;
this.kernels_vert = kernels_vert;
this.kernels_hor = quadCLT.getCLTKernels()[0][0][0].length; // kernels_hor;
this.kernels_vert = quadCLT.getCLTKernels()[0][0].length; // kernels_vert;
this.kern_tiles = kernels_hor * kernels_vert * num_colors;
this.kern_size = kern_tiles * 4 * 64;
......@@ -797,10 +740,30 @@ public class GPUTileProcessor {
texture_stride_rgba = (int)(device_stride[0] / Sizeof.FLOAT);
}
public void resetGeometryCorrection() {
geometry_correction_set = false;
geometry_correction_vector_set = false;
}
public void resetGeometryCorrectionVector() {
geometry_correction_vector_set = false;
}
public int getImageWidth() {return this.img_width;}
public int getImageHeight() {return this.img_height;}
public int getDttSize() {return DTT_SIZE;}
public int getNumCams() {return NUM_CAMS;}
public void setGeometryCorrection() { // will reset geometry_correction_set when running GPU kernel
// if (geometry_correction_set) return;
setGeometryCorrection(
quadCLT.getGeometryCorrection(),
false);
}
public void setGeometryCorrectionVector() { // will reset geometry_correction_vector_set when running GPU kernel
// if (geometry_correction_vector_set) return;
setExtrinsicsVector(
quadCLT.getGeometryCorrection().getCorrVector());
}
public void setGeometryCorrection(GeometryCorrection gc,
boolean use_java_rByRDist) { // false - use newer GPU execCalcReverseDistortions
......@@ -816,7 +779,10 @@ 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
}
/**
* Copy extrinsic correction vector to the GPU memory
* @param cv correction vector
*/
public void setExtrinsicsVector(GeometryCorrection.CorrVector cv) {
double [] dcv = cv.toFullRollArray();
float [] fcv = new float [dcv.length];
......@@ -826,7 +792,11 @@ public class GPUTileProcessor {
cuMemcpyHtoD(gpu_correction_vector, Pointer.to(fcv), fcv.length * Sizeof.FLOAT);
}
/**
* Copy array of CPU-prepared tasks to the GPU memory
* @param tile_tasks array of TpTask prepared by the CPU (before geometry correction is appllied)
* @param use_aux Use second (aux) camera
*/
public void setTasks(TpTask [] tile_tasks, boolean use_aux) // while is it in class member? - just to be able to free
{
num_task_tiles = tile_tasks.length;
......@@ -847,7 +817,7 @@ public class GPUTileProcessor {
cuMemcpyHtoD(gpu_corr_indices, Pointer.to(fcorr_indices), num_corr_tiles * Sizeof.FLOAT);
}
public void setTextureIndices(int [] texture_indices)
public void setTextureIndices(int [] texture_indices) // never used
{
num_texture_tiles = texture_indices.length;
float [] ftexture_indices = new float [texture_indices.length];
......@@ -868,6 +838,7 @@ public class GPUTileProcessor {
for (int i = 0; i < num_tiles; i++) {
texture_indices[i] = Float.floatToIntBits(ftexture_indices[i]);
}
num_texture_tiles = num_tiles; //TODO: was it missing?
return texture_indices;
}
......@@ -880,6 +851,25 @@ public class GPUTileProcessor {
cuMemcpyHtoD(gpu_kernel_offsets_h[ncam], Pointer.to(kernel_offsets), kern_tiles * CLTEXTRA_SIZE * Sizeof.FLOAT);
}
/**
* Set CLT kernels in GPU memory reading them using quadCLT instance
* @param clt_kernels - [num_cameras==4][num_colors][tilesY][tilesX][quadrants==4][elements==8*8==64]
* @param force re-set kernels even if they were already set
*/
public void setConvolutionKernels(
boolean force)
{
if (kernels_set && ! force) {
return;
}
setConvolutionKernels(quadCLT.getCLTKernels(), true);
}
/**
* Set CLT kernels in GPU memory
* @param clt_kernels - [num_cameras==4][num_colors][tilesY][tilesX][quadrants==4][elements==8*8==64]
* @param force re-set kernels even if they were already set
*/
public void setConvolutionKernels(
double [][][][][][] clt_kernels,
boolean force)
......@@ -950,7 +940,28 @@ public class GPUTileProcessor {
copyH2D.Height = img_height; // /4;
cuMemcpy2D(copyH2D);
}
// combines 3 bayer channels into one and transfers to GPU memory
/**
* Copy a set of images to the GPU (if they are new)
* @param force set even if there is no new Bayer data available
*/
public void setBayerImages(
boolean force) {
if (!force && bayer_set && !quadCLT.hasNewImageData()) {
return;
}
double [][][] bayer_data = quadCLT.getImageData(); // resets hasNewImageData()
setBayerImages(
bayer_data,
true);
}
/**
* Copy a set of Bayer images to the GPU
* @param bayer_data per camera a set of split-color images [4][3][w*h], may be [4][1][w*h]
* @param force copy regardless of if it was copied before
*/
// combines 3 bayer channels into one and transfers to GPU memory
public void setBayerImages(
double [][][] bayer_data,
boolean force) {
......@@ -960,7 +971,10 @@ public class GPUTileProcessor {
float [] fbayer = new float [bayer_data[0][0].length];
for (int ncam = 0; ncam < bayer_data.length; ncam++) {
for (int i = 0; i < bayer_data[ncam][0].length; i++) {
fbayer[i] = (float) (bayer_data[ncam][0][i] + bayer_data[ncam][1][i] + bayer_data[ncam][2][i]);
fbayer[i] = (float) (bayer_data[ncam][0][i]); // + bayer_data[ncam][1][i] + bayer_data[ncam][2][i]);
for (int j = 1; j < bayer_data[ncam].length; j++) {
fbayer[i] += (float) (bayer_data[ncam][j][i]);
}
}
setBayerImage(
fbayer, // float [] bayer_image,
......@@ -971,19 +985,26 @@ public class GPUTileProcessor {
// prepare tasks for full frame, same dispaity.
// need to run setTasks(TpTask [] tile_tasks, boolean use_aux) to format/transfer to GPU memory
/**
* CPU-side preparation of per-tile tasks , will need GPU kernel run to finalize
* @param woi - rectangle WOI in tiles (x,y,w,h) to process, normally full window (may provide null for full frame, faster)
* @param round_woi - debug feature - use round (elliptic) WOI, ignored for woi == null
* @param target_disparity - same disparity for all tiles
* @param disparity_corr - add to all disparities (minus disparity at infinity)
* @param out_image from which camera channels to generate image (currently 0/1)
* @param corr_mask from which correlation pairs to generate (0x3f - all 6)
* @param threadsMax - historic
* @param debugLevel - not yet used
* @return per-tile array of TpTask instances (need to be finalized by the GPU kernel)
*/
public TpTask [] setFullFrameImages(
boolean calc_offsets, // old way, now not needed with GPU calculation
Rectangle woi,
boolean round_woi,
float target_disparity, // apply same disparity to all tiles
float disparity_corr, // add to disparity (at infinity)
int out_image, // from which tiles to generate image (currently 0/1)
int corr_mask, // which correlation pairs to generate (maybe later - reduce size from 15x15)
boolean use_master,
boolean use_aux,
final GeometryCorrection geometryCorrection_main,
final GeometryCorrection geometryCorrection_aux, // if null, will only calculate offsets fro the main camera
final double [][][] ers_delay, // if not null - fill with tile center acquisition delay
final int threadsMax, // maximal number of threads to launch
// final int threadsMax, // maximal number of threads to launch
final int debugLevel) {
int tilesX = img_width / DTT_SIZE;
int tilesY = img_height / DTT_SIZE;
......@@ -998,22 +1019,92 @@ public class GPUTileProcessor {
corr_masks[i] = corr_mask; // 0x3f; // all 6 correlations
}
return setFullFrameImages(
calc_offsets, // boolean calc_offsets, // old way, now not needed with GPU calculation
woi, // Rectangle woi,
round_woi, // boolean round_woi,
target_disparities, // should be tilesX*tilesY long
disparity_corr, // float disparity_corr, // add to disparity (at infinity)
out_images, // int [] out_images, // from which tiles to generate image (currently 0/1)
corr_masks, // int [] corr_mask, // which correlation pairs to generate (maybe later - reduce size from 15x15)
use_master,
use_aux,
geometryCorrection_main,
geometryCorrection_aux, // if null, will only calculate offsets fro the main camera
ers_delay, // if not null - fill with tile center acquisition delay
threadsMax, // maximal number of threads to launch
// threadsMax, // maximal number of threads to launch
debugLevel);
}
/**
* CPU-side preparation of per-tile tasks, will need GPU kernel run to finalize
* @param woi - rectangle woi in tiles (x,y,w,h) to process, normally full window (may provide null for full frame, faster)
* @param round_woi - debug feature - use round (elliptic) WOI, ignored for woi == null
* @param target_disparities 1D array [tilesX*tilesY] of target disparities
* @param disparity_corr - add to all disparities (minus disparity at infinity)
* @param out_images per-tile array [tilesX*tilesY] from which tiles to generate image (currently 0/1)
* @param corr_mask per-tile array [tilesX*tilesY] from which correlation pairs to generate (0x3f - all 6)
* @param threadsMax - historic
* @param debugLevel - not yet used
* @return per-tile array of TpTask instances (need to be finalized by the GPU kernel)
*/
public TpTask [] setFullFrameImages(
Rectangle woi, // or null
boolean round_woi,
float [] target_disparities, // should be tilesX*tilesY long
float disparity_corr, // add to disparity (at infinity)
int [] out_images, // from which tiles to generate image (currently 0/1)
int [] corr_mask, // which correlation pairs to generate (maybe later - reduce size from 15x15)
final int debugLevel)
{
int tilesX = img_width / DTT_SIZE;
int tilesY = img_height / DTT_SIZE;
boolean [] mask = new boolean[tilesX*tilesY];
int num_tiles = 0;
if (woi == null) {
round_woi = false;
woi = new Rectangle(0,0,tilesX,tilesY);
num_tiles = tilesX*tilesY;
for (int i = 0; i < num_tiles; i++) {
mask[i] = true;
}
} else { // mostly for debug
if (woi.x < 0) woi.x = 0;
if (woi.y < 0) woi.y = 0;
if (woi.x > (tilesX - 2)) woi.x = tilesX - 2;
if (woi.y > (tilesY - 2)) woi.y = tilesY - 2;
if ((woi.x + woi.width) > tilesX) woi.width = tilesX - woi.x;
if ((woi.y + woi.height) > tilesY) woi.height = tilesY - woi.y;
double rx = 0.5*woi.width;
double ry = 0.5*woi.height;
double xc = woi.x + rx - 0.5;
double yc = woi.y + ry - 0.5;
for (int ty = woi.y; ty < (woi.y +woi.height); ty++) {
double ry2 = (ty - yc) / ry;
ry2*=ry2;
for (int tx = woi.x; tx < (woi.x +woi.width); tx++) {
double rx2 = (tx - xc) / rx;
rx2*=rx2;
if (!round_woi || ((rx2+ry2) < 1.0)) {
mask[ty * tilesX + tx] = true;
num_tiles ++;
}
}
}
}
TpTask [] tp_tasks = new TpTask[num_tiles];
int indx = 0;
for (int ty = 0; ty < tilesY; ty++) {
for (int tx = 0; tx < tilesX; tx++) {
int nt = ty * tilesX + tx;
if (mask[nt]) {
// Only generate for non-empty tasks, use 1 empty empty as a terminator?
tp_tasks[indx] = new TpTask(tx,ty, target_disparities[indx]+disparity_corr,
((out_images[nt] & 0x0f) << 0) |
((corr_mask [nt] & 0x3f) << 4)
); // task == 1 for now
indx++;
}
}
}
return tp_tasks;
}
public TpTask [] setFullFrameImages_dbg( // keeping debug features
boolean calc_offsets, // old way, now not needed with GPU calculation
Rectangle woi, // or null
boolean round_woi,
......@@ -1119,7 +1210,8 @@ public class GPUTileProcessor {
}
return tp_tasks;
}
/**
* Prepare contents pointers for calculation of the correlation pairs
* @param tp_tasks array of tasks that contain masks of the required pairs
......@@ -1178,14 +1270,18 @@ public class GPUTileProcessor {
return iarr;
}
// All data is already copied to GPU memory
/**
* Calculate rotation matrices and their derivatives
* Needed after extrinsics changed
*/
public void execRotDerivs() {
if (geometry_correction_vector_set) return;
if (GPU_ROT_DERIV_kernel == null)
{
IJ.showMessage("Error", "No GPU kernel: GPU_ROT_DERIV_kernel");
return;
}
setGeometryCorrectionVector();
// kernel parameters: pointer to pointers
int [] GridFullWarps = {num_cams, 1, 1}; // round up
int [] ThreadsFullWarps = {3, 3, 3};
......@@ -1202,14 +1298,21 @@ 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
geometry_correction_vector_set = true;
}
/**
* Calculate reverse radial distortion table from gpu_geometry_correction
* Needed once during initialization of GeometryCorrection
*/
public void execCalcReverseDistortions() {
if (geometry_correction_set) return;
if (GPU_CALC_REVERSE_DISTORTION_kernel == null)
{
IJ.showMessage("Error", "No GPU kernel: GPU_CALC_REVERSE_DISTORTION_kernel");
return;
}
setGeometryCorrection();
// kernel parameters: pointer to pointers
int [] GridFullWarps = {num_cams, 1, 1}; // round up
int [] ThreadsFullWarps = {3, 3, 3};
......@@ -1224,9 +1327,15 @@ 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
geometry_correction_set = true;
}
/**
* Calculate tiles offsets (before each direct conversion run)
*/
public void execSetTilesOffsets() {
execCalcReverseDistortions(); // will check if it is needed first
execRotDerivs(); // will check if it is needed first
if (GPU_SET_TILES_OFFSETS_kernel == null)
{
IJ.showMessage("Error", "No GPU kernel: GPU_SET_TILES_OFFSETS_kernel");
......@@ -1251,12 +1360,17 @@ public class GPUTileProcessor {
cuCtxSynchronize(); // remove later
}
/**
* Direct CLT conversion and aberration correction
*/
public void execConvertDirect() {
if (GPU_CONVERT_DIRECT_kernel == null)
{
IJ.showMessage("Error", "No GPU kernel: GPU_CONVERT_DIRECT_kernel");
return;
}
setConvolutionKernels(false); // set kernels if they are not set already
setBayerImages(false); // set Bayer images if this.quadCLT instance has new ones
// kernel parameters: pointer to pointers
int tilesX = img_width / DTT_SIZE;
int [] GridFullWarps = {1, 1, 1};
......@@ -1291,7 +1405,10 @@ public class GPUTileProcessor {
}
/**
* Generate corrected image(s) from the CLT representation created by execConvertDirect()
* @param is_mono monochrome image (false - RGB)
*/
public void execImcltRbgAll(
boolean is_mono
) {
......@@ -1324,6 +1441,12 @@ public class GPUTileProcessor {
cuCtxSynchronize();
}
/**
* Generate 2D phase correlations from the CLT representation
* @param scales R,G,B weights
* @param fat_zero - absolute fat zero - add to correlations before normalization
* @param corr_radius - correlation result size (maximal 7 for 15x15)
*/
public void execCorr2D(
double [] scales,
double fat_zero,
......@@ -1366,7 +1489,18 @@ public class GPUTileProcessor {
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize();
}
/**
* Generate combined (overlapping) texture
* @param color_weights - [3] (RGB) or [1] (mono) color weights for matching
* @param is_lwir ignore shot noise normalization
* @param min_shot - minimal value to use sqrt() for shot noise normalization (default - 10)
* @param scale_shot scale shot noise (3.0)
* @param diff_sigma pixel value/pixel change (1.5) - normalize channel values difference by the offsets of the cameras
* @param diff_threshold - never used?
* @param min_agree minimal number of channels to agree on a point (real number to work with fuzzy averages
* @param dust_remove do not reduce average weight when only one image differs much from the average
*/
public void execRBGA(
double [] color_weights,
boolean is_lwir,
......@@ -1376,8 +1510,8 @@ public class GPUTileProcessor {
double diff_threshold, // pixel value/pixel change
double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
boolean dust_remove) {
if (GPU_RBGA_kernel == null)
{
execCalcReverseDistortions(); // will check if it is needed first
if (GPU_RBGA_kernel == null) {
IJ.showMessage("Error", "No GPU kernel: GPU_TEXTURES_kernel");
return;
}
......@@ -1387,45 +1521,45 @@ public class GPUTileProcessor {
fcolor_weights[0] = (float) color_weights[0];
fcolor_weights[1] = (num_colors >1)?((float) color_weights[1]):0.0f;
fcolor_weights[2] = (num_colors >2)?((float) color_weights[2]):0.0f;
cuMemcpyHtoD(gpu_color_weights, Pointer.to(fcolor_weights), fcolor_weights.length * Sizeof.FLOAT);
cuMemcpyHtoD(gpu_color_weights, Pointer.to(fcolor_weights), fcolor_weights.length * Sizeof.FLOAT);
float [] generate_RBGA_params = {
float [] generate_RBGA_params = {
(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)
};
cuMemcpyHtoD(gpu_generate_RBGA_params, Pointer.to(generate_RBGA_params), generate_RBGA_params.length * Sizeof.FLOAT);
cuMemcpyHtoD(gpu_generate_RBGA_params, Pointer.to(generate_RBGA_params), generate_RBGA_params.length * Sizeof.FLOAT);
int iis_lwir = (is_lwir)? 1:0;
int idust_remove = (dust_remove)? 1 : 0;
// uses dynamic parallelization, top kernel is a single-thread one
int [] GridFullWarps = {1, 1, 1};
int [] GridFullWarps = {1, 1, 1};
int [] ThreadsFullWarps = {1, 1, 1};
Pointer kernelParameters = Pointer.to(
Pointer.to(gpu_tasks), // struct tp_task * gpu_tasks,
Pointer.to(new int[] { num_task_tiles }), // int num_tiles, // number of tiles in task list
// declare arrays in device code?
Pointer.to(gpu_texture_indices_ovlp), // int * gpu_texture_indices_ovlp,// packed tile + bits (now only (1 << 7)
Pointer.to(gpu_tasks), // struct tp_task * gpu_tasks,
Pointer.to(new int[] { num_task_tiles }), // int num_tiles, // number of tiles in task list
// declare arrays in device code?
Pointer.to(gpu_texture_indices_ovlp), // int * gpu_texture_indices_ovlp,// packed tile + bits (now only (1 << 7)
Pointer.to(gpu_num_texture_ovlp), // int * num_texture_tiles, // number of texture tiles to process (8 elements)
Pointer.to(gpu_woi), // int * woi, // x,y,width,height of the woi
// set smaller for LWIR - it is used to reduce work aread
Pointer.to(new int[] {img_width / DTT_SIZE}), // int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
Pointer.to(new int[] {img_height / DTT_SIZE}), // int height); // <= TILESY, use for faster processing of LWIR images
// Parameters for the texture generation
// Parameters for the texture generation
Pointer.to(gpu_clt), // float ** gpu_clt, // [num_cams] ->[TILESY][TILESX][num_colors][DTT_SIZE*DTT_SIZE]
Pointer.to(gpu_geometry_correction), // struct gc * gpu_geometry_correction,
Pointer.to(new int[] {num_colors}), // int colors, // number of colors (3/1)
Pointer.to(new int[] {iis_lwir}), // int is_lwir, // do not perform shot correction
Pointer.to(gpu_generate_RBGA_params), // float generate_RBGA_params[5],
Pointer.to(gpu_color_weights), // float weights[3], // scale for R,B,G
Pointer.to(new int[] { idust_remove }), // int dust_remove, // Do not reduce average weight when only one image differes much from the average
Pointer.to(new int[] {0}), // int keep_weights, // return channel weights after A in RGBA
Pointer.to(new int[] { texture_stride_rgba }), // const size_t texture_rbga_stride, // in floats
Pointer.to(gpu_textures_rgba)); // float * gpu_texture_tiles) // (number of colors +1 + ?)*16*16 rgba texture tiles
Pointer.to(gpu_geometry_correction), // struct gc * gpu_geometry_correction,
Pointer.to(new int[] {num_colors}), // int colors, // number of colors (3/1)
Pointer.to(new int[] {iis_lwir}), // int is_lwir, // do not perform shot correction
Pointer.to(gpu_generate_RBGA_params), // float generate_RBGA_params[5],
Pointer.to(gpu_color_weights), // float weights[3], // scale for R,B,G
Pointer.to(new int[] { idust_remove }), // int dust_remove, // Do not reduce average weight when only one image differes much from the average
Pointer.to(new int[] {0}), // int keep_weights, // return channel weights after A in RGBA
Pointer.to(new int[] { texture_stride_rgba }), // const size_t texture_rbga_stride, // in floats
Pointer.to(gpu_textures_rgba)); // float * gpu_texture_tiles) // (number of colors +1 + ?)*16*16 rgba texture tiles
cuCtxSynchronize();
// Call the kernel function
......@@ -1436,8 +1570,20 @@ public class GPUTileProcessor {
kernelParameters, null); // Kernel- and extra parameters
cuCtxSynchronize();
}
public void execTextures(
/**
* Generate non-overlapping (16x16) texture tiles
* @param color_weights - [3] (RGB) or [1] (mono) color weights for matching
* @param is_lwir ignore shot noise normalization
* @param min_shot - minimal value to use sqrt() for shot noise normalization (default - 10)
* @param scale_shot scale shot noise (3.0)
* @param diff_sigma pixel value/pixel change (1.5) - normalize channel values difference by the offsets of the cameras
* @param diff_threshold - never used?
* @param min_agree minimal number of channels to agree on a point (real number to work with fuzzy averages
* @param dust_remove do not reduce average weight when only one image differs much from the average
* @param calc_textures calculate textures (false for macro-only output)
* @param calc_extra calculate extra output - low-res for macro
*/
public void execTextures(
double [] color_weights,
boolean is_lwir,
double min_shot, // 10.0
......@@ -1445,7 +1591,10 @@ public class GPUTileProcessor {
double diff_sigma, // pixel value/pixel change
double diff_threshold, // pixel value/pixel change
double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
boolean dust_remove) {
boolean dust_remove, // Do not reduce average weight when only one image differs much from the average
boolean calc_textures,
boolean calc_extra)
{
if (GPU_TEXTURES_kernel == null)
{
IJ.showMessage("Error", "No GPU kernel: GPU_TEXTURES_kernel");
......@@ -1458,41 +1607,42 @@ public class GPUTileProcessor {
fcolor_weights[0] = (float) color_weights[0];
fcolor_weights[1] = (num_colors >1)?((float) color_weights[1]):0.0f;
fcolor_weights[2] = (num_colors >2)?((float) color_weights[2]):0.0f;
cuMemcpyHtoD(gpu_color_weights, Pointer.to(fcolor_weights), fcolor_weights.length * Sizeof.FLOAT);
cuMemcpyHtoD(gpu_color_weights, Pointer.to(fcolor_weights), fcolor_weights.length * Sizeof.FLOAT);
float [] generate_RBGA_params = {
float [] generate_RBGA_params = {
(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)
};
cuMemcpyHtoD(gpu_generate_RBGA_params, Pointer.to(generate_RBGA_params), generate_RBGA_params.length * Sizeof.FLOAT);
cuMemcpyHtoD(gpu_generate_RBGA_params, Pointer.to(generate_RBGA_params), generate_RBGA_params.length * Sizeof.FLOAT);
int iis_lwir = (is_lwir)? 1:0;
int idust_remove = (dust_remove)? 1 : 0;
int [] GridFullWarps = {1, 1, 1};
int [] GridFullWarps = {1, 1, 1};
int [] ThreadsFullWarps = {1, 1, 1};
// CUdeviceptr gpu_diff_rgb_combo_local = calc_extra ? gpu_diff_rgb_combo : null;
Pointer kernelParameters = Pointer.to(
Pointer.to(gpu_tasks), // struct tp_task * gpu_tasks,
Pointer.to(new int[] { num_task_tiles }), // int num_tiles, // number of tiles in task list
// declare arrays in device code?
Pointer.to(gpu_texture_indices), // int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
Pointer.to(gpu_tasks), // struct tp_task * gpu_tasks,
Pointer.to(new int[] { num_task_tiles }), // int num_tiles, // number of tiles in task list
// declare arrays in device code?
Pointer.to(gpu_texture_indices), // int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
Pointer.to(gpu_texture_indices_len),
// Parameters for the texture generation
// Parameters for the texture generation
Pointer.to(gpu_clt), // float ** gpu_clt, // [num_cams] ->[TILESY][TILESX][num_colors][DTT_SIZE*DTT_SIZE]
Pointer.to(gpu_geometry_correction), // struct gc * gpu_geometry_correction,
Pointer.to(gpu_geometry_correction), // struct gc * gpu_geometry_correction,
Pointer.to(new int[] { num_colors }),
Pointer.to(new int[] { iis_lwir }),
Pointer.to(gpu_generate_RBGA_params), // float generate_RBGA_params[5],
Pointer.to(gpu_color_weights), // float weights[3], // scale for R,B,G
Pointer.to(gpu_generate_RBGA_params), // float generate_RBGA_params[5],
Pointer.to(gpu_color_weights), // float weights[3], // scale for R,B,G
Pointer.to(new int[] { idust_remove }),
Pointer.to(new int[] {texture_stride}), // can be a null pointer - will not be used! float * gpu_texture_rbg, // (number of colors +1 + ?)*16*16 rgba texture tiles
Pointer.to(new int[] {calc_textures? texture_stride : 0}),
Pointer.to(gpu_textures),
Pointer.to(gpu_diff_rgb_combo), // ); // float * gpu_diff_rgb_combo); // diff[num_cams], R[num_cams], B[num_cams],G[num_cams]
Pointer.to(new int[] { tilesX }));
calc_extra ? Pointer.to(gpu_diff_rgb_combo) : Pointer.to(new int[] { 0 }),
Pointer.to(new int[] { tilesX }));
cuCtxSynchronize();
// Call the kernel function
cuLaunchKernel(GPU_TEXTURES_kernel,
......@@ -1504,29 +1654,29 @@ public class GPUTileProcessor {
}
public float [][] getCorr2D(int corr_rad){
int corr_size = (2 * corr_rad + 1) * (2 * corr_rad + 1);
float [] cpu_corrs = new float [ num_corr_tiles * corr_size];
CUDA_MEMCPY2D copyD2H = new CUDA_MEMCPY2D();
copyD2H.srcMemoryType = CUmemorytype.CU_MEMORYTYPE_DEVICE;
copyD2H.srcDevice = gpu_corrs;
copyD2H.srcPitch = corr_stride * Sizeof.FLOAT;
int corr_size = (2 * corr_rad + 1) * (2 * corr_rad + 1);
float [] cpu_corrs = new float [ num_corr_tiles * corr_size];
CUDA_MEMCPY2D copyD2H = new CUDA_MEMCPY2D();
copyD2H.srcMemoryType = CUmemorytype.CU_MEMORYTYPE_DEVICE;
copyD2H.srcDevice = gpu_corrs;
copyD2H.srcPitch = corr_stride * Sizeof.FLOAT;
copyD2H.dstMemoryType = CUmemorytype.CU_MEMORYTYPE_HOST;
copyD2H.dstHost = Pointer.to(cpu_corrs);
copyD2H.dstPitch = corr_size * Sizeof.FLOAT;
copyD2H.dstMemoryType = CUmemorytype.CU_MEMORYTYPE_HOST;
copyD2H.dstHost = Pointer.to(cpu_corrs);
copyD2H.dstPitch = corr_size * Sizeof.FLOAT;
copyD2H.WidthInBytes = corr_size * Sizeof.FLOAT;
copyD2H.Height = num_corr_tiles;
copyD2H.WidthInBytes = corr_size * Sizeof.FLOAT;
copyD2H.Height = num_corr_tiles;
cuMemcpy2D(copyD2H); // run copy
cuMemcpy2D(copyD2H); // run copy
float [][] corrs = new float [num_corr_tiles][ corr_size];
for (int ncorr = 0; ncorr < num_corr_tiles; ncorr++) {
System.arraycopy(cpu_corrs, ncorr*corr_size, corrs[ncorr], 0, corr_size);
}
return corrs;
float [][] corrs = new float [num_corr_tiles][ corr_size];
for (int ncorr = 0; ncorr < num_corr_tiles; ncorr++) {
System.arraycopy(cpu_corrs, ncorr*corr_size, corrs[ncorr], 0, corr_size);
}
return corrs;
}
public int [] getCorrIndices() {
float [] fnum_corrs = new float[1];
cuMemcpyDtoH(Pointer.to(fnum_corrs), gpu_num_corr_tiles, 1 * Sizeof.FLOAT);
......@@ -1542,7 +1692,12 @@ public class GPUTileProcessor {
}
// read extra data for macro generation: 4 DIFFs, 4 of R, 4 of B, 4 of G
//
/**
* Read extra data for macro generation: 4 DIFFs, 4 of R, 4 of B, 4 of G
* Available after texture
* @return [ num_cams*(num_colors+1)][tilesX*tilesY] array for macro generation
*/
public float [][] getExtra(){
int [] texture_indices = getTextureIndices();
int num_tile_extra = num_cams*(num_colors+1);
......@@ -1554,9 +1709,6 @@ public class GPUTileProcessor {
for (int i = 0; i < texture_indices.length; i++) {
if (((texture_indices[i] >> CORR_TEXTURE_BIT) & 1) != 0) {
int ntile = texture_indices[i] >> CORR_NTILE_SHIFT;
// if (ntile == 22507) {
// System.out.println("i="+i+", ntile="+ntile);
// }
for (int l = 0; l < num_tile_extra; l++) {
extra[l][ntile] = diff_rgb_combo[i * num_tile_extra + l];
}
......@@ -1839,7 +1991,8 @@ public class GPUTileProcessor {
}
public void setLpfRbg(
float [][] lpf_rbg) // 4 64-el. arrays: r,b,g,m
float [][] lpf_rbg, // 4 64-el. arrays: r,b,g,m
boolean debug)
{
int l = lpf_rbg[0].length; // 64
......@@ -1854,24 +2007,26 @@ public class GPUTileProcessor {
long constantMemorySizeArray[] = { 0 };
cuModuleGetGlobal(constantMemoryPointer, constantMemorySizeArray, module, "lpf_data");
int constantMemorySize = (int)constantMemorySizeArray[0];
System.out.println("constantMemoryPointer: " + constantMemoryPointer);
System.out.println("constantMemorySize: " + constantMemorySize);
if (debug) System.out.println("constantMemoryPointer: " + constantMemoryPointer);
if (debug) System.out.println("constantMemorySize: " + constantMemorySize);
cuMemcpyHtoD(constantMemoryPointer, Pointer.to(lpf_flat), constantMemorySize);
System.out.println();
if (debug) System.out.println();
}
public void setLpfCorr(
String const_name, // "lpf_corr"
float [] lpf_flat)
String const_name, // "lpf_corr"
float [] lpf_flat,
boolean debug)
{
CUdeviceptr constantMemoryPointer = new CUdeviceptr();
long constantMemorySizeArray[] = { 0 };
cuModuleGetGlobal(constantMemoryPointer, constantMemorySizeArray, module, const_name);
int constantMemorySize = (int)constantMemorySizeArray[0];
System.out.println("constantMemoryPointer: " + constantMemoryPointer);
System.out.println("constantMemorySize: " + constantMemorySize);
if (debug) System.out.println("constantMemoryPointer: " + constantMemoryPointer);
if (debug) System.out.println("constantMemorySize: " + constantMemorySize);
cuMemcpyHtoD(constantMemoryPointer, Pointer.to(lpf_flat), constantMemorySize);
System.out.println();
if (debug) System.out.println();
}
public float [] floatSetCltLpfFd(
......
src/main/java/com/elphel/imagej/tileprocessor/AlignmentCorrection.java
View file @ ac185fe7
......@@ -48,7 +48,8 @@ public class AlignmentCorrection {
static final int INDEX_14_DISPARITY = 0;
static final double DISP_SCALE = 2.0;
QuadCLT qc;
QuadCLT qc_gpu; // maybe will be used later?
QuadCLTCPU qc;
public class Sample{
public int series;
......@@ -265,9 +266,12 @@ public class AlignmentCorrection {
}
//System.arraycopy(dpixels, (tileY*width+tileX)*dct_size + i*width, tile_in, i*n2, n2);
AlignmentCorrection (QuadCLT qc){
AlignmentCorrection (QuadCLTCPU qc){
this.qc = qc;
}
AlignmentCorrection (QuadCLT qc){
this.qc_gpu = qc;
}
public double [][][] infinityCorrection(
final boolean use_poly,
......
src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java
View file @ ac185fe7
......@@ -2364,7 +2364,6 @@ public class ImageDtt {
final double [][][][] clt_corr_combo, // [type][tilesY][tilesX][(2*transform_size-1)*(2*transform_size-1)] // if null - will not calculate
// [type][tilesY][tilesX] should be set by caller
// types: 0 - selected correlation (product+offset), 1 - sum
final double [][][][][] clt_corr_partial,// [tilesY][tilesX][quad]color][(2*transform_size-1)*(2*transform_size-1)] // if null - will not calculate
// [tilesY][tilesX] should be set by caller
// When clt_mismatch is non-zero, no far objects extraction will be attempted
......@@ -2374,7 +2373,6 @@ public class ImageDtt {
final double [][] disparity_map, // [8][tilesY][tilesX], only [6][] is needed on input or null - do not calculate
// last 2 - contrast, avg/ "geometric average)
final double [][][][] texture_tiles, // [tilesY][tilesX]["RGBA".length()][]; null - will skip images combining
final int width,
final double corr_fat_zero, // add to denominator to modify phase correlation (same units as data1, data2). <0 - pure sum
final boolean corr_sym,
......@@ -3230,7 +3228,7 @@ public class ImageDtt {
disparity_map[DISPARITY_INDEX_POLY] [tIndex] = lma_disparity_strength[0];
// if enabled overwrite - replace DISPARITY_INDEX_CM and DISPARITY_STRENGTH_INDEX
if (imgdtt_params.mix_corr_poly) {
if (imgdtt_params.mix_corr_poly) { //true
disp_str[0] = lma_disparity_strength[0];
disp_str[1] = lma_disparity_strength[1];
disparity_map[DISPARITY_INDEX_CM] [tIndex] = disp_str[0];
......
src/main/java/com/elphel/imagej/tileprocessor/ImageDttParameters.java
View file @ ac185fe7
......@@ -30,14 +30,14 @@ import com.elphel.imagej.common.GenericJTabbedDialog;
public class ImageDttParameters {
public boolean corr_mode_debug = true;
public boolean mix_corr_poly = true;
public double min_poly_strength = 0.2;
public double min_poly_strength = 0.2; /// 0.1
public double max_poly_hwidth = 2.5; // Maximal polynomial approximation half-width (in both directions)
public double poly_corr_scale = 2.0; // Shift value if correlation maximum is wide in X than in Y to detect near objects (negative - far ones)
public double poly_corr_scale = 2.0; /// 0.0 // Shift value if correlation maximum is wide in X than in Y to detect near objects (negative - far ones)
public double poly_pwr = 1.0;
public double poly_vasw_pwr = 2.0; // raise value to this power and apply as weight (0 - disable)
public double corr_magic_scale_cm = 1.0; //0.85; // reported correlation offset vs. actual one (not yet understood)
public double corr_magic_scale_poly = 1.0; // 0.95; // reported correlation offset vs. actual one (not yet understood)
public double corr_magic_scale_cm = 1.0; /// 1.0 //0.85; // reported correlation offset vs. actual one (not yet understood)
public double corr_magic_scale_poly = 1.0; /// 1.0 // 0.95; // reported correlation offset vs. actual one (not yet understood)
public int ortho_height = 7; // height of non-zero weights for hor/vert correlation to compensate borders
public double ortho_eff_height = 2.58; // effective correlation stripe height to match strengths
......@@ -50,7 +50,7 @@ public class ImageDttParameters {
private double enhortho_scale = 0.0; // 0.2; // multiply center correlation pixels (inside enhortho_width)
private double enhortho_scale_aux = 0.0; // 0.2; // multiply center correlation pixels (inside enhortho_width)
public boolean ly_poly = false; // Use polynomial when measuring mismatch (false - use center of mass)
public double ly_crazy_poly = 1.0; // Maximal allowed mismatch difference calculated as polynomial maximum
public double ly_crazy_poly = 1.0; ///2.0 // Maximal allowed mismatch difference calculated as polynomial maximum
public boolean ly_poly_backup = true; // Use CM offset measuremets if poly failed
public boolean fo_correct = true; // correct far objects by comparing orthogonal correlations
......@@ -111,7 +111,7 @@ public class ImageDttParameters {
public double lmas_lambda_initial = 0.03; //
public double lmas_rms_diff = 0.0003; //
public int lmas_num_iter = 20; //
public int lmas_num_iter = 20; ///10
// Filtering and strength calculation
public double lmas_max_rel_rms = 0.3; // maximal relative (to average max/min amplitude LMA RMS) // May be up to 0.3)
public double lmas_min_strength = 0.7; // minimal composite strength (sqrt(average amp squared over absolute RMS)
......@@ -141,8 +141,8 @@ public class ImageDttParameters {
public double lma_half_width = 2.0; //
public double lma_cost_wy = 0.0; // cost of parallel-to-disparity correction
public double lma_cost_wxy = 0.0; // cost of ortho-to-disparity correction
public double lma_cost_wy = 0.0; /// 0.00050 // cost of parallel-to-disparity correction
public double lma_cost_wxy = 0.0; /// 0.00100 // cost of ortho-to-disparity correction
public double lma_lambda_initial = 0.03; //
public double lma_lambda_scale_good = 0.5; //
......
src/main/java/com/elphel/imagej/tileprocessor/QuadCLT.java
View file @ ac185fe7
This source diff could not be displayed because it is too large. You can view the blob instead.
src/main/java/com/elphel/imagej/tileprocessor/QuadCLTCPU.java 0 → 100644
View file @ ac185fe7
This source diff could not be displayed because it is too large. You can view the blob instead.
src/main/java/com/elphel/imagej/tileprocessor/TwoQuadCLT.java
View file @ ac185fe7
......@@ -125,8 +125,8 @@ public class TwoQuadCLT {
}
return biCamDSI_persistent.biScans.get(indx);
}
public void processCLTQuadCorrs(
/*
public void processCLTTwoQuadCorrs( // not referenced?
QuadCLT quadCLT_main,
QuadCLT quadCLT_aux,
CLTParameters clt_parameters,
......@@ -190,7 +190,7 @@ public class TwoQuadCLT {
saturation_imp_aux, //output // boolean [][] saturation_imp,
debugLevel); // int debugLevel);
// Tempporarily processing individaully with the old code
// Temporarily processing individaully with the old code
quadCLT_main.processCLTQuadCorr( // returns ImagePlus, but it already should be saved/shown
imp_srcs_main, // [srcChannel], // should have properties "name"(base for saving results), "channel","path"
saturation_imp_main, // boolean [][] saturation_imp, // (near) saturated pixels or null
......@@ -236,12 +236,12 @@ public class TwoQuadCLT {
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("processCLTTwoQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
*/
......@@ -293,6 +293,7 @@ public class TwoQuadCLT {
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
ImagePlus [] imp_srcs_aux = quadCLT_aux.conditionImageSet(
......@@ -304,6 +305,7 @@ public class TwoQuadCLT {
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
// Tempporarily processing individaully with the old code
......@@ -342,7 +344,7 @@ public class TwoQuadCLT {
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("processCLTQuadCorrPairs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
......@@ -396,6 +398,7 @@ public class TwoQuadCLT {
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
ImagePlus [] imp_srcs_aux = quadCLT_aux.conditionImageSet(
......@@ -407,6 +410,7 @@ public class TwoQuadCLT {
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
// Tempporarily processing individually with the old code
......@@ -446,13 +450,12 @@ public class TwoQuadCLT {
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("prepareFilesForGPUDebug(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
public void processCLTQuadCorrPairsGpu(
// GPUTileProcessor gPUTileProcessor,
GPUTileProcessor.GpuQuad gpuQuad_main,
GPUTileProcessor.GpuQuad gpuQuad_aux,
QuadCLT quadCLT_main,
......@@ -467,7 +470,7 @@ public class TwoQuadCLT {
final boolean updateStatus,
final int debugLevel) throws Exception
{
// TwoQuadGPU twoQuadGPU = new TwoQuadGPU(this);
this.startTime=System.nanoTime();
String [] sourceFiles=quadCLT_main.correctionsParameters.getSourcePaths();
QuadCLT.SetChannels [] set_channels_main = quadCLT_main.setChannels(debugLevel);
......@@ -503,6 +506,7 @@ public class TwoQuadCLT {
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
ImagePlus [] imp_srcs_aux = quadCLT_aux.conditionImageSet(
......@@ -514,13 +518,13 @@ public class TwoQuadCLT {
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
// Tempporarily processing individaully with the old code
processCLTQuadCorrPairGpu(
// gPUTileProcessor, // GPUTileProcessor gPUTileProcessor,
gpuQuad_main, // GPUTileProcessor.GpuQuad gpuQuad_main,
gpuQuad_aux, // GPUTileProcessor.GpuQuad gpuQuad_aux,
QuadCLT.processCLTQuadCorrPairGpu(
// gpuQuad_main, // GPUTileProcessor.GpuQuad gpuQuad_main,
// gpuQuad_aux, // GPUTileProcessor.GpuQuad gpuQuad_aux,
quadCLT_main, // QuadCLT quadCLT_main,
quadCLT_aux, // QuadCLT quadCLT_aux,
imp_srcs_main, // ImagePlus [] imp_quad_main,
......@@ -556,10 +560,11 @@ public class TwoQuadCLT {
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("processCLTQuadCorrPairsGpu(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
public ImagePlus [] processCLTQuadCorrPair(
QuadCLT quadCLT_main,
QuadCLT quadCLT_aux,
......@@ -599,6 +604,7 @@ public class TwoQuadCLT {
results[i].setTitle(results[i].getTitle()+"RAW");
}
if (debugLevel>1) System.out.println("processing: "+path);
/* // 08/12/2020 Moved to conditionImageSet
getRigImageStacks(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
quadCLT_main, // QuadCLT quadCLT_main,
......@@ -609,7 +615,7 @@ public class TwoQuadCLT {
saturation_aux, // boolean [][] saturation_aux, // (near) saturated pixels or null
threadsMax, // maximal number of threads to launch
debugLevel); // final int debugLevel);
*/
// temporary setting up tile task file (one integer per tile, bitmask
// for testing defined for a window, later the tiles to process will be calculated based on previous passes results
......@@ -617,9 +623,7 @@ public class TwoQuadCLT {
// int [][] tile_op_aux = quadCLT_aux.tp.setSameTileOp (clt_parameters, clt_parameters.tile_task_op, debugLevel);
double [][] disparity_array_main = quadCLT_main.tp.setSameDisparity(clt_parameters.disparity); // [tp.tilesY][tp.tilesX] - individual per-tile expected disparity
//TODO: Add array of default disparity - use for combining images in force disparity mode (no correlation), when disparity is predicted from other tiles
// start with all old arrays, remove some later
double [][][][] clt_corr_combo = null;
double [][][][] texture_tiles_main = null; // [tp.tilesY][tp.tilesX]["RGBA".length()][]; // tiles will be 16x16, 2 visualization mode full 16 or overlapped
......@@ -642,7 +646,6 @@ public class TwoQuadCLT {
texture_tiles_aux[i][j] = null;
}
}
}
double [][] disparity_bimap = new double [ImageDtt.BIDISPARITY_TITLES.length][]; //[0] -residual disparity, [1] - orthogonal (just for debugging) last 4 - max pixel differences
......@@ -1325,6 +1328,7 @@ public class TwoQuadCLT {
results[i].setTitle(results[i].getTitle()+"RAW");
}
if (debugLevel>1) System.out.println("processing: "+path);
/* // 08/12/2020 Moved to conditionImageSet
getRigImageStacks(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
quadCLT_main, // QuadCLT quadCLT_main,
......@@ -1335,9 +1339,7 @@ public class TwoQuadCLT {
saturation_aux, // boolean [][] saturation_aux, // (near) saturated pixels or null
threadsMax, // maximal number of threads to launch
debugLevel); // final int debugLevel);
*/
// temporary setting up tile task file (one integer per tile, bitmask
// for testing defined for a window, later the tiles to process will be calculated based on previous passes results
......@@ -1936,602 +1938,7 @@ public class TwoQuadCLT {
return results;
}
public ImagePlus [] processCLTQuadCorrPairGpu(
// GPUTileProcessor gPUTileProcessor,
GPUTileProcessor.GpuQuad gpuQuad_main,
GPUTileProcessor.GpuQuad gpuQuad_aux,
QuadCLT quadCLT_main,
QuadCLT quadCLT_aux,
ImagePlus [] imp_quad_main,
ImagePlus [] imp_quad_aux,
boolean [][] saturation_main, // (near) saturated pixels or null
boolean [][] saturation_aux, // (near) saturated pixels or null
CLTParameters clt_parameters,
EyesisCorrectionParameters.CorrectionParameters ecp,
EyesisCorrectionParameters.DebayerParameters debayerParameters,
ColorProcParameters colorProcParameters,
ColorProcParameters colorProcParameters_aux,
EyesisCorrectionParameters.RGBParameters rgbParameters,
double [] scaleExposures_main, // probably not needed here - restores brightness of the final image
double [] scaleExposures_aux, // probably not needed here - restores brightness of the final image
boolean notch_mode, // use pole-detection mode for inter-camera correlation
final int lt_rad, // low texture mode - inter-correlation is averaged between the neighbors before argmax-ing, using
final int threadsMax, // maximal number of threads to launch
final boolean updateStatus,
final int debugLevel){
// get fat_zero (absolute) and color scales
boolean is_mono = quadCLT_main.isMonochrome();
boolean is_lwir = quadCLT_main.isLwir();
double fat_zero = clt_parameters.getGpuFatZero(is_mono); // 30.0;
double [] scales = (is_mono) ? (new double [] {1.0}) :(new double [] {
clt_parameters.gpu_weight_r, // 0.25
clt_parameters.gpu_weight_b, // 0.25
1.0 - clt_parameters.gpu_weight_r - clt_parameters.gpu_weight_b}); // 0.5
double cwgreen = 1.0/(1.0 + clt_parameters.corr_red + clt_parameters.corr_blue); // green color
double [] col_weights= (is_mono) ? (new double [] {1.0}) :(new double [] {
clt_parameters.corr_red * cwgreen,
clt_parameters.corr_blue * cwgreen,
cwgreen});
ImageDtt image_dtt = new ImageDtt(
clt_parameters.transform_size,
is_mono,
is_lwir,
1.0);
float [][] lpf_rgb = new float[][] {
image_dtt.floatGetCltLpfFd(clt_parameters.gpu_sigma_r),
image_dtt.floatGetCltLpfFd(clt_parameters.gpu_sigma_b),
image_dtt.floatGetCltLpfFd(clt_parameters.gpu_sigma_g),
image_dtt.floatGetCltLpfFd(clt_parameters.gpu_sigma_m)
};
gpuQuad_main.setLpfRbg(
lpf_rgb);
float [] lpf_flat = image_dtt.floatGetCltLpfFd(clt_parameters.getGpuCorrSigma(is_mono));
gpuQuad_main.setLpfCorr(
"lpf_corr", // String const_name, // "lpf_corr"
lpf_flat);
float [] lpf_rb_flat = image_dtt.floatGetCltLpfFd(clt_parameters.getGpuCorrRBSigma(is_mono));
gpuQuad_main.setLpfCorr(
"lpf_rb_corr", // String const_name, // "lpf_corr"
lpf_rb_flat);
final boolean use_aux = false; // currently GPU is configured for a single quad camera
final boolean batch_mode = clt_parameters.batch_run; //disable any debug images
boolean toRGB= quadCLT_main.correctionsParameters.toRGB;
// may use this.StartTime to report intermediate steps execution times
String name=quadCLT_main.correctionsParameters.getModelName((String) imp_quad_main[0].getProperty("name"));
String path= (String) imp_quad_main[0].getProperty("path"); // Only for debug output
// now set to only 4 !
ImagePlus [] results = new ImagePlus[imp_quad_main.length]; // + imp_quad_aux.length];
for (int i = 0; i < results.length; i++) {
if (i< imp_quad_main.length) {
results[i] = imp_quad_main[i];
} else {
results[i] = imp_quad_aux[i-imp_quad_main.length];
}
results[i].setTitle(results[i].getTitle()+"RAW");
}
if (debugLevel>1) System.out.println("processing: "+path);
getRigImageStacks(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
quadCLT_main, // QuadCLT quadCLT_main,
quadCLT_aux, // QuadCLT quadCLT_aux,
imp_quad_main, // ImagePlus [] imp_quad_main,
imp_quad_aux, // ImagePlus [] imp_quad_aux,
saturation_main, // boolean [][] saturation_main, // (near) saturated pixels or null
saturation_aux, // boolean [][] saturation_aux, // (near) saturated pixels or null
threadsMax, // maximal number of threads to launch
debugLevel); // final int debugLevel);
gpuQuad_main.setConvolutionKernels(
(use_aux?quadCLT_aux.getCLTKernels() : quadCLT_main.getCLTKernels()), // double [][][][][][] clt_kernels,
false); // boolean force)
gpuQuad_main.setBayerImages(
(use_aux? quadCLT_aux.image_data: quadCLT_main.image_data), // double [][][] bayer_data,
true); // boolean force);
// Set task clt_parameters.disparity
Rectangle twoi = new Rectangle (
clt_parameters.gpu_woi_tx,
clt_parameters.gpu_woi_ty,
clt_parameters.gpu_woi_twidth,
clt_parameters.gpu_woi_theight);
GPUTileProcessor.TpTask [] tp_tasks = gpuQuad_main.setFullFrameImages(
false, // boolean calc_offsets, // old way, now not needed with GPU calculation
twoi, // Rectangle woi,
clt_parameters.gpu_woi_round, // boolean round_woi,
(float) clt_parameters.disparity, // float target_disparity, // apply same disparity to all tiles
0xf, // int out_image, // from which tiles to generate image (currently 0/1)
0x3f, // int corr_mask, // which correlation pairs to generate (maybe later - reduce size from 15x15)
!use_aux, // boolean use_master,
use_aux, // boolean use_aux,
quadCLT_main.getGeometryCorrection(), // final GeometryCorrection geometryCorrection_main,
quadCLT_aux.getGeometryCorrection(), // final GeometryCorrection geometryCorrection_aux, // if null, will only calculate offsets fro the main camera
null, // final double [][][] ers_delay, // if not null - fill with tile center acquisition delay
threadsMax, // final int threadsMax, // maximal number of threads to launch
debugLevel); // final int debugLevel)
// Optionally save offsets here?
// EyesisCorrectionParameters.CorrectionParameters ecp,
boolean save_ports_xy = false; // true; Same files as saved with the kernels
if ((ecp.tile_processor_gpu != null) && !ecp.tile_processor_gpu.isEmpty() && save_ports_xy) {
int quad = 4;
String file_prefix = ecp.tile_processor_gpu +"clt/main";
for (int chn = 0; chn < quad; chn++) {
String img_path = file_prefix+"_chn"+chn+".portsxy";
FileOutputStream fos;
try {
fos = new FileOutputStream(img_path);
} catch (FileNotFoundException e) {
// TODO Auto-generated catch block
System.out.println("Could not write to "+img_path+" (file not found) port offsets");
break;
}
DataOutputStream dos = new DataOutputStream(fos);
WritableByteChannel channel = Channels.newChannel(dos);
ByteBuffer bb = ByteBuffer.allocate(tp_tasks.length * 2 * 4);
bb.order(ByteOrder.LITTLE_ENDIAN);
bb.clear();
for (int i = 0; i < tp_tasks.length; i++) {
bb.putFloat((tp_tasks[i].xy[chn][0])); // x-offset
bb.putFloat((tp_tasks[i].xy[chn][1])); // y-offset
}
bb.flip();
try {
channel.write(bb);
} catch (IOException e) {
System.out.println("Could not write to "+img_path+" port offsets");
break;
}
try {
dos.close();
} catch (IOException e) {
System.out.println("Could not close DataOutputStream for "+img_path+" port offsets");
}
System.out.println("Wrote port offsets to "+img_path+".");
}
}
gpuQuad_main.setTasks(
tp_tasks, // TpTask [] tile_tasks,
use_aux); // boolean use_aux)
gpuQuad_main.setGeometryCorrection(
quadCLT_main.getGeometryCorrection(),
false); // boolean use_java_rByRDist) { // false - use newer GPU execCalcReverseDistortions); // once
gpuQuad_main.setExtrinsicsVector(quadCLT_main.getGeometryCorrection().getCorrVector()); // for each new image
/* // TODO: calculate from the camera geometry?
double[][] port_offsets = { // used only in textures to scale differences
{-0.5, -0.5},
{ 0.5, -0.5},
{-0.5, 0.5},
{ 0.5, 0.5}};
*/
// All set, run kernel (correct and convert)
int NREPEAT = 1; // 00;
System.out.println("\n------------ Running GPU "+NREPEAT+" times ----------------");
long startGPU=System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) {
gpuQuad_main.execCalcReverseDistortions();
}
long startRotDerivs=System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) {
gpuQuad_main.execRotDerivs();
}
long startTasksSetup=System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) {
gpuQuad_main.execSetTilesOffsets();
}
long startDirectConvert=System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) {
gpuQuad_main.execConvertDirect();
}
// run imclt;
long startIMCLT=System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) {
gpuQuad_main.execImcltRbgAll(quadCLT_main.isMonochrome());
}
long endImcltTime = System.nanoTime();
// run correlation
long startCorr2d=System.nanoTime(); // System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) gpuQuad_main.execCorr2D(
scales,// double [] scales,
fat_zero, // double fat_zero);
clt_parameters.gpu_corr_rad); // int corr_radius
long endCorr2d = System.nanoTime();
// run textures
long startTextures = System.nanoTime(); // System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) gpuQuad_main.execTextures(
col_weights, // double [] color_weights,
quadCLT_main.isLwir(), // boolean is_lwir,
clt_parameters.min_shot, // double min_shot, // 10.0
clt_parameters.scale_shot, // double scale_shot, // 3.0
clt_parameters.diff_sigma, // double diff_sigma, // pixel value/pixel change
clt_parameters.diff_threshold, // double diff_threshold, // pixel value/pixel change
clt_parameters.min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
clt_parameters.dust_remove); // boolean dust_remove,
long endTextures = System.nanoTime();
// run texturesRBGA
long startTexturesRBGA = System.nanoTime(); // System.nanoTime();
for (int i = 0; i < NREPEAT; i++ ) gpuQuad_main.execRBGA(
col_weights, // double [] color_weights,
quadCLT_main.isLwir(), // boolean is_lwir,
clt_parameters.min_shot, // double min_shot, // 10.0
clt_parameters.scale_shot, // double scale_shot, // 3.0
clt_parameters.diff_sigma, // double diff_sigma, // pixel value/pixel change
clt_parameters.diff_threshold, // double diff_threshold, // pixel value/pixel change
clt_parameters.min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
clt_parameters.dust_remove); // boolean dust_remove,
long endTexturesRBGA = System.nanoTime();
long endGPUTime = System.nanoTime();
long calcReverseTime= (startRotDerivs- startGPU) /NREPEAT;
long rotDerivsTime= (startTasksSetup- startRotDerivs) /NREPEAT;
long tasksSetupTime= (startDirectConvert- startTasksSetup) /NREPEAT;
long firstGPUTime= (startIMCLT- startDirectConvert) /NREPEAT;
long runImcltTime = (endImcltTime - startIMCLT) /NREPEAT;
long runCorr2DTime = (endCorr2d - startCorr2d) /NREPEAT;
long runTexturesTime = (endTextures - startTextures) /NREPEAT;
long runTexturesRBGATime = (endTexturesRBGA - startTexturesRBGA) /NREPEAT;
long runGPUTime = (endGPUTime - startGPU) /NREPEAT;
// run corr2d
//RotDerivs
System.out.println("\n------------ End of running GPU "+NREPEAT+" times ----------------");
System.out.println("GPU run time ="+ (runGPUTime * 1.0e-6)+"ms");
System.out.println(" - calc reverse dist.: "+(calcReverseTime*1.0e-6)+"ms");
System.out.println(" - rot/derivs: "+(rotDerivsTime*1.0e-6)+"ms");
System.out.println(" - tasks setup: "+(tasksSetupTime*1.0e-6)+"ms");
System.out.println(" - direct conversion: "+(firstGPUTime*1.0e-6)+"ms");
System.out.println(" - imclt: "+(runImcltTime*1.0e-6)+"ms");
System.out.println(" - corr2D: "+(runCorr2DTime*1.0e-6)+"ms");
System.out.println(" - textures: "+(runTexturesTime*1.0e-6)+"ms");
System.out.println(" - RGBA: "+(runTexturesRBGATime*1.0e-6)+"ms");
// get data back from GPU
float [][][] iclt_fimg = new float [gpuQuad_main.getNumCams()][][];
for (int ncam = 0; ncam < iclt_fimg.length; ncam++) {
iclt_fimg[ncam] = gpuQuad_main.getRBG(ncam);
}
// gpuQuad_main
int out_width = gpuQuad_main.getImageWidth() + gpuQuad_main.getDttSize();
int out_height = gpuQuad_main.getImageHeight() + gpuQuad_main.getDttSize();
int tilesX = gpuQuad_main.getImageWidth() / gpuQuad_main.getDttSize();
int tilesY = gpuQuad_main.getImageHeight() / gpuQuad_main.getDttSize();
// show extra
/* */
String [] extra_group_titles = {"DIFF","Red","Blue","Green"};
String [] extra_titles = new String [extra_group_titles.length*gpuQuad_main.getNumCams()];
for (int g = 0; g < extra_group_titles.length;g++) {
for (int ncam=0; ncam < gpuQuad_main.getNumCams();ncam++) {
extra_titles[g * gpuQuad_main.getNumCams() + ncam]= extra_group_titles[g]+"-"+ncam;
}
}
float [][] extra = gpuQuad_main.getExtra();
(new ShowDoubleFloatArrays()).showArrays(
extra,
tilesX,
tilesY,
true,
name+"-EXTRA-D"+clt_parameters.disparity,
extra_titles);
/* */
ImagePlus [] imps_RGB = new ImagePlus[iclt_fimg.length];
for (int ncam = 0; ncam < iclt_fimg.length; ncam++) {
String title=name+"-"+String.format("%02d", ncam);
imps_RGB[ncam] = quadCLT_main.linearStackToColor( // probably no need to separate and process the second half with quadCLT_aux
clt_parameters,
colorProcParameters,
rgbParameters,
title, // String name,
"-D"+clt_parameters.disparity, //String suffix, // such as disparity=...
toRGB,
!quadCLT_main.correctionsParameters.jpeg, // boolean bpp16, // 16-bit per channel color mode for result
!batch_mode, // true, // boolean saveShowIntermediate, // save/show if set globally
false, // boolean saveShowFinal, // save/show result (color image?)
iclt_fimg[ncam],
out_width,
out_height,
1.0, // scaleExposures[iAux][iSubCam], // double scaleExposure, // is it needed?
debugLevel );
}
//show_corr
int [] wh = new int[2];
if (clt_parameters.show_corr) {
int [] corr_indices = gpuQuad_main.getCorrIndices();
float [][] corr2D = gpuQuad_main.getCorr2D(
clt_parameters.gpu_corr_rad); // int corr_rad);
// convert to 6-layer image using tasks
double [][] dbg_corr = GPUTileProcessor.getCorr2DView(
tilesX,
tilesY,
corr_indices,
corr2D,
wh);
(new ShowDoubleFloatArrays()).showArrays(
dbg_corr,
wh[0],
wh[1],
true,
name+"-CORR2D-D"+clt_parameters.disparity,
GPUTileProcessor.getCorrTitles());
}
// convert to overlapping and show
if (clt_parameters.gen_chn_img) {
// combine to a sliced color image
// assuming total number of images to be multiple of 4
// int [] slice_seq = {0,1,3,2}; //clockwise
int [] slice_seq = new int[results.length];
for (int i = 0; i < slice_seq.length; i++) {
slice_seq[i] = i ^ ((i >> 1) & 1); // 0,1,3,2,4,5,7,6, ...
}
int width = imps_RGB[0].getWidth();
int height = imps_RGB[0].getHeight();
ImageStack array_stack=new ImageStack(width,height);
for (int i = 0; i<slice_seq.length; i++){
if (imps_RGB[slice_seq[i]] != null) {
array_stack.addSlice("port_"+slice_seq[i], imps_RGB[slice_seq[i]].getProcessor().getPixels());
} else {
array_stack.addSlice("port_"+slice_seq[i], results[slice_seq[i]].getProcessor().getPixels());
}
}
ImagePlus imp_stack = new ImagePlus(name+"-SHIFTED-D"+clt_parameters.disparity, array_stack);
imp_stack.getProcessor().resetMinAndMax();
if (!batch_mode) {
imp_stack.updateAndDraw();
}
//imp_stack.getProcessor().resetMinAndMax();
//imp_stack.show();
// eyesisCorrections.saveAndShow(imp_stack, this.correctionsParameters);
quadCLT_main.eyesisCorrections.saveAndShowEnable(
imp_stack, // ImagePlus imp,
quadCLT_main.correctionsParameters, // EyesisCorrectionParameters.CorrectionParameters correctionsParameters,
true, // boolean enableSave,
!batch_mode) ;// boolean enableShow);
}
if (clt_parameters.gen_4_img) {
// Save as individual JPEG images in the model directory
String x3d_path= quadCLT_main.correctionsParameters.selectX3dDirectory(
name, // quad timestamp. Will be ignored if correctionsParameters.use_x3d_subdirs is false
quadCLT_main.correctionsParameters.x3dModelVersion,
true, // smart,
true); //newAllowed, // save
for (int sub_img = 0; sub_img < imps_RGB.length; sub_img++){
quadCLT_main.eyesisCorrections.saveAndShow(
imps_RGB[sub_img],
x3d_path,
quadCLT_main.correctionsParameters.png && !clt_parameters.black_back,
!batch_mode && clt_parameters.show_textures,
quadCLT_main.correctionsParameters.JPEG_quality, // jpegQuality); // jpegQuality){// <0 - keep current, 0 - force Tiff, >0 use for JPEG
(debugLevel > 0) ? debugLevel : 1); // int debugLevel (print what it saves)
}
String model_path= quadCLT_main.correctionsParameters.selectX3dDirectory(
name, // quad timestamp. Will be ignored if correctionsParameters.use_x3d_subdirs is false
null,
true, // smart,
true); //newAllowed, // save
quadCLT_main.createThumbNailImage(
imps_RGB[0],
model_path,
"thumb",
debugLevel);
}
// Use GPU prepared RBGA
if (clt_parameters.show_rgba_color) {
Rectangle woi = new Rectangle();
float [][] rbga = gpuQuad_main.getRBGA(
(is_mono?1:3), // int num_colors,
woi);
(new ShowDoubleFloatArrays()).showArrays(
rbga,
woi.width,
woi.height,
true,
name+"-RGBA-STACK-D"+clt_parameters.disparity+
":"+clt_parameters.gpu_woi_tx+":"+clt_parameters.gpu_woi_ty+
":"+clt_parameters.gpu_woi_twidth+":"+clt_parameters.gpu_woi_theight+
":"+(clt_parameters.gpu_woi_round?"C":"R"),
new String[] {"R","B","G","A"}
);
// for now - use just RGB. Later add option for RGBA
float [][] rgb_main = {rbga[0],rbga[1],rbga[2]};
float [][] rgba_main = {rbga[0],rbga[1],rbga[2],rbga[3]};
ImagePlus imp_rgba_main = quadCLT_main.linearStackToColor(
clt_parameters,
colorProcParameters,
rgbParameters,
name+"-texture", // String name,
"-D"+clt_parameters.disparity+"-MAINGPU", //String suffix, // such as disparity=...
toRGB,
!quadCLT_main.correctionsParameters.jpeg, // boolean bpp16, // 16-bit per channel color mode for result
false, // true, // boolean saveShowIntermediate, // save/show if set globally
false, // true, // boolean saveShowFinal, // save/show result (color image?)
((clt_parameters.alpha1 > 0)? rgba_main: rgb_main),
woi.width, // clt_parameters.gpu_woi_twidth * image_dtt.transform_size, // tilesX * image_dtt.transform_size,
woi.height, // clt_parameters.gpu_woi_theight *image_dtt.transform_size, // tilesY * image_dtt.transform_size,
1.0, // double scaleExposure, // is it needed?
debugLevel );
int width = imp_rgba_main.getWidth();
int height =imp_rgba_main.getHeight();
ImageStack texture_stack=new ImageStack(width,height);
texture_stack.addSlice("main", imp_rgba_main.getProcessor().getPixels()); // single slice
// ImagePlus imp_texture_stack = new ImagePlus(name+"-RGBA-D"+clt_parameters.disparity, texture_stack);
ImagePlus imp_texture_stack = new ImagePlus(
name+"-RGBA-D"+clt_parameters.disparity+
":"+clt_parameters.gpu_woi_tx+":"+clt_parameters.gpu_woi_ty+
":"+clt_parameters.gpu_woi_twidth+":"+clt_parameters.gpu_woi_theight+
":"+(clt_parameters.gpu_woi_round?"C":"R"),
texture_stack);
imp_texture_stack.getProcessor().resetMinAndMax();
// imp_texture_stack.show();
String results_path= quadCLT_main.correctionsParameters.selectResultsDirectory( // selectX3dDirectory(
// name, // quad timestamp. Will be ignored if correctionsParameters.use_x3d_subdirs is false
// quadCLT_main.correctionsParameters.x3dModelVersion,
true, // smart,
true); //newAllowed, // save
quadCLT_main.eyesisCorrections.saveAndShow(
imp_texture_stack,
results_path,
true, // quadCLT_main.correctionsParameters.png && !clt_parameters.black_back,
true, // !batch_mode && clt_parameters.show_textures,
0, // quadCLT_main.correctionsParameters.JPEG_quality, // jpegQuality); // jpegQuality){// <0 - keep current, 0 - force Tiff, >0 use for JPEG
(debugLevel > 0) ? debugLevel : 1); // int debugLevel (print what it saves)
}
// convert textures to RGBA in Java
if (clt_parameters.show_rgba_color && (debugLevel > 100)) { // disabling
int numcol = quadCLT_main.isMonochrome()?1:3;
int ports = imp_quad_main.length;
int [] texture_indices = gpuQuad_main.getTextureIndices();
int num_src_slices = numcol + 1 + (clt_parameters.keep_weights?(ports + numcol + 1):0); // 12 ; // calculate
// float [][][] ftextures = gPUTileProcessor.getTextures(
// (is_mono?1:3), // int num_colors,
// clt_parameters.keep_weights); // boolean keep_weights);
float [] flat_textures = gpuQuad_main.getFlatTextures(
texture_indices.length,
(is_mono?1:3), // int num_colors,
clt_parameters.keep_weights); // boolean keep_weights);
int texture_slice_size = (2 * gpuQuad_main.getDttSize())* (2 * gpuQuad_main.getDttSize());
int texture_tile_size = texture_slice_size * num_src_slices ;
if (debugLevel > -1) {
for (int indx = 0; indx < texture_indices.length; indx++) if ((texture_indices[indx] & (1 << GPUTileProcessor.LIST_TEXTURE_BIT)) != 0){
int tile = texture_indices[indx] >> GPUTileProcessor.CORR_NTILE_SHIFT;
int tileX = tile % tilesX;
int tileY = tile / tilesX;
if ((tileY == clt_parameters.tileY) && (tileX == clt_parameters.tileX)) {
System.out.println("=== tileX= "+tileX+" tileY= "+tileY+" tile="+tile+" ===");
for (int slice =0; slice < num_src_slices; slice++) {
System.out.println("=== Slice="+slice+" ===");
for (int i = 0; i < 2 * gpuQuad_main.getDttSize(); i++) {
for (int j = 0; j < 2 * gpuQuad_main.getDttSize(); j++) {
System.out.print(String.format("%10.4f ",
flat_textures[indx*texture_tile_size + slice* texture_slice_size + 2 * gpuQuad_main.getDttSize() * i + j]));
}
System.out.println();
}
}
}
}
}
double [][][][] texture_tiles = gpuQuad_main.doubleTextures(
new Rectangle(0, 0, tilesX, tilesY), // Rectangle woi,
texture_indices, // int [] indices,
flat_textures, // float [][][] ftextures,
tilesX, // int full_width,
4, // rbga only /int num_slices
num_src_slices // int num_src_slices
);
if ((debugLevel > -1) && (clt_parameters.tileX >= 0) && (clt_parameters.tileY >= 0) && (clt_parameters.tileX < tilesX) && (clt_parameters.tileY < tilesY)) {
String [] rgba_titles = {"red","blue","green","alpha"};
String [] rgba_weights_titles = {"red","blue","green","alpha","port0","port1","port2","port3","r-rms","b-rms","g-rms","w-rms"};
double [][] texture_tile = texture_tiles[clt_parameters.tileY][clt_parameters.tileX];
int tile = +clt_parameters.tileY * tilesX +clt_parameters.tileX;
System.out.println("=== tileX= "+clt_parameters.tileX+" tileY= "+clt_parameters.tileY+" tile="+tile+" ===");
for (int slice =0; slice < texture_tile.length; slice++) {
System.out.println("\n=== Slice="+slice+" ===");
for (int i = 0; i < 2 * gpuQuad_main.getDttSize(); i++) {
for (int j = 0; j < 2 * gpuQuad_main.getDttSize(); j++) {
System.out.print(String.format("%10.4f ",
texture_tile[slice][2 * gpuQuad_main.getDttSize() * i + j]));
}
System.out.println();
}
}
(new ShowDoubleFloatArrays()).showArrays(
texture_tile,
2 * image_dtt.transform_size,
2 * image_dtt.transform_size,
true,
name + "-TXTNOL-GPU-D"+clt_parameters.disparity+"-X"+clt_parameters.tileX+"-Y"+clt_parameters.tileY,
(clt_parameters.keep_weights?rgba_weights_titles:rgba_titles));
}
int alpha_index = 3;
// in monochrome mode only MONO_CHN == GREEN_CHN is used, R and B are null
double [][] texture_overlap_main = image_dtt.combineRBGATiles(
texture_tiles, // array [tp.tilesY][tp.tilesX][4][4*transform_size] or [tp.tilesY][tp.tilesX]{null}
// image_dtt.transform_size,
true, // when false - output each tile as 16x16, true - overlap to make 8x8
clt_parameters.sharp_alpha, // combining mode for alpha channel: false - treat as RGB, true - apply center 8x8 only
threadsMax, // maximal number of threads to launch
debugLevel);
if (clt_parameters.alpha1 > 0){ // negative or 0 - keep alpha as it was
double scale = (clt_parameters.alpha1 > clt_parameters.alpha0) ? (1.0/(clt_parameters.alpha1 - clt_parameters.alpha0)) : 0.0;
for (int i = 0; i < texture_overlap_main[alpha_index].length; i++){
double d = texture_overlap_main[alpha_index][i];
if (d >=clt_parameters.alpha1) d = 1.0;
else if (d <=clt_parameters.alpha0) d = 0.0;
else d = scale * (d- clt_parameters.alpha0);
texture_overlap_main[alpha_index][i] = d;
}
}
// for now - use just RGB. Later add option for RGBA
double [][] texture_rgb_main = {texture_overlap_main[0],texture_overlap_main[1],texture_overlap_main[2]};
double [][] texture_rgba_main = {texture_overlap_main[0],texture_overlap_main[1],texture_overlap_main[2],texture_overlap_main[3]};
ImagePlus imp_texture_main = quadCLT_main.linearStackToColor(
clt_parameters,
colorProcParameters,
rgbParameters,
name+"-texture", // String name,
"-D"+clt_parameters.disparity+"-MAINGPU", //String suffix, // such as disparity=...
toRGB,
!quadCLT_main.correctionsParameters.jpeg, // boolean bpp16, // 16-bit per channel color mode for result
false, // true, // boolean saveShowIntermediate, // save/show if set globally
false, // true, // boolean saveShowFinal, // save/show result (color image?)
((clt_parameters.alpha1 > 0)? texture_rgba_main: texture_rgb_main),
tilesX * image_dtt.transform_size,
tilesY * image_dtt.transform_size,
1.0, // double scaleExposure, // is it needed?
debugLevel );
int width = imp_texture_main.getWidth();
int height =imp_texture_main.getHeight();
ImageStack texture_stack=new ImageStack(width,height);
texture_stack.addSlice("main", imp_texture_main.getProcessor().getPixels()); // single slice
ImagePlus imp_texture_stack = new ImagePlus(name+"-TEXTURES-D"+clt_parameters.disparity, texture_stack);
imp_texture_stack.getProcessor().resetMinAndMax();
imp_texture_stack.show();
}
return results;
}
public void showERSDelay(double [][][] ers_delay)
{
int tilesX = quadCLT_main.tp.getTilesX();
......@@ -2562,9 +1969,10 @@ public class TwoQuadCLT {
}
// public double [][][][] getRigImageStacks(
public void getRigImageStacks(
// 08/12/2020 Moved to conditionImageSet
/*
public void getRigImageStacks( // Removed all references
CLTParameters clt_parameters,
QuadCLT quadCLT_main,
QuadCLT quadCLT_aux,
......@@ -2621,7 +2029,7 @@ public class TwoQuadCLT {
quadCLT_main.tp.setTrustedCorrelation(clt_parameters.grow_disp_trust);
quadCLT_aux.tp.setTrustedCorrelation(clt_parameters.grow_disp_trust);
}
*/
public void processInfinityRigs( // actually there is no sense to process multiple image sets. Combine with other processing?
QuadCLT quadCLT_main,
......@@ -2669,6 +2077,7 @@ public class TwoQuadCLT {
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
ImagePlus [] imp_srcs_aux = quadCLT_aux.conditionImageSet(
......@@ -2680,6 +2089,7 @@ public class TwoQuadCLT {
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
// Temporary processing individually with the old code
......@@ -2707,7 +2117,7 @@ public class TwoQuadCLT {
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("processInfinityRigs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
......@@ -2843,6 +2253,7 @@ public class TwoQuadCLT {
final int debugLevel0){
final int debugLevel = debugLevel0 + (clt_parameters.rig.rig_mode_debug?2:0);
// double [][][][] double_stacks =
/* // 08/12/2020 Moved to condifuinImageSet
getRigImageStacks(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
quadCLT_main, // QuadCLT quadCLT_main,
......@@ -2856,6 +2267,7 @@ public class TwoQuadCLT {
quadCLT_main.tp.resetCLTPasses();
quadCLT_aux.tp.resetCLTPasses();
*/
/*
FIXME - make it work
prealignInfinityRig(
......@@ -8300,6 +7712,7 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevelInner); // int debugLevel);
if (updateStatus) IJ.showStatus("Conditioning aux camera image set for "+quadCLT_main.image_name);
ImagePlus [] imp_srcs_aux = quadCLT_aux.conditionImageSet(
......@@ -8311,6 +7724,7 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevelInner); // int debugLevel);
// optionally adjust main, aux, rig here
......@@ -8769,13 +8183,15 @@ if (debugLevel > -100) return true; // temporarily !
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("batchRig(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
public void batchLwirRig(
QuadCLT quadCLT_main, // tiles should be set
// GPUTileProcessor.GpuQuad gpuQuad_main, // may be null if GPU for MAIN is not use3d
// GPUTileProcessor.GpuQuad gpuQuad_aux, // may be null if GPU for AUX is not use3d
QuadCLT quadCLT_main, // tiles should be set
QuadCLT quadCLT_aux,
CLTParameters clt_parameters,
EyesisCorrectionParameters.DebayerParameters debayerParameters,
......@@ -8789,6 +8205,15 @@ if (debugLevel > -100) return true; // temporarily !
final boolean updateStatus,
final int debugLevel) throws Exception
{
if ((quadCLT_main != null) && (quadCLT_main.getGPU() != null)) {
quadCLT_main.getGPU().resetGeometryCorrection();
quadCLT_main.getGPU().resetGeometryCorrectionVector();
}
if ((quadCLT_aux != null) && (quadCLT_aux.getGPU() != null)) {
quadCLT_aux.getGPU().resetGeometryCorrection();
quadCLT_aux.getGPU().resetGeometryCorrectionVector();
}
// final boolean batch_mode = clt_parameters.batch_run;
// Reset dsi data (only 2 slices will be used)
this.dsi = new double [DSI_SLICES.length][];
......@@ -8833,6 +8258,7 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_imp_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevelInner); // int debugLevel);
if (updateStatus) IJ.showStatus("Conditioning aux camera image set for "+quadCLT_main.image_name);
......@@ -8883,22 +8309,38 @@ if (debugLevel > -100) return true; // temporarily !
}
// Generate 4 main camera images and thumbnail
if (quadCLT_main.correctionsParameters.clt_batch_4img){
if (updateStatus) IJ.showStatus("Rendering 4 image set (disparity = 0) for "+quadCLT_main.image_name);
if (clt_parameters.gpu_use_main) {
if (updateStatus) IJ.showStatus("GPU: Rendering 4 image set (disparity = 0) for "+quadCLT_main.image_name+ "and a thumb nail");
quadCLT_main.processCLTQuadCorrGPU(
imp_srcs_main, // ImagePlus [] imp_quad,
saturation_imp_main, // boolean [][] saturation_imp, // (near) saturated pixels or null
clt_parameters, // CLTParameters clt_parameters,
debayerParameters, // EyesisCorrectionParameters.DebayerParameters debayerParameters,
colorProcParameters, // ColorProcParameters colorProcParameters,
channelGainParameters,
rgbParameters, // EyesisCorrectionParameters.RGBParameters rgbParameters,
scaleExposures_main, // double [] scaleExposures, // probably not needed here - restores brightness of the final image
threadsMax, // final int threadsMax, // maximal number of threads to launch
updateStatus, // final boolean updateStatus,
debugLevel); // final int debugLevel);
} else {
if (updateStatus) IJ.showStatus("CPU: Rendering 4 image set (disparity = 0) for "+quadCLT_main.image_name+ "and a thumb nail");
quadCLT_main.processCLTQuadCorr( // returns ImagePlus, but it already should be saved/shown
imp_srcs_main, // [srcChannel], // should have properties "name"(base for saving results), "channel","path"
saturation_imp_main, // boolean [][] saturation_imp, // (near) saturated pixels or null
clt_parameters,
debayerParameters,
colorProcParameters,
channelGainParameters,
rgbParameters,
scaleExposures_main,
false, // calculate and apply additional fine geometry correction
false, // calculate and apply geometry correction at infinity
threadsMax, // maximal number of threads to launch
updateStatus,
debugLevel);
quadCLT_main.processCLTQuadCorrCPU( // returns ImagePlus, but it already should be saved/shown
imp_srcs_main, // [srcChannel], // should have properties "name"(base for saving results), "channel","path"
saturation_imp_main, // boolean [][] saturation_imp, // (near) saturated pixels or null
clt_parameters,
debayerParameters,
colorProcParameters,
channelGainParameters,
rgbParameters,
scaleExposures_main,
false, // calculate and apply additional fine geometry correction
false, // calculate and apply geometry correction at infinity
threadsMax, // maximal number of threads to launch
updateStatus,
debugLevel);
}
quadCLT_main.tp.resetCLTPasses();
}
......@@ -9050,6 +8492,7 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_imp_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevelInner); // int debugLevel);
// optionally adjust AUX extrinsics (using quadCLT_aux.ds_from_main )
......@@ -9105,7 +8548,7 @@ if (debugLevel > -100) return true; // temporarily !
if (quadCLT_main.correctionsParameters.clt_batch_4img_aux){
if (updateStatus) IJ.showStatus("Rendering 4 AUX image set (disparity = 0) for "+quadCLT_aux.image_name);
quadCLT_aux.processCLTQuadCorr( // returns ImagePlus, but it already should be saved/shown
quadCLT_aux.processCLTQuadCorrCPU( // returns ImagePlus, but it already should be saved/shown
imp_srcs_aux, // [srcChannel], // should have properties "name"(base for saving results), "channel","path"
saturation_imp_aux, // boolean [][] saturation_imp, // (near) saturated pixels or null
clt_parameters,
......@@ -9193,7 +8636,7 @@ if (debugLevel > -100) return true; // temporarily !
return;
}
}
System.out.println("processCLTQuadCorrs(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
System.out.println("batchLwirRig(): processing "+(quadCLT_main.getTotalFiles(set_channels_main)+quadCLT_aux.getTotalFiles(set_channels_aux))+" files ("+set_channels_main.length+" file sets) finished at "+
IJ.d2s(0.000000001*(System.nanoTime()-this.startTime),3)+" sec, --- Free memory="+Runtime.getRuntime().freeMemory()+" (of "+Runtime.getRuntime().totalMemory()+")");
}
......@@ -9609,6 +9052,7 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_main, // int [] channelFiles,
scaleExposures_main, //output // double [] scaleExposures
saturation_main, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
ImagePlus [] imp_quad_aux = quadCLT_aux.conditionImageSet(
......@@ -9620,8 +9064,9 @@ if (debugLevel > -100) return true; // temporarily !
channelFiles_aux, // int [] channelFiles,
scaleExposures_aux, //output // double [] scaleExposures
saturation_aux, //output // boolean [][] saturation_imp,
threadsMax, // int threadsMax,
debugLevel); // int debugLevel);
/* // 08/12/2020 Moved to conditionImageSet
getRigImageStacks(
clt_parameters, // EyesisCorrectionParameters.CLTParameters clt_parameters,
quadCLT_main, // QuadCLT quadCLT_main,
......@@ -9633,6 +9078,8 @@ if (debugLevel > -100) return true; // temporarily !
threadsMax, // maximal number of threads to launch
debugLevel); // final int debugLevel);
// now tp is defined
*/
/*
main_dsi = enhanceMainDSI(
main_dsi, // double [][] main_dsi,
......
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