Loading src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java +6 −1 Original line number Diff line number Diff line Loading @@ -102,6 +102,7 @@ public class GPUTileProcessor { static String GPU_CLEAR_TEXTURE_RBGA_NAME = "clear_texture_rbga"; static String GPU_TEXTURES_ACCUMULATE_NAME = "textures_accumulate"; static String GPU_CREATE_NONOVERLAP_LIST_NAME ="create_nonoverlap_list"; static String GPU_ERASE_CLT_TILES_NAME = "erase_clt_tiles"; Loading Loading @@ -180,6 +181,7 @@ public class GPUTileProcessor { CUfunction GPU_CLEAR_TEXTURE_RBGA_kernel = null; // "clear_texture_rbga" CUfunction GPU_TEXTURES_ACCUMULATE_kernel = null; // "textures_accumulate" CUfunction GPU_CREATE_NONOVERLAP_LIST_kernel = null; // "create_nonoverlap_list" CUfunction GPU_ERASE_CLT_TILES_kernel = null; // "erase_clt_tiles" Loading Loading @@ -331,7 +333,8 @@ public class GPUTileProcessor { GPU_GEN_TEXTURE_LIST_NAME, GPU_CLEAR_TEXTURE_RBGA_NAME, GPU_TEXTURES_ACCUMULATE_NAME, GPU_CREATE_NONOVERLAP_LIST_NAME GPU_CREATE_NONOVERLAP_LIST_NAME, GPU_ERASE_CLT_TILES_NAME }; CUfunction[] functions = createFunctions(kernelSources, func_names, Loading @@ -356,6 +359,7 @@ public class GPUTileProcessor { GPU_CLEAR_TEXTURE_RBGA_kernel = functions[15]; GPU_TEXTURES_ACCUMULATE_kernel = functions[16]; GPU_CREATE_NONOVERLAP_LIST_kernel = functions[17]; GPU_ERASE_CLT_TILES_kernel = functions[18]; System.out.println("GPU kernel functions initialized"); Loading @@ -378,6 +382,7 @@ public class GPUTileProcessor { System.out.println(GPU_CLEAR_TEXTURE_RBGA_kernel.toString()); System.out.println(GPU_TEXTURES_ACCUMULATE_kernel.toString()); System.out.println(GPU_CREATE_NONOVERLAP_LIST_kernel.toString()); System.out.println(GPU_ERASE_CLT_TILES_kernel.toString()); // GPU data structures are now initialized through GpuQuad instances } Loading src/main/java/com/elphel/imagej/gpu/GpuQuad.java +40 −29 Original line number Diff line number Diff line Loading @@ -1464,33 +1464,44 @@ public class GpuQuad{ // quad camera description /** * Direct CLT conversion and aberration correction */ public void execConvertDirect() { public void execConvertDirect(int erase_clt) { execConvertDirect( false, null); null, erase_clt); } /** * Convert and save TD representation in either normal or reference scene. Reference scene TD representation * is used for interscene correlation (for "IMU") * @param ref_scene save result into a separate buffer for interscene correlation when true. * @param wh window width, height (or null) * @param erase_clt erase CLT data. Only needed before execImcltRbgAll() if not all the * tiles are converted. <0 - do not erase, 0 - erase to 0, 1 - erase to NaN */ public void execConvertDirect( boolean ref_scene, int [] wh) { int [] wh, int erase_clt) { if (this.gpuTileProcessor.GPU_CONVERT_DIRECT_kernel == null) { IJ.showMessage("Error", "No GPU kernel: GPU_CONVERT_DIRECT_kernel"); return; } if (this.gpuTileProcessor.GPU_ERASE_CLT_TILES_kernel == null) { IJ.showMessage("Error", "No GPU kernel: GPU_ERASE_CLT_TILES_kernel"); return; } if (wh == null) { wh = new int[] {img_width, img_height}; } 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 / GPUTileProcessor.DTT_SIZE; int tilesX = wh[0] / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; // De-allocate if size mismatch, allocate if needed. Now it is the only place where clt is allocated if (ref_scene) { if ((gpu_clt_ref_wh != null) && ((gpu_clt_ref_wh[0] != wh[0]) || (gpu_clt_ref_wh[1] != wh[1]))) { Loading @@ -1503,8 +1514,6 @@ public class GpuQuad{ // quad camera description } if (gpu_clt_ref == null) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_ref_l = new long [num_cams]; // int tilesY = img_height / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; gpu_clt_ref_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_h[ncam] = new CUdeviceptr(); Loading @@ -1530,8 +1539,6 @@ public class GpuQuad{ // quad camera description } if (gpu_clt == null) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_l = new long [num_cams]; // int tilesY = img_height / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; gpu_clt_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_h[ncam] = new CUdeviceptr(); Loading @@ -1547,28 +1554,32 @@ public class GpuQuad{ // quad camera description gpu_clt_wh = wh.clone(); } } /* if (ref_scene && (gpu_clt_ref == null)) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_ref_l = new long [num_cams]; int tilesY = img_height / GPUTileProcessor.DTT_SIZE; gpu_clt_ref_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_h[ncam] = new CUdeviceptr(); cuMemAlloc(gpu_clt_ref_h[ncam], tilesY * tilesX * num_colors * 4 * GPUTileProcessor.DTT_SIZE * GPUTileProcessor.DTT_SIZE * Sizeof.FLOAT ); } gpu_clt_ref = new CUdeviceptr(); cuMemAlloc(gpu_clt_ref, num_cams * Sizeof.POINTER); for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_l[ncam] = GPUTileProcessor.getPointerAddress(gpu_clt_ref_h[ncam]); } cuMemcpyHtoD(gpu_clt_ref, Pointer.to(gpu_clt_ref_l), num_cams * Sizeof.POINTER); gpu_clt_ref_wh = wh.clone(); } */ CUdeviceptr gpu_clt_selected = ref_scene ? gpu_clt_ref : gpu_clt; int [] GridFullWarps = {1, 1, 1}; int [] ThreadsFullWarps = {1, 1, 1}; if (erase_clt >= 0) { float fill_data = (erase_clt > 0) ? Float.NaN : 0.0f; Pointer kernelParametersEraseClt = Pointer.to( Pointer.to(new int[] { num_cams}), // int num_cams, Pointer.to(new int[] { num_colors}), // int num_colors, Pointer.to(new int[] { tilesX}), // int tiles_x, Pointer.to(new int[] { tilesY}), // int tiles_y, Pointer.to(gpu_clt_selected), // float ** gpu_clt,// [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE] Pointer.to(new float[] { fill_data})); // float fill_data cuCtxSynchronize(); // Call the kernel function cuLaunchKernel(this.gpuTileProcessor.GPU_ERASE_CLT_TILES_kernel, GridFullWarps[0], GridFullWarps[1], GridFullWarps[2], // Grid dimension ThreadsFullWarps[0], ThreadsFullWarps[1],ThreadsFullWarps[2],// Block dimension 0, null, // Shared memory size and stream (shared - only dynamic, static is in code) kernelParametersEraseClt, null); // Kernel- and extra parameters cuCtxSynchronize(); // remove later CUDA_ERROR_ILLEGAL_ADDRESS if (getGpu_debug_level() > -1) { System.out.println("======execConvertDirect(): erased CLT"); } } Pointer kernelParameters = Pointer.to( Pointer.to(new int[] { num_cams}), // int num_cams, Pointer.to(new int[] { num_colors}), // int num_colors, Loading Loading @@ -1598,7 +1609,7 @@ public class GpuQuad{ // quad camera description kernelParameters, null); // Kernel- and extra parameters cuCtxSynchronize(); // remove later CUDA_ERROR_ILLEGAL_ADDRESS if (getGpu_debug_level() > -1) { System.out.println("======execConvertDirect("+ref_scene+")"); System.out.println("======execConvertDirect("+ref_scene+", "+erase_clt+")"); } } Loading src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java +11 −6 Original line number Diff line number Diff line Loading @@ -339,7 +339,7 @@ public class ImageDtt extends ImageDttCPU { } gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (iclt_fimg != null) { gpuQuad.execImcltRbgAll(isMonochrome()); // execute GPU kernel for (int ncam = 0; ncam < iclt_fimg.length; ncam++) { Loading Loading @@ -958,7 +958,7 @@ public class ImageDtt extends ImageDttCPU { // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (mcorr_sel == 0) { // no correlation at all return; } Loading Loading @@ -1124,7 +1124,7 @@ public class ImageDtt extends ImageDttCPU { // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (sensor_mask_inter == 0) { // no correlation at all return; } Loading @@ -1143,6 +1143,9 @@ public class ImageDtt extends ImageDttCPU { /** * Convert reference scene to FD and save result in extra GPU array for the future interscene correlation * Geometry correction and images will come from gpuQuad instance - * @param erase_clt erase CLT (<0 - do not erase, 0 - erase to 0.0, >0 - erase to NaN). Needed only for later IMCLT * end rendering images. NaN produces sharp, distinct borders; 0f - blended * @param wh if null, will uses sensor dimensions. Otherwise {width, height} in pixels * @param imgdtt_params * @param use_reference_buffer true - use extra GPU array, false - use main one * @param tp_tasks Loading @@ -1154,6 +1157,7 @@ public class ImageDtt extends ImageDttCPU { * @param globalDebugLevel */ public void setReferenceTD( final int erase_clt, final int [] wh, // null (use sensor dimensions) or pair {width, height} in pixels final ImageDttParameters imgdtt_params, // Now just extra correlation parameters, later will include, most others final boolean use_reference_buffer, Loading Loading @@ -1186,7 +1190,8 @@ public class ImageDtt extends ImageDttCPU { false); // boolean use_aux // while is it in class member? - just to be able to free // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... // int [] wh = null; gpuQuad.execConvertDirect(use_reference_buffer, wh); // put results into a "reference" buffer // int erase_clt = 1; // NaN; gpuQuad.execConvertDirect(use_reference_buffer, wh, erase_clt); // put results into a "reference" buffer } Loading Loading @@ -1297,7 +1302,7 @@ public class ImageDtt extends ImageDttCPU { /// gpuQuad.execSetTilesOffsets(); // prepare tiles offsets in GPU memory // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt //Generate 2D phase correlations from the CLT representation int mcorr_sel = Correlation2d.corrSelEncode(imgdtt_params,numSensors); Loading Loading @@ -1688,7 +1693,7 @@ public class ImageDtt extends ImageDttCPU { // GPUTileProcessor.TpTask[] tp_tasks_full2) gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (iclt_fimg != null) { gpuQuad.execImcltRbgAll(isMonochrome()); // execute GPU kernel for (int ncam = 0; ncam < iclt_fimg.length; ncam++) { Loading src/main/java/com/elphel/imagej/tileprocessor/IntersceneMatchParameters.java +15 −3 Original line number Diff line number Diff line Loading @@ -37,6 +37,8 @@ public class IntersceneMatchParameters { public boolean show_images = false; // color, infinity public boolean show_images_bgfg = false; // bg and fg public boolean show_images_mono = false; // float, monochrome 16-slice images (same disparity, COMBO_DSN_INDX_DISP_FG and COMBO_DSN_INDX_DISP_BG_ALL, public boolean show_color_nan = true; // use NAN background for color images (sharp, but distinct black) public boolean show_mono_nan = false; // use NAN background for monochrome images (sharp, but distinct black) public boolean show_ranges = true; public double range_disparity_offset = -0.08; Loading Loading @@ -120,9 +122,13 @@ public class IntersceneMatchParameters { "Show foreground and background exported images"); gd.addCheckbox ("Show floating-point monochrome images", this.show_images_mono, "Display generated/saved monochrome images"); gd.addCheckbox ("Color NaN background", this.show_color_nan, "Use NaN for undefined tiles (false - 0.0f). NaN produces sharp distinct result, 0.0f - blended"); gd.addCheckbox ("Mono NaN background", this.show_mono_nan, "Use NaN for undefined tiles (false - 0.0f). NaN produces sharp distinct result, 0.0f - blended"); gd.addCheckbox ("Show distances in meters", this.show_ranges, "Calculate strength, distance, X, and Y in meters"); gd.addNumericField("Disparity at infinity", this.range_disparity_offset, 5,7,"pix", "Disparity at infinity - subtract from measured disparity when converting to ranges."); gd.addNumericField("Minimal strength for range calculation", this.range_min_strength, 5,7,"", Loading Loading @@ -234,6 +240,8 @@ public class IntersceneMatchParameters { this.show_images = gd.getNextBoolean(); this.show_images_bgfg = gd.getNextBoolean(); this.show_images_mono = gd.getNextBoolean(); this.show_color_nan = gd.getNextBoolean(); this.show_mono_nan = gd.getNextBoolean(); this.show_ranges = gd.getNextBoolean(); this.range_disparity_offset = gd.getNextNumber(); this.range_min_strength = gd.getNextNumber(); Loading Loading @@ -292,12 +300,12 @@ public class IntersceneMatchParameters { properties.setProperty(prefix+"show_images", this.show_images + ""); // boolean properties.setProperty(prefix+"show_images_bgfg", this.show_images_bgfg + ""); // boolean properties.setProperty(prefix+"show_images_mono", this.show_images_mono + ""); // boolean properties.setProperty(prefix+"show_color_nan", this.show_color_nan + ""); // boolean properties.setProperty(prefix+"show_mono_nan", this.show_mono_nan + ""); // boolean properties.setProperty(prefix+"show_ranges", this.show_ranges + ""); // boolean properties.setProperty(prefix+"range_disparity_offset",this.range_disparity_offset+""); // double properties.setProperty(prefix+"range_min_strength", this.range_min_strength+""); // double properties.setProperty(prefix+"range_max", this.range_max+""); // double properties.setProperty(prefix+"margin", this.margin+""); // int properties.setProperty(prefix+"sensor_mask_inter", this.sensor_mask_inter+""); // int properties.setProperty(prefix+"use_partial", this.use_partial+""); // boolean Loading Loading @@ -349,6 +357,8 @@ public class IntersceneMatchParameters { if (properties.getProperty(prefix+"show_images")!=null) this.show_images=Boolean.parseBoolean(properties.getProperty(prefix+"show_images")); if (properties.getProperty(prefix+"show_images_bgfg")!=null) this.show_images_bgfg=Boolean.parseBoolean(properties.getProperty(prefix+"show_images_bgfg")); if (properties.getProperty(prefix+"show_images_mono")!=null) this.show_images_mono=Boolean.parseBoolean(properties.getProperty(prefix+"show_images_mono")); if (properties.getProperty(prefix+"show_color_nan")!=null) this.show_color_nan=Boolean.parseBoolean(properties.getProperty(prefix+"show_color_nan")); if (properties.getProperty(prefix+"show_mono_nan")!=null) this.show_mono_nan=Boolean.parseBoolean(properties.getProperty(prefix+"show_mono_nan")); if (properties.getProperty(prefix+"show_ranges")!=null) this.show_images=Boolean.parseBoolean(properties.getProperty(prefix+"show_ranges")); if (properties.getProperty(prefix+"range_disparity_offset")!=null) this.range_disparity_offset=Double.parseDouble(properties.getProperty(prefix+"range_disparity_offset")); if (properties.getProperty(prefix+"range_min_strength")!=null) this.range_min_strength=Double.parseDouble(properties.getProperty(prefix+"range_min_strength")); Loading Loading @@ -406,6 +416,8 @@ public class IntersceneMatchParameters { imp.show_images = this.show_images; imp.show_images_bgfg = this.show_images_bgfg; imp.show_images_mono = this.show_images_mono; imp.show_color_nan = this.show_color_nan; imp.show_mono_nan = this.show_mono_nan; imp.show_ranges = this.show_ranges; imp.range_disparity_offset = this.range_disparity_offset; imp.range_min_strength = this.range_min_strength; Loading src/main/java/com/elphel/imagej/tileprocessor/OpticalFlow.java +18 −3 Original line number Diff line number Diff line Loading @@ -3963,6 +3963,9 @@ public class OpticalFlow { boolean export_images = clt_parameters.imp.export_images; boolean export_dsi_image = clt_parameters.imp.show_ranges; boolean show_images = clt_parameters.imp.show_images; boolean show_color_nan = clt_parameters.imp.show_color_nan; boolean show_mono_nan = clt_parameters.imp.show_mono_nan; boolean show_images_bgfg = clt_parameters.imp.show_images_bgfg; boolean show_images_mono = clt_parameters.imp.show_images_mono; Loading Loading @@ -4279,11 +4282,12 @@ public class OpticalFlow { Arrays.fill(constant_disparity,clt_parameters.disparity); Rectangle testr = new Rectangle(10, 8, 100,80); ImagePlus imp_constant = QuadCLT.renderGPUFromDSI( testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) -1, // final int sensor_mask, null, // testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, constant_disparity, // double [] disparity_ref, ZERO3, // final double [] scene_xyz, // camera center in world coordinates new double[] {.1,0.1,.1}, // ZERO3, // final double [] scene_atr, // camera orientation relative to world frame ZERO3, // new double[] {.1,0.1,.1}, // ZERO3, // final double [] scene_atr, // camera orientation relative to world frame quadCLTs[ref_index], // final QuadCLT scene, true, // toRGB, // final boolean toRGB, "GPU-SHIFTED-D"+clt_parameters.disparity, // String suffix, Loading @@ -4293,7 +4297,8 @@ public class OpticalFlow { null, // "GPU-SHIFTED-D"+clt_parameters.disparity, // String suffix, imp_constant); // ImagePlus imp) ImagePlus imp_constant_mono = QuadCLT.renderGPUFromDSI( testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) -1, // final int sensor_mask, null, // testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, constant_disparity, // double [] disparity_ref, ZERO3, // final double [] scene_xyz, // camera center in world coordinates Loading @@ -4313,6 +4318,7 @@ public class OpticalFlow { } } ImagePlus imp_fg = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, fg_disparity, // double [] disparity_ref, Loading @@ -4327,6 +4333,7 @@ public class OpticalFlow { null, // "GPU-SHIFTED-FOREGROUND", // String suffix, imp_fg); // ImagePlus imp) ImagePlus imp_fg_mono = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, fg_disparity, // double [] disparity_ref, Loading @@ -4347,6 +4354,7 @@ public class OpticalFlow { } } ImagePlus imp_bg = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, bg_disparity, // double [] disparity_ref, Loading @@ -4361,6 +4369,7 @@ public class OpticalFlow { null, // "GPU-SHIFTED-BACKGROUND", // String suffix, imp_bg); // ImagePlus imp) ImagePlus imp_bg_mono = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, bg_disparity, // double [] disparity_ref, Loading Loading @@ -9812,6 +9821,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad final float [][][] accum_2d_corr, // if [1][][] - return accumulated 2d correlations (all pairs)final float [][][] accum_2d_corr, // if [1][][] - return accumulated 2d correlations (all pairs) int debug_level) { TileProcessor tp = ref_scene.getTileProcessor(); // Temporary reusing same ref scene ****** boolean scene_is_ref_test = clt_parameters.imp.scene_is_ref_test; // false; // true; Loading @@ -9821,6 +9831,8 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad boolean show_render_scene = clt_parameters.imp.renderScene(); // false; // true; boolean toRGB = clt_parameters.imp.toRGB ; // true; boolean show_coord_motion = clt_parameters.imp.showCorrMotion(); // mae its own int erase_clt = (toRGB? clt_parameters.imp.show_color_nan : clt_parameters.imp.show_mono_nan) ? 1:0; if (scene_is_ref_test) { scene_xyz = ZERO3.clone(); scene_atr = ZERO3.clone(); Loading Loading @@ -9907,6 +9919,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad float [][][][] fcorr_td = null; // no accumulation, use data in GPU ref_scene.saveQuadClt(); // to re-load new set of Bayer images to the GPU (do nothing for CPU) and Geometry image_dtt.setReferenceTD( erase_clt, null, // final int [] wh, // null (use sensor dimensions) or pair {width, height} in pixels clt_parameters.img_dtt, // final ImageDttParameters imgdtt_params, // Now just extra correlation parameters, later will include, most others true, // final boolean use_reference_buffer, Loading @@ -9919,6 +9932,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad debug_level); // final int globalDebugLevel); if (show_render_ref) { ImagePlus imp_render_ref = ref_scene.renderFromTD ( -1, // final int sensor_mask, clt_parameters, // CLTParameters clt_parameters, clt_parameters.getColorProcParameters(ref_scene.isAux()), //ColorProcParameters colorProcParameters, clt_parameters.getRGBParameters(), //EyesisCorrectionParameters.RGBParameters rgbParameters, Loading Loading @@ -9948,6 +9962,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad debug_level); // final int globalDebugLevel); if (show_render_scene) { ImagePlus imp_render_scene = scene.renderFromTD ( -1, // final int sensor_mask, clt_parameters, // CLTParameters clt_parameters, clt_parameters.getColorProcParameters(ref_scene.isAux()), //ColorProcParameters colorProcParameters, clt_parameters.getRGBParameters(), //EyesisCorrectionParameters.RGBParameters rgbParameters, Loading Loading
src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java +6 −1 Original line number Diff line number Diff line Loading @@ -102,6 +102,7 @@ public class GPUTileProcessor { static String GPU_CLEAR_TEXTURE_RBGA_NAME = "clear_texture_rbga"; static String GPU_TEXTURES_ACCUMULATE_NAME = "textures_accumulate"; static String GPU_CREATE_NONOVERLAP_LIST_NAME ="create_nonoverlap_list"; static String GPU_ERASE_CLT_TILES_NAME = "erase_clt_tiles"; Loading Loading @@ -180,6 +181,7 @@ public class GPUTileProcessor { CUfunction GPU_CLEAR_TEXTURE_RBGA_kernel = null; // "clear_texture_rbga" CUfunction GPU_TEXTURES_ACCUMULATE_kernel = null; // "textures_accumulate" CUfunction GPU_CREATE_NONOVERLAP_LIST_kernel = null; // "create_nonoverlap_list" CUfunction GPU_ERASE_CLT_TILES_kernel = null; // "erase_clt_tiles" Loading Loading @@ -331,7 +333,8 @@ public class GPUTileProcessor { GPU_GEN_TEXTURE_LIST_NAME, GPU_CLEAR_TEXTURE_RBGA_NAME, GPU_TEXTURES_ACCUMULATE_NAME, GPU_CREATE_NONOVERLAP_LIST_NAME GPU_CREATE_NONOVERLAP_LIST_NAME, GPU_ERASE_CLT_TILES_NAME }; CUfunction[] functions = createFunctions(kernelSources, func_names, Loading @@ -356,6 +359,7 @@ public class GPUTileProcessor { GPU_CLEAR_TEXTURE_RBGA_kernel = functions[15]; GPU_TEXTURES_ACCUMULATE_kernel = functions[16]; GPU_CREATE_NONOVERLAP_LIST_kernel = functions[17]; GPU_ERASE_CLT_TILES_kernel = functions[18]; System.out.println("GPU kernel functions initialized"); Loading @@ -378,6 +382,7 @@ public class GPUTileProcessor { System.out.println(GPU_CLEAR_TEXTURE_RBGA_kernel.toString()); System.out.println(GPU_TEXTURES_ACCUMULATE_kernel.toString()); System.out.println(GPU_CREATE_NONOVERLAP_LIST_kernel.toString()); System.out.println(GPU_ERASE_CLT_TILES_kernel.toString()); // GPU data structures are now initialized through GpuQuad instances } Loading
src/main/java/com/elphel/imagej/gpu/GpuQuad.java +40 −29 Original line number Diff line number Diff line Loading @@ -1464,33 +1464,44 @@ public class GpuQuad{ // quad camera description /** * Direct CLT conversion and aberration correction */ public void execConvertDirect() { public void execConvertDirect(int erase_clt) { execConvertDirect( false, null); null, erase_clt); } /** * Convert and save TD representation in either normal or reference scene. Reference scene TD representation * is used for interscene correlation (for "IMU") * @param ref_scene save result into a separate buffer for interscene correlation when true. * @param wh window width, height (or null) * @param erase_clt erase CLT data. Only needed before execImcltRbgAll() if not all the * tiles are converted. <0 - do not erase, 0 - erase to 0, 1 - erase to NaN */ public void execConvertDirect( boolean ref_scene, int [] wh) { int [] wh, int erase_clt) { if (this.gpuTileProcessor.GPU_CONVERT_DIRECT_kernel == null) { IJ.showMessage("Error", "No GPU kernel: GPU_CONVERT_DIRECT_kernel"); return; } if (this.gpuTileProcessor.GPU_ERASE_CLT_TILES_kernel == null) { IJ.showMessage("Error", "No GPU kernel: GPU_ERASE_CLT_TILES_kernel"); return; } if (wh == null) { wh = new int[] {img_width, img_height}; } 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 / GPUTileProcessor.DTT_SIZE; int tilesX = wh[0] / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; // De-allocate if size mismatch, allocate if needed. Now it is the only place where clt is allocated if (ref_scene) { if ((gpu_clt_ref_wh != null) && ((gpu_clt_ref_wh[0] != wh[0]) || (gpu_clt_ref_wh[1] != wh[1]))) { Loading @@ -1503,8 +1514,6 @@ public class GpuQuad{ // quad camera description } if (gpu_clt_ref == null) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_ref_l = new long [num_cams]; // int tilesY = img_height / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; gpu_clt_ref_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_h[ncam] = new CUdeviceptr(); Loading @@ -1530,8 +1539,6 @@ public class GpuQuad{ // quad camera description } if (gpu_clt == null) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_l = new long [num_cams]; // int tilesY = img_height / GPUTileProcessor.DTT_SIZE; int tilesY = wh[1] / GPUTileProcessor.DTT_SIZE; gpu_clt_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_h[ncam] = new CUdeviceptr(); Loading @@ -1547,28 +1554,32 @@ public class GpuQuad{ // quad camera description gpu_clt_wh = wh.clone(); } } /* if (ref_scene && (gpu_clt_ref == null)) { // Allocate memory, create pointers for reference scene TD representation long [] gpu_clt_ref_l = new long [num_cams]; int tilesY = img_height / GPUTileProcessor.DTT_SIZE; gpu_clt_ref_h = new CUdeviceptr[num_cams]; for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_h[ncam] = new CUdeviceptr(); cuMemAlloc(gpu_clt_ref_h[ncam], tilesY * tilesX * num_colors * 4 * GPUTileProcessor.DTT_SIZE * GPUTileProcessor.DTT_SIZE * Sizeof.FLOAT ); } gpu_clt_ref = new CUdeviceptr(); cuMemAlloc(gpu_clt_ref, num_cams * Sizeof.POINTER); for (int ncam = 0; ncam < num_cams; ncam++) { gpu_clt_ref_l[ncam] = GPUTileProcessor.getPointerAddress(gpu_clt_ref_h[ncam]); } cuMemcpyHtoD(gpu_clt_ref, Pointer.to(gpu_clt_ref_l), num_cams * Sizeof.POINTER); gpu_clt_ref_wh = wh.clone(); } */ CUdeviceptr gpu_clt_selected = ref_scene ? gpu_clt_ref : gpu_clt; int [] GridFullWarps = {1, 1, 1}; int [] ThreadsFullWarps = {1, 1, 1}; if (erase_clt >= 0) { float fill_data = (erase_clt > 0) ? Float.NaN : 0.0f; Pointer kernelParametersEraseClt = Pointer.to( Pointer.to(new int[] { num_cams}), // int num_cams, Pointer.to(new int[] { num_colors}), // int num_colors, Pointer.to(new int[] { tilesX}), // int tiles_x, Pointer.to(new int[] { tilesY}), // int tiles_y, Pointer.to(gpu_clt_selected), // float ** gpu_clt,// [num_cams][tiles_y][tiles_x][num_colors][4*DTT_SIZE*DTT_SIZE] Pointer.to(new float[] { fill_data})); // float fill_data cuCtxSynchronize(); // Call the kernel function cuLaunchKernel(this.gpuTileProcessor.GPU_ERASE_CLT_TILES_kernel, GridFullWarps[0], GridFullWarps[1], GridFullWarps[2], // Grid dimension ThreadsFullWarps[0], ThreadsFullWarps[1],ThreadsFullWarps[2],// Block dimension 0, null, // Shared memory size and stream (shared - only dynamic, static is in code) kernelParametersEraseClt, null); // Kernel- and extra parameters cuCtxSynchronize(); // remove later CUDA_ERROR_ILLEGAL_ADDRESS if (getGpu_debug_level() > -1) { System.out.println("======execConvertDirect(): erased CLT"); } } Pointer kernelParameters = Pointer.to( Pointer.to(new int[] { num_cams}), // int num_cams, Pointer.to(new int[] { num_colors}), // int num_colors, Loading Loading @@ -1598,7 +1609,7 @@ public class GpuQuad{ // quad camera description kernelParameters, null); // Kernel- and extra parameters cuCtxSynchronize(); // remove later CUDA_ERROR_ILLEGAL_ADDRESS if (getGpu_debug_level() > -1) { System.out.println("======execConvertDirect("+ref_scene+")"); System.out.println("======execConvertDirect("+ref_scene+", "+erase_clt+")"); } } Loading
src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java +11 −6 Original line number Diff line number Diff line Loading @@ -339,7 +339,7 @@ public class ImageDtt extends ImageDttCPU { } gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (iclt_fimg != null) { gpuQuad.execImcltRbgAll(isMonochrome()); // execute GPU kernel for (int ncam = 0; ncam < iclt_fimg.length; ncam++) { Loading Loading @@ -958,7 +958,7 @@ public class ImageDtt extends ImageDttCPU { // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (mcorr_sel == 0) { // no correlation at all return; } Loading Loading @@ -1124,7 +1124,7 @@ public class ImageDtt extends ImageDttCPU { // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (sensor_mask_inter == 0) { // no correlation at all return; } Loading @@ -1143,6 +1143,9 @@ public class ImageDtt extends ImageDttCPU { /** * Convert reference scene to FD and save result in extra GPU array for the future interscene correlation * Geometry correction and images will come from gpuQuad instance - * @param erase_clt erase CLT (<0 - do not erase, 0 - erase to 0.0, >0 - erase to NaN). Needed only for later IMCLT * end rendering images. NaN produces sharp, distinct borders; 0f - blended * @param wh if null, will uses sensor dimensions. Otherwise {width, height} in pixels * @param imgdtt_params * @param use_reference_buffer true - use extra GPU array, false - use main one * @param tp_tasks Loading @@ -1154,6 +1157,7 @@ public class ImageDtt extends ImageDttCPU { * @param globalDebugLevel */ public void setReferenceTD( final int erase_clt, final int [] wh, // null (use sensor dimensions) or pair {width, height} in pixels final ImageDttParameters imgdtt_params, // Now just extra correlation parameters, later will include, most others final boolean use_reference_buffer, Loading Loading @@ -1186,7 +1190,8 @@ public class ImageDtt extends ImageDttCPU { false); // boolean use_aux // while is it in class member? - just to be able to free // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... // int [] wh = null; gpuQuad.execConvertDirect(use_reference_buffer, wh); // put results into a "reference" buffer // int erase_clt = 1; // NaN; gpuQuad.execConvertDirect(use_reference_buffer, wh, erase_clt); // put results into a "reference" buffer } Loading Loading @@ -1297,7 +1302,7 @@ public class ImageDtt extends ImageDttCPU { /// gpuQuad.execSetTilesOffsets(); // prepare tiles offsets in GPU memory // Skipping if ((fdisp_dist != null) || (fpxpy != null)) {... gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt //Generate 2D phase correlations from the CLT representation int mcorr_sel = Correlation2d.corrSelEncode(imgdtt_params,numSensors); Loading Loading @@ -1688,7 +1693,7 @@ public class ImageDtt extends ImageDttCPU { // GPUTileProcessor.TpTask[] tp_tasks_full2) gpuQuad.execConvertDirect(); gpuQuad.execConvertDirect(-1); // boolean erase_clt if (iclt_fimg != null) { gpuQuad.execImcltRbgAll(isMonochrome()); // execute GPU kernel for (int ncam = 0; ncam < iclt_fimg.length; ncam++) { Loading
src/main/java/com/elphel/imagej/tileprocessor/IntersceneMatchParameters.java +15 −3 Original line number Diff line number Diff line Loading @@ -37,6 +37,8 @@ public class IntersceneMatchParameters { public boolean show_images = false; // color, infinity public boolean show_images_bgfg = false; // bg and fg public boolean show_images_mono = false; // float, monochrome 16-slice images (same disparity, COMBO_DSN_INDX_DISP_FG and COMBO_DSN_INDX_DISP_BG_ALL, public boolean show_color_nan = true; // use NAN background for color images (sharp, but distinct black) public boolean show_mono_nan = false; // use NAN background for monochrome images (sharp, but distinct black) public boolean show_ranges = true; public double range_disparity_offset = -0.08; Loading Loading @@ -120,9 +122,13 @@ public class IntersceneMatchParameters { "Show foreground and background exported images"); gd.addCheckbox ("Show floating-point monochrome images", this.show_images_mono, "Display generated/saved monochrome images"); gd.addCheckbox ("Color NaN background", this.show_color_nan, "Use NaN for undefined tiles (false - 0.0f). NaN produces sharp distinct result, 0.0f - blended"); gd.addCheckbox ("Mono NaN background", this.show_mono_nan, "Use NaN for undefined tiles (false - 0.0f). NaN produces sharp distinct result, 0.0f - blended"); gd.addCheckbox ("Show distances in meters", this.show_ranges, "Calculate strength, distance, X, and Y in meters"); gd.addNumericField("Disparity at infinity", this.range_disparity_offset, 5,7,"pix", "Disparity at infinity - subtract from measured disparity when converting to ranges."); gd.addNumericField("Minimal strength for range calculation", this.range_min_strength, 5,7,"", Loading Loading @@ -234,6 +240,8 @@ public class IntersceneMatchParameters { this.show_images = gd.getNextBoolean(); this.show_images_bgfg = gd.getNextBoolean(); this.show_images_mono = gd.getNextBoolean(); this.show_color_nan = gd.getNextBoolean(); this.show_mono_nan = gd.getNextBoolean(); this.show_ranges = gd.getNextBoolean(); this.range_disparity_offset = gd.getNextNumber(); this.range_min_strength = gd.getNextNumber(); Loading Loading @@ -292,12 +300,12 @@ public class IntersceneMatchParameters { properties.setProperty(prefix+"show_images", this.show_images + ""); // boolean properties.setProperty(prefix+"show_images_bgfg", this.show_images_bgfg + ""); // boolean properties.setProperty(prefix+"show_images_mono", this.show_images_mono + ""); // boolean properties.setProperty(prefix+"show_color_nan", this.show_color_nan + ""); // boolean properties.setProperty(prefix+"show_mono_nan", this.show_mono_nan + ""); // boolean properties.setProperty(prefix+"show_ranges", this.show_ranges + ""); // boolean properties.setProperty(prefix+"range_disparity_offset",this.range_disparity_offset+""); // double properties.setProperty(prefix+"range_min_strength", this.range_min_strength+""); // double properties.setProperty(prefix+"range_max", this.range_max+""); // double properties.setProperty(prefix+"margin", this.margin+""); // int properties.setProperty(prefix+"sensor_mask_inter", this.sensor_mask_inter+""); // int properties.setProperty(prefix+"use_partial", this.use_partial+""); // boolean Loading Loading @@ -349,6 +357,8 @@ public class IntersceneMatchParameters { if (properties.getProperty(prefix+"show_images")!=null) this.show_images=Boolean.parseBoolean(properties.getProperty(prefix+"show_images")); if (properties.getProperty(prefix+"show_images_bgfg")!=null) this.show_images_bgfg=Boolean.parseBoolean(properties.getProperty(prefix+"show_images_bgfg")); if (properties.getProperty(prefix+"show_images_mono")!=null) this.show_images_mono=Boolean.parseBoolean(properties.getProperty(prefix+"show_images_mono")); if (properties.getProperty(prefix+"show_color_nan")!=null) this.show_color_nan=Boolean.parseBoolean(properties.getProperty(prefix+"show_color_nan")); if (properties.getProperty(prefix+"show_mono_nan")!=null) this.show_mono_nan=Boolean.parseBoolean(properties.getProperty(prefix+"show_mono_nan")); if (properties.getProperty(prefix+"show_ranges")!=null) this.show_images=Boolean.parseBoolean(properties.getProperty(prefix+"show_ranges")); if (properties.getProperty(prefix+"range_disparity_offset")!=null) this.range_disparity_offset=Double.parseDouble(properties.getProperty(prefix+"range_disparity_offset")); if (properties.getProperty(prefix+"range_min_strength")!=null) this.range_min_strength=Double.parseDouble(properties.getProperty(prefix+"range_min_strength")); Loading Loading @@ -406,6 +416,8 @@ public class IntersceneMatchParameters { imp.show_images = this.show_images; imp.show_images_bgfg = this.show_images_bgfg; imp.show_images_mono = this.show_images_mono; imp.show_color_nan = this.show_color_nan; imp.show_mono_nan = this.show_mono_nan; imp.show_ranges = this.show_ranges; imp.range_disparity_offset = this.range_disparity_offset; imp.range_min_strength = this.range_min_strength; Loading
src/main/java/com/elphel/imagej/tileprocessor/OpticalFlow.java +18 −3 Original line number Diff line number Diff line Loading @@ -3963,6 +3963,9 @@ public class OpticalFlow { boolean export_images = clt_parameters.imp.export_images; boolean export_dsi_image = clt_parameters.imp.show_ranges; boolean show_images = clt_parameters.imp.show_images; boolean show_color_nan = clt_parameters.imp.show_color_nan; boolean show_mono_nan = clt_parameters.imp.show_mono_nan; boolean show_images_bgfg = clt_parameters.imp.show_images_bgfg; boolean show_images_mono = clt_parameters.imp.show_images_mono; Loading Loading @@ -4279,11 +4282,12 @@ public class OpticalFlow { Arrays.fill(constant_disparity,clt_parameters.disparity); Rectangle testr = new Rectangle(10, 8, 100,80); ImagePlus imp_constant = QuadCLT.renderGPUFromDSI( testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) -1, // final int sensor_mask, null, // testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, constant_disparity, // double [] disparity_ref, ZERO3, // final double [] scene_xyz, // camera center in world coordinates new double[] {.1,0.1,.1}, // ZERO3, // final double [] scene_atr, // camera orientation relative to world frame ZERO3, // new double[] {.1,0.1,.1}, // ZERO3, // final double [] scene_atr, // camera orientation relative to world frame quadCLTs[ref_index], // final QuadCLT scene, true, // toRGB, // final boolean toRGB, "GPU-SHIFTED-D"+clt_parameters.disparity, // String suffix, Loading @@ -4293,7 +4297,8 @@ public class OpticalFlow { null, // "GPU-SHIFTED-D"+clt_parameters.disparity, // String suffix, imp_constant); // ImagePlus imp) ImagePlus imp_constant_mono = QuadCLT.renderGPUFromDSI( testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) -1, // final int sensor_mask, null, // testr, // null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, constant_disparity, // double [] disparity_ref, ZERO3, // final double [] scene_xyz, // camera center in world coordinates Loading @@ -4313,6 +4318,7 @@ public class OpticalFlow { } } ImagePlus imp_fg = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, fg_disparity, // double [] disparity_ref, Loading @@ -4327,6 +4333,7 @@ public class OpticalFlow { null, // "GPU-SHIFTED-FOREGROUND", // String suffix, imp_fg); // ImagePlus imp) ImagePlus imp_fg_mono = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, fg_disparity, // double [] disparity_ref, Loading @@ -4347,6 +4354,7 @@ public class OpticalFlow { } } ImagePlus imp_bg = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, bg_disparity, // double [] disparity_ref, Loading @@ -4361,6 +4369,7 @@ public class OpticalFlow { null, // "GPU-SHIFTED-BACKGROUND", // String suffix, imp_bg); // ImagePlus imp) ImagePlus imp_bg_mono = QuadCLT.renderGPUFromDSI( -1, // final int sensor_mask, null, // final Rectangle full_woi_in, // show larger than sensor WOI (or null) clt_parameters, // CLTParameters clt_parameters, bg_disparity, // double [] disparity_ref, Loading Loading @@ -9812,6 +9821,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad final float [][][] accum_2d_corr, // if [1][][] - return accumulated 2d correlations (all pairs)final float [][][] accum_2d_corr, // if [1][][] - return accumulated 2d correlations (all pairs) int debug_level) { TileProcessor tp = ref_scene.getTileProcessor(); // Temporary reusing same ref scene ****** boolean scene_is_ref_test = clt_parameters.imp.scene_is_ref_test; // false; // true; Loading @@ -9821,6 +9831,8 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad boolean show_render_scene = clt_parameters.imp.renderScene(); // false; // true; boolean toRGB = clt_parameters.imp.toRGB ; // true; boolean show_coord_motion = clt_parameters.imp.showCorrMotion(); // mae its own int erase_clt = (toRGB? clt_parameters.imp.show_color_nan : clt_parameters.imp.show_mono_nan) ? 1:0; if (scene_is_ref_test) { scene_xyz = ZERO3.clone(); scene_atr = ZERO3.clone(); Loading Loading @@ -9907,6 +9919,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad float [][][][] fcorr_td = null; // no accumulation, use data in GPU ref_scene.saveQuadClt(); // to re-load new set of Bayer images to the GPU (do nothing for CPU) and Geometry image_dtt.setReferenceTD( erase_clt, null, // final int [] wh, // null (use sensor dimensions) or pair {width, height} in pixels clt_parameters.img_dtt, // final ImageDttParameters imgdtt_params, // Now just extra correlation parameters, later will include, most others true, // final boolean use_reference_buffer, Loading @@ -9919,6 +9932,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad debug_level); // final int globalDebugLevel); if (show_render_ref) { ImagePlus imp_render_ref = ref_scene.renderFromTD ( -1, // final int sensor_mask, clt_parameters, // CLTParameters clt_parameters, clt_parameters.getColorProcParameters(ref_scene.isAux()), //ColorProcParameters colorProcParameters, clt_parameters.getRGBParameters(), //EyesisCorrectionParameters.RGBParameters rgbParameters, Loading Loading @@ -9948,6 +9962,7 @@ public double[][] correlateIntersceneDebug( // only uses GPU and quad debug_level); // final int globalDebugLevel); if (show_render_scene) { ImagePlus imp_render_scene = scene.renderFromTD ( -1, // final int sensor_mask, clt_parameters, // CLTParameters clt_parameters, clt_parameters.getColorProcParameters(ref_scene.isAux()), //ColorProcParameters colorProcParameters, clt_parameters.getRGBParameters(), //EyesisCorrectionParameters.RGBParameters rgbParameters, Loading