src/TileProcessor.h

+/**
+ **
+ ** TileProcessor.h
+ **
+ ** Copyright (C) 2020 Elphel, Inc.
+ **
+ ** -----------------------------------------------------------------------------**
+ **
+ **  TileProcessor.h is free software: you can redistribute it and/or modify
+ **  it under the terms of the GNU General Public License as published by
+ **  the Free Software Foundation, either version 3 of the License, or
+ **  (at your option) any later version.
+ **
+ **  This program is distributed in the hope that it will be useful,
+ **  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ **  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ **  GNU General Public License for more details.
+ **
+ **  You should have received a copy of the GNU General Public License
+ **  along with this program.  If not, see <http://www.gnu.org/licenses/>.
+ **
+ **  Additional permission under GNU GPL version 3 section 7
+ **
+ **  If you modify this Program, or any covered work, by linking or
+ **  combining it with NVIDIA Corporation's CUDA libraries from the
+ **  NVIDIA CUDA Toolkit (or a modified version of those libraries),
+ **  containing parts covered by the terms of NVIDIA CUDA Toolkit
+ **  EULA, the licensors of this Program grant you additional
+ **  permission to convey the resulting work.
+ ** -----------------------------------------------------------------------------**
+ */
+
+/**
+ **************************************************************************
+ * \file TileProcessor.h
+ * \brief header file for  the Tile Processor for frequency domain
+
+ */
+#pragma once
+#ifndef NUM_CAMS
+#include "tp_defines.h"
+#endif
+
+
+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
+		float           ** gpu_kernel_offsets, // [NUM_CAMS],
+		float           ** gpu_kernels,        // [NUM_CAMS],
+		float           ** gpu_images,         // [NUM_CAMS],
+		struct tp_task   * gpu_tasks,
+		float           ** gpu_clt,            // [NUM_CAMS][TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
+		size_t             dstride,            // in floats (pixels)
+		int                num_tiles,          // number of tiles in task
+		int                lpf_mask,           // apply lpf to colors : bit 0 - red, bit 1 - blue, bit2 - green. Now - always 0 !
+		int                woi_width,
+		int                woi_height,
+		int                kernels_hor,
+		int                kernels_vert,
+		int *              gpu_active_tiles,      // pointer to the calculated number of non-zero tiles
+		int *              pnum_active_tiles);  //  indices to gpu_tasks
+
+extern "C" __global__ void correlate2D(
+		float          ** gpu_clt,            // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
+		int               colors,             // number of colors (3/1)
+		float             scale0,             // scale for R
+		float             scale1,             // scale for B
+		float             scale2,             // scale for G
+		float             fat_zero,           // here - absolute
+		struct tp_task  * gpu_tasks,          // array of per-tile tasks (now bits 4..9 - correlation pairs)
+		int               num_tiles,          // number of tiles in task
+		int             * gpu_corr_indices,   // packed tile+pair
+		int             * pnum_corr_tiles,    // pointer to a number of correlation tiles to process
+		const size_t      corr_stride,        // in floats
+		int               corr_radius,        // radius of the output correlation (7 for 15x15)
+		float           * gpu_corrs);          // correlation output data
+
+
+extern "C" __global__ void textures_nonoverlap(
+		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             * pnum_texture_tiles,  // returns total number of elements in gpu_texture_indices array
+		float          ** gpu_clt,            // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
+		// 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             params[5],
+//		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)
+		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_stride,     // in floats (now 256*4 = 1024)  // may be 0 if not needed
+		float           * gpu_texture_tiles,  // (number of colors +1 + ?)*16*16 rgba texture tiles    // may be 0 if not needed
+		float           * gpu_diff_rgb_combo); // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS] // may be 0 if not needed
+
+extern "C"
+__global__ void imclt_rbg_all(
+		float           ** gpu_clt,            // [NUM_CAMS][TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
+		float           ** gpu_corr_images,    // [NUM_CAMS][WIDTH, 3 * HEIGHT]
+		int                apply_lpf,
+		int                colors,
+		int                woi_twidth,
+		int                woi_theight,
+		const size_t       dstride);            // in floats (pixels)
+
+extern "C" __global__ void imclt_rbg(
+		float           * gpu_clt,            // [TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
+		float           * gpu_rbg,            // WIDTH, 3 * HEIGHT
+		int               apply_lpf,
+		int               mono,               // defines lpf filter
+		int               color,              // defines location of clt data
+		int               v_offset,
+		int               h_offset,
+		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,
+		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]
+		// 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             params[5],          // mitigating CUDA_ERROR_INVALID_PTX
+		/*
+			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)
+		 */
+		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
+		float           * gpu_texture_tiles);  // (number of colors +1 + ?)*16*16 rgba texture tiles
+

