Skip to content

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

intermediate, not broken

parent 14ed8e38
Show whitespace changes
Inline Side-by-side
Showing with 106 additions and 24 deletions
src/main/java/com/elphel/imagej/cameras/CLTParameters.java
View file @ 95e514a4
...@@ -121,7 +121,7 @@ public class CLTParameters { ...@@ -121,7 +121,7 @@ public class CLTParameters {
// pixel location by quadratic approximation // pixel location by quadratic approximation
public double corr_border_contrast = 0.01; // contrast of dotted border on correlation results public double corr_border_contrast = 0.01; // contrast of dotted border on correlation results
public int tile_task_op = 0xff; // bitmask of operation modes applied to tiles (0 - nothing), bits TBD later public int tile_task_op = 0x7ff; // bitmask of operation modes applied to tiles (0 - nothing), bits TBD later
// +(0..f) - images, +(00.f0) - process pairs + 256 - force disparity when combining images // +(0..f) - images, +(00.f0) - process pairs + 256 - force disparity when combining images
// window to process tiles (later arbitrary masks will be generated to follow particular stages); // window to process tiles (later arbitrary masks will be generated to follow particular stages);
public int tile_task_wl = 0; // public int tile_task_wl = 0; //
......
src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java
View file @ 95e514a4
...@@ -135,15 +135,28 @@ public class GPUTileProcessor { ...@@ -135,15 +135,28 @@ public class GPUTileProcessor {
static int CLTEXTRA_SIZE = 8; static int CLTEXTRA_SIZE = 8;
static int CORR_SIZE = (2* DTT_SIZE - 1) * (2* DTT_SIZE - 1); // 15x15 static int CORR_SIZE = (2* DTT_SIZE - 1) * (2* DTT_SIZE - 1); // 15x15
public static int CORR_NTILE_SHIFT = 8; // also for texture tiles list - not anymore 11/18/2022 public static int CORR_NTILE_SHIFT = 8; // also for texture tiles list - not anymore 11/18/2022
// FIXME: CORR_PAIRS_MASK will not work !!! // FIXME: CORR_PAIRS_MASK will not work !!! It is already removed from the GPU code
public static int CORR_PAIRS_MASK = 0x3f; // lower bits used to address correlation pair for the selected tile public static int CORR_PAIRS_MASK = 0x3f; // lower bits used to address correlation pair for the selected tile
// public static int CORR_TEXTURE_BIT = 7; // bit 7 used to request texture for the tile GET RID !!! public static int TASK_INTER_EN = 10; // Task bit to enable interscene correlation
public static int TASK_CORR_EN = 9; // Task bit to enable intrascene correlation (pairs defined separately)
public static int TASK_TEXT_EN = 8; // task bit to enable texture generation
public static int TASK_CORR_BITS = 4; // start of pair mask public static int TASK_CORR_BITS = 4; // start of pair mask
public static int TASK_TEXT_N_BIT = 0; // Texture with North neighbor
public static int TASK_TEXT_NE_BIT = 1; // Texture with North-East neighbor
public static int TASK_TEXT_E_BIT = 2; // Texture with East neighbor
public static int TASK_TEXT_SE_BIT = 3; // Texture with South-East neighbor
public static int TASK_TEXT_S_BIT = 4; // Texture with South neighbor
public static int TASK_TEXT_SW_BIT = 5; // Texture with South-West neighbor
public static int TASK_TEXT_W_BIT = 6; // Texture with West neighbor
public static int TASK_TEXT_NW_BIT = 7; // Texture with North-West neighbor
public static int TASK_TEXTURE_N_BIT = 0; // Texture with North neighbor public static int TASK_TEXTURE_N_BIT = 0; // Texture with North neighbor
public static int TASK_TEXTURE_E_BIT = 1; // Texture with East neighbor public static int TASK_TEXTURE_E_BIT = 1; // Texture with East neighbor
public static int TASK_TEXTURE_S_BIT = 2; // Texture with South neighbor public static int TASK_TEXTURE_S_BIT = 2; // Texture with South neighbor
public static int TASK_TEXTURE_W_BIT = 3; // Texture with West neighbor public static int TASK_TEXTURE_W_BIT = 3; // Texture with West neighbor
// public static int TASK_TEXTURE_BIT = 3; // bit to request texture calculation int task field of struct tp_task
public static int LIST_TEXTURE_BIT = 8; // 7; // bit to request texture calculation public static int LIST_TEXTURE_BIT = 8; // 7; // bit to request texture calculation
public static int TEXT_NTILE_SHIFT = 9; // 8; // split from CORR_NTILE_SHIFT public static int TEXT_NTILE_SHIFT = 9; // 8; // split from CORR_NTILE_SHIFT
...@@ -155,8 +168,9 @@ public class GPUTileProcessor { ...@@ -155,8 +168,9 @@ public class GPUTileProcessor {
public static int RBYRDIST_LEN = 5001; //for double, 10001 - float; // length of rByRDist to allocate shared memory public static int RBYRDIST_LEN = 5001; //for double, 10001 - float; // length of rByRDist to allocate shared memory
public static double RBYRDIST_STEP = 0.0004; // for double, 0.0002 - for float; // to fit into GPU shared memory (was 0.001); public static double RBYRDIST_STEP = 0.0004; // for double, 0.0002 - for float; // to fit into GPU shared memory (was 0.001);
public static int TILES_PER_BLOCK_GEOM = 32/MAX_NUM_CAMS; // blockDim.x = NUM_CAMS; blockDim.x = TILES_PER_BLOCK_GEOM public static int TILES_PER_BLOCK_GEOM = 32/MAX_NUM_CAMS; // blockDim.x = NUM_CAMS; blockDim.x = TILES_PER_BLOCK_GEOM
public static int TASK_TEXTURE_BITS = ((1 << TASK_TEXTURE_N_BIT) | (1 << TASK_TEXTURE_E_BIT) | (1 << TASK_TEXTURE_S_BIT) | (1 << TASK_TEXTURE_W_BIT)); public static int TASK_TEXTURE_BITS =
( (1 << TASK_TEXT_N_BIT) | (1 << TASK_TEXT_NE_BIT) | (1 << TASK_TEXT_E_BIT) | (1 << TASK_TEXT_SE_BIT) |
(1 << TASK_TEXT_S_BIT) | (1 << TASK_TEXT_SW_BIT) | (1 << TASK_TEXT_W_BIT) | (1 << TASK_TEXT_NW_BIT));
int DTTTEST_BLOCK_WIDTH = 32; // may be read from the source code int DTTTEST_BLOCK_WIDTH = 32; // may be read from the source code
int DTTTEST_BLOCK_HEIGHT = 16; // may be read from the source code int DTTTEST_BLOCK_HEIGHT = 16; // may be read from the source code
...@@ -227,16 +241,22 @@ public class GPUTileProcessor { ...@@ -227,16 +241,22 @@ public class GPUTileProcessor {
"#define IMCLT_THREADS_PER_TILE " + IMCLT_THREADS_PER_TILE+"\n"+ "#define IMCLT_THREADS_PER_TILE " + IMCLT_THREADS_PER_TILE+"\n"+
"#define IMCLT_TILES_PER_BLOCK " + IMCLT_TILES_PER_BLOCK+"\n"+ "#define IMCLT_TILES_PER_BLOCK " + IMCLT_TILES_PER_BLOCK+"\n"+
"#define CORR_NTILE_SHIFT " + CORR_NTILE_SHIFT+"\n"+ "#define CORR_NTILE_SHIFT " + CORR_NTILE_SHIFT+"\n"+
// "#define CORR_PAIRS_MASK " + CORR_PAIRS_MASK+"\n"+ "#define TASK_INTER_EN " + TASK_INTER_EN+"\n"+
// "#define CORR_TEXTURE_BIT " + CORR_TEXTURE_BIT+"\n"+ "#define TASK_CORR_EN " + TASK_CORR_EN+"\n"+
"#define TASK_CORR_BITS " + TASK_CORR_BITS+"\n"+ "#define TASK_TEXT_EN " + TASK_TEXT_EN+"\n"+
"#define TASK_TEXTURE_N_BIT " + TASK_TEXTURE_N_BIT+"\n"+
"#define TASK_TEXTURE_E_BIT " + TASK_TEXTURE_E_BIT+"\n"+ "#define TASK_TEXT_N_BIT " + TASK_TEXT_N_BIT+"\n"+
"#define TASK_TEXTURE_S_BIT " + TASK_TEXTURE_S_BIT+"\n"+ "#define TASK_TEXT_NE_BIT " + TASK_TEXT_NE_BIT+"\n"+
"#define TASK_TEXTURE_W_BIT " + TASK_TEXTURE_W_BIT+"\n"+ "#define TASK_TEXT_E_BIT " + TASK_TEXT_E_BIT+"\n"+
"#define TASK_TEXT_SE_BIT " + TASK_TEXT_SE_BIT+"\n"+
"#define TASK_TEXT_S_BIT " + TASK_TEXT_S_BIT+"\n"+
"#define TASK_TEXT_SW_BIT " + TASK_TEXT_SW_BIT+"\n"+
"#define TASK_TEXT_W_BIT " + TASK_TEXT_W_BIT+"\n"+
"#define TASK_TEXT_NW_BIT " + TASK_TEXT_NW_BIT+"\n"+
"#define LIST_TEXTURE_BIT " + LIST_TEXTURE_BIT+"\n"+ "#define LIST_TEXTURE_BIT " + LIST_TEXTURE_BIT+"\n"+
"#define TEXT_NTILE_SHIFT " + TEXT_NTILE_SHIFT+"\n"+ "#define TEXT_NTILE_SHIFT " + TEXT_NTILE_SHIFT+"\n"+
// "#define CORR_OUT_RAD " + CORR_OUT_RAD+"\n" +
"#define FAT_ZERO_WEIGHT " + FAT_ZERO_WEIGHT+"\n"+ "#define FAT_ZERO_WEIGHT " + FAT_ZERO_WEIGHT+"\n"+
"#define THREADS_DYNAMIC_BITS " + THREADS_DYNAMIC_BITS+"\n"+ "#define THREADS_DYNAMIC_BITS " + THREADS_DYNAMIC_BITS+"\n"+
"#define RBYRDIST_LEN " + RBYRDIST_LEN+"\n"+ "#define RBYRDIST_LEN " + RBYRDIST_LEN+"\n"+
......
src/main/java/com/elphel/imagej/gpu/GpuQuad.java
View file @ 95e514a4
...@@ -991,6 +991,10 @@ public class GpuQuad{ // quad camera description ...@@ -991,6 +991,10 @@ public class GpuQuad{ // quad camera description
((out_images[nt] & 0x0f) << 0) | ((out_images[nt] & 0x0f) << 0) |
((corr_mask [nt] & 0x3f) << 4) ((corr_mask [nt] & 0x3f) << 4)
); // task == 1 for now ); // task == 1 for now
// Old code, enabling all options
tp_tasks[indx]. setTextureEnable(true);
tp_tasks[indx]. setIntraCorrelationEnable(true);
tp_tasks[indx]. setInterCorrelationEnable(true);
indx++; indx++;
} }
} }
...@@ -1091,6 +1095,10 @@ public class GpuQuad{ // quad camera description ...@@ -1091,6 +1095,10 @@ public class GpuQuad{ // quad camera description
((out_images[indx] & 0x0f) << 0) | ((out_images[indx] & 0x0f) << 0) |
((corr_mask [indx] & 0x3f) << 4) ((corr_mask [indx] & 0x3f) << 4)
); // task == 1 for now ); // task == 1 for now
// Old code, enabling all options
tp_tasks[indx]. setTextureEnable(true);
tp_tasks[indx]. setIntraCorrelationEnable(true);
tp_tasks[indx]. setInterCorrelationEnable(true);
indx++; indx++;
} }
} }
...@@ -1150,14 +1158,20 @@ public class GpuQuad{ // quad camera description ...@@ -1150,14 +1158,20 @@ public class GpuQuad{ // quad camera description
acorrs.set(true); acorrs.set(true);
} }
} }
task_list.add(new TpTask( TpTask tptask = new TpTask(
num_cams, num_cams,
tileX, tileX,
tileY, tileY,
(float) (disparity_array[tileY][tileX] + disparity_corr), (float) (disparity_array[tileY][tileX] + disparity_corr),
((img_mask & 0x0f) << 0) | ((img_mask & 0x0f) << 0) |
((corr_mask_tp & 0x3f) << 4) ((corr_mask_tp & 0x3f) << 4)
)); // task == 1 for now ); // task == 1 for now
// Old code, guessing options
if (img_mask != 0) tptask. setTextureEnable(true);
if (corr_mask_tp != 0) tptask. setIntraCorrelationEnable(true);
if (corr_mask_tp != 0) tptask. setInterCorrelationEnable(true);
task_list.add(tptask);
// mask out pairs that use missing channels // mask out pairs that use missing channels
} }
...@@ -1252,6 +1266,8 @@ public class GpuQuad{ // quad camera description ...@@ -1252,6 +1266,8 @@ public class GpuQuad{ // quad camera description
int op = ImageDtt.setImgMask(0, 0xf); // use if tile_op is not provided int op = ImageDtt.setImgMask(0, 0xf); // use if tile_op is not provided
op = ImageDtt.setPairMask(op,0xf); op = ImageDtt.setPairMask(op,0xf);
op = ImageDtt.setForcedDisparity(op,true); op = ImageDtt.setForcedDisparity(op,true);
// setting new (11/18/2022) bits
op |= (1 << GPUTileProcessor.TASK_CORR_EN) | (1 << GPUTileProcessor.TASK_INTER_EN) | (1 << GPUTileProcessor.TASK_TEXT_EN);
final int fop = op; final int fop = op;
int tx = -1; int tx = -1;
for (int i = 0; i < disparity_array.length; i++) if (disparity_array[i] != null) { for (int i = 0; i < disparity_array.length; i++) if (disparity_array[i] != null) {
...@@ -3594,7 +3610,10 @@ public class GpuQuad{ // quad camera description ...@@ -3594,7 +3610,10 @@ public class GpuQuad{ // quad camera description
{ {
int num_pairs = Correlation2d.getNumPairs(num_cams); int num_pairs = Correlation2d.getNumPairs(num_cams);
//change to fixed 511? //change to fixed 511?
final int task_code = ((1 << num_pairs)-1) << GPUTileProcessor.TASK_CORR_BITS; // correlation only // final int task_code = ((1 << num_pairs)-1) << GPUTileProcessor.TASK_CORR_BITS; // correlation only
final int task_code = (1 << GPUTileProcessor.TASK_CORR_EN) | (1 << GPUTileProcessor.TASK_INTER_EN);
final double min_px = margin; final double min_px = margin;
final double max_px = geometryCorrection.getSensorWH()[0] - 1 - margin; // sensor width here, not window width final double max_px = geometryCorrection.getSensorWH()[0] - 1 - margin; // sensor width here, not window width
final double [] min_py = new double[num_cams] ; final double [] min_py = new double[num_cams] ;
...@@ -3703,7 +3722,9 @@ public class GpuQuad{ // quad camera description ...@@ -3703,7 +3722,9 @@ public class GpuQuad{ // quad camera description
final double min_len = 0.1; // pix final double min_len = 0.1; // pix
int num_pairs = Correlation2d.getNumPairs(num_cams); int num_pairs = Correlation2d.getNumPairs(num_cams);
//change to fixed 511? //change to fixed 511?
final int task_code = ((1 << num_pairs)-1) << GPUTileProcessor.TASK_CORR_BITS; // correlation only // final int task_code = ((1 << num_pairs)-1) << GPUTileProcessor.TASK_CORR_BITS; // correlation only
final int task_code = (1 << GPUTileProcessor.TASK_CORR_EN) | (1 << GPUTileProcessor.TASK_INTER_EN);
final double min_px = margin; final double min_px = margin;
final double max_px = geometryCorrection.getSensorWH()[0] - 1 - margin; // sensor width here, not window width final double max_px = geometryCorrection.getSensorWH()[0] - 1 - margin; // sensor width here, not window width
final double [] min_py = new double[num_cams] ; final double [] min_py = new double[num_cams] ;
......
src/main/java/com/elphel/imagej/gpu/TpTask.java
View file @ 95e514a4
package com.elphel.imagej.gpu; package com.elphel.imagej.gpu;
public class TpTask { public class TpTask {
public int task; // [0](+1) - generate 4 images, [4..9]+16..+512 - correlation pairs, 2 - generate texture tiles public int task;
// task bits 0..7 - texture neighbors (0 - N, 1 - NE, ..., 7 - NW)
// bit 8 (GPUTileProcessor.TASK_TEXT_EN) - enable texture generation
// bit 9 (GPUTileProcessor.TASK_CORR_EN) - enable intrascene correlations
// bit 10 (GPUTileProcessor.TASK_INTER_EN) - enable interscene correlations
// Old (still not updated for CPU): [0](+1) - generate 4 images, [4..9]+16..+512 - correlation pairs, 2 - generate texture tiles
public float target_disparity; public float target_disparity;
public int num_sensors = 4; public int num_sensors = 4;
public int ty; public int ty;
...@@ -90,6 +95,36 @@ public class TpTask { ...@@ -90,6 +95,36 @@ public class TpTask {
} }
} }
} }
public void setTextureEnable(boolean en) {
if (en) task |= (1 << GPUTileProcessor.TASK_TEXT_EN);
else task &= ~(1 << GPUTileProcessor.TASK_TEXT_EN);
}
public boolean getTextureEnable() {
return (task & (1 << GPUTileProcessor.TASK_TEXT_EN)) != 0;
}
public void setIntraCorrelationEnable(boolean en) {
if (en) task |= (1 << GPUTileProcessor.TASK_CORR_EN);
else task &= ~(1 << GPUTileProcessor.TASK_CORR_EN);
}
public boolean getIntraCorrelationEnable() {
return (task & (1 << GPUTileProcessor.TASK_CORR_EN)) != 0;
}
public void setInterCorrelationEnable(boolean en) {
if (en) task |= (1 << GPUTileProcessor.TASK_INTER_EN);
else task &= ~(1 << GPUTileProcessor.TASK_INTER_EN);
}
public boolean getInterCorrelationEnable() {
return (task & (1 << GPUTileProcessor.TASK_INTER_EN)) != 0;
}
public float [][] getDispDist(){ public float [][] getDispDist(){
return disp_dist; return disp_dist;
} }
......
src/main/java/com/elphel/imagej/tileprocessor/CLTPass3d.java
View file @ 95e514a4
...@@ -28,6 +28,8 @@ import java.awt.Rectangle; ...@@ -28,6 +28,8 @@ import java.awt.Rectangle;
import java.util.Arrays; import java.util.Arrays;
import java.util.concurrent.atomic.AtomicInteger; import java.util.concurrent.atomic.AtomicInteger;
import com.elphel.imagej.gpu.GPUTileProcessor;
public class CLTPass3d{ public class CLTPass3d{
// static double max_overexposed = 0.8; // TODO: make parameter // static double max_overexposed = 0.8; // TODO: make parameter
public double [][] disparity; // per-tile disparity set for the pass[tileY][tileX] public double [][] disparity; // per-tile disparity set for the pass[tileY][tileX]
...@@ -1079,6 +1081,8 @@ public class CLTPass3d{ ...@@ -1079,6 +1081,8 @@ public class CLTPass3d{
int op = ImageDtt.setImgMask(0, 0xf); int op = ImageDtt.setImgMask(0, 0xf);
op = ImageDtt.setPairMask(op,0xf); op = ImageDtt.setPairMask(op,0xf);
op = ImageDtt.setForcedDisparity(op,true); op = ImageDtt.setForcedDisparity(op,true);
// New as for 11/18/2022 - no CPU support yet
op |= (1 << GPUTileProcessor.TASK_TEXT_EN) | (1 << GPUTileProcessor.TASK_CORR_EN) | (1 << GPUTileProcessor.TASK_INTER_EN);
return setTileOpDisparity( return setTileOpDisparity(
op, // int tile_op, op, // int tile_op,
selection, // boolean [] selection, selection, // boolean [] selection,
...@@ -1120,7 +1124,7 @@ public class CLTPass3d{ ...@@ -1120,7 +1124,7 @@ public class CLTPass3d{
} }
return num_op_tiles; return num_op_tiles;
} }
/*
public void setTileOp( public void setTileOp(
int tile_op) int tile_op)
{ {
...@@ -1132,7 +1136,7 @@ public class CLTPass3d{ ...@@ -1132,7 +1136,7 @@ public class CLTPass3d{
this.tile_op[ty][tx] = tile_op; this.tile_op[ty][tx] = tile_op;
} }
} }
*/
public double [] getDA() { public double [] getDA() {
final int tilesX = tileProcessor.getTilesX(); final int tilesX = tileProcessor.getTilesX();
final int tilesY = tileProcessor.getTilesY(); final int tilesY = tileProcessor.getTilesY();
......
src/main/java/com/elphel/imagej/tileprocessor/QuadCLT.java
View file @ 95e514a4
...@@ -3982,6 +3982,8 @@ public class QuadCLT extends QuadCLTCPU { ...@@ -3982,6 +3982,8 @@ public class QuadCLT extends QuadCLTCPU {
if ((scan.disparity == null) || (scan.tile_op == null)) { if ((scan.disparity == null) || (scan.tile_op == null)) {
int d = ImageDtt.setImgMask(0, 0xf); // no correlations int d = ImageDtt.setImgMask(0, 0xf); // no correlations
d = ImageDtt.setForcedDisparity(d,true); d = ImageDtt.setForcedDisparity(d,true);
// New as for 11/18/2022 - no CPU support yet
d |= (1 << GPUTileProcessor.TASK_TEXT_EN);
scan.setTileOpDisparity( scan.setTileOpDisparity(
d, d,
scan.getSelected(), // boolean [] selection, scan.getSelected(), // boolean [] selection,
......
src/main/java/com/elphel/imagej/tileprocessor/TexturedModel.java
View file @ 95e514a4
...@@ -2137,7 +2137,7 @@ public class TexturedModel { ...@@ -2137,7 +2137,7 @@ public class TexturedModel {
(tileY * transform_size + row) * width + (tileX * transform_size), (tileY * transform_size + row) * width + (tileX * transform_size),
transform_size); transform_size);
} }
int navg = inter_textures_wd[fnslice][tileY][tileX].length - num_colors + y_color; int navg = y_color; // inter_textures_wd[fnslice][tileY][tileX].length - num_colors + y_color;
System.arraycopy( System.arraycopy(
inter_textures_wd[fnslice][tileY][tileX][navg], inter_textures_wd[fnslice][tileY][tileX][navg],
row*transform_size, row*transform_size,
......
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