Skip to content

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

Adding channels to overlapped textures

parent 7780b796
Show whitespace changes
Inline Side-by-side
Showing with 170 additions and 51 deletions
src/main/java/com/elphel/imagej/cameras/CLTParameters.java
View file @ f6006466
......@@ -135,10 +135,11 @@ public class CLTParameters {
public double diff_threshold = 1.5; // RMS difference from average to discard channel (~ 1.0 - 1/255 full scale image)
public boolean diff_gauss = true; // when averaging images, use Gaussian around average as weight (false - sharp all/nothing)
public double min_agree = 3.0; // minimal number of channels to agree on a point (real number to work with fuzzy averages)
public boolean dust_remove = true; // Do not reduce average weight when only one image differes much from the average
public boolean dust_remove = true; // Do not reduce average weight when only one image differs much from the average
public boolean black_back = true; // use Black for backdrop outside of the FOV
public boolean keep_weights = true; // add port weights to RGBA stack (debug feature)
public boolean replace_weights = false; //replace port weights with ports textures (16x16)
public boolean sharp_alpha = false; // combining mode for alpha channel: false - treat as RGB, true - apply center 8x8 only
public double alpha0 = 0.6; // > .525 Alpha channel 0.0 threshold (lower - transparent) (watch for far objects)
public double alpha1 = 0.8; // Alpha channel 1.0 threshold (higher - opaque) (watch for window dust)
......@@ -1195,6 +1196,7 @@ public class CLTParameters {
properties.setProperty(prefix+"dust_remove", this.dust_remove+"");
properties.setProperty(prefix+"black_back", this.black_back+"");
properties.setProperty(prefix+"keep_weights", this.keep_weights+"");
properties.setProperty(prefix+"replace_weights", this.replace_weights+"");
properties.setProperty(prefix+"sharp_alpha", this.sharp_alpha+"");
properties.setProperty(prefix+"alpha0", this.alpha0 +"");
......@@ -2095,6 +2097,7 @@ public class CLTParameters {
if (properties.getProperty(prefix+"dust_remove")!=null) this.dust_remove=Boolean.parseBoolean(properties.getProperty(prefix+"dust_remove"));
if (properties.getProperty(prefix+"black_back")!=null) this.black_back=Boolean.parseBoolean(properties.getProperty(prefix+"black_back"));
if (properties.getProperty(prefix+"keep_weights")!=null) this.keep_weights=Boolean.parseBoolean(properties.getProperty(prefix+"keep_weights"));
if (properties.getProperty(prefix+"replace_weights")!=null) this.replace_weights=Boolean.parseBoolean(properties.getProperty(prefix+"replace_weights"));
if (properties.getProperty(prefix+"sharp_alpha")!=null) this.sharp_alpha=Boolean.parseBoolean(properties.getProperty(prefix+"sharp_alpha"));
if (properties.getProperty(prefix+"alpha0")!=null) this.alpha0=Double.parseDouble(properties.getProperty(prefix+"alpha0"));
if (properties.getProperty(prefix+"alpha1")!=null) this.alpha1=Double.parseDouble(properties.getProperty(prefix+"alpha1"));
......@@ -3033,6 +3036,7 @@ public class CLTParameters {
gd.addCheckbox ("Do not reduce average weight when only one image differes much from the average", this.dust_remove);
gd.addCheckbox ("Use black for backdrop outside of the FOV", this.black_back);
gd.addCheckbox ("Add port weights to RGBA stack (debug feature)", this.keep_weights);
gd.addCheckbox ("Replace port weights with per-channel MIDCT (16x16) - were relevant", this.replace_weights);
gd.addCheckbox ("Alpha channel: use center 8x8 (unchecked - treat same as RGB)", this.sharp_alpha);
gd.addNumericField("Alpha channel 0.0 thereshold (lower - transparent)", this.alpha0, 3);
gd.addNumericField("Alpha channel 1.0 threshold (higher - opaque)", this.alpha1, 3);
......@@ -4238,6 +4242,7 @@ public class CLTParameters {
this.dust_remove= gd.getNextBoolean();
this.black_back= gd.getNextBoolean();
this.keep_weights= gd.getNextBoolean();
this.replace_weights= gd.getNextBoolean();
this.sharp_alpha= gd.getNextBoolean();
this.alpha0= gd.getNextNumber();
this.alpha1= gd.getNextNumber();
......
src/main/java/com/elphel/imagej/gpu/GpuQuad.java
View file @ f6006466
......@@ -190,7 +190,7 @@ public class GpuQuad{ // quad camera description
int host_get_textures_shared_size( // in bytes
static int host_get_textures_shared_size( // in bytes
//__device__ int get_textures_shared_size( // in bytes
int num_cams, // actual number of cameras
int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
......@@ -440,7 +440,8 @@ public class GpuQuad{ // quad camera description
texture_stride = (int)(device_stride[0] / Sizeof.FLOAT);
int max_rgba_width = (tilesX + 1) * GPUTileProcessor.DTT_SIZE;
int max_rgba_height = (tilesY + 1) * GPUTileProcessor.DTT_SIZE;
int max_rbga_slices = num_colors + 1;
// int max_rbga_slices = num_colors + 1; // num_cams
int max_rbga_slices = num_colors + 1 + (num_colors * num_cams); // colors, alpha, colors per channel
cuMemAllocPitch (
gpu_textures_rgba, // CUdeviceptr dptr,
device_stride, // long[] pPitch,
......@@ -2046,7 +2047,8 @@ public class GpuQuad{ // quad camera description
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,
int keep_weights) {
if (GPUTileProcessor.USE_DS_DP) {
execRBGA_DP(
color_weights, // double [] color_weights,
......@@ -2056,7 +2058,8 @@ public class GpuQuad{ // quad camera description
diff_sigma, // double diff_sigma, // pixel value/pixel change
diff_threshold, // double diff_threshold, // pixel value/pixel change
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove); // boolean dust_remove);
dust_remove, // boolean dust_remove);
keep_weights); // int keep_weights)
} else {
execRBGA_noDP(
color_weights, // double [] color_weights,
......@@ -2066,7 +2069,8 @@ public class GpuQuad{ // quad camera description
diff_sigma, // double diff_sigma, // pixel value/pixel change
diff_threshold, // double diff_threshold, // pixel value/pixel change
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove); // boolean dust_remove);
dust_remove, // boolean dust_remove);
keep_weights); // int keep_weights)
}
}
/**
......@@ -2079,6 +2083,8 @@ public class GpuQuad{ // quad camera description
* @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 keep_weights (since 11/13/2022). Now 2 separate bits: +1 - generate channel weights and combo metrics.
* +2 replace port_weights with raw per-channel (+1 is not used here for overlapping)
*/
public void execRBGA_DP(
double [] color_weights,
......@@ -2088,7 +2094,8 @@ public class GpuQuad{ // quad camera description
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,
int keep_weights) {
execCalcReverseDistortions(); // will check if it is needed first
if (this.gpuTileProcessor.GPU_RBGA_kernel == null) {
IJ.showMessage("Error", "No GPU kernel: GPU_TEXTURES_kernel");
......@@ -2139,7 +2146,7 @@ public class GpuQuad{ // quad camera description
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[] {keep_weights}), // 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
......@@ -2165,6 +2172,8 @@ public class GpuQuad{ // quad camera description
* @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 keep_weights (since 11/13/2022). Now 2 separate bits: +1 - generate channel weights and combo metrics.
* +2 replace port_weights with raw per-channel (+1 is not used here for overlapping)
*/
public void execRBGA_noDP(
......@@ -2175,7 +2184,8 @@ public class GpuQuad{ // quad camera description
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,
int keep_weights) {
execCalcReverseDistortions(); // will check if it is needed first
if ( (this.gpuTileProcessor.GPU_CLEAR_TEXTURE_LIST_kernel == null) &&
(this.gpuTileProcessor.GPU_MARK_TEXTURE_LIST_kernel == null) &&
......@@ -2478,7 +2488,7 @@ public class GpuQuad{ // quad camera description
Pointer.to(new float[] {(float) min_agree}), // float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
Pointer.to(gpu_color_weights), // float weights[3], // scale for R,B,G (or {1.0,0.0,0.0}
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[] {keep_weights}), // int keep_weights, // return channel weights after A in RGBA
// combining both non-overlap and overlap (each calculated if pointer is not null )
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
......@@ -2512,6 +2522,8 @@ public class GpuQuad{ // quad camera description
* @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 keep_weights (since 11/13/2022). Now 2 separate bits: +1 - generate channel weights and combo metrics.
* +2 replace port_weights with raw per-channel
* @param calc_textures calculate textures (false for macro-only output)
* @param calc_extra calculate extra output - low-res for macro
*/
......@@ -2524,6 +2536,7 @@ public class GpuQuad{ // quad camera description
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, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // 2 bits now, move to parameters
boolean calc_textures,
boolean calc_extra,
boolean linescan_order
......@@ -2538,6 +2551,7 @@ public class GpuQuad{ // quad camera description
diff_threshold, // pixel value/pixel change
min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // Do not reduce average weight when only one image differs much from the average
keep_weights, // int keep_weights, // 2 bits now, move to parameters
calc_textures,
calc_extra,
linescan_order);
......@@ -2551,6 +2565,7 @@ public class GpuQuad{ // quad camera description
diff_threshold, // pixel value/pixel change
min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // Do not reduce average weight when only one image differs much from the average
keep_weights, // int keep_weights, // 2 bits now, move to parameters
calc_textures,
calc_extra,
linescan_order);
......@@ -2566,6 +2581,7 @@ public class GpuQuad{ // quad camera description
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, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // 2 bits now, move to parameters
boolean calc_textures,
boolean calc_extra,
boolean linescan_order)
......@@ -2597,7 +2613,7 @@ public class GpuQuad{ // quad camera description
int iis_lwir = (is_lwir)? 1:0;
int ilinescan_order = linescan_order? 1 : 0;
int idust_remove = (dust_remove)? 1 : 0;
// int keep_weights = 0; // 2 bits now, move to parameters
int [] GridFullWarps = {1, 1, 1};
int [] ThreadsFullWarps = {1, 1, 1};
......@@ -2618,6 +2634,7 @@ public class GpuQuad{ // quad camera description
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[] { keep_weights }), // +1 : "keep_weights", +2 - replace port_weights with channel output
Pointer.to(new int[] {calc_textures? texture_stride : 0}),
Pointer.to(gpu_textures),
Pointer.to(new int[] {ilinescan_order}), // 1, // int linescan_order, // if !=0 then output gpu_diff_rgb_combo in linescan order, else - in gpu_texture_indices order
......@@ -2642,6 +2659,7 @@ public class GpuQuad{ // quad camera description
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, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // 2 bits now, move to parameters
boolean calc_textures,
boolean calc_extra,
boolean linescan_order)
......@@ -2656,7 +2674,7 @@ public class GpuQuad{ // quad camera description
IJ.showMessage("Error", "No GPU kernel(s)");
return;
}
int keep_texture_weights = 0; // pass as parameter?
int keep_texture_weights = keep_weights; // pass as parameter?
int tilesX = img_width / GPUTileProcessor.DTT_SIZE;
int num_colors = is_lwir? 1 : color_weights.length;
if (num_colors > 3) num_colors = 3;
......@@ -3364,7 +3382,7 @@ public class GpuQuad{ // quad camera description
return textures;
}
public double [][][][] doubleTextures(
public static double [][][][] doubleTextures( // not used
Rectangle woi,
int [] indices,
float [][][] ftextures,
......@@ -3390,7 +3408,7 @@ public class GpuQuad{ // quad camera description
return textures;
}
public double [][][][] doubleTextures( // may be accelerated with multithreading if needed.
public static double [][][][] doubleTextures( // may be accelerated with multithreading if needed.
Rectangle woi, // null or width and height match texture_tiles
double [][][][] texture_tiles, // null or [tilesY][tilesX]
int [] indices,
......
src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java
View file @ f6006466
......@@ -356,7 +356,8 @@ public class ImageDtt extends ImageDttCPU {
diff_sigma, // double diff_sigma, // pixel value/pixel change
diff_threshold, // double diff_threshold, // pixel value/pixel change
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove); // boolean dust_remove,
dust_remove, // boolean dust_remove,
0); // int keep_weights)
float [][] rbga = gpuQuad.getRBGA(
(isMonochrome() ? 1 : 3), // int num_colors,
(texture_woi_pix != null)? texture_woi_pix : woi);
......@@ -376,6 +377,7 @@ public class ImageDtt extends ImageDttCPU {
diff_threshold, // double diff_threshold, // pixel value/pixel change - never used in GPU ?
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // boolean dust_remove, // Do not reduce average weight when only one image differs much from the average
0, // int keep_weights, // 2 bits now, move to parameters
false, // boolean calc_textures,
true, // boolean calc_extra
false); // boolean linescan_order) // TODO: use true to avoid reordering of the low-res output
......@@ -427,6 +429,7 @@ public class ImageDtt extends ImageDttCPU {
diff_threshold, // double diff_threshold, // pixel value/pixel change - never used in GPU ?
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // boolean dust_remove, // Do not reduce average weight when only one image differs much from the average
0, // int keep_weights, // 2 bits now, move to parameters
true, // boolean calc_textures,
false, // boolean calc_extra
false); // boolean linescan_order)
......@@ -438,7 +441,7 @@ public class ImageDtt extends ImageDttCPU {
numcol, // int num_colors,
false); // clt_parameters.keep_weights); // boolean keep_weights);
gpuQuad.doubleTextures(
GpuQuad.doubleTextures(
new Rectangle(0, 0, tilesX, tilesY), // Rectangle woi,
texture_tiles, // double [][][][] texture_tiles, // null or [tilesY][tilesX]
texture_indices, // int [] indices,
......@@ -663,7 +666,8 @@ public class ImageDtt extends ImageDttCPU {
double diff_sigma, // pixel value/pixel change Used much larger sigma = 10.0 instead of 1.5
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,
int keep_weights // 2 bits now, move to parameters
){
int numcol = isMonochrome()? 1 : 3;
double [] col_weights = new double[numcol];
......@@ -684,6 +688,7 @@ public class ImageDtt extends ImageDttCPU {
diff_threshold, // double diff_threshold, // pixel value/pixel change
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // boolean dust_remove,
keep_weights, // int keep_weights, // 2 bits now, move to parameters
true, // boolean calc_textures,
false, // boolean calc_extra
false); // boolean linescan_order)
......@@ -696,7 +701,7 @@ public class ImageDtt extends ImageDttCPU {
int tilesX = gpuQuad.img_width / GPUTileProcessor.DTT_SIZE;
int tilesY = gpuQuad.img_height / GPUTileProcessor.DTT_SIZE;
double [][][][] texture_tiles = new double [tilesY][tilesX][][];
gpuQuad.doubleTextures(
GpuQuad.doubleTextures(
new Rectangle(0, 0, tilesX, tilesY), // Rectangle woi,
texture_tiles, // double [][][][] texture_tiles, // null or [tilesY][tilesX]
texture_indices, // int [] indices,
......@@ -744,6 +749,7 @@ public class ImageDtt extends ImageDttCPU {
diff_threshold, // double diff_threshold, // pixel value/pixel change
min_agree, // double min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
dust_remove, // boolean dust_remove,
0, // int keep_weights, // 2 bits now, move to parameters
false, // boolean calc_textures,
true, // boolean calc_extra
false); // boolean linescan_order)
......
src/main/java/com/elphel/imagej/tileprocessor/QuadCLT.java
View file @ f6006466
......@@ -1655,7 +1655,9 @@ public class QuadCLT extends QuadCLTCPU {
clt_parameters.diff_sigma, // double diff_sigma, // pixel value/pixel change Used much larger sigma = 10.0 instead of 1.5
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
clt_parameters.dust_remove, // boolean dust_remove
0); // int keep_weights // 2 bits now, move to parameters
if (max_distortion > 0) {// remove distorted tiles
double max_distortion2 = max_distortion * max_distortion;
......@@ -2225,6 +2227,10 @@ public class QuadCLT extends QuadCLTCPU {
col_weights[1] = clt_parameters.corr_blue * col_weights[2];
}
Rectangle woi = new Rectangle(); // will be filled out to match actual available image
// int keep_weights = (clt_parameters.keep_weights? 1 : 0) + (clt_parameters.replace_weights? 2 : 0);
int keep_weights = (clt_parameters.replace_weights? 2 : 0); // just 0/2
int text_colors = (isMonochrome() ? 1 : 3);
int texture_layers = (text_colors + 1)+((keep_weights != 0)?(text_colors * getNumSensors()):0);
gpuQuad.execRBGA(
col_weights, // double [] color_weights,
isLwir(), // boolean is_lwir,
......@@ -2233,9 +2239,10 @@ public class QuadCLT extends QuadCLTCPU {
clt_parameters.diff_sigma, // double diff_sigma, // pixel value/pixel change Used much larger sigma = 10.0 instead of 1.5
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,
clt_parameters.dust_remove, // boolean dust_remove,
keep_weights); // int keep_weights)
float [][] rbga = gpuQuad.getRBGA(
(isMonochrome() ? 1 : 3), // int num_colors,
texture_layers - 1, // (isMonochrome() ? 1 : 3), // int num_colors,
woi);
for (int ncol = 0; ncol < texture_img.length; ncol++) if (ncol < rbga.length) {
texture_img[ncol] = rbga[ncol];
......@@ -2301,6 +2308,7 @@ public class QuadCLT extends QuadCLTCPU {
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,
0, // int keep_weights, // 2 bits now, move to parameters
false, // boolean calc_textures,
true, // boolean calc_extra
false); // boolean linescan_order) // TODO: use true to avoid reordering of the low-res output
......@@ -2320,6 +2328,7 @@ public class QuadCLT extends QuadCLTCPU {
}
if (try_textures) {
int keep_weights = (clt_parameters.keep_weights? 1 : 0) + (clt_parameters.replace_weights? 2 : 0);
//Generate non-overlapping (16x16) texture tiles, prepare
gpuQuad.execTextures(
col_weights, // double [] color_weights,
......@@ -2330,27 +2339,33 @@ public class QuadCLT extends QuadCLTCPU {
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,
keep_weights, // int keep_weights, // 2 bits now, move to parameters
true, // boolean calc_textures,
false, // boolean calc_extra
false); // boolean linescan_order)
int [] texture_indices = gpuQuad.getTextureIndices();
int numcol = isMonochrome()? 1 : 3;
int num_src_slices = numcol + 1; // + (clt_parameters.keep_weights?(ports + numcol + 1):0); // 12 ; // calculate
// int num_src_slices = numcol + 1; // + (clt_parameters.keep_weights?(ports + numcol + 1):0); // 12 ; // calculate
int num_src_slices = numcol + 1 + (clt_parameters.keep_weights?(getNumSensors() + numcol + 1):0); // 12 ; // calculate
float [] flat_textures = gpuQuad.getFlatTextures( // fatal error has been detected by the Java Runtime Environment:
texture_indices.length,
numcol, // int num_colors,
false); // clt_parameters.keep_weights); // boolean keep_weights);
(keep_weights != 0)); // false); // clt_parameters.keep_weights); // boolean keep_weights);
if (keep_weights != 0) {
}
// numcol + 1 + (clt_parameters.keep_weights?(getNumSensors() + numcol + 1):0); // 12 ; // calculate
int tilesX = gpuQuad.img_width / GPUTileProcessor.DTT_SIZE;
int tilesY = gpuQuad.img_height / GPUTileProcessor.DTT_SIZE;
double [][][][] texture_tiles = new double [tilesY][tilesX][][];
gpuQuad.doubleTextures(
GpuQuad.doubleTextures(
new Rectangle(0, 0, tilesX, tilesY), // Rectangle woi,
texture_tiles, // double [][][][] texture_tiles, // null or [tilesY][tilesX]
texture_indices, // int [] indices,
flat_textures, // float [][][] ftextures,
tilesX, // int full_width,
isMonochrome()? 2: 4, // rbga only /int num_slices
num_src_slices, // isMonochrome()? 2: 4, // rbga only /int num_slices
num_src_slices // int num_src_slices
);
int num_out_slices = 0;
......@@ -2361,6 +2376,29 @@ public class QuadCLT extends QuadCLTCPU {
}
}
if (num_out_slices > 0) {
String [] dbg_titles = new String[num_out_slices];
fill_titles:
{
int indx = 0;
for (int i = 0; i < numcol; i++) {
if (indx >= num_out_slices) break fill_titles;
else dbg_titles[indx++] = "c"+i;
}
if (indx >= num_out_slices) break fill_titles;
else dbg_titles[indx++] = "alpha";
for (int i = 0; i < getNumSensors(); i++) {
if (indx >= num_out_slices) break fill_titles;
else dbg_titles[indx++] = (((keep_weights & 2) != 0)?"t":"w")+i;
}
for (int i = 0; i < numcol; i++) {
if (indx >= num_out_slices) break fill_titles;
else dbg_titles[indx++] = "crms"+i;
}
if (indx >= num_out_slices) break fill_titles;
else dbg_titles[indx++] = "crms";
}
//crms
int ssize = 2*GPUTileProcessor.DTT_SIZE;
int width = tilesX * ssize;
int height = tilesY * ssize;
......@@ -2390,8 +2428,8 @@ public class QuadCLT extends QuadCLTCPU {
width,
height,
true,
getImageName()+"-textures"
);
getImageName()+"-textures-"+keep_weights,
dbg_titles);
}
System.out.println("try_textures DONE");
}
......@@ -2661,6 +2699,7 @@ public class QuadCLT extends QuadCLTCPU {
clt_parameters.diff_threshold, // double diff_threshold, // pixel value/pixel change - never used in GPU ?
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, // Do not reduce average weight when only one image differs much from the average
0, // int keep_weights, // 2 bits now, move to parameters
calc_textures, // boolean calc_textures,
calc_extra, // boolean calc_extra)
false); // boolean linescan_order) // TODO: use true to avoid reordering of the low-res output
......@@ -2680,7 +2719,9 @@ public class QuadCLT extends QuadCLTCPU {
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,
clt_parameters.dust_remove, // boolean dust_remove,
0); // int keep_weights)
long endTexturesRBGA = System.nanoTime();
long endGPUTime = System.nanoTime();
......@@ -2973,7 +3014,7 @@ public class QuadCLT extends QuadCLTCPU {
String results_path= quadCLT_main.correctionsParameters.selectResultsDirectory( // selectX3dDirectory(
true, // smart,
true); //newAllowed, // save
quadCLT_main.eyesisCorrections.saveAndShow( // save and show color RGBA texture
EyesisCorrections.saveAndShow( // save and show color RGBA texture
imp_texture_stack,
results_path,
true, // quadCLT_main.correctionsParameters.png && !clt_parameters.black_back,
......@@ -3021,7 +3062,8 @@ public class QuadCLT extends QuadCLTCPU {
}
}
double [][][][] texture_tiles = new double [tilesY][tilesX][][];
quadCLT_main.getGPU().doubleTextures(
quadCLT_main.getGPU();
GpuQuad.doubleTextures(
new Rectangle(0, 0, tilesX, tilesY), // Rectangle woi,
texture_tiles, // double [][][][] texture_tiles, // null or [tilesY][tilesX]
texture_indices, // int [] indices,
......
src/main/java/com/elphel/imagej/tileprocessor/QuadCLTCPU.java
View file @ f6006466
......@@ -13730,7 +13730,8 @@ public class QuadCLTCPU {
clt_parameters.diff_sigma, // double diff_sigma, // pixel value/pixel change Used much larger sigma = 10.0 instead of 1.5
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
clt_parameters.dust_remove, // boolean dust_remove
0); // int keep_weights // 2 bits now, move to parameters
}
if (save_diff || save_lowres) {
image_dtt.get_diffs_lowres( // CUDA_ERROR_INVALID_VALUE (because no tiles)
......
src/main/resources/kernels/TileProcessor.cuh
View file @ f6006466
......@@ -2350,13 +2350,16 @@ extern "C" __global__ void generate_RBGA(
int texture_width = (*(woi + 2) + 1)* DTT_SIZE;
int texture_tiles_height = (*(woi + 3) + 1) * DTT_SIZE;
int texture_slices = colors + 1;
if (keep_weights & 2){
texture_slices += colors * num_cams;
}
if (threadIdx.x == 0) {
dim3 threads2((1 << THREADS_DYNAMIC_BITS), 1, 1);
int blocks_x = (texture_width + ((1 << (THREADS_DYNAMIC_BITS + DTT_SIZE_LOG2 )) - 1)) >> (THREADS_DYNAMIC_BITS + DTT_SIZE_LOG2);
dim3 blocks2 (blocks_x, texture_tiles_height * texture_slices, 1); // each thread - 8 vertical
clear_texture_rbga<<<blocks2,threads2>>>( // illegal value error
clear_texture_rbga<<<blocks2,threads2>>>( // add clearing of multi-sensor output (keep_weights & 2 !=0)
texture_width,
texture_tiles_height * texture_slices, // int texture_slice_height,
texture_rbga_stride, // const size_t texture_rbga_stride, // in floats 8*stride
......@@ -2378,7 +2381,7 @@ extern "C" __global__ void generate_RBGA(
dim3 grid_texture((ntt + TEXTURE_TILES_PER_BLOCK-1) / TEXTURE_TILES_PER_BLOCK,1,1); // TEXTURE_TILES_PER_BLOCK = 1
int ti_offset = (pass & 3) * (width * (tilesya >> 2)); // (TILES-X * (TILES-YA >> 2)); // 1/4
if (border_tile){
ti_offset += width * (tilesya >> 2) - ntt;; // TILES-X * (TILES-YA >> 2) - ntt;
ti_offset += width * (tilesya >> 2) - ntt; // TILES-X * (TILES-YA >> 2) - ntt;
}
#ifdef DEBUG12
printf("\ngenerate_RBGA() pass= %d, border_tile= %d, ti_offset= %d, ntt=%d\n",
......@@ -2413,7 +2416,7 @@ extern "C" __global__ void generate_RBGA(
min_agree, // float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
weights, // float weights[3], // scale for R,B,G
dust_remove, // int dust_remove, // Do not reduce average weight when only one image differs much from the average
0, // int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
keep_weights, // int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
// combining both non-overlap and overlap (each calculated if pointer is not null )
texture_rbga_stride, // size_t texture_rbg_stride, // in floats
gpu_texture_tiles, // float * gpu_texture_rbg, // (number of colors +1 + ?)*16*16 rgba texture tiles
......@@ -3218,6 +3221,7 @@ __global__ void convert_correct_tiles(
* min_agree minimal number of channels to agree on a point (real number to work with fuzzy averages) (3.0)
* @param weights scales for R,B,G {0.294118, 0.117647, 0.588235}
* @param dust_remove do not reduce average weight when only one image differs much from the average (true)
* @param keep_weights Was not here before 10/12/2022. return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
* @param texture_stride output stride in floats (now 256*4 = 1024)
* @param gpu_texture_tiles output array (number of colors +1 + ?)*16*16 rgba texture tiles) float values. Will not be calculated if null
* @param inescan_order 0 low-res tiles have the same order, as gpu_texture_indices, 1 - in linescan order
......@@ -3228,7 +3232,6 @@ extern "C" __global__ void textures_nonoverlap(
int num_cams, // number of cameras
float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats
int num_tiles, // number of tiles in task list
// int num_tilesx, // number of tiles in a row
// declare arrays in device code?
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int * pnum_texture_tiles, // returns total number of elements in gpu_texture_indices array
......@@ -3240,6 +3243,7 @@ extern "C" __global__ void textures_nonoverlap(
float params[5],
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // Was not here before 10/12/2022. return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
// combining both non-overlap and overlap (each calculated if pointer is not null )
size_t texture_stride, // in floats (now 256*4 = 1024)
float * gpu_texture_tiles, // (number of colors +1 + ?)*16*16 rgba texture tiles
......@@ -3302,7 +3306,7 @@ extern "C" __global__ void textures_nonoverlap(
min_agree, // float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
weights, // float weights[3], // scale for R,B,G
dust_remove, // int dust_remove, // Do not reduce average weight when only one image differs much from the average
0, // int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
keep_weights, // int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
// combining both non-overlap and overlap (each calculated if pointer is not null )
0, // size_t texture_rbg_stride, // in floats
(float *) 0, // float * gpu_texture_rbg, // (number of colors +1 + ?)*16*16 rgba texture tiles
......@@ -3339,7 +3343,7 @@ extern "C" __global__ void textures_nonoverlap(
* @param min_agree minimal number of channels to agree on a point (real number to work with fuzzy averages) (3.0)
* @param weights scales for R,B,G {0.294118, 0.117647, 0.588235}
* @param dust_remove do not reduce average weight when only one image differs much from the average (true)
* @param keep_weights return channel weights after A in RGBA (was removed)
* @param keep_weights return channel weights after A in RGBA (was removed). Now (11/12/2022): +1 - old meaning, +2 - replace port_weights with channel imclt
* @param texture_rbg_stride output stride for overlapped texture in floats, or 0 to skip
* @param gpu_texture_rbg output array (number of colors +1 + ?) * woi.height * output stride(first woi.width valid) float values (or 0)
* @param texture_stride output stride for non-overlapping texture tile output in floats (or 0 to skip)
......@@ -3366,7 +3370,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)? Now +2 - output raw channels
// combining both non-overlap and overlap (each calculated if pointer is not null )
size_t texture_rbg_stride, // in floats
float * gpu_texture_rbg, // (number of colors +1 + ?)*16*16 rgba texture tiles
......@@ -3433,7 +3437,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
float * mclt_tiles = &all_shared[offsets[0]] ; // [num_cams][colors][2*DTT_SIZE][DTT_SIZE21]; // 16*1*16*17=0x1100 | 4*3*16*17=0xcc0
float * clt_tiles = &all_shared[offsets[1]] ; // [num_cams][colors][4][DTT_SIZE][DTT_SIZE1]; // 16 * 1 * 4 * 8 * 9 = 0x1200 | 4 * 3 * 4 * 8 * 9 = 0xd80
float * mclt_debayer = &all_shared[offsets[1]] ; // [num_cams][colors][MCLT_UNION_LEN]; // 16 * 1 * 16 * 18 = 0x1200 | 4 * 3 * 16 * 18 = 0xd80 | to align with clt_tiles
float * mclt_tmps = &all_shared[offsets[2]] ; // [num_cams][colors][DTT_SIZE2][DTT_SIZE21]; // 16*1*16*17=0x1100 | 4*3*16*17=0xcc0
float * mclt_tmps = &all_shared[offsets[2]] ; // [num_cams][colors][DTT_SIZE2][DTT_SIZE21]; // 16*1*16*17=0x1100 | 4*3*16*17=0xcc0. Used only with Bayer, not with mono
float * rgbaw = &all_shared[offsets[2]] ; // [colors + 1 + num_cams + colors + 1][DTT_SIZE2][DTT_SIZE21];
float * port_offsets = &all_shared[offsets[3]] ; // [num_cams][2]; // 16 * 2 = 0x20 | 4*2 = 0x8
......@@ -3691,7 +3695,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
min_agree, // float min_agree, NOT USED? // minimal number of channels to agree on a point (real number to work with fuzzy averages)
weights, // float * chn_weights, // color channel weights, sum == 1.0
dust_remove, // int dust_remove, // Do not reduce average weight when only one image differs much from the average
keep_weights, // int keep_weights, // return channel weights and rms after A in RGBA (weight are always calculated)
(keep_weights & 1), // int keep_weights, // return channel weights and rms after A in RGBA (weight are always calculated)
debug ); // int debug );
__syncthreads(); // _syncthreads();1
......@@ -3712,7 +3716,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
float * rgba_i = rgbaw + i;
// always copy 3 (1) colors + alpha
if (colors == 3){
if (keep_weights) {
if (keep_weights & 1) {
for (int ncol = 0; ncol < colors + 1 + num_cams + colors + 1 ; ncol++) { // 12
*(gpu_texture_tile_gi + ncol * (DTT_SIZE2 * DTT_SIZE2)) = *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
......@@ -3721,8 +3725,15 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
*(gpu_texture_tile_gi + ncol * (DTT_SIZE2 * DTT_SIZE2)) = *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
}
if (keep_weights & 2) {
float * mclt_dst_i = mclt_debayer + i;
float * gpu_texture_tile_raw_gi = gpu_texture_tile_gi + (colors + 1) * (DTT_SIZE2 * DTT_SIZE2); // skip colors + alpha
for (int ncam = 0; ncam < num_cams ; ncam++) { // 8
*(gpu_texture_tile_raw_gi + ncam * (DTT_SIZE2 * DTT_SIZE2)) = *(mclt_dst_i + (ncam * 3 + 2) * (MCLT_UNION_LEN)); // green colors
}
}
} else { // assuming colors = 1
if (keep_weights) {
if (keep_weights & 1) {
for (int ncol = 0; ncol < 1 + 1 + num_cams + 1 + 1 ; ncol++) { // 8
*(gpu_texture_tile_gi + ncol * (DTT_SIZE2 * DTT_SIZE2)) = *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
......@@ -3731,8 +3742,14 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
*(gpu_texture_tile_gi + ncol * (DTT_SIZE2 * DTT_SIZE2)) = *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
}
if (keep_weights & 2) {
float * mclt_dst_i = mclt_debayer + i;
float * gpu_texture_tile_raw_gi = gpu_texture_tile_gi + (colors + 1) * (DTT_SIZE2 * DTT_SIZE2); // skip colors + alpha
for (int ncam = 0; ncam < num_cams ; ncam++) { // 8
*(gpu_texture_tile_raw_gi + ncam * (DTT_SIZE2 * DTT_SIZE2)) = *(mclt_dst_i + (ncam * 1 + 0) * (MCLT_UNION_LEN));
}
}
}
}
#ifdef DEBUG7A
if ((tile_num == DBG_TILE) && (threadIdx.x == 0) && (threadIdx.y == 0)){
......@@ -3747,8 +3764,8 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
if (!tile_code){
return; // should not happen
}
if (gpu_texture_rbg && (texture_rbg_stride != 0)) { // generate RGBA
// if no extra and no overlap -> nothing remains, return
if (gpu_texture_rbg && (texture_rbg_stride != 0)) { // generate RGBA (overlapped) // keep_weights
#ifdef DEBUG7A
if ((tile_num == DBG_TILE) && (threadIdx.x == 0) && (threadIdx.y == 0)){
// printf("\ntextures_accumulate accumulating tile = %d, tile_code= %d, border_tile=%d\n",
......@@ -3766,7 +3783,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
__syncthreads();// __syncwarp();
#endif // DEBUG12
int alpha_mode = alphaIndex[tile_code]; // only 4 lowest bits
if (!alpha_mode){ // only multiply if needed, alpha_mode == 0 - keep as is.
if (!alpha_mode){ // only multiply if needed, alpha_mode == 0 - keep as is. FIXME: alpha_mode ???
for (int pass = 0; pass < 8; pass ++) {
int row = pass * 2 + (threadIdx.y >> 1);
int col = ((threadIdx.y & 1) << 3) + threadIdx.x;
......@@ -3833,7 +3850,7 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
// always copy 3 (1) colors + alpha
if (colors == 3){
#pragma unroll
for (int ncol = 0; ncol < colors + 1; ncol++) { // 4
for (int ncol = 0; ncol < 3 + 1; ncol++) { // 4
*(gpu_texture_rbg_gi + ncol * slice_stride) += *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
} else { // assuming colors = 1
......@@ -3842,7 +3859,36 @@ extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
*(gpu_texture_rbg_gi + ncol * slice_stride) += *(rgba_i + ncol * (DTT_SIZE2 * DTT_SIZE21));
}
}
/// }
}
if (keep_weights & 2){ // copy individual sensors output
for (int ncam = 0; ncam < num_cams; ncam++) {
float * mclt_dst_ncam = mclt_debayer + (ncam * colors ) * (MCLT_UNION_LEN);
if (!alpha_mode){ // only multiply if needed, alpha_mode == 0 - keep as is. FIXME: alpha_mode ???
for (int pass = 0; pass < 8; pass ++) {
int row = pass * 2 + (threadIdx.y >> 1);
int col = ((threadIdx.y & 1) << 3) + threadIdx.x;
int i = row * DTT_SIZE21 + col;
int gi = row * DTT_SIZE2 + col;
float * mclt_dst_i = mclt_dst_ncam + i;
for (int ncol = 0; ncol < colors; ncol++) {
*(mclt_dst_i + ncol * (MCLT_UNION_LEN)) *= alphaFade[alpha_mode][gi]; // reduce [tile_code] by LUT
}
}
}
for (int pass = 0; pass < 8; pass ++) {
int row = pass * 2 + (threadIdx.y >> 1); // row inside a tile (0..15)
int col = ((threadIdx.y & 1) << 3) + threadIdx.x; // column inside a tile (0..15)
int g_row = row + tile_y0;
int g_col = col + tile_x0;
int i = row * DTT_SIZE21 + col;
int gi = g_row * texture_rbg_stride + g_col; // offset to the top left corner
float * gpu_texture_rbg_gi = gpu_texture_rbg + gi + (colors + 1 + colors * ncam) * slice_stride;
float * mclt_dst_i = mclt_dst_ncam + i;
for (int ncol = 0; ncol < colors; ncol++) { // 4
*(gpu_texture_rbg_gi + ncol * slice_stride) += *(mclt_dst_i + ncol * (MCLT_UNION_LEN));
}
}
}
}
} // if (gpu_texture_rbg) { // generate RGBA
if (calc_extra){ // gpu_diff_rgb_combo
......
src/main/resources/kernels/TileProcessor.h
View file @ f6006466
......@@ -144,6 +144,7 @@ extern "C" __global__ void textures_nonoverlap(
float params[5],
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // Was not here before 10/12/2022. return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
// combining both non-overlap and overlap (each calculated if pointer is not null )
size_t texture_stride, // in floats (now 256*4 = 1024) // may be 0 if not needed
float * gpu_texture_tiles, // (number of colors +1 + ?)*16*16 rgba texture tiles // may be 0 if not needed
......
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