troubleshooting (fcd21d6e) · Commits · Elphel / tile_processor_gpu

src/TileProcessor.cuh

+1 −4

Original line number	Original line	Diff line number	Diff line
	@@ -1902,8 +1902,7 @@ extern "C" __global__ void textures_nonoverlap(


	//#undef USE_textures_gen		//#undef USE_textures_gen
	extern "C"		extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
	__global__ void textures_accumulate( // (8,4,1) (N,1,1)
	int * woi, // x, y, width,height		int * woi, // x, y, width,height
	float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]		float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
	size_t num_texture_tiles, // number of texture tiles to process		size_t num_texture_tiles, // number of texture tiles to process
	@@ -3952,7 +3951,6 @@ __device__ void tile_combine_rgba(
	}		}
	max_diff_shared[cam] = sqrtf(mx);		max_diff_shared[cam] = sqrtf(mx);
	}		}
	__syncthreads(); //?
	#ifdef DEBUG22		#ifdef DEBUG22
	if (debug && (threadIdx.x == 0) && (threadIdx.y == 0)){		if (debug && (threadIdx.x == 0) && (threadIdx.y == 0)){
	printf("\n 1. max_diff\n");		printf("\n 1. max_diff\n");
	@@ -4022,7 +4020,6 @@ __device__ void tile_combine_rgba(
	ports_rgb_shared[ncol][cam] /= DTT_SIZE2*DTT_SIZE2; // correct for window?		ports_rgb_shared[ncol][cam] /= DTT_SIZE2*DTT_SIZE2; // correct for window?
	}		}
	}		}
	__syncthreads(); //?
	#ifdef DEBUG22		#ifdef DEBUG22
	if (debug && (threadIdx.x == 0) && (threadIdx.y == 0)){		if (debug && (threadIdx.x == 0) && (threadIdx.y == 0)){
	printf("\n 2. max_diff\n");		printf("\n 2. max_diff\n");

src/TileProcessor.h

+13 −20

Original line number	Original line	Diff line number	Diff line
	@@ -120,22 +120,13 @@ extern "C" __global__ void imclt_rbg(
	int woi_twidth,		int woi_twidth,
	int woi_theight,		int woi_theight,
	const size_t dstride); // in floats (pixels)		const size_t dstride); // in floats (pixels)
	/*		extern "C" __global__ void textures_accumulate( // (8,4,1) (N,1,1)
	extern "C" __global__ void generate_RBGA(		int * woi, // x, y, width,height
	// Parameters to generate texture tasks
	struct tp_task * gpu_tasks,
	int num_tiles, // number of tiles in task list
	// declare arrays in device code?
	int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
	int * num_texture_tiles, // number of texture tiles to process (8 separate elements for accumulation)
	int * woi, // x,y,width,height of the woi
	int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
	int height, // <= TILESY, use for faster processing of LWIR images
	// Parameters for the texture generation
	float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]		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)
	// TODO: use geometry_correction rXY !		// TODO: use geometry_correction rXY !
	struct gc * gpu_geometry_correction,		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 colors, // number of colors (3/1)
	int is_lwir, // do not perform shot correction		int is_lwir, // do not perform shot correction
	float min_shot, // 10.0		float min_shot, // 10.0
	@@ -143,11 +134,13 @@ extern "C" __global__ void generate_RBGA(
	float diff_sigma, // pixel value/pixel change		float diff_sigma, // pixel value/pixel change
	float diff_threshold, // 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 min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
	float weight0, // scale for R		float weights[3], // scale for R,B,G
	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 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)		int keep_weights, // return channel weights after A in RGBA (was removed) (should be 0 if gpu_texture_rbg)?
	const size_t texture_rbga_stride, // in floats		// combining both non-overlap and overlap (each calculated if pointer is not null )
	float * gpu_texture_tiles); // (number of colors +1 + ?)1616 rgba texture tiles		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
			float * gpu_diff_rgb_combo); // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]

src/test_tp.cu

+2 −2

Original line number	Original line	Diff line number	Diff line
	@@ -1276,8 +1276,8 @@ int main(int argc, char **argv)
	1, // int dust_remove, // Do not reduce average weight when only one image differes much from the average		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		0, // int keep_weights, // return channel weights after A in RGBA
	dstride_textures_rbga/sizeof(float), // const size_t texture_rbga_stride, // in floats		dstride_textures_rbga/sizeof(float), // const size_t texture_rbga_stride, // in floats
	gpu_textures_rbga, // float * gpu_texture_tiles) // (number of colors +1 + ?)1616 rgba texture tiles		gpu_textures_rbga); // float * gpu_texture_tiles) // (number of colors +1 + ?)1616 rgba texture tiles
	(float ) 0 ); // gpu_diff_rgb_combo); // float gpu_diff_rgb_combo) // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]		// (float ) 0 ); // gpu_diff_rgb_combo); // float gpu_diff_rgb_combo) // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]

	getLastCudaError("Kernel failure");		getLastCudaError("Kernel failure");
	checkCudaErrors(cudaDeviceSynchronize());		checkCudaErrors(cudaDeviceSynchronize());

src/tp_defines.h

+1 −1

Original line number	Original line	Diff line number	Diff line
	@@ -118,7 +118,7 @@
	//#define DEBUG21 1 // Geometry Correction		//#define DEBUG21 1 // Geometry Correction

	//#define DEBUG22 1		//#define DEBUG22 1
	#define DEBUG23 1		//#define DEBUG23 1

	#endif //#if (DBG_TILE_X >= 0) && (DBG_TILE_Y >= 0)		#endif //#if (DBG_TILE_X >= 0) && (DBG_TILE_Y >= 0)

Admin message