Implemented variable fat zero, depending on per-tile number of averaging (4398007c) · Commits · Elphel / tile_processor_gpu

src/TileProcessor.cuh

+38 −22

Original line number	Diff line number	Diff line
		@@ -921,7 +921,7 @@ __device__ void resetCorrelation(

		__device__ void normalizeTileAmplitude(
		float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports
		float fat_zero); // fat zero is absolute, scale it outside
		float fat_zero2); // fat zero is absolute, scale it outside

		__device__ void imclt8threads(// for 8 threads
		int do_acc, // 1 - add to previous value, 0 - overwrite
		@@ -1076,7 +1076,7 @@ extern "C" __global__ void correlate2D_inner(
		float scale0, // scale for R
		float scale1, // scale for B
		float scale2, // scale for G
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute
		size_t num_corr_tiles, // number of correlation tiles to process
		int * gpu_corr_indices, // packed tile+pair
		const size_t corr_stride, // in floats
		@@ -1087,9 +1087,10 @@ extern "C" __global__ void corr2D_normalize_inner(
		int num_corr_tiles, // number of correlation tiles to process
		const size_t corr_stride_td, // (in floats) stride for the input TD correlations
		float * gpu_corrs_td, // correlation tiles in transform domain
		float * corr_weights, // null or per-tile weight (fat_zero2 will be divided by it)
		const size_t corr_stride, // in floats
		float * gpu_corrs, // correlation output data (either pixel domain or transform domain
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute
		int corr_radius); // radius of the output correlation (7 for 15x15)

		extern "C" __global__ void corr2D_combine_inner(
		@@ -1151,7 +1152,7 @@ __device__ int get_textures_shared_size( // in bytes
		* @param scale0 scale red (or mono) component before mixing
		* @param scale1 scale blue (if colors = 3) component before mixing
		* @param scale2 scale green (if colors = 3) component before mixing
		* @param fat_zero add this value squared to the sum of squared components before normalization\
		* @param fat_zero2 add this value squared to the sum of squared components before normalization (squared)
		* @param gpu_ftasks flattened tasks, 29 floats for quad EO, 101 floats for LWIR16
		// * @param gpu_tasks array of per-tile tasks (now bits 4..9 - correlation pairs)
		* @param num_tiles number of tiles int gpu_tasks array prepared for processing
		@@ -1174,7 +1175,7 @@ extern "C" __global__ void correlate2D(
		float scale0, // scale for R
		float scale1, // scale for B
		float scale2, // scale for G
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		// struct tp_task * gpu_tasks, // array of per-tile tasks (now bits 4..9 - correlation pairs)
		int num_tiles, // number of tiles in task
		@@ -1211,7 +1212,7 @@ extern "C" __global__ void correlate2D(
		scale0, // float scale0, // scale for R
		scale1, // float scale1, // scale for B
		scale2, // float scale2, // scale for G
		fat_zero, // float fat_zero, // here - absolute
		fat_zero2, // float fat_zero2, // here - absolute
		*pnum_corr_tiles, // size_t num_corr_tiles, // number of correlation tiles to process
		gpu_corr_indices, // int * gpu_corr_indices, // packed tile+pair
		corr_stride, // const size_t corr_stride, // in floats
		@@ -1231,7 +1232,7 @@ extern "C" __global__ void correlate2D(
		* @param scale0 scale red (or mono) component before mixing
		* @param scale1 scale blue (if colors = 3) component before mixing
		* @param scale2 scale green (if colors = 3) component before mixing
		* @param fat_zero add this value squared to the sum of squared components before normalization
		* @param fat_zero2 add this value squared to the sum of squared components before normalization
		* @param num_corr_tiles number of correlation tiles to process
		* @param gpu_corr_indices packed array (each element, integer contains tile+pair) of correlation tasks
		* @param corr_stride, stride (in floats) for correlation outputs.
		@@ -1245,7 +1246,7 @@ extern "C" __global__ void correlate2D_inner(
		float scale0, // scale for R
		float scale1, // scale for B
		float scale2, // scale for G
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute
		size_t num_corr_tiles, // number of correlation tiles to process
		int * gpu_corr_indices, // packed tile+pair
		const size_t corr_stride, // in floats
		@@ -1384,13 +1385,13 @@ extern "C" __global__ void correlate2D_inner(
		if (corr_radius > 0) {
		normalizeTileAmplitude(
		clt_corr, // float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports
		fat_zero); // float fat_zero ) // fat zero is absolute, scale it outside
		fat_zero2); // float fat_zero2 ) // fat zero is absolute, scale it outside
		// Low Pass Filter from constant area (is it possible to replace?)

		#ifdef DBG_TILE
		#ifdef DEBUG6
		if ((tile_num == DBG_TILE) && (corr_pair == 0) && (threadIdx.x == 0)){
		printf("\ncorrelate2D CORRELATION NORMALIZED, fat_zero=%f\n",fat_zero);
		printf("\ncorrelate2D CORRELATION NORMALIZED, fat_zero2=%f\n",fat_zero2);
		debug_print_clt1(clt_corr, -1, 0xf);
		}
		__syncthreads();// __syncwarp();
		@@ -1695,18 +1696,20 @@ extern "C" __global__ void corr2D_combine_inner(
		* @param num_tiles number of correlation tiles to process
		* @param corr_stride_td, stride (in floats) for correlation input (transform domain).
		* @param gpu_corrs_td correlation data in transform domain
		* @param corr_weights null or per-tile weight (fat_zero2 will be divided by it)
		* @param corr_stride, stride (in floats) for correlation pixel-domain outputs.
		* @param gpu_corrs allocated array for the correlation output data (each element stride, payload: (2*corr_radius+1)^2
		* @param fat_zero add this value squared to the sum of squared components before normalization
		* @param fat_zero2 add this value squared to the sum of squared components before normalization (squared)
		* @param corr_radius, radius of the output correlation (maximal 7 for 15x15)
		*/
		extern "C" __global__ void corr2D_normalize(
		int num_corr_tiles, // number of correlation tiles to process
		const size_t corr_stride_td, // in floats
		float * gpu_corrs_td, // correlation tiles in transform domain
		float * corr_weights, // null or per correlation tile weight (fat_zero2 will be divided by it)
		const size_t corr_stride, // in floats
		float * gpu_corrs, // correlation output data (either pixel domain or transform domain
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute, squared
		int corr_radius) // radius of the output correlation (7 for 15x15)
		{
		if (threadIdx.x == 0) { // only 1 thread, 1 block
		@@ -1716,9 +1719,10 @@ extern "C" __global__ void corr2D_normalize(
		num_corr_tiles, // int num_corr_tiles, // number of correlation tiles to process
		corr_stride_td, // const size_t corr_stride, // in floats
		gpu_corrs_td, // float * gpu_corrs_td, // correlation tiles in transform domain
		corr_weights, // float * corr_weights, // null or per-tile weight (fat_zero2 will be divided by it)
		corr_stride, // const size_t corr_stride, // in floats
		gpu_corrs, // float * gpu_corrs, // correlation output data (either pixel domain or transform domain
		fat_zero, // float fat_zero, // here - absolute
		fat_zero2, // float fat_zero2, // here - absolute
		corr_radius); // int corr_radius, // radius of the output correlation (7 for 15x15)
		}
		}
		@@ -1730,9 +1734,10 @@ extern "C" __global__ void corr2D_normalize(
		* @param num_tiles number of correlation tiles to process
		* @param corr_stride_td, stride (in floats) for correlation input (transform domain).
		* @param gpu_corrs_td correlation data in transform domain
		* @param corr_weights null or per-tile weight (fat_zero2 will be divided by it)
		* @param corr_stride, stride (in floats) for correlation pixel-domain outputs.
		* @param gpu_corrs allocated array for the correlation output data (each element stride, payload: (2*corr_radius+1)^2
		* @param fat_zero add this value squared to the sum of squared components before normalization
		* @param fat_zero2 add this value squared to the sum of squared components before normalization
		* @param corr_radius, radius of the output correlation (maximal 7 for 15x15)
		*/

		@@ -1740,9 +1745,10 @@ extern "C" __global__ void corr2D_normalize_inner(
		int num_corr_tiles, // number of correlation tiles to process
		const size_t corr_stride_td, // (in floats) stride for the input TD correlations
		float * gpu_corrs_td, // correlation tiles in transform domain
		float * corr_weights, // null or per-tile weight (fat_zero2 will be divided by it)
		const size_t corr_stride, // in floats
		float * gpu_corrs, // correlation output data (either pixel domain or transform domain
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute, squared
		int corr_radius) // radius of the output correlation (7 for 15x15)
		{
		int corr_in_block = threadIdx.y;
		@@ -1753,6 +1759,7 @@ extern "C" __global__ void corr2D_normalize_inner(
		__syncthreads();// __syncwarp();
		__shared__ float clt_corrs [CORR_TILES_PER_BLOCK_NORMALIZE][4][DTT_SIZE][DTT_SIZE1];
		__shared__ float mlt_corrs [CORR_TILES_PER_BLOCK_NORMALIZE][DTT_SIZE2M1][DTT_SIZE2M1]; // result correlation
		__shared__ float norm_fat_zero [CORR_TILES_PER_BLOCK_NORMALIZE];
		// set clt_corr to all zeros
		float * clt_corr = ((float ) clt_corrs) + corr_in_block (4 * DTT_SIZE * DTT_SIZE1); // top left quadrant0
		float * mclt_corr = ((float ) mlt_corrs) + corr_in_block (DTT_SIZE2M1*DTT_SIZE2M1);
		@@ -1768,16 +1775,25 @@ extern "C" __global__ void corr2D_normalize_inner(
		}
		__syncthreads();// __syncwarp();

		if (threadIdx.x == 0){
		norm_fat_zero[corr_in_block] = fat_zero2;
		if (corr_weights) { // same for all
		norm_fat_zero[corr_in_block] /= * (corr_weights + corr_num);
		}
		}
		__syncthreads();// __syncwarp();


		// normalize Amplitude
		normalizeTileAmplitude(
		clt_corr, // float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports
		fat_zero); // float fat_zero ) // fat zero is absolute, scale it outside
		norm_fat_zero[corr_in_block]); // fat_zero2); // float fat_zero2 ) // fat zero is absolute, scale it outside
		// Low Pass Filter from constant area (is it possible to replace?)

		#ifdef DBG_TILE
		#ifdef DEBUG6
		if ((tile_num == DBG_TILE) && (corr_pair == 0) && (threadIdx.x == 0)){
		printf("\ncorrelate2D CORRELATION NORMALIZED, fat_zero=%f\n",fat_zero);
		printf("\ncorrelate2D CORRELATION NORMALIZED, fat_zero2=%f\n",fat_zero2);
		debug_print_clt1(clt_corr, -1, 0xf);
		}
		__syncthreads();// __syncwarp();
		@@ -3976,12 +3992,12 @@ __device__ void resetCorrelation(
		* Called from correlate2D()->correlate2D_inner()
		*
		* @param clt_tile pointer to a correlation result tile [4][8][8+1] to be normalized
		* @param fat_zero value to add to amplitudes for regularization. Absolute value,
		* @param fat_zero2 value to add to amplitudes for regularization. Absolute value,
		* scale if needed outside.
		*/
		__device__ void normalizeTileAmplitude(
		float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports
		float fat_zero ) // fat zero is absolute, scale it outside
		float fat_zero2 ) // fat zero is absolute, scale it outside
		{
		int joffs = threadIdx.x * DTT_SIZE1;
		float * clt_tile_j0 = clt_tile + joffs; // ==&clt_tile[0][j][0]
		@@ -3990,7 +4006,7 @@ __device__ void normalizeTileAmplitude(
		float * clt_tile_j3 = clt_tile_j2 + (DTT_SIZE1*DTT_SIZE); // ==&clt_tile[3][j][0]
		#pragma unroll
		for (int i = 0; i < DTT_SIZE; i++) {
		float s2 = fat_zero * fat_zero +
		float s2 = fat_zero2 +
		(clt_tile_j0) *(clt_tile_j0) +
		(clt_tile_j1) *(clt_tile_j1) +
		(clt_tile_j2) *(clt_tile_j2) +

src/TileProcessor.h

+3 −2

Original line number	Diff line number	Diff line
		@@ -75,7 +75,7 @@ extern "C" __global__ void correlate2D(
		float scale0, // scale for R
		float scale1, // scale for B
		float scale2, // scale for G
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute, squared
		float * gpu_ftasks, // flattened tasks, 29 floats for quad EO, 101 floats for LWIR16
		// struct tp_task * gpu_tasks, // array of per-tile tasks (now bits 4..9 - correlation pairs)
		int num_tiles, // number of tiles in task
		@@ -90,9 +90,10 @@ extern "C" __global__ void corr2D_normalize(
		int num_corr_tiles, // number of correlation tiles to process
		const size_t corr_stride_td, // in floats
		float * gpu_corrs_td, // correlation tiles in transform domain
		float * corr_weights, // null or per-tile weight (fat_zero2 will be divided by it)
		const size_t corr_stride, // in floats
		float * gpu_corrs, // correlation output data (either pixel domain or transform domain
		float fat_zero, // here - absolute
		float fat_zero2, // here - absolute, squared
		int corr_radius); // radius of the output correlation (7 for 15x15)

		extern "C" __global__ void corr2D_combine(