Commit 907cda8d authored by Andrey Filippov's avatar Andrey Filippov
Browse files

Using Marco's help to make jcuda work with nvrtc with cdp

parent 0d087c0b

src/main/java/com/elphel/imagej/cameras/CLTParameters.java

+1 −1
Original line number Original line Diff line number Diff line
@@ -333,7 +333,7 @@ public class CLTParameters { 
	public double     infinityDistance  = 10000;  // Distance to generate backdrop (0 - use regular backdrop)

	public double     infinityDistance  = 10000;  // Distance to generate backdrop (0 - use regular backdrop)

	public int        min_bgnd_tiles    = 10;     // Minimal number of background tiles to generate background

	public int        min_bgnd_tiles    = 10;     // Minimal number of background tiles to generate background

	public boolean    shUseFlaps        = true;  // Split into shells with flaps

	public boolean    shUseFlaps        = true;  // Split into shells with flaps

	public boolean    shAggrFade        = true;  // Aggressive fade alpha (whole boundary)

	public boolean    shAggrFade        = false; // true;  // Aggressive fade alpha (whole boundary)

	public int        shMinArea         = 1;     // Minimal shell area (not counting flaps

	public int        shMinArea         = 1;     // Minimal shell area (not counting flaps

	public double     shMinStrength    = 0.2;   // Minimal value of the shell maximum strength

	public double     shMinStrength    = 0.2;   // Minimal value of the shell maximum strength

src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java

+99 −59
Original line number Original line Diff line number Diff line
@@ -26,19 +26,29 @@ package com.elphel.imagej.gpu; 
**

**

*/

*/





import static jcuda.driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR;

import static jcuda.driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR;

// Uses code by Marco Hutter - http://www.jcuda.org

// Uses code by Marco Hutter - http://www.jcuda.org

import static jcuda.driver.CUjitInputType.CU_JIT_INPUT_LIBRARY;

import static jcuda.driver.CUjitInputType.CU_JIT_INPUT_PTX;

import static jcuda.driver.JCudaDriver.cuCtxCreate;

import static jcuda.driver.JCudaDriver.cuCtxCreate;

import static jcuda.driver.JCudaDriver.cuCtxSynchronize;

import static jcuda.driver.JCudaDriver.cuCtxSynchronize;

import static jcuda.driver.JCudaDriver.cuDeviceGet;

import static jcuda.driver.JCudaDriver.cuDeviceGet;

import static jcuda.driver.JCudaDriver.cuDeviceGetAttribute;

import static jcuda.driver.JCudaDriver.cuInit;

import static jcuda.driver.JCudaDriver.cuInit;

import static jcuda.driver.JCudaDriver.cuLaunchKernel;

import static jcuda.driver.JCudaDriver.cuLaunchKernel;

import static jcuda.driver.JCudaDriver.cuLinkAddData;

import static jcuda.driver.JCudaDriver.cuLinkAddFile;

import static jcuda.driver.JCudaDriver.cuLinkComplete;

import static jcuda.driver.JCudaDriver.cuLinkCreate;

import static jcuda.driver.JCudaDriver.cuLinkDestroy;

import static jcuda.driver.JCudaDriver.cuMemAlloc;

import static jcuda.driver.JCudaDriver.cuMemAlloc;

import static jcuda.driver.JCudaDriver.cuMemAllocPitch;

import static jcuda.driver.JCudaDriver.cuMemAllocPitch;

import static jcuda.driver.JCudaDriver.cuMemcpy2D;

import static jcuda.driver.JCudaDriver.cuMemcpy2D;

import static jcuda.driver.JCudaDriver.cuMemcpyHtoD;

import static jcuda.driver.JCudaDriver.cuMemcpyHtoD;

import static jcuda.driver.JCudaDriver.cuModuleGetFunction;

import static jcuda.driver.JCudaDriver.cuModuleGetFunction;

import static jcuda.driver.JCudaDriver.cuModuleGetGlobal;

import static jcuda.driver.JCudaDriver.cuModuleGetGlobal;

import static jcuda.driver.JCudaDriver.cuModuleLoadData;

import static jcuda.driver.JCudaDriver.cuModuleLoadDataEx;

import static jcuda.nvrtc.JNvrtc.nvrtcCompileProgram;

import static jcuda.nvrtc.JNvrtc.nvrtcCompileProgram;

import static jcuda.nvrtc.JNvrtc.nvrtcCreateProgram;

import static jcuda.nvrtc.JNvrtc.nvrtcCreateProgram;

import static jcuda.nvrtc.JNvrtc.nvrtcDestroyProgram;

import static jcuda.nvrtc.JNvrtc.nvrtcDestroyProgram;
@@ -66,14 +76,17 @@ import jcuda.driver.CUcontext; 
import jcuda.driver.CUdevice;

import jcuda.driver.CUdevice;

import jcuda.driver.CUdeviceptr;

import jcuda.driver.CUdeviceptr;

import jcuda.driver.CUfunction;

import jcuda.driver.CUfunction;

import jcuda.driver.CUlinkState;

import jcuda.driver.CUmemorytype;

import jcuda.driver.CUmemorytype;

import jcuda.driver.CUmodule;

import jcuda.driver.CUmodule;

import jcuda.driver.JCudaDriver;

import jcuda.driver.JCudaDriver;

import jcuda.driver.JITOptions;

import jcuda.nvrtc.JNvrtc;

import jcuda.nvrtc.JNvrtc;

import jcuda.nvrtc.nvrtcProgram;

import jcuda.nvrtc.nvrtcProgram;





public class GPUTileProcessor {

public class GPUTileProcessor {

	static String GPU_KERNEL_FILE = "dtt8x8.cuh";

	String LIBRARY_PATH = "/usr/local/cuda/targets/x86_64-linux/lib/libcudadevrt.a"; // linux

	static String GPU_RESOURCE_DIR =              "kernels";

	static String [] GPU_KERNEL_FILES = {"dtt8x8.cuh","TileProcessor.cuh"};

	static String [] GPU_KERNEL_FILES = {"dtt8x8.cuh","TileProcessor.cuh"};

	static String GPU_CONVERT_CORRECT_TILES_NAME = "convert_correct_tiles"; // name in C code

	static String GPU_CONVERT_CORRECT_TILES_NAME = "convert_correct_tiles"; // name in C code

	static String GPU_IMCLT_RBG_NAME =             "imclt_rbg"; // name in C code

	static String GPU_IMCLT_RBG_NAME =             "imclt_rbg"; // name in C code
@@ -112,6 +125,7 @@ public class GPUTileProcessor { 
	public static int CORR_OUT_RAD =              4;  // output radius of the correlations (implemented)

	public static int CORR_OUT_RAD =              4;  // output radius of the correlations (implemented)

	public static double FAT_ZERO_WEIGHT =        0.0001; // add to port weights to avoid nan

	public static double FAT_ZERO_WEIGHT =        0.0001; // add to port weights to avoid nan





	public static int THREADS_DYNAMIC_BITS =      5; // treads in block for CDP creation of the texture list





    int DTTTEST_BLOCK_WIDTH =        32; // may be read from the source code

    int DTTTEST_BLOCK_WIDTH =        32; // may be read from the source code

    int DTTTEST_BLOCK_HEIGHT =       16; // may be read from the source code

    int DTTTEST_BLOCK_HEIGHT =       16; // may be read from the source code
@@ -129,7 +143,6 @@ public class GPUTileProcessor { 
    private CUdeviceptr [] gpu_kernel_offsets_h = new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_kernel_offsets_h = new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_bayer_h =          new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_bayer_h =          new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_clt_h =            new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_clt_h =            new CUdeviceptr[NUM_CAMS];

//    private CUdeviceptr [] gpu_lpf_h =            new CUdeviceptr[NUM_COLORS];

    private CUdeviceptr [] gpu_corr_images_h=     new CUdeviceptr[NUM_CAMS];

    private CUdeviceptr [] gpu_corr_images_h=     new CUdeviceptr[NUM_CAMS];





    // GPU pointers to array of GPU pointers

    // GPU pointers to array of GPU pointers
@@ -143,9 +156,7 @@ public class GPUTileProcessor { 
    private CUdeviceptr gpu_corr_indices =        new CUdeviceptr(); //  allocate tilesX * tilesY * 6 * Sizeof.POINTER

    private CUdeviceptr gpu_corr_indices =        new CUdeviceptr(); //  allocate tilesX * tilesY * 6 * Sizeof.POINTER

    private CUdeviceptr gpu_texture_indices =     new CUdeviceptr(); //  allocate tilesX * tilesY * 6 * Sizeof.POINTER

    private CUdeviceptr gpu_texture_indices =     new CUdeviceptr(); //  allocate tilesX * tilesY * 6 * Sizeof.POINTER

    private CUdeviceptr gpu_port_offsets =        new CUdeviceptr(); //  allocate Quad * 2 * Sizeof.POINTER

    private CUdeviceptr gpu_port_offsets =        new CUdeviceptr(); //  allocate Quad * 2 * Sizeof.POINTER

    //    private

    CUmodule    module; // to access constants memory

    CUmodule    module; // to access constants memory

//    private CUdeviceptr gpu_lpf =            new CUdeviceptr();

    private int mclt_stride;

    private int mclt_stride;

    private int corr_stride;

    private int corr_stride;

    private int imclt_stride;

    private int imclt_stride;
@@ -153,7 +164,6 @@ public class GPUTileProcessor { 
    public int num_task_tiles;

    public int num_task_tiles;

    public int num_corr_tiles;

    public int num_corr_tiles;

    public int num_texture_tiles;

    public int num_texture_tiles;



    public class TpTask {

    public class TpTask {

    	public int   task; // [0](+1) - generate 4 images, [4..9]+16..+512 - correlation pairs, 2 - generate texture tiles

    	public int   task; // [0](+1) - generate 4 images, [4..9]+16..+512 - correlation pairs, 2 - generate texture tiles

    	public float target_disparity;

    	public float target_disparity;
@@ -267,6 +277,7 @@ public class GPUTileProcessor { 




    public GPUTileProcessor(String cuda_project_directory) throws IOException

    public GPUTileProcessor(String cuda_project_directory) throws IOException

    {

    {



    	// From code by Marco Hutter - http://www.jcuda.org

    	// From code by Marco Hutter - http://www.jcuda.org

        // Enable exceptions and omit all subsequent error checks

        // Enable exceptions and omit all subsequent error checks

        JCudaDriver.setExceptionsEnabled(true);

        JCudaDriver.setExceptionsEnabled(true);
@@ -274,13 +285,24 @@ public class GPUTileProcessor { 




        // Initialize the driver and create a context for the first device.

        // Initialize the driver and create a context for the first device.

        cuInit(0);

        cuInit(0);

        CUdevice device = new CUdevice();

        //2020 - making them global

        CUdevice

        device = new CUdevice();

        cuDeviceGet(device, 0);

        cuDeviceGet(device, 0);

        CUcontext context = new CUcontext();

        CUcontext

        context = new CUcontext();

        cuCtxCreate(context, 0, device);

        cuCtxCreate(context, 0, device);





        int majorArray[] = { 0 };

        int minorArray[] = { 0 };

        cuDeviceGetAttribute(majorArray,  CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device);

        cuDeviceGetAttribute(minorArray,  CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device);

        int major = majorArray[0];

        int minor = minorArray[0];

        int capability = major * 10 + minor;



        // Obtain the CUDA source code from the CUDA file

        // Obtain the CUDA source code from the CUDA file

        // Get absolute path to the file in resource foldder, then read it as a normal file.

        // Get absolute path to the file in resource folder, then read it as a normal file.

        // When using just Eclipse resources - it does not notice that the file

        // When using just Eclipse resources - it does not notice that the file

        // was edited (happens frequently during kernel development).

        // was edited (happens frequently during kernel development).

        ClassLoader classLoader = getClass().getClassLoader();

        ClassLoader classLoader = getClass().getClassLoader();
@@ -311,12 +333,14 @@ public class GPUTileProcessor { 
        				"#define TASK_TEXTURE_BIT " +         TASK_TEXTURE_BIT+"\n"+

        				"#define TASK_TEXTURE_BIT " +         TASK_TEXTURE_BIT+"\n"+

        				"#define LIST_TEXTURE_BIT " +         LIST_TEXTURE_BIT+"\n"+

        				"#define LIST_TEXTURE_BIT " +         LIST_TEXTURE_BIT+"\n"+

        				"#define CORR_OUT_RAD " +             CORR_OUT_RAD+"\n" +

        				"#define CORR_OUT_RAD " +             CORR_OUT_RAD+"\n" +

        				"#define FAT_ZERO_WEIGHT " +          FAT_ZERO_WEIGHT+"\n";

        				"#define FAT_ZERO_WEIGHT " +          FAT_ZERO_WEIGHT+"\n"+

        				"#define THREADS_DYNAMIC_BITS " +     THREADS_DYNAMIC_BITS+"\n";







        for (String src_file:GPU_KERNEL_FILES) {

        for (String src_file:GPU_KERNEL_FILES) {

        	File file = null;

        	File file = null;

        	if ((cuda_project_directory == null) || (cuda_project_directory == "")) {

        	if ((cuda_project_directory == null) || cuda_project_directory.isEmpty()) {

        		file = new File(classLoader.getResource(src_file).getFile());

        		file = new File(classLoader.getResource(GPU_RESOURCE_DIR+"/"+src_file).getFile());

        		System.out.println("Loading resource "+file);

        		System.out.println("Loading resource "+file);

        	} else {

        	} else {

        		File src_dir = new File(cuda_project_directory, "src");

        		File src_dir = new File(cuda_project_directory, "src");
@@ -336,11 +360,15 @@ public class GPUTileProcessor { 
        }

        }

        // Create the kernel functions (first - just test)

        // Create the kernel functions (first - just test)

        String [] func_names = {GPU_CONVERT_CORRECT_TILES_NAME, GPU_IMCLT_RBG_NAME, GPU_CORRELATE2D_NAME, GPU_TEXTURES_NAME};

        String [] func_names = {GPU_CONVERT_CORRECT_TILES_NAME, GPU_IMCLT_RBG_NAME, GPU_CORRELATE2D_NAME, GPU_TEXTURES_NAME};

        CUfunction[] functions = createFunctions(kernelSource, func_names);

        CUfunction[] functions = createFunctions(kernelSource,

        		                                 func_names,

        		                                 capability); // on my - 75



        this.GPU_CONVERT_CORRECT_TILES_kernel = functions[0];

        this.GPU_CONVERT_CORRECT_TILES_kernel = functions[0];

        this.GPU_IMCLT_RBG_kernel =             functions[1];

        this.GPU_IMCLT_RBG_kernel =             functions[1];

        this.GPU_CORRELATE2D_kernel =           functions[2];

        this.GPU_CORRELATE2D_kernel =           functions[2];

        this.GPU_TEXTURES_kernel=               functions[3];

        this.GPU_TEXTURES_kernel=               functions[3];



        System.out.println("GPU kernel functions initialized");

        System.out.println("GPU kernel functions initialized");

        System.out.println(GPU_CONVERT_CORRECT_TILES_kernel.toString());

        System.out.println(GPU_CONVERT_CORRECT_TILES_kernel.toString());

        System.out.println(GPU_IMCLT_RBG_kernel.toString());

        System.out.println(GPU_IMCLT_RBG_kernel.toString());
@@ -968,11 +996,12 @@ public class GPUTileProcessor { 
    		int     num_colors,

    		int     num_colors,

    		boolean keep_weights){

    		boolean keep_weights){





    	int texture_slices =     (num_colors + 1 + (keep_weights?(NUM_CAMS + num_colors + 1):0));

    	int texture_slices =     (num_colors + 1 + (keep_weights?(NUM_CAMS + num_colors + 1):0)); // number of texture slices

    	int texture_slice_size = (2 * DTT_SIZE)* (2 * DTT_SIZE);

    	int texture_slice_size = (2 * DTT_SIZE)* (2 * DTT_SIZE);        // number of (float) elements in a single slice of a tile

    	int texture_tile_size =  texture_slices * texture_slice_size;

    	int texture_tile_size =  texture_slices * texture_slice_size;   // number of (float) elements in a multi-slice tile

    	int texture_size =       texture_tile_size * num_texture_tiles;

    	int texture_size =       texture_tile_size * num_texture_tiles; // number of (float) elements in the whole texture

        float [] cpu_textures = new float [ num_texture_tiles * texture_size];

//        float [] cpu_textures = new float [ num_texture_tiles * texture_size];

        float [] cpu_textures = new float [texture_size];

        CUDA_MEMCPY2D copyD2H =   new CUDA_MEMCPY2D();

        CUDA_MEMCPY2D copyD2H =   new CUDA_MEMCPY2D();

        copyD2H.srcMemoryType =   CUmemorytype.CU_MEMORYTYPE_DEVICE;

        copyD2H.srcMemoryType =   CUmemorytype.CU_MEMORYTYPE_DEVICE;

        copyD2H.srcDevice =       gpu_textures;

        copyD2H.srcDevice =       gpu_textures;
@@ -1105,21 +1134,25 @@ public class GPUTileProcessor { 




//    private static CUfunction [] createFunctions(

//    private static CUfunction [] createFunctions(

    private CUfunction [] createFunctions(

    private CUfunction [] createFunctions(

            String sourceCode, String [] kernelNames) throws IOException

    		String     sourceCode,

    		String []  kernelNames,

    		int        capability

    		) throws IOException

    {

    {

    	CUfunction [] functions = new CUfunction [kernelNames.length];

    	CUfunction [] functions = new CUfunction [kernelNames.length];

    	boolean OK = false;

    	boolean OK = false;

    	// Use the NVRTC to create a program by compiling the source code

    	// Use the NVRTC to create a program by compiling the source code

    	nvrtcProgram program = new nvrtcProgram();

    	nvrtcProgram program = new nvrtcProgram();

            nvrtcCreateProgram(

    	nvrtcCreateProgram(	program, sourceCode, null, 0, null, null);

                program, sourceCode, null, 0, null, null);

    	String options[] = {"--gpu-architecture=compute_"+capability};



    	try {

    	try {

            	nvrtcCompileProgram(program, 0, null);

    		nvrtcCompileProgram(program, options.length, options);

    		OK = true;

    		OK = true;

    	} catch (Exception e) {

    	} catch (Exception e) {

    		System.out.println("nvrtcCompileProgram() FAILED");

    		System.out.println("nvrtcCompileProgram() FAILED");

    	}

    	}

            // Compilation log with errors/warnongs

    	// Compilation log with errors/warnings

    	String programLog[] = new String[1];

    	String programLog[] = new String[1];

    	nvrtcGetProgramLog(program, programLog);

    	nvrtcGetProgramLog(program, programLog);

    	String log = programLog[0].trim();

    	String log = programLog[0].trim();
@@ -1135,11 +1168,18 @@ public class GPUTileProcessor { 
    	String[] ptx = new String[1];

    	String[] ptx = new String[1];

    	nvrtcGetPTX(program, ptx);

    	nvrtcGetPTX(program, ptx);

    	nvrtcDestroyProgram(program);

    	nvrtcDestroyProgram(program);



    	byte[] ptxData = ptx[0].getBytes();

            // Create a CUDA module from the PTX code

    	JITOptions jitOptions = new JITOptions();

//            CUmodule

    	CUlinkState state = new CUlinkState();

    	cuLinkCreate(jitOptions, state);

    	cuLinkAddFile(state, CU_JIT_INPUT_LIBRARY, LIBRARY_PATH, jitOptions);

    	cuLinkAddData(state, CU_JIT_INPUT_PTX,     Pointer.to(ptxData), ptxData.length, "input.ptx", jitOptions);

    	long size[] = { 0 };

    	Pointer image = new Pointer();

    	cuLinkComplete(state, image, size);

    	module = new CUmodule();

    	module = new CUmodule();

            cuModuleLoadData(module, ptx[0]);

    	cuModuleLoadDataEx(module, image, 0, new int[0], Pointer.to(new int[0]));

    	cuLinkDestroy(state);





    	for (int i = 0; i < kernelNames.length; i++) {

    	for (int i = 0; i < kernelNames.length; i++) {

    		// Find the function in the source by name, get its pointer

    		// Find the function in the source by name, get its pointer

src/main/java/com/elphel/imagej/tileprocessor/ImageDtt.java

+21 −7
Original line number Original line Diff line number Diff line
@@ -3867,7 +3867,9 @@ public class ImageDtt { 










		double [][] port_weights = new double[ports][tile_len];




		double [][] port_weights = new double[ports][tile_len];








		double [][] color_avg =    new double[numcol][tile_len];




		double [][] color_avg =    new double[numcol][tile_len];








		double [][] rgba = new double[numcol + 1 + (keep_weights?(ports + numcol + 1):0)][];




//		double [][] rgba = new double[numcol + 1 + (keep_weights?(ports + numcol + 1):0)][];










// need to pass keep_weights to the caller










		double [][] rgba = new double[numcol + 1 + ports + (keep_weights?(numcol + 1):0)][];








		int rms_start = numcol + 1 + ports;




		int rms_start = numcol + 1 + ports;








		if (keep_weights){




		if (keep_weights){








			for (int ncol = 0; ncol <= numcol ; ncol++) if ((ncol == numcol) || (iclt_tile[0][ncol] != null)) {




			for (int ncol = 0; ncol <= numcol ; ncol++) if ((ncol == numcol) || (iclt_tile[0][ncol] != null)) {










@@ -4069,8 +4071,11 @@ public class ImageDtt {






			rgba[ncol] = color_avg[ncol];




			rgba[ncol] = color_avg[ncol];








		}




		}








		rgba[numcol] = alpha;




		rgba[numcol] = alpha;











//		if (keep_weights){








		for (int i = 0; i < ports; i++)  rgba[numcol + 1 + i] = port_weights[i];




			for (int i = 0; i < ports; i++) {










				rgba[numcol + 1 + i] = port_weights[i];










			}










//		}








		if (max_diff != null){




		if (max_diff != null){








			for (int ip = 0; ip < ports; ip++){




			for (int ip = 0; ip < ports; ip++){








				max_diff[ip] = 0;




				max_diff[ip] = 0;










@@ -5026,7 +5031,7 @@ public class ImageDtt {






			System.out.println("iclt_2d():sharp_alpha=   "+sharp_alpha);




			System.out.println("iclt_2d():sharp_alpha=   "+sharp_alpha);








		}




		}








		boolean has_weights = false;




		boolean has_weights = false;








		boolean set_has_weight = false;




		boolean set_has_weight = false; // not used








		for (int i = 0; (i < tilesY) && !set_has_weight; i++){




		for (int i = 0; (i < tilesY) && !set_has_weight; i++){








			for (int j = 0; (j < tilesX) && !set_has_weight; j++){




			for (int j = 0; (j < tilesX) && !set_has_weight; j++){








				if (texture_tiles[i][j] != null) {




				if (texture_tiles[i][j] != null) {










@@ -5064,14 +5069,14 @@ public class ImageDtt {






						int n_half = transform_size / 2;




						int n_half = transform_size / 2;








						int lastY = tilesY-1;




						int lastY = tilesY-1;








						int lastX = tilesX-1;




						int lastX = tilesX-1;








						int offset = n_half * (transform_size * tilesX) + n_half;




						int offset = n_half * (transform_size * tilesX) + n_half; // 4 pixels left and down (right/up when subtracted below)








						for (int nTile = ai.getAndIncrement(); nTile < tiles_list[nser.get()].length; nTile = ai.getAndIncrement()) {




						for (int nTile = ai.getAndIncrement(); nTile < tiles_list[nser.get()].length; nTile = ai.getAndIncrement()) {








							tileX = tiles_list[nser.get()][nTile][0];




							tileX = tiles_list[nser.get()][nTile][0];








							tileY = tiles_list[nser.get()][nTile][1];




							tileY = tiles_list[nser.get()][nTile][1];








							double [][] texture_tile =texture_tiles[tileY][tileX];




							double [][] texture_tile =texture_tiles[tileY][tileX];








							if (texture_tile != null) {




							if (texture_tile != null) {








								if (overlap) {




								if (overlap) {








									if ((tileY >0) && (tileX > 0) && (tileY < lastY) && (tileX < lastX)) { // fast, no extra checks




									if ((tileY >0) && (tileX > 0) && (tileY < lastY) && (tileX < lastX)) { // fast, no extra checks - ignore first/last rows and columns








										for (int i = 0; i < n2; i++){




										for (int i = 0; i < n2; i++){








											int start_line = ((tileY*transform_size + i) * tilesX + tileX)*transform_size - offset;




											int start_line = ((tileY*transform_size + i) * tilesX + tileX)*transform_size - offset;








											for (int chn = 0; chn < texture_tile.length; chn++) {




											for (int chn = 0; chn < texture_tile.length; chn++) {










@@ -5082,6 +5087,7 @@ public class ImageDtt {






												if (texture_tile[schn] == null) {




												if (texture_tile[schn] == null) {








													dpixels[chn] = null;




													dpixels[chn] = null;








												} else {




												} else {










													// should it be better to multiply each color by alpha before accumulating? No, it is already windowed!








													if ((chn != 3) || !sharp_alpha) {




													if ((chn != 3) || !sharp_alpha) {








														for (int j = 0; j<n2;j++) {




														for (int j = 0; j<n2;j++) {








															dpixels[chn][start_line + j] += texture_tile[schn][n2 * i + j];




															dpixels[chn][start_line + j] += texture_tile[schn][n2 * i + j];










@@ -9278,6 +9284,14 @@ public class ImageDtt {






				}




				}








			}




			}








		}




		}










		// fix: removing extra slices










		if (!clt_parameters.keep_weights && (texture_tiles[tileY][tileX]!=null)) {










			if (numcol == 3 ) {










				texture_tiles[tileY][tileX] = new double[][] {texture_tiles[tileY][tileX][0],texture_tiles[tileY][tileX][1],texture_tiles[tileY][tileX][2],texture_tiles[tileY][tileX][3]};










			} else {










				texture_tiles[tileY][tileX] = new double[][] {texture_tiles[tileY][tileX][0],texture_tiles[tileY][tileX][1]};










			}










		}








	}




	}


















	public double [][] get2DCorrs(




	public double [][] get2DCorrs(


































src/main/java/com/elphel/imagej/tileprocessor/QuadCLT.java



+3
−1






























Original line number
Original line
Diff line number
Diff line







@@ -8985,6 +8985,7 @@ public class QuadCLT {






			  tp.clt_3d_passes.add(latest_scan); // put it back




			  tp.clt_3d_passes.add(latest_scan); // put it back








		  }




		  }








		  int next_pass = tp.clt_3d_passes.size(); //




		  int next_pass = tp.clt_3d_passes.size(); //










		  // Create tasks to scan, have tasks, disparity and border tiles in tp.clt_3d_passes








		  tp.thirdPassSetupSurf( // prepare tile tasks for the second pass based on the previous one(s) // needs last scan




		  tp.thirdPassSetupSurf( // prepare tile tasks for the second pass based on the previous one(s) // needs last scan








				  clt_parameters,




				  clt_parameters,








				  //FIXME: make a special parameter?




				  //FIXME: make a special parameter?










@@ -9564,6 +9565,7 @@ public class QuadCLT {






		  double [] alpha_zero = new double [4*image_dtt.transform_size*image_dtt.transform_size];




		  double [] alpha_zero = new double [4*image_dtt.transform_size*image_dtt.transform_size];








		  int alpha_index = 3;




		  int alpha_index = 3;








		  for (int i = 0; i < alpha_zero.length; i++) alpha_zero[i]=0.0;




		  for (int i = 0; i < alpha_zero.length; i++) alpha_zero[i]=0.0;










		  // border tiles are copied, alpha from alphaFade (not multiplied?)








		  for (int tileY = 0; tileY < tilesY; tileY++){




		  for (int tileY = 0; tileY < tilesY; tileY++){








			  for (int tileX = 0; tileX < tilesX; tileX++){




			  for (int tileX = 0; tileX < tilesX; tileX++){








				  texture_tiles_cluster[tileY][tileX]= null;




				  texture_tiles_cluster[tileY][tileX]= null;










@@ -9607,7 +9609,7 @@ public class QuadCLT {






				  texture_overlap[alpha_index][i] = d;




				  texture_overlap[alpha_index][i] = d;








			  }




			  }








		  }




		  }








		  // for now - use just RGB. Later add option for RGBA




		  // for now - use just RGB. Later add option for RGBA (?)








		  double [][] texture_rgb = {texture_overlap[0],texture_overlap[1],texture_overlap[2]};




		  double [][] texture_rgb = {texture_overlap[0],texture_overlap[1],texture_overlap[2]};








		  double [][] texture_rgba = {texture_overlap[0],texture_overlap[1],texture_overlap[2],texture_overlap[3]};




		  double [][] texture_rgba = {texture_overlap[0],texture_overlap[1],texture_overlap[2],texture_overlap[3]};








		  double [][] texture_rgbx = ((clt_parameters.alpha1 > 0)? texture_rgba: texture_rgb);




		  double [][] texture_rgbx = ((clt_parameters.alpha1 > 0)? texture_rgba: texture_rgb);


































src/main/java/com/elphel/imagej/tileprocessor/TwoQuadCLT.java



+4
−2






























Original line number
Original line
Diff line number
Diff line







@@ -2191,7 +2191,9 @@ public class TwoQuadCLT {
















		}




		}








		if (clt_parameters.show_rgba_color) {




		if (clt_parameters.show_rgba_color) {








			int          num_src_slices = 12 ; // calculate




			int numcol = quadCLT_main.isMonochrome()?1:3;










			int ports = imp_quad_main.length;










			int          num_src_slices = numcol + 1 + (clt_parameters.keep_weights?(ports + numcol + 1):0); // 12 ; // calculate








//			float [][][] ftextures = gPUTileProcessor.getTextures(




//			float [][][] ftextures = gPUTileProcessor.getTextures(








//		    		(is_mono?1:3), // int     num_colors,




//		    		(is_mono?1:3), // int     num_colors,








//		    		clt_parameters.keep_weights); // boolean keep_weights);




//		    		clt_parameters.keep_weights); // boolean keep_weights);










@@ -3533,7 +3535,7 @@ if (debugLevel > -100) return true; // temporarily !






			}




			}








		}




		}


















		quadCLT_main.writeKml(debugLevel ); // alos generated with x3d model




		quadCLT_main.writeKml(debugLevel ); // also generated with x3d model


















		String jp4_copy_path= quadCLT_main.correctionsParameters.selectX3dDirectory(




		String jp4_copy_path= quadCLT_main.correctionsParameters.selectX3dDirectory(








				set_name, // quad timestamp. Will be ignored if correctionsParameters.use_x3d_subdirs is false




				set_name, // quad timestamp. Will be ignored if correctionsParameters.use_x3d_subdirs is false