Loading src/TileProcessor.cuh +7 −1 Original line number Diff line number Diff line Loading @@ -41,6 +41,8 @@ #ifndef JCUDA #include "tp_defines.h" #include "dtt8x8.h" #include "geometry_correction.h" #include "TileProcessor.h" #endif // #ifndef JCUDA #define TASK_TEXTURE_BITS ((1 << TASK_TEXTURE_N_BIT) | (1 << TASK_TEXTURE_E_BIT) | (1 << TASK_TEXTURE_S_BIT) | (1 << TASK_TEXTURE_W_BIT)) Loading Loading @@ -106,11 +108,12 @@ #define DBG_TILE_Y 111 // 66 #define DBG_TILE (DBG_TILE_Y * 324 + DBG_TILE_X) #undef DBG_MARK_DBG_TILE 1 #undef DBG_MARK_DBG_TILE //56494 // struct tp_task //#define TASK_SIZE 12 #if 0 struct tp_task { int task; union { Loading @@ -119,6 +122,7 @@ struct tp_task { }; float xy[NUM_CAMS][2]; }; #endif struct CltExtra{ float data_x; // kernel data is relative to this displacement X (0.5 pixel increments) float data_y; // kernel data is relative to this displacement Y (0.5 pixel increments) Loading Loading @@ -826,6 +830,7 @@ __device__ void imclt_plane( // not implemented, not used float * gpu_rbg, // WIDTH, HEIGHT const size_t dstride); // in floats (pixels) #if 0 extern "C" __global__ void clear_texture_list( int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) Loading Loading @@ -892,6 +897,7 @@ __global__ void imclt_rbg( int h_offset, const size_t dstride); // in floats (pixels) //=========================== #endif extern "C" __global__ void correlate2D( Loading src/TileProcessor.h 0 → 100644 +106 −0 Original line number Diff line number Diff line /** ** ** TileProcessor.h ** ** Copyright (C) 2020 Elphel, Inc. ** ** -----------------------------------------------------------------------------** ** ** TileProcessor.h is free software: you can redistribute it and/or modify ** it under the terms of the GNU General Public License as published by ** the Free Software Foundation, either version 3 of the License, or ** (at your option) any later version. ** ** This program is distributed in the hope that it will be useful, ** but WITHOUT ANY WARRANTY; without even the implied warranty of ** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the ** GNU General Public License for more details. ** ** You should have received a copy of the GNU General Public License ** along with this program. If not, see <http://www.gnu.org/licenses/>. ** ** Additional permission under GNU GPL version 3 section 7 ** ** If you modify this Program, or any covered work, by linking or ** combining it with NVIDIA Corporation's CUDA libraries from the ** NVIDIA CUDA Toolkit (or a modified version of those libraries), ** containing parts covered by the terms of NVIDIA CUDA Toolkit ** EULA, the licensors of this Program grant you additional ** permission to convey the resulting work. ** -----------------------------------------------------------------------------** */ /** ************************************************************************** * \file TileProcessor.h * \brief header file for the Tile Processor for frequency domain */ #pragma once #ifndef NUM_CAMS #include "tp_defines.h" #endif extern "C" __global__ void clear_texture_list( int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) int width, // <= TILESX, use for faster processing of LWIR images int height); // <= TILESY, use for faster processing of LWIR images extern "C" __global__ void mark_texture_tiles( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices); // packed tile + bits (now only (1 << 7) extern "C" __global__ void mark_texture_neighbor_tiles( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices, // packed tile + bits (now only (1 << 7) int * woi); // x,y,width,height of the woi extern "C" __global__ void gen_texture_list( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices, // packed tile + bits (now only (1 << 7) int * num_texture_tiles, // number of texture tiles to process int * woi); // x,y,width,height of the woi extern "C" __global__ void clear_texture_rbga( int texture_width, int texture_slice_height, const size_t texture_rbga_stride, // in floats 8*stride float * gpu_texture_tiles); // (number of colors +1 + ?)*16*16 rgba texture tiles extern "C" __global__ void textures_accumulate( // int border_tile, // if 1 - watch for border int * woi, // x, y, width,height float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE] size_t num_texture_tiles, // number of texture tiles to process int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate int colors, // number of colors (3/1) int is_lwir, // do not perform shot correction float min_shot, // 10.0 float scale_shot, // 3.0 float diff_sigma, // pixel value/pixel change float diff_threshold, // pixel value/pixel change float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages) float weight0, // scale for R float weight1, // scale for B float weight2, // scale for 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)? // 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 size_t texture_stride, // in floats (now 256*4 = 1024) float * gpu_texture_tiles); // (number of colors +1 + ?)*16*16 rgba texture tiles extern "C" __global__ void imclt_rbg( float * gpu_clt, // [TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE] float * gpu_rbg, // WIDTH, 3 * HEIGHT int apply_lpf, int mono, // defines lpf filter int color, // defines location of clt data int v_offset, int h_offset, const size_t dstride); // in floats (pixels) src/geometry_correction.h 0 → 100644 +97 −0 Original line number Diff line number Diff line /** ** ** geometry_correction.h ** ** Copyright (C) 2020 Elphel, Inc. ** ** -----------------------------------------------------------------------------** ** ** geometry_correction.h is free software: you can redistribute it and/or modify ** it under the terms of the GNU General Public License as published by ** the Free Software Foundation, either version 3 of the License, or ** (at your option) any later version. ** ** This program is distributed in the hope that it will be useful, ** but WITHOUT ANY WARRANTY; without even the implied warranty of ** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the ** GNU General Public License for more details. ** ** You should have received a copy of the GNU General Public License ** along with this program. If not, see <http://www.gnu.org/licenses/>. ** ** Additional permission under GNU GPL version 3 section 7 ** ** If you modify this Program, or any covered work, by linking or ** combining it with NVIDIA Corporation's CUDA libraries from the ** NVIDIA CUDA Toolkit (or a modified version of those libraries), ** containing parts covered by the terms of NVIDIA CUDA Toolkit ** EULA, the licensors of this Program grant you additional ** permission to convey the resulting work. ** -----------------------------------------------------------------------------** */ /** ************************************************************************** * \file geometry_correction.h * \brief header file for geometry correction - per-tile/per camera calculation of the tile offset */ #pragma once #ifndef NUM_CAMS #include "tp_defines.h" #endif struct tp_task { int task; union { int txy; unsigned short sxy[2]; }; float xy[NUM_CAMS][2]; }; struct corr_vector{ float tilt [NUM_CAMS-1]; // 0..2 float azimuth [NUM_CAMS-1]; // 3..5 float roll [NUM_CAMS]; // 6..9 float zoom [NUM_CAMS-1]; // 10..12 // for ERS correction: float imu_rot [3]; // d_tilt/dt (rad/s), d_az/dt, d_roll/dt 13..15 float imu_move[3]; // dx/dt, dy/dt, dz/dt 16..19 }; struct gc { float focalLength; // =FOCAL_LENGTH; float pixelSize; // = PIXEL_SIZE; //um float distortionRadius; // = DISTORTION_RADIUS; // mm - half width of the sensor float distortionA8; //r^8 (normalized to focal length or to sensor half width?) float distortionA7; //r^7 (normalized to focal length or to sensor half width?) float distortionA6; //r^6 (normalized to focal length or to sensor half width?) float distortionA5; //r^5 (normalized to focal length or to sensor half width?) float distortionA; // r^4 (normalized to focal length or to sensor half width?) float distortionB; // r^3 float distortionC; // r^2 // parameters, common for all sensors float elevation; // degrees, up - positive; float heading; // degrees, CW (from top) - positive float forward [NUM_CAMS]; float right [NUM_CAMS]; float height [NUM_CAMS]; float roll [NUM_CAMS]; // degrees, CW (to target) - positive float common_right; // mm right, camera center float common_forward; // mm forward (to target), camera center float common_height; // mm up, camera center float common_roll; // degrees CW (to target) camera as a whole // float [][] XYZ_he; // all cameras coordinates transformed to eliminate heading and elevation (rolls preserved) // float [][] XYZ_her = null; // XYZ of the lenses in a corrected CCS (adjusted for to elevation, heading, common_roll) float rXY [NUM_CAMS][3]; // XY pairs of the in a normal plane, relative to disparityRadius // float [][] rXY_ideal = {{-0.5, -0.5}, {0.5,-0.5}, {-0.5, 0.5}, {0.5,0.5}}; // only used for the multi-quad systems float cameraRadius; // =0; // average distance from the "mass center" of the sensors to the sensors float disparityRadius; // =150.0; // distance between cameras to normalize disparity units to. sqrt(2)*disparityRadius for quad }; Loading
src/TileProcessor.cuh +7 −1 Original line number Diff line number Diff line Loading @@ -41,6 +41,8 @@ #ifndef JCUDA #include "tp_defines.h" #include "dtt8x8.h" #include "geometry_correction.h" #include "TileProcessor.h" #endif // #ifndef JCUDA #define TASK_TEXTURE_BITS ((1 << TASK_TEXTURE_N_BIT) | (1 << TASK_TEXTURE_E_BIT) | (1 << TASK_TEXTURE_S_BIT) | (1 << TASK_TEXTURE_W_BIT)) Loading Loading @@ -106,11 +108,12 @@ #define DBG_TILE_Y 111 // 66 #define DBG_TILE (DBG_TILE_Y * 324 + DBG_TILE_X) #undef DBG_MARK_DBG_TILE 1 #undef DBG_MARK_DBG_TILE //56494 // struct tp_task //#define TASK_SIZE 12 #if 0 struct tp_task { int task; union { Loading @@ -119,6 +122,7 @@ struct tp_task { }; float xy[NUM_CAMS][2]; }; #endif struct CltExtra{ float data_x; // kernel data is relative to this displacement X (0.5 pixel increments) float data_y; // kernel data is relative to this displacement Y (0.5 pixel increments) Loading Loading @@ -826,6 +830,7 @@ __device__ void imclt_plane( // not implemented, not used float * gpu_rbg, // WIDTH, HEIGHT const size_t dstride); // in floats (pixels) #if 0 extern "C" __global__ void clear_texture_list( int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) Loading Loading @@ -892,6 +897,7 @@ __global__ void imclt_rbg( int h_offset, const size_t dstride); // in floats (pixels) //=========================== #endif extern "C" __global__ void correlate2D( Loading
src/TileProcessor.h 0 → 100644 +106 −0 Original line number Diff line number Diff line /** ** ** TileProcessor.h ** ** Copyright (C) 2020 Elphel, Inc. ** ** -----------------------------------------------------------------------------** ** ** TileProcessor.h is free software: you can redistribute it and/or modify ** it under the terms of the GNU General Public License as published by ** the Free Software Foundation, either version 3 of the License, or ** (at your option) any later version. ** ** This program is distributed in the hope that it will be useful, ** but WITHOUT ANY WARRANTY; without even the implied warranty of ** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the ** GNU General Public License for more details. ** ** You should have received a copy of the GNU General Public License ** along with this program. If not, see <http://www.gnu.org/licenses/>. ** ** Additional permission under GNU GPL version 3 section 7 ** ** If you modify this Program, or any covered work, by linking or ** combining it with NVIDIA Corporation's CUDA libraries from the ** NVIDIA CUDA Toolkit (or a modified version of those libraries), ** containing parts covered by the terms of NVIDIA CUDA Toolkit ** EULA, the licensors of this Program grant you additional ** permission to convey the resulting work. ** -----------------------------------------------------------------------------** */ /** ************************************************************************** * \file TileProcessor.h * \brief header file for the Tile Processor for frequency domain */ #pragma once #ifndef NUM_CAMS #include "tp_defines.h" #endif extern "C" __global__ void clear_texture_list( int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) int width, // <= TILESX, use for faster processing of LWIR images int height); // <= TILESY, use for faster processing of LWIR images extern "C" __global__ void mark_texture_tiles( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices); // packed tile + bits (now only (1 << 7) extern "C" __global__ void mark_texture_neighbor_tiles( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices, // packed tile + bits (now only (1 << 7) int * woi); // x,y,width,height of the woi extern "C" __global__ void gen_texture_list( struct tp_task * gpu_tasks, int num_tiles, // number of tiles in task list int * gpu_texture_indices, // packed tile + bits (now only (1 << 7) int * num_texture_tiles, // number of texture tiles to process int * woi); // x,y,width,height of the woi extern "C" __global__ void clear_texture_rbga( int texture_width, int texture_slice_height, const size_t texture_rbga_stride, // in floats 8*stride float * gpu_texture_tiles); // (number of colors +1 + ?)*16*16 rgba texture tiles extern "C" __global__ void textures_accumulate( // int border_tile, // if 1 - watch for border int * woi, // x, y, width,height float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE] size_t num_texture_tiles, // number of texture tiles to process int * gpu_texture_indices,// packed tile + bits (now only (1 << 7) float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate int colors, // number of colors (3/1) int is_lwir, // do not perform shot correction float min_shot, // 10.0 float scale_shot, // 3.0 float diff_sigma, // pixel value/pixel change float diff_threshold, // pixel value/pixel change float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages) float weight0, // scale for R float weight1, // scale for B float weight2, // scale for 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)? // 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 size_t texture_stride, // in floats (now 256*4 = 1024) float * gpu_texture_tiles); // (number of colors +1 + ?)*16*16 rgba texture tiles extern "C" __global__ void imclt_rbg( float * gpu_clt, // [TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE] float * gpu_rbg, // WIDTH, 3 * HEIGHT int apply_lpf, int mono, // defines lpf filter int color, // defines location of clt data int v_offset, int h_offset, const size_t dstride); // in floats (pixels)
src/geometry_correction.h 0 → 100644 +97 −0 Original line number Diff line number Diff line /** ** ** geometry_correction.h ** ** Copyright (C) 2020 Elphel, Inc. ** ** -----------------------------------------------------------------------------** ** ** geometry_correction.h is free software: you can redistribute it and/or modify ** it under the terms of the GNU General Public License as published by ** the Free Software Foundation, either version 3 of the License, or ** (at your option) any later version. ** ** This program is distributed in the hope that it will be useful, ** but WITHOUT ANY WARRANTY; without even the implied warranty of ** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the ** GNU General Public License for more details. ** ** You should have received a copy of the GNU General Public License ** along with this program. If not, see <http://www.gnu.org/licenses/>. ** ** Additional permission under GNU GPL version 3 section 7 ** ** If you modify this Program, or any covered work, by linking or ** combining it with NVIDIA Corporation's CUDA libraries from the ** NVIDIA CUDA Toolkit (or a modified version of those libraries), ** containing parts covered by the terms of NVIDIA CUDA Toolkit ** EULA, the licensors of this Program grant you additional ** permission to convey the resulting work. ** -----------------------------------------------------------------------------** */ /** ************************************************************************** * \file geometry_correction.h * \brief header file for geometry correction - per-tile/per camera calculation of the tile offset */ #pragma once #ifndef NUM_CAMS #include "tp_defines.h" #endif struct tp_task { int task; union { int txy; unsigned short sxy[2]; }; float xy[NUM_CAMS][2]; }; struct corr_vector{ float tilt [NUM_CAMS-1]; // 0..2 float azimuth [NUM_CAMS-1]; // 3..5 float roll [NUM_CAMS]; // 6..9 float zoom [NUM_CAMS-1]; // 10..12 // for ERS correction: float imu_rot [3]; // d_tilt/dt (rad/s), d_az/dt, d_roll/dt 13..15 float imu_move[3]; // dx/dt, dy/dt, dz/dt 16..19 }; struct gc { float focalLength; // =FOCAL_LENGTH; float pixelSize; // = PIXEL_SIZE; //um float distortionRadius; // = DISTORTION_RADIUS; // mm - half width of the sensor float distortionA8; //r^8 (normalized to focal length or to sensor half width?) float distortionA7; //r^7 (normalized to focal length or to sensor half width?) float distortionA6; //r^6 (normalized to focal length or to sensor half width?) float distortionA5; //r^5 (normalized to focal length or to sensor half width?) float distortionA; // r^4 (normalized to focal length or to sensor half width?) float distortionB; // r^3 float distortionC; // r^2 // parameters, common for all sensors float elevation; // degrees, up - positive; float heading; // degrees, CW (from top) - positive float forward [NUM_CAMS]; float right [NUM_CAMS]; float height [NUM_CAMS]; float roll [NUM_CAMS]; // degrees, CW (to target) - positive float common_right; // mm right, camera center float common_forward; // mm forward (to target), camera center float common_height; // mm up, camera center float common_roll; // degrees CW (to target) camera as a whole // float [][] XYZ_he; // all cameras coordinates transformed to eliminate heading and elevation (rolls preserved) // float [][] XYZ_her = null; // XYZ of the lenses in a corrected CCS (adjusted for to elevation, heading, common_roll) float rXY [NUM_CAMS][3]; // XY pairs of the in a normal plane, relative to disparityRadius // float [][] rXY_ideal = {{-0.5, -0.5}, {0.5,-0.5}, {-0.5, 0.5}, {0.5,0.5}}; // only used for the multi-quad systems float cameraRadius; // =0; // average distance from the "mass center" of the sensors to the sensors float disparityRadius; // =150.0; // distance between cameras to normalize disparity units to. sqrt(2)*disparityRadius for quad };