src/dtt8x8.h

+/**
+ **
+ ** dtt8x8.h
+ **
+ ** Copyright (C) 2018 Elphel, Inc.
+ **
+ ** -----------------------------------------------------------------------------**
+ **
+ **  dtt8x8.cuh is free software: you can redistribute it and/or modify
+ **  it under the terms of the GNU General Public License as published by
+ **  the Free Software Foundation, either version 3 of the License, or
+ **  (at your option) any later version.
+ **
+ **  This program is distributed in the hope that it will be useful,
+ **  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ **  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ **  GNU General Public License for more details.
+ **
+ **  You should have received a copy of the GNU General Public License
+ **  along with this program.  If not, see <http://www.gnu.org/licenses/>.
+ **
+ **  Additional permission under GNU GPL version 3 section 7
+ **
+ **  If you modify this Program, or any covered work, by linking or
+ **  combining it with NVIDIA Corporation's CUDA libraries from the
+ **  NVIDIA CUDA Toolkit (or a modified version of those libraries),
+ **  containing parts covered by the terms of NVIDIA CUDA Toolkit
+ **  EULA, the licensors of this Program grant you additional
+ **  permission to convey the resulting work.
+ ** -----------------------------------------------------------------------------**
+ */
+
+/**
+**************************************************************************
+* \file dtt8x8.h
+* \brief DCT-II, DST-II, DCT-IV and DST-IV for Complex Lapped Transform of 16x16 (stride 8)
+*        in GPU
+* This file contains building blocks for the 16x16 stride 8 COmplex Lapped Transform (CLT)
+* implementation. DTT-IV are used for forward and inverse 2D CLT, DTT-II - to convert correlation
+* results from the frequency to pixel domain. DTT-III (inverse of DTT-II) is not implemented
+* here it is used to convert convolution kernels and LPF to the frequency domain - done in
+* software.
+*
+* This file is cpompatible with both runtime and driver API, runtime is used for development
+* with Nvidia Nsight, driver API when calling these kernels from Java
+*/
+#ifndef JCUDA
+#define DTT_SIZE_LOG2                 3
+#endif
+
+#pragma once
+#define DTT_SIZE                     (1 << DTT_SIZE_LOG2)
+#define DTT_SIZE1        (DTT_SIZE + 1)
+#define DTT_SIZE2        (2 * DTT_SIZE)
+#define DTT_SIZE21       (DTT_SIZE2 + 1)
+#define DTT_SIZE4        (4 * DTT_SIZE)
+#define DTT_SIZE2M1      (DTT_SIZE2 - 1)
+#define BAYER_RED   0
+#define BAYER_BLUE  1
+#define BAYER_GREEN 2
+// assuming GR/BG as now
+#define BAYER_RED_ROW 0
+#define BAYER_RED_COL 1
+
+#define DTTTEST_BLOCK_WIDTH          32
+#define DTTTEST_BLOCK_HEIGHT         16
+#define DTTTEST_BLK_STRIDE     (DTTTEST_BLOCK_WIDTH+1)
+
+//extern __constant__ float idct_signs[4][4][4];
+//extern __constant__ int imclt_indx9[16];
+//extern __constant__ float HWINDOW2[];
+
+
+// kernels (not used so far)
+#if 0
+extern "C" __global__ void GPU_DTT24_DRV(float *dst, float *src, int src_stride, int dtt_mode);
+#endif// #if 0
+
+//=========================== 2D functions ===============
+extern __device__ void corrUnfoldTile(
+		int corr_radius,
+		float* qdata0, //    [4][DTT_SIZE][DTT_SIZE1], // 4 quadrants of the clt data, rows extended to optimize shared ports
+		float* rslt);  //   [DTT_SIZE2M1][DTT_SIZE2M1]) // 15x15
+
+extern __device__ void dttii_2d(
+		float * clt_corr); // shared memory, [4][DTT_SIZE1][DTT_SIZE]
+
+extern __device__ void dttiv_color_2d(
+		float * clt_tile,
+		int color);
+extern __device__ void imclt(
+		float * clt_tile,   //        [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports [4][8][9]
+		float * mclt_tile );
+
+extern __device__ void imclt8threads(
+		int     do_acc,     // 1 - add to previous value, 0 - overwrite
+		float * clt_tile,   //        [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports [4][8][9]
+		float * mclt_tile,  //           [2* DTT_SIZE][DTT_SIZE1+ DTT_SIZE], // +1 to alternate column ports[16][17]
+		int     debug);

src/geometry_correction.h

+/**
+ **
+ ** geometry_correction.h
+ **
+ ** Copyright (C) 2020 Elphel, Inc.
+ **
+ ** -----------------------------------------------------------------------------**
+ **
+ **  geometry_correction.h is free software: you can redistribute it and/or modify
+ **  it under the terms of the GNU General Public License as published by
+ **  the Free Software Foundation, either version 3 of the License, or
+ **  (at your option) any later version.
+ **
+ **  This program is distributed in the hope that it will be useful,
+ **  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ **  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ **  GNU General Public License for more details.
+ **
+ **  You should have received a copy of the GNU General Public License
+ **  along with this program.  If not, see <http://www.gnu.org/licenses/>.
+ **
+ **  Additional permission under GNU GPL version 3 section 7
+ **
+ **  If you modify this Program, or any covered work, by linking or
+ **  combining it with NVIDIA Corporation's CUDA libraries from the
+ **  NVIDIA CUDA Toolkit (or a modified version of those libraries),
+ **  containing parts covered by the terms of NVIDIA CUDA Toolkit
+ **  EULA, the licensors of this Program grant you additional
+ **  permission to convey the resulting work.
+ ** -----------------------------------------------------------------------------**
+ */
+
+/**
+**************************************************************************
+* \file geometry_correction.h
+* \brief header file for geometry correction - per-tile/per camera calculation of the tile offset
+
+*/
+#pragma once
+#ifndef NUM_CAMS
+#include "tp_defines.h"
+#endif
+
+#define NVRTC_BUG 1
+#ifndef M_PI
+#define M_PI  3.14159265358979323846 /* pi */
+#endif
+#ifndef offsetof
+#define offsetof(st, m) \
+    ((size_t)&(((st *)0)->m))
+//#define offsetof(TYPE, MEMBER) __builtin_offsetof (TYPE, MEMBER)
+#endif
+
+
+#define SCENE_UNITS_SCALE  0.001 // meters from mm
+#define MIN_DISPARITY      0.01  // minimal disparity to try to convert to world coordinates
+struct tp_task {
+	int   task;
+	union {
+		int      txy;
+		unsigned short sxy[2];
+	};
+	float xy[NUM_CAMS][2];
+	float target_disparity;
+	float disp_dist[NUM_CAMS][4]; // calculated with getPortsCoordinates()
+};
+
+struct corr_vector{
+	float tilt    [NUM_CAMS-1]; // 0..2
+	float azimuth [NUM_CAMS-1]; // 3..5
+	float roll    [NUM_CAMS];   // 6..9
+	float zoom    [NUM_CAMS-1]; // 10..12
+	// for ERS correction:
+	float imu_rot [3]; // d_tilt/dt (rad/s), d_az/dt, d_roll/dt 13..15
+	float imu_move[3]; // dx/dt, dy/dt, dz/dt 16..19
+};
+#ifdef NVRTC_BUG
+struct trot_deriv{
+	float rots    [NUM_CAMS][3][3];
+	float d_daz   [NUM_CAMS][3][3];
+	float d_tilt  [NUM_CAMS][3][3];
+	float d_roll  [NUM_CAMS][3][3];
+	float d_zoom  [NUM_CAMS][3][3];
+};
+#else
+union trot_deriv{
+	struct {
+		float rots    [NUM_CAMS][3][3];
+		float d_daz   [NUM_CAMS][3][3];
+		float d_tilt  [NUM_CAMS][3][3];
+		float d_roll  [NUM_CAMS][3][3];
+		float d_zoom  [NUM_CAMS][3][3];
+	};
+	float matrices [5][NUM_CAMS][3][3];
+};
+#endif
+
+struct gc {
+	float pixelCorrectionWidth; //  =2592;   // virtual camera center is at (pixelCorrectionWidth/2, pixelCorrectionHeight/2)
+	float pixelCorrectionHeight; // =1936;
+	float line_time;        // duration of one scan line readout (for ERS)
+	float focalLength;      // =FOCAL_LENGTH;
+	float pixelSize;        // =  PIXEL_SIZE; //um
+	float distortionRadius; // =  DISTORTION_RADIUS; // mm - half width of the sensor
+#ifndef	NVRTC_BUG
+	union {
+		struct {
+#endif
+			float distortionC;      // r^2
+			float distortionB;      // r^3
+			float distortionA;      // r^4 (normalized to focal length or to sensor half width?)
+			float distortionA5;     //r^5 (normalized to focal length or to sensor half width?)
+			float distortionA6;     //r^6 (normalized to focal length or to sensor half width?)
+			float distortionA7;     //r^7 (normalized to focal length or to sensor half width?)
+			float distortionA8;     //r^8 (normalized to focal length or to sensor half width?)
+#ifndef	NVRTC_BUG
+		};
+		float rad_coeff [7];
+	};
+#endif
+	// parameters, common for all sensors
+	float    elevation;     // degrees, up - positive;
+	float    heading;       // degrees, CW (from top) - positive
+
+	float forward    [NUM_CAMS];
+	float right      [NUM_CAMS];
+	float height     [NUM_CAMS];
+	float roll       [NUM_CAMS];    // degrees, CW (to target) - positive
+	float pXY0       [NUM_CAMS][2];
+	float common_right;    // mm right, camera center
+	float common_forward;  // mm forward (to target), camera center
+	float common_height;   // mm up, camera center
+	float common_roll;     // degrees CW (to target) camera as a whole
+//	float [][] XYZ_he;     // all cameras coordinates transformed to eliminate heading and elevation (rolls preserved)
+//	float [][] XYZ_her = null; // XYZ of the lenses in a corrected CCS (adjusted for to elevation, heading,  common_roll)
+	float rXY        [NUM_CAMS][2]; // XY pairs of the in a normal plane, relative to disparityRadius
+//	float [][] rXY_ideal = {{-0.5, -0.5}, {0.5,-0.5}, {-0.5, 0.5}, {0.5,0.5}};
+// only used for the multi-quad systems
+	float cameraRadius; // =0; // average distance from the "mass center" of the sensors to the sensors
+	float disparityRadius; // =150.0; // distance between cameras to normalize disparity units to. sqrt(2)*disparityRadius for quad
+};
+#define RAD_COEFF_LEN 7
+extern "C" __global__ void get_tiles_offsets(
+		struct tp_task     * gpu_tasks,
+		int                  num_tiles,          // number of tiles in task
+		struct gc          * gpu_geometry_correction,
+		struct corr_vector * gpu_correction_vector,
+		float *              gpu_rByRDist, // length should match RBYRDIST_LEN
+		trot_deriv   * gpu_rot_deriv);
+
+// uses NUM_CAMS blocks, (3,3,3) threads
+extern "C" __global__ void calc_rot_deriv(
+		struct corr_vector * gpu_correction_vector,
+		trot_deriv   * gpu_rot_deriv);
+
+#define CALC_REVERSE_TABLE_BLOCK_THREADS (NUM_CAMS * 3 * 3 * 3) // fixed blockDim
+// Use same blocks/threads as with calc_rot_deriv() - NUM_CAMS blocks, (3,3,3) threads
+extern "C" __global__ void calcReverseDistortionTable(
+		struct gc * geometry_correction,
+		float * rByRDist);
+
+

src/tp_defines.h

+/**
+ **
+ ** tp_defines.h
+ **
+ ** Copyright (C) 2020 Elphel, Inc.
+ **
+ ** -----------------------------------------------------------------------------**
+ **
+ **  tp_defines.h is free software: you can redistribute it and/or modify
+ **  it under the terms of the GNU General Public License as published by
+ **  the Free Software Foundation, either version 3 of the License, or
+ **  (at your option) any later version.
+ **
+ **  This program is distributed in the hope that it will be useful,
+ **  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ **  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ **  GNU General Public License for more details.
+ **
+ **  You should have received a copy of the GNU General Public License
+ **  along with this program.  If not, see <http://www.gnu.org/licenses/>.
+ **
+ **  Additional permission under GNU GPL version 3 section 7
+ **
+ **  If you modify this Program, or any covered work, by linking or
+ **  combining it with NVIDIA Corporation's CUDA libraries from the
+ **  NVIDIA CUDA Toolkit (or a modified version of those libraries),
+ **  containing parts covered by the terms of NVIDIA CUDA Toolkit
+ **  EULA, the licensors of this Program grant you additional
+ **  permission to convey the resulting work.
+ ** -----------------------------------------------------------------------------**
+ */
+
+/**
+**************************************************************************
+* \file tp_defines.h
+* \brief Defines for running in C++ environment, replaced when called from Java
+
+*/
+// Avoiding includes in jcuda, all source files will be merged
+#pragma once
+#ifndef JCUDA
+#include <stdio.h>
+#define THREADSX         (DTT_SIZE)
+#define NUM_CAMS                  4
+#define NUM_PAIRS                 6
+#define NUM_COLORS                3
+#define IMG_WIDTH              2592
+#define IMG_HEIGHT             1936
+#define KERNELS_HOR             164
+#define KERNELS_VERT            123
+#define KERNELS_LSTEP             4
+#define THREADS_PER_TILE          8
+#define TILES_PER_BLOCK           4
+#define CORR_THREADS_PER_TILE     8
+#define CORR_TILES_PER_BLOCK      4
+#define TEXTURE_THREADS_PER_TILE  8
+#define TEXTURE_TILES_PER_BLOCK   1
+#define IMCLT_THREADS_PER_TILE   16
+#define IMCLT_TILES_PER_BLOCK     4
+#define CORR_NTILE_SHIFT          8 // higher bits - number of a pair, other bits tile number
+#define CORR_PAIRS_MASK        0x3f// lower bits used to address correlation pair for the selected tile
+#define CORR_TEXTURE_BIT          7 // bit 7 used to request texture for the tile
+#define TASK_CORR_BITS            4
+#define TASK_TEXTURE_N_BIT        0 // Texture with North neighbor
+#define TASK_TEXTURE_E_BIT        1 // Texture with East  neighbor
+#define TASK_TEXTURE_S_BIT        2 // Texture with South neighbor
+#define TASK_TEXTURE_W_BIT        3 // Texture with West  neighbor
+#define TASK_TEXTURE_BIT          3 // bit to request texture calculation int task field of struct tp_task
+#define LIST_TEXTURE_BIT          7 // bit to request texture calculation
+#define CORR_OUT_RAD              4
+#define FAT_ZERO_WEIGHT           0.0001 // add to port weights to avoid nan
+
+#define THREADS_DYNAMIC_BITS      5 // treads in block for CDP creation of the texture list
+
+#define DBG_DISPARITY            0.0 // 56.0 // disparity for which to calculate offsets (not needed in Java)
+#define RBYRDIST_LEN           5001   // for doubles 10001 - floats   // length of rByRDist to allocate shared memory
+#define RBYRDIST_STEP             0.0004 // for doubles, 0.0002 - floats // to fit into GPU shared memory (was 0.001);
+#define TILES_PER_BLOCK_GEOM     (32/NUM_CAMS)   // each tile has NUM_CAMS threads
+
+#define DEBUG_OOB1 1
+
+// Use CORR_OUT_RAD for the correlation output
+
+//#define DBG_TILE_X     40
+//#define DBG_TILE_Y     80
+#define DBG_TILE_X     162 // 151 // 161 // 49
+#define DBG_TILE_Y     121 // 69  // 111 // 66
+
+#define DBG_TILE    (DBG_TILE_Y * 324 + DBG_TILE_X)
+#undef DBG_MARK_DBG_TILE
+
+
+
+//#undef HAS_PRINTF
+#define HAS_PRINTF
+//7
+//#define DEBUG1 1
+//#define DEBUG2 1
+//#define DEBUG3 1
+//#define DEBUG4 1
+//#define DEBUG5 1
+//#define DEBUG6 1
+/*
+#define DEBUG7 1
+#define DEBUG8 1
+#define DEBUG9 1
+*/
+//textures
+//#define DEBUG10 1
+//#define DEBUG11 1
+//#define DEBUG12 1
+//#define USE_textures_gen
+//#define DEBUG_OOB1 1
+// geom
+//#define DEBUG20 1
+
+#if (DBG_TILE_X >= 0) && (DBG_TILE_Y >= 0)
+//#define DEBUG21 1 // Geometry Correction
+
+//#define DEBUG22 1
+//#define DEBUG23 1
+
+#endif //#if (DBG_TILE_X >= 0) && (DBG_TILE_Y >= 0)
+
+#endif //#ifndef JCUDA
+