Skip to content

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

fixed leak between colors in textures

parent 10c327ae
Hide whitespace changes
Inline Side-by-side
Showing with 135 additions and 18 deletions
src/main/java/com/elphel/imagej/cameras/CLTParameters.java
View file @ 3f0b0bc0
......@@ -779,6 +779,12 @@ 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 replaceWeakOutliers = true; // false;
public boolean debug_initial_discriminate = false;
......@@ -1552,6 +1558,12 @@ 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_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+"");
......@@ -2320,6 +2332,12 @@ 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_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"));
......@@ -3243,11 +3261,20 @@ public class CLTParameters {
"Add squared fat zero to the sum of squared amplitudes, color images");
gd.addNumericField("Fat zero (absolute) for phase correlation of monochrome images", this.gpu_fatz_m, 4, 6,"",
"Add squared fat zero to the sum of squared amplitudes, monochrome images");
gd.addMessage ("--- GPU WOI selection ---");
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",
"Top WOI margin, in tiles (0..241");
gd.addNumericField("WOI width", this.gpu_woi_twidth, 0, 6,"tiles",
"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.addTab ("LWIR", "parameters for LWIR/EO 8-camera rig");
this.lwir.dialogQuestions(gd);
gd.addTab ("Debug", "Other debug images");
gd.addMessage ("--- Other debug images ---");
// clt_parameters.debug_initial_discriminate, // final boolean debug_initial_discriminate,
......@@ -3987,6 +4014,12 @@ public class CLTParameters {
this.gpu_fatz = gd.getNextNumber();
this.gpu_fatz_m = gd.getNextNumber();
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.lwir.dialogAnswers(gd);
this.debug_initial_discriminate= gd.getNextBoolean();
......
src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java
View file @ 3f0b0bc0
......@@ -641,6 +641,8 @@ 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
public TpTask [] setFullFrameImages(
Rectangle woi,
boolean round_woi,
float target_disparity, // apply same disparity to all tiles
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)
......@@ -664,19 +666,23 @@ public class GPUTileProcessor {
corr_masks[i] = corr_mask; // 0x3f; // all 6 correlations
}
return setFullFrameImages(
woi, // Rectangle woi,
round_woi, // boolean round_woi,
target_disparities, // should be tilesX*tilesY long
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)
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
ers_delay, // if not null - fill with tile center acquisition delay
threadsMax, // maximal number of threads to launch
debugLevel);
}
public TpTask [] setFullFrameImages(
Rectangle woi, // or null
boolean round_woi,
float [] target_disparities, // should be tilesX*tilesY long
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)
......@@ -690,11 +696,49 @@ public class GPUTileProcessor {
{
int tilesX = IMG_WIDTH / DTT_SIZE;
int tilesY = IMG_HEIGHT / DTT_SIZE;
if (woi == null) {
woi = new Rectangle(0,0,tilesX,tilesY);
}
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;
boolean dbg1 = false; // true;
boolean [] mask = new boolean[tilesX*tilesY];
int num_tiles = 0;
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 ++;
}
}
}
if (dbg1) {
// mask[(woi.y-1) * tilesX + (woi.x-1)] = true;
mask[(woi.y+woi.height) * tilesX + (woi.x+woi.width)] = true;
num_tiles += 1; // 2;
}
// TpTask [] tp_tasks = new TpTask[tilesX*tilesY];
TpTask [] tp_tasks = new TpTask[num_tiles];
TpTask [] tp_tasks = new TpTask[tilesX*tilesY];
int indx = 0;
for (int ty = 0; ty < tilesY; ty++) {
for (int tx = 0; tx < tilesX; tx++) {
for (int tx = 0; tx < tilesX; tx++) if (mask[ty * tilesX + tx]) {
// tp_tasks[indx] = new TpTask(tx,ty, target_disparities[indx], 1); // task == 1 for now
// Only generate for non-empty tasks, use 1 empty empty as a terminator?
tp_tasks[indx] = new TpTask(tx,ty, target_disparities[indx],
......
src/main/java/com/elphel/imagej/tileprocessor/TwoQuadCLT.java
View file @ 3f0b0bc0
......@@ -1996,7 +1996,14 @@ public class TwoQuadCLT {
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 = gPUTileProcessor.setFullFrameImages(
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)
......@@ -2008,6 +2015,7 @@ public class TwoQuadCLT {
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
......@@ -2253,6 +2261,20 @@ public class TwoQuadCLT {
float [][] rbga = gPUTileProcessor.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]};
......@@ -2267,8 +2289,8 @@ public class TwoQuadCLT {
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),
tilesX * image_dtt.transform_size,
tilesY * image_dtt.transform_size,
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 );
......@@ -2276,7 +2298,13 @@ public class TwoQuadCLT {
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, 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(
......
src/main/resources/kernels/TileProcessor.cuh
View file @ 3f0b0bc0
......@@ -1503,6 +1503,15 @@ __global__ void generate_RBGA(
num_tiles, // number of tiles in task list
gpu_texture_indices, // packed tile + bits (now only (1 << 7)
woi); // min_x, min_y, max_x, max_y
// REMOVE when done!
/*
*(woi + 0) -= 1;
*(woi + 1) -= 1;
*(woi + 2) += 1;
*(woi + 3) += 1;
*/
cudaDeviceSynchronize();
// Generate tile indices list, upper 24 bits - tile index, lower 4 bits: n/e/s/w neighbors, bit 7 - set to 1
*(num_texture_tiles+0) = 0;
......@@ -1807,7 +1816,8 @@ __global__ void gen_texture_list(
// int is_border = (x == woi[0]) || (y == woi[1]) || (x == woi[2]) || (y == woi[3]);
// don't care if calculate extra pixels that still fit into memory
int is_border = (x == woi[0]) || (y == woi[1]) || (x == (TILESX - 1)) || (y == (TILESY - 1));
// int is_border = (x == woi[0]) || (y == woi[1]) || (x == (TILESX - 1)) || (y == (TILESY - 1));
int is_border = (x == woi[0]) || (y == woi[1]) || (x == (TILESX - 1)) || (y == woi[3]);
int buff_head = 0;
int num_offset = 0;
if (x & 1) {
......@@ -2512,6 +2522,7 @@ __global__ void textures_accumulate(
int tileX = tile_num - tileY * TILESX;
int tile_x0 = (tileX - *(woi + 0)) * DTT_SIZE - (DTT_SIZE/2); // may be negative == -4
int tile_y0 = (tileY - *(woi + 1)) * DTT_SIZE - (DTT_SIZE/2); // may be negative == -4
int height = *(woi + 3) << DTT_SIZE_LOG2;
#ifdef DEBUG12
if ((tile_num == DBG_TILE) && (threadIdx.x == 0) && (threadIdx.y == 0)){
......@@ -2537,16 +2548,17 @@ __global__ void textures_accumulate(
float * gpu_texture_rbg_gi = gpu_texture_rbg + gi;
float * rgba_i = ((float *) shr1.rgbaw) + i;
#ifdef DEBUG12
if ((tile_num == DBG_TILE) && (threadIdx.x == 0) && (threadIdx.y == 0)){
printf("\ntextures_accumulate () pass=%d, row=%d, col=%d, g_row=%d, g_col=%d, i=%d, gi=%d\n",
pass, row, col, g_row, g_col, i, gi);
if ((tile_num == DBG_TILE) && (threadIdx.x == 0) && (threadIdx.y == 0)){
printf("\ntextures_accumulate () pass=%d, row=%d, col=%d, g_row=%d, g_col=%d, i=%d, gi=%d\n",
pass, row, col, g_row, g_col, i, gi);
}
__syncthreads();// __syncwarp();
}
__syncthreads();// __syncwarp();
#endif // DEBUG12
if (!border_tile ||
((g_row >= 0) && (g_col >= 0) && (g_row < (DTT_SIZE * TILESX)) && (g_col < (DTT_SIZE * TILESY)))){
/// ((g_row >= 0) && (g_col >= 0) && (g_row < (DTT_SIZE * TILESY)) && (g_col < (DTT_SIZE * TILESX)))){
((g_row >= 0) && (g_col >= 0) && (g_row < height) && (g_col < (DTT_SIZE * TILESX)))){
// always copy 3 (1) colors + alpha
if (colors == 3){
#pragma unroll
......
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