Refactoring, working state (1bf5e8bf) · Commits · Elphel / tile_processor_gpu

src/TileProcessor.cuh→src/TileProcessor.cu

+11 −131

Original line number	Diff line number	Diff line
		@@ -36,12 +36,17 @@
		* \brief Top level of the Tile Processor for frequency domain
		*/
		// Avoiding includes in jcuda, all source files will be merged
		#pragma once
		//#pragma once
		#ifndef JCUDA
		#include "tp_defines.h"
		#include "dtt8x8.h"
		#include "geometry_correction.h"
		#include "TileProcessor.h"

		#include <cuda_runtime.h>
		// #include <helper_cuda.h>
		// #include <helper_functions.h>

		#endif // #ifndef JCUDA

		// CUDA fast math is slower!
		@@ -55,67 +60,16 @@ GPU run time =523.451927ms, (direct conversion: 24.080189999999998ms, imclt: 17.

		//#define TASK_TEXTURE_BITS ((1 << TASK_TEXTURE_N_BIT) \| (1 << TASK_TEXTURE_E_BIT) \| (1 << TASK_TEXTURE_S_BIT) \| (1 << TASK_TEXTURE_W_BIT))

		#define 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))


		//#define IMCLT14
		//#define NOICLT 1
		//#define TEST_IMCLT
		//#define SAVE_CLT
		// Not enough shared memory to have more threads per block,even just for the result clt tiles
		// What to do:
		// 1) make single image aberration correction: 1/4 of the result tiles
		// With 4 cameras = calculate correlations (9x9), reusing kernel or just clt ones after color reducing, then output them to device memory
		//Average run time =1308.124146 ms
		//#define TILES_PER_BLOCK 2
		//Average run time =12502.638672 - with 2 tiles/block it is longer!
		///12129.268555 ms
		//Average run time =4704.506348 ms (syncwarp)
		//Average run time =4705.612305 ms (syncthreads)
		//Average run time =1051.411255 ms
		//Average run time =861.866577 ms
		//Average run time =850.871277 ms had bugs
		//Average run time =857.947632 ms fixed bugs


		// Something broke, even w/o LPF: Average run time =1093.115112 ms
		// without clt copying to device memory - Average run time =965.342407 ms - still worse
		//Average run time =965.880554 ms
		// combined tx and ty into a single int : Average run time =871.017944 ms
		//Average run time =873.386597 ms (reduced number of registers)
		//__umul24 : Average run time =879.125122 ms
		// without __umul24 - back to Average run time =871.315552 ms
		// Added copying clt to device memory - Average run time =942.071960 ms
		// Removed rest of NOICLT : Average run time =943.456177 ms
		// Added lpf: Average run time =1046.101318 ms (0.1 sec, 10%) - can be combined with the PSF kernel
		//#define USE_UMUL24
		////#define TILES_PER_BLOCK 4
		//Average run time =5155.922852 ms
		//Average run time =1166.388306 ms
		//Average run time =988.750977 ms
		//#define TILES_PER_BLOCK 8
		//Average run time =9656.743164 ms
		// Average run time =9422.057617 ms (reducing divergence)
		//#define TILES_PER_BLOCK 1

		//#define THREADS_PER_TILE 8
		//#define IMCLT_THREADS_PER_TILE 16
		//#define IMCLT_TILES_PER_BLOCK 4


		/// #define KERNELS_STEP (1 << KERNELS_LSTEP)
		//#define TILES-X (IMG-WIDTH / DTT_SIZE)
		//#define TILES-Y (IMG-HEIGHT / DTT_SIZE)
		/*
		#define CONVERT_DIRECT_INDEXING_THREADS_LOG2 5
		#define CONVERT_DIRECT_INDEXING_THREADS (1 << CONVERT_DIRECT_INDEXING_THREADS_LOG2) // 32

		*/
		// Make TILES-YA >= TILES-X and a multiple of 4
		//#define TILES-YA ((TILES-Y +3) & (~3))

		// increase row length by 1 so vertical passes will use different ports
		#define MCLT_UNION_LEN (DTT_SIZE2 * (DTT_SIZE2 + 2))

		//#define MCLT_UNION_LEN (DTT_SIZE2 * (DTT_SIZE2 + 2))
		/*
		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)
		@@ -126,7 +80,7 @@ struct CltExtra{
		float dyc_dx;
		float dyc_dy;
		};

		*/
		/*
		Python code to generate constant coefficients:
		def setup_hwindow(n=8, l=4):
		@@ -735,48 +689,11 @@ __device__ void imclt_plane( // not implemented, not used
		float * gpu_rbg, // WIDTH, HEIGHT
		const size_t dstride); // in floats (pixels)

		extern "C" __global__ void clear_texture_list(
		int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
		int width, // <= TILES-X, use for faster processing of LWIR images
		int height); // <= TILES-Y, use for faster processing of LWIR images

		extern "C" __global__ void mark_texture_tiles(
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int * gpu_texture_indices);// packed tile + bits (now only (1 << 7)

		extern "C" __global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int height, // number of tiles rows
		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(
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int height, // number of tiles rows
		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
		int * num_texture_tiles, // number of texture tiles to process
		int * woi); // min_x, min_y, max_x, max_y input

		extern "C" __global__ void update_woi(
		int texture_slices,
		int * woi, // min_x, min_y, max_x, max_y input, not modified, max_x - not used
		int * twh); // 2-element in device global memory

		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 + ?)1616 rgba texture tiles

		extern "C" __global__ void clear_texture_rbga2( // version for CDP2
		int * twh, // texture_width, // aligned to DTT_SIZE
		const size_t texture_rbga_stride, // in floats 8*stride
		@@ -818,15 +735,6 @@ __global__ void index_inter_correlate(
		int * gpu_corr_indices, // array of correlation tasks
		int * pnum_corr_tiles); // pointer to the length of correlation tasks array

		extern "C" __global__ void create_nonoverlap_list(
		int num_cams,
		float * gpu_ftasks , // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task
		int width, // number of tiles in a row
		int * nonoverlap_list, // pointer to the calculated number of non-zero tiles
		int * pnonoverlap_length); // indices to gpu_tasks // should be initialized to zero


		__global__ void convert_correct_tiles(
		int num_cams, // actual number of cameras
		int num_colors, // actual number of colors: 3 for RGB, 1 for LWIR/mono
		@@ -906,34 +814,6 @@ extern "C" __global__ void corr2D_combine_inner(
		const size_t corr_stride_combo, // (in floats) stride for the output TD correlations (same as input)
		float * gpu_corrs_combo); // combined correlation output (one per tile)

		extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
		int num_cams, // number of cameras used
		int * woi, // x, y, width,height
		float ** gpu_clt, // [num_cams] ->[TILES-Y][TILES-X][colors][DTT_SIZE*DTT_SIZE]
		/// size_t num_texture_tiles, // number of texture tiles to process
		int * pnum_texture_tiles, // pointer to a number of texture tiles to process
		int gpu_texture_indices_offset,// add to gpu_texture_indices
		int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
		// TODO: use geometry_correction rXY !
		struct gc * gpu_geometry_correction,
		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)
		const float weights[3], // scale for R,B,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 + ?)1616 rgba texture tiles
		size_t texture_stride, // in floats (now 256*4 = 1024)
		float * gpu_texture_tiles, // (number of colors +1 + ?)1616 rgba texture tiles
		int linescan_order, // if !=0 then output gpu_diff_rgb_combo in linescan order, else - in gpu_texture_indices order
		float * gpu_diff_rgb_combo, //) // diff[num_cams], R[num_cams], B[num_cams],G[num_cams]
		int tilesx);

		__device__ int get_textures_shared_size( // in bytes
		int num_cams, // actual number of cameras

src/TileProcessor.h

+101 −0

Original line number	Diff line number	Diff line
		@@ -37,10 +37,34 @@

		*/
		#pragma once
		#ifndef TILE_PROCESSOR_H_
		#define TILE_PROCESSOR_H_


		#ifndef NUM_CAMS
		#include "tp_defines.h"
		#endif

		#define 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))

		#define CONVERT_DIRECT_INDEXING_THREADS_LOG2 5
		#define CONVERT_DIRECT_INDEXING_THREADS (1 << CONVERT_DIRECT_INDEXING_THREADS_LOG2) // 32

		#define MCLT_UNION_LEN (DTT_SIZE2 * (DTT_SIZE2 + 2))

		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)
		float center_x; // actual center X (use to find derivatives)
		float center_y; // actual center X (use to find derivatives)
		float dxc_dx; // add this to data_x per each pixel X-shift relative to the kernel center location
		float dxc_dy; // same per each Y-shift pixel
		float dyc_dx;
		float dyc_dy;
		};



		extern "C" __global__ void convert_direct( // called with a single block, single thread
		// struct CltExtra ** gpu_kernel_offsets, // [NUM_CAMS], // changed for jcuda to avoid struct parameters
		@@ -219,4 +243,81 @@ extern "C" __global__ void generate_RBGA(
		float * gpu_texture_tiles, // (number of colors +1 + ?)1616 rgba texture tiles
		int * twh);

		extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
		int num_cams, // number of cameras used
		int * woi, // x, y, width,height
		float ** gpu_clt, // [num_cams] ->[TILES-Y][TILES-X][colors][DTT_SIZE*DTT_SIZE]
		/// size_t num_texture_tiles, // number of texture tiles to process
		int * pnum_texture_tiles, // pointer to a number of texture tiles to process
		int gpu_texture_indices_offset,// add to gpu_texture_indices
		int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
		// TODO: use geometry_correction rXY !
		struct gc * gpu_geometry_correction,
		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)
		const float weights[3], // scale for R,B,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 + ?)1616 rgba texture tiles
		size_t texture_stride, // in floats (now 256*4 = 1024)
		float * gpu_texture_tiles, // (number of colors +1 + ?)1616 rgba texture tiles
		int linescan_order, // if !=0 then output gpu_diff_rgb_combo in linescan order, else - in gpu_texture_indices order
		float * gpu_diff_rgb_combo, //) // diff[num_cams], R[num_cams], B[num_cams],G[num_cams]
		int tilesx);

		extern "C" __global__ void clear_texture_list(
		int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
		int width, // <= TILES-X, use for faster processing of LWIR images
		int height); // <= TILES-Y, use for faster processing of LWIR images

		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 + ?)1616 rgba texture tiles

		extern "C" __global__ void create_nonoverlap_list(
		int num_cams,
		float * gpu_ftasks , // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task
		int width, // number of tiles in a row
		int * nonoverlap_list, // pointer to the calculated number of non-zero tiles
		int * pnonoverlap_length); // indices to gpu_tasks // should be initialized to zero

		extern "C" __global__ void mark_texture_tiles(
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int * gpu_texture_indices);// packed tile + bits (now only (1 << 7)

		extern "C" __global__ void mark_texture_neighbor_tiles( // TODO: remove __global__?
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int height, // number of tiles rows
		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(
		int num_cams,
		float * gpu_ftasks, // flattened tasks, 27 floats for quad EO, 99 floats for LWIR16
		int num_tiles, // number of tiles in task list
		int width, // number of tiles in a row
		int height, // number of tiles rows
		int * gpu_texture_indices, // packed tile + bits (now only (1 << 7)
		int * num_texture_tiles, // number of texture tiles to process
		int * woi); // min_x, min_y, max_x, max_y input




		#endif

src/geometry_correction.cu

+5 −0

Original line number	Diff line number	Diff line
		@@ -43,6 +43,11 @@

		// #include "TileProcessor.h"

		#include <cuda_runtime.h>
		// #include <helper_cuda.h>
		// #include <helper_functions.h>


		#endif // #ifndef JCUDA

		#ifndef get_task_size

src/test_tp.cu

+9 −4

Original line number	Diff line number	Diff line
		@@ -33,9 +33,9 @@
		// all of the next 5 were disabled
		#define NOCORR
		#define NOCORR_TD
		//#define NOTEXTURES
		//#define NOTEXTURE_RGBA
		//#define NOTEXTURE_RGBAXXX
		#define NOTEXTURES //
		#define NOTEXTURE_RGBA //
		//#define NOTEXTURE_RGBAXXX //


		#define SAVE_CLT
		@@ -55,11 +55,15 @@
		#include <iterator>
		#include <vector>


		#include "tp_defines.h" // was not here
		#include "dtt8x8.h"
		#include "geometry_correction.h"
		#include "TileProcessor.cuh"
		//#include "TileProcessor.cuh"
		#include "TileProcessor.h"
		#include "tp_utils.h"
		#include "tp_files.h"
		//#include "tp_paths.cuh"
		#include "tp_paths.h"

		#if TEST_LWIR
		@@ -2130,6 +2134,7 @@ int main(int argc, char **argv)
		cudaMemcpyDeviceToHost));

		#ifndef NSAVE_TEXTURES
		printf("Import raw, real, little-endian, 18 images 648x520\n");
		#ifdef NO_DP
		printf("Writing RBGA texture slices to %s\n", tpPaths.result_textures_rgba_file);
		writeFloatsToFile(

src/tp_files.cu

+45 −2

Original line number	Diff line number	Diff line
		@@ -17,7 +17,7 @@


		#include "tp_files.h"
		int get_file_size(std::string filename) // path to file
		int get_file_size(const std::string filename) // path to file
		{
		FILE *p_file = NULL;
		p_file = fopen(filename.c_str(),"rb");
		@@ -26,6 +26,7 @@ int get_file_size(std::string filename) // path to file
		fclose(p_file);
		return size;
		}

		int readFloatsFromFile(float * data, // allocated array
		const char * path) // file path
		{
		@@ -42,6 +43,8 @@ int readFloatsFromFile(float * data, // allocated array
		return 0;
		}



		float * readAllFloatsFromFile(const char * path,
		int * len_in_floats) //
		{
		@@ -55,7 +58,6 @@ float * readAllFloatsFromFile(const char * path,
		printf("---- Bytes read: %d from %s\n", fsize, path);
		* len_in_floats = fsize/sizeof(float);
		return data;

		}

		int writeFloatsToFile(float * data, // allocated array
		@@ -68,5 +70,46 @@ int writeFloatsToFile(float * data, // allocated array
		return 0;
		}

		/*
		int readFloatsFromFile(float * data, // allocated array
		const std::string path) // file path
		{
		printf("readFloatsFromFile(%s)\n", path);

		int fsize = get_file_size(path);
		std::ifstream input(path, std::ios::binary );
		// copies all data into buffer
		std::vector<char> buffer((
		std::istreambuf_iterator<char>(input)),
		(std::istreambuf_iterator<char>()));
		std::copy( buffer.begin(), buffer.end(), (char *) data);
		printf("---- Bytes read: %d from %s\n", fsize, path);
		return 0;
		}

		float * readAllFloatsFromFile(const std::string path,
		int * len_in_floats) //
		{
		int fsize = get_file_size(path);
		float * data = (float *) malloc(fsize);
		std::ifstream input(path, std::ios::binary );
		std::vector<char> buffer((
		std::istreambuf_iterator<char>(input)),
		(std::istreambuf_iterator<char>()));
		std::copy( buffer.begin(), buffer.end(), (char *) data);
		printf("---- Bytes read: %d from %s\n", fsize, path);
		* len_in_floats = fsize/sizeof(float);
		return data;
		}

		int writeFloatsToFile(float * data, // allocated array
		int size, // length in elements
		const std::string path) // file path
		{
		std::ofstream ofile(path, std::ios::binary);
		ofile.write((char ) data, size sizeof(float));
		return 0;
		}

		*/

Admin message