testing GC direct+inverse for 16 mono. Need to remove debayer (d26457f8) · Commits · Elphel / tile_processor_gpu

src/geometry_correction.cu

+43 −24

Original line number	Diff line number	Diff line
		@@ -57,8 +57,8 @@

		#define DBG_CAM 3

		__device__ void printGeometryCorrection(struct gc * g);
		__device__ void printExtrinsicCorrection(corr_vector * cv);
		__device__ void printGeometryCorrection(struct gc * g, int num_cams);
		__device__ void printExtrinsicCorrection(corr_vector * cv, int num_cams);


		/**
		@@ -436,8 +436,8 @@ extern "C" __global__ void get_tiles_offsets(
		if ((ncam == DBG_CAM) && (task_num == DBG_TILE)){
		printf("\nTile = %d, camera= %d\n", task_num, ncam);
		printf("\nget_tiles_offsets() threadIdx.x = %d, threadIdx.y = %d,blockIdx.x= %d\n", (int)threadIdx.x, (int)threadIdx.y, (int) blockIdx.x);
		printGeometryCorrection(&geometry_correction);
		printExtrinsicCorrection(&extrinsic_corr);
		printGeometryCorrection(&geometry_correction, num_cams);
		printExtrinsicCorrection(&extrinsic_corr,num_cams);
		}
		__syncthreads();// __syncwarp();
		#endif // DEBUG21
		@@ -852,7 +852,7 @@ inline __device__ float getRByRDist(float rDist,
		return result;
		}

		__device__ void printGeometryCorrection(struct gc * g){
		__device__ void printGeometryCorrection(struct gc * g, int num_cams){
		#ifndef JCUDA
		printf("\nGeometry Correction\n------------------\n");
		printf("%22s: %f\n","pixelCorrectionWidth", g->pixelCorrectionWidth);
		@@ -874,39 +874,58 @@ __device__ void printGeometryCorrection(struct gc * g){
		printf("%22s: %f\n","elevation", g->elevation);
		printf("%22s: %f\n","heading", g->heading);

		printf("%22s: %f, %f, %f, %f \n","forward", g->forward[0], g->forward[1], g->forward[2], g->forward[3]);
		printf("%22s: %f, %f, %f, %f \n","right", g->right[0], g->right[1], g->right[2], g->right[3]);
		printf("%22s: %f, %f, %f, %f \n","height", g->height[0], g->height[1], g->height[2], g->height[3]);
		printf("%22s: %f, %f, %f, %f \n","roll", g->roll[0], g->roll[1], g->roll[2], g->roll[3]);
		printf("%22s: %f, %f \n", "pXY0[0]", g->pXY0[0][0], g->pXY0[0][1]);
		printf("%22s: %f, %f \n", "pXY0[1]", g->pXY0[1][0], g->pXY0[1][1]);
		printf("%22s: %f, %f \n", "pXY0[2]", g->pXY0[2][0], g->pXY0[2][1]);
		printf("%22s: %f, %f \n", "pXY0[3]", g->pXY0[3][0], g->pXY0[3][1]);
		// printf("%22s: %f, %f, %f, %f \n","forward", g->forward[0], g->forward[1], g->forward[2], g->forward[3]);
		// printf("%22s: %f, %f, %f, %f \n","right", g->right[0], g->right[1], g->right[2], g->right[3]);
		// printf("%22s: %f, %f, %f, %f \n","height", g->height[0], g->height[1], g->height[2], g->height[3]);
		// printf("%22s: %f, %f, %f, %f \n","roll", g->roll[0], g->roll[1], g->roll[2], g->roll[3]);
		// printf("%22s: %f, %f \n", "pXY0[0]", g->pXY0[0][0], g->pXY0[0][1]);
		// printf("%22s: %f, %f \n", "pXY0[1]", g->pXY0[1][0], g->pXY0[1][1]);
		// printf("%22s: %f, %f \n", "pXY0[2]", g->pXY0[2][0], g->pXY0[2][1]);
		// printf("%22s: %f, %f \n", "pXY0[3]", g->pXY0[3][0], g->pXY0[3][1]);
		printf("%22s:","forward"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", g->forward[ncam]); printf("\n");
		printf("%22s:","right"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", g->right [ncam]); printf("\n");
		printf("%22s:","height"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", g->height [ncam]); printf("\n");
		printf("%22s:","roll"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", g->roll [ncam]); printf("\n");
		for (int ncam = 0; ncam < num_cams; ncam++) {
		printf("%19s%2d]: %f, %f \n", "pXY0[",ncam, g->pXY0[ncam][0], g->pXY0[ncam][1]);
		}

		printf("%22s: %f\n","common_right", g->common_right);
		printf("%22s: %f\n","common_forward", g->common_forward);
		printf("%22s: %f\n","common_height", g->common_height);
		printf("%22s: %f\n","common_roll", g->common_roll);

		printf("%22s: x=%f, y=%f\n","rXY[0]", g->rXY[0][0], g->rXY[0][1]);
		printf("%22s: x=%f, y=%f\n","rXY[1]", g->rXY[1][0], g->rXY[1][1]);
		printf("%22s: x=%f, y=%f\n","rXY[2]", g->rXY[2][0], g->rXY[2][1]);
		printf("%22s: x=%f, y=%f\n","rXY[3]", g->rXY[3][0], g->rXY[3][1]);

		// printf("%22s: x=%f, y=%f\n","rXY[0]", g->rXY[0][0], g->rXY[0][1]);
		// printf("%22s: x=%f, y=%f\n","rXY[1]", g->rXY[1][0], g->rXY[1][1]);
		// printf("%22s: x=%f, y=%f\n","rXY[2]", g->rXY[2][0], g->rXY[2][1]);
		// printf("%22s: x=%f, y=%f\n","rXY[3]", g->rXY[3][0], g->rXY[3][1]);
		for (int ncam = 0; ncam < num_cams; ncam++) {
		printf("%19s%2d]: %f, %f \n", "rXY[", ncam, g->rXY[ncam][0], g->rXY[ncam][1]);
		}
		printf("%22s: %f\n","cameraRadius", g->cameraRadius);
		printf("%22s: %f\n","disparityRadius", g->disparityRadius);
		printf("%22s: %f, %f, %f, %f \n","woi_tops", g->woi_tops[0], g->woi_tops[1], g->woi_tops[2], g->woi_tops[3]);

		// printf("%22s: %f, %f, %f, %f \n","woi_tops", g->woi_tops[0], g->woi_tops[1], g->woi_tops[2], g->woi_tops[3]);
		printf("%22s:","woi_tops"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", g->woi_tops[ncam]); printf("\n");

		#endif //ifndef JCUDA
		}

		__device__ void printExtrinsicCorrection(corr_vector * cv)
		__device__ void printExtrinsicCorrection(corr_vector * cv, int num_cams)
		{
		#ifndef JCUDA
		printf("\nExtrinsic Correction Vector\n---------------------------\n");
		printf("%22s: %f, %f, %f\n", "tilt", cv->tilt[0], cv->tilt[1], cv->tilt[2]);
		printf("%22s: %f, %f, %f\n", "azimuth", cv->azimuth[0], cv->azimuth[1], cv->azimuth[2]);
		printf("%22s: %f, %f, %f, %f\n", "roll", cv->roll[0], cv->roll[1], cv->roll[2], cv->roll[3]);
		printf("%22s: %f, %f, %f\n", "zoom", cv->zoom[0], cv->zoom[1], cv->zoom[2]);
		// printf("%22s: %f, %f, %f\n", "tilt", cv->tilt[0], cv->tilt[1], cv->tilt[2]);
		// printf("%22s: %f, %f, %f\n", "azimuth", cv->azimuth[0], cv->azimuth[1], cv->azimuth[2]);
		// printf("%22s: %f, %f, %f, %f\n", "roll", cv->roll[0], cv->roll[1], cv->roll[2], cv->roll[3]);
		// printf("%22s: %f, %f, %f\n", "zoom", cv->zoom[0], cv->zoom[1], cv->zoom[2]);

		printf("%22s:","tilt"); for (int ncam = 0; ncam < (num_cams-1); ncam++) printf(" %f,", cv->tilt[ncam]); printf("\n");
		printf("%22s:","azimuth"); for (int ncam = 0; ncam < (num_cams-1); ncam++) printf(" %f,", cv->azimuth[ncam]); printf("\n");
		printf("%22s:","roll"); for (int ncam = 0; ncam < num_cams; ncam++) printf(" %f,", cv->roll[ncam]); printf("\n");
		printf("%22s:","zoom"); for (int ncam = 0; ncam < (num_cams-1); ncam++) printf(" %f,", cv->zoom[ncam]); printf("\n");



		printf("%22s: %f(t), %f(a), %f(r)\n", "imu_rot", cv->imu_rot[0], cv->imu_rot[1], cv->imu_rot[2]);
		printf("%22s: %f(x), %f(y), %f(z)\n", "imu_move", cv->imu_move[0], cv->imu_move[1], cv->imu_move[2]);

src/test_tp.cu

+131 −3

Original line number	Diff line number	Diff line
		@@ -156,6 +156,8 @@ int get_file_size(std::string filename) // path to file
		int readFloatsFromFile(float * data, // allocated array
		const char * 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
		@@ -254,8 +256,128 @@ int main(int argc, char **argv)
		//initialize CUDA
		findCudaDevice(argc, (const char **)argv);

		// CLT testing
		#if TEST_LWIR
		const char* kernel_file[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15_transposed.kernel"};

		const char* kernel_offs_file[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14_transposed.kernel_offsets",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15_transposed.kernel_offsets"};

		const char* image_files[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14.bayer",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15.bayer"};

		const char* ports_offs_xy_file[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14.portsxy",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15.portsxy"};

		//#ifndef DBG_TILE
		/*
		const char* ports_clt_file[] = { // never referenced
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14.clt",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15.clt"};
		*/
		const char* result_rbg_file[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn0.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn1.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn2.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn3.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn4.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn5.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn6.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn7.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn8.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn9.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn10.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn11.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn12.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn13.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn14.rbg",
		"/home/eyesis/git/tile_processor_gpu/clt/aux_chn15.rbg"};
		//#endif
		const char* result_corr_file = "/home/eyesis/git/tile_processor_gpu/clt/aux_corr.corr";
		const char* result_corr_quad_file = "/home/eyesis/git/tile_processor_gpu/clt/aux_corr-quad.corr";
		/// const char* result_corr_cross_file = "/home/eyesis/git/tile_processor_gpu/clt/aux_corr-cross.corr";
		const char* result_textures_file = "/home/eyesis/git/tile_processor_gpu/clt/texture_aux.rgba";
		const char* result_textures_rgba_file = "/home/eyesis/git/tile_processor_gpu/clt/texture_rgba_aux.rgba";

		const char* rByRDist_file = "/home/eyesis/git/tile_processor_gpu/clt/aux.rbyrdist";
		const char* correction_vector_file = "/home/eyesis/git/tile_processor_gpu/clt/aux.correction_vector";
		const char* geometry_correction_file = "/home/eyesis/git/tile_processor_gpu/clt/aux.geometry_correction";

		#else
		const char* kernel_file[] = {
		"/home/eyesis/git/tile_processor_gpu/clt/main_chn0_transposed.kernel",
		"/home/eyesis/git/tile_processor_gpu/clt/main_chn1_transposed.kernel",
		@@ -303,6 +425,10 @@ int main(int argc, char **argv)
		const char* rByRDist_file = "/home/eyesis/git/tile_processor_gpu/clt/main.rbyrdist";
		const char* correction_vector_file = "/home/eyesis/git/tile_processor_gpu/clt/main.correction_vector";
		const char* geometry_correction_file = "/home/eyesis/git/tile_processor_gpu/clt/main.geometry_correction";
		#endif



		int sel_pairs[4];

		#if TEST_LWIR
		@@ -490,7 +616,8 @@ int main(int argc, char **argv)
		gpu_corr_images_h[ncam] = alloc_image_gpu(
		&dstride_rslt, // size_t* dstride, // in bytes!!
		IMG_WIDTH + DTT_SIZE, // int width,
		3*(IMG_HEIGHT + DTT_SIZE)); // int height);
		// 3*(IMG_HEIGHT + DTT_SIZE)); // int height);
		num_colors*(IMG_HEIGHT + DTT_SIZE)); // int height);
		}
		// allocates one correlation kernel per line (15x15 floats), number of rows - number of tiles * number of pairs
		gpu_corrs = alloc_image_gpu(
		@@ -1083,7 +1210,8 @@ int main(int argc, char **argv)
		gpu_corr_images_h[ncam],
		dstride_rslt,
		(IMG_WIDTH + DTT_SIZE) * sizeof(float),
		3* (IMG_HEIGHT + DTT_SIZE),
		// 3* (IMG_HEIGHT + DTT_SIZE),
		num_colors* (IMG_HEIGHT + DTT_SIZE),
		cudaMemcpyDeviceToHost));

		///#ifndef DBG_TILE

src/tp_defines.h

+12 −7

Original line number	Diff line number	Diff line
		@@ -41,7 +41,7 @@
		#ifndef JCUDA
		#include <stdio.h>
		#define THREADSX (DTT_SIZE)
		#define TEST_LWIR 0
		#define TEST_LWIR 1
		#define NUM_CAMS 16 // now maximal number of cameras
		//#define NUM_PAIRS 6
		#define NUM_COLORS 1 //3
		@@ -105,12 +105,16 @@

		//#define DBG_TILE_X 40
		//#define DBG_TILE_Y 80
		#if TEST_LWIR
		#define DBG_TILE_X 32 // 162 // 151 // 161 // 49
		#define DBG_TILE_Y 36 // 88 // 121 // 69 // 111 // 66
		#define DBG_TILE (DBG_TILE_Y * 80 + DBG_TILE_X)
		#else
		#define DBG_TILE_X 32 // 162 // 151 // 161 // 49
		#define DBG_TILE_Y 88 // 121 // 69 // 111 // 66

		#define DBG_TILE (DBG_TILE_Y * 324 + DBG_TILE_X)
		#endif
		#undef DBG_MARK_DBG_TILE

		//#undef DBG_TILE

		//#undef HAS_PRINTF
		@@ -137,6 +141,7 @@
		//#define DEBUG20 1

		#if (DBG_TILE_X >= 0) && (DBG_TILE_Y >= 0)
		#define DEBUG20 1 // Geometry Correction
		#define DEBUG21 1 // Geometry Correction
		//#define DEBUG210 1

Admin message