using geometry correction for port offsets (d8f9defc) · Commits · Elphel / tile_processor_gpu

src/TileProcessor.cuh

+11 −9

Original line number	Diff line number	Diff line
		@@ -1182,7 +1182,9 @@ __global__ void generate_RBGA(
		// Parameters for the texture generation
		float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
		// TODO: use geometry_correction rXY !
		float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
		struct gc * gpu_geometry_correction,
		// float * gpu_geometry_correction,
		// 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
		@@ -1190,9 +1192,7 @@ __global__ void generate_RBGA(
		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
		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)
		const size_t texture_rbga_stride, // in floats
		@@ -1200,6 +1200,7 @@ __global__ void generate_RBGA(
		{
		// TODO use atomic_add to increment num_texture_tiles
		// TODO calculate woi

		dim3 threads0((1 << THREADS_DYNAMIC_BITS), 1, 1);
		int blocks_x = (width + ((1 << THREADS_DYNAMIC_BITS) - 1)) >> THREADS_DYNAMIC_BITS;
		dim3 blocks0 (blocks_x, height, 1);
		@@ -1311,7 +1312,8 @@ __global__ void generate_RBGA(
		gpu_clt, // float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
		ntt, // size_t num_texture_tiles, // number of texture tiles to process
		gpu_texture_indices + ti_offset, // int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
		gpu_port_offsets, // float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
		// gpu_port_offsets, // float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
		(float *) gpu_geometry_correction ->pXY0,
		colors, // int colors, // number of colors (3/1)
		is_lwir, // int is_lwir, // do not perform shot correction
		min_shot, // float min_shot, // 10.0
		@@ -1319,9 +1321,9 @@ __global__ void generate_RBGA(
		diff_sigma, // float diff_sigma, // pixel value/pixel change
		diff_threshold, // float diff_threshold, // pixel value/pixel change
		min_agree, // float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
		weight0, // float weight0, // scale for R
		weight1, // float weight1, // scale for B
		weight2, // float weight2, // scale for G
		weights[0], // float weight0, // scale for R
		weights[1], // float weight1, // scale for B
		weights[2], // float weight2, // scale for G
		dust_remove, // int dust_remove, // Do not reduce average weight when only one image differs much from the average
		0, // 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 )

src/TileProcessor.h

+3 −2

Original line number	Diff line number	Diff line
		@@ -121,7 +121,7 @@ extern "C" __global__ void imclt_rbg(
		int woi_twidth,
		int woi_theight,
		const size_t dstride); // in floats (pixels)

		/*
		extern "C" __global__ void generate_RBGA(
		// Parameters to generate texture tasks
		struct tp_task * gpu_tasks,
		@@ -135,6 +135,7 @@ extern "C" __global__ void generate_RBGA(
		// Parameters for the texture generation
		float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
		// TODO: use geometry_correction rXY !
		struct gc * gpu_geometry_correction,
		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
		@@ -150,4 +151,4 @@ extern "C" __global__ void generate_RBGA(
		int keep_weights, // return channel weights after A in RGBA (was removed)
		const size_t texture_rbga_stride, // in floats
		float * gpu_texture_tiles); // (number of colors +1 + ?)1616 rgba texture tiles

		*/

src/test_tp.cu

+15 −4

Original line number	Diff line number	Diff line
		@@ -346,6 +346,7 @@ int main(int argc, char **argv)
		int * gpu_woi;
		int * gpu_num_texture_tiles;
		float * gpu_port_offsets;
		float * gpu_color_weights;
		int num_corrs;
		int num_textures;
		int num_ports = NUM_CAMS;
		@@ -556,7 +557,13 @@ int main(int argc, char **argv)
		// number of border tiles

		// copy port indices to gpu
		float color_weights [] = {
		0.294118, // float weight0, // scale for R
		0.117647, // float weight1, // scale for B
		0.588235}; // float weight2, // scale for G

		gpu_port_offsets = (float ) copyalloc_kernel_gpu((float ) port_offsets, num_ports * 2);
		gpu_color_weights = (float ) copyalloc_kernel_gpu((float ) color_weights, sizeof(color_weights));



		@@ -1270,7 +1277,9 @@ int main(int argc, char **argv)
		TILESY, // int height); // <= TILESY, use for faster processing of LWIR images
		// Parameters for the texture generation
		gpu_clt , // float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
		gpu_port_offsets, // float * port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
		// (float *)
		gpu_geometry_correction, // struct gc * gpu_geometry_correction,
		// gpu_port_offsets, // float * port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
		texture_colors, // int colors, // number of colors (3/1)
		(texture_colors == 1), // int is_lwir, // do not perform shot correction
		10.0, // float min_shot, // 10.0
		@@ -1278,9 +1287,10 @@ int main(int argc, char **argv)
		1.5f, // float diff_sigma, // pixel value/pixel change
		10.0f, // float diff_threshold, // pixel value/pixel change
		3.0, // float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
		0.294118, // float weight0, // scale for R
		0.117647, // float weight1, // scale for B
		0.588235, // float weight2, // scale for G
		gpu_color_weights, // float weights[3], // scale for R
		// 0.294118, // float weight0, // scale for R
		// 0.117647, // float weight1, // scale for B
		// 0.588235, // float weight2, // scale for G
		1, // int dust_remove, // Do not reduce average weight when only one image differes much from the average
		0, // int keep_weights, // return channel weights after A in RGBA
		dstride_textures_rbga/sizeof(float), // const size_t texture_rbga_stride, // in floats
		@@ -1372,6 +1382,7 @@ int main(int argc, char **argv)
		checkCudaErrors(cudaFree(gpu_num_corr_tiles));
		checkCudaErrors(cudaFree(gpu_texture_indices));
		checkCudaErrors(cudaFree(gpu_port_offsets));
		checkCudaErrors(cudaFree(gpu_color_weights));
		checkCudaErrors(cudaFree(gpu_textures));
		checkCudaErrors(cudaFree(gpu_textures_rbga));
		checkCudaErrors(cudaFree(gpu_woi));

Admin message