Got GPU-processed LWIR images

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

@@ -42,8 +42,20 @@ public class ThermalColor {
 		double k = out_range/PALETTE_RANGE;
 		double value = (v-min)/(max-min) * (this.palette.length - 1);
 		int ivalue = (int) (value);
-		if (ivalue < 0)                         return getRGB((float) min); // this.palette[0];
-		if (ivalue >= (this.palette.length -1)) return getRGB((float) max); // this.palette[this.palette.length -1];
+		if (ivalue < 0) {
+//			return getRGB((float) min); // this.palette[0];
+			return new float[] {
+					(float) this.palette[0][0],
+					(float) this.palette[0][1],
+					(float) this.palette[0][2]};
+		}
+		if (ivalue >= (this.palette.length -1)) {
+//			return getRGB((float) max); // this.palette[this.palette.length -1];
+			return new float[] {
+				(float) this.palette[this.palette.length -1][0],
+				(float) this.palette[this.palette.length -1][1],
+				(float) this.palette[this.palette.length -1][2]};
+		}
 		double a = (value-ivalue); // 0..1
 		float [] rslt = {
 				(float) (k*((1 - a) * this.palette[ivalue][0] + a * this.palette[ivalue+1][0])),

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

@@ -1346,7 +1346,7 @@ public class GpuQuad{ // quad camera description
 		cuCtxSynchronize();
 		
 		// Call the kernel function
-		cuLaunchKernel(this.gpuTileProcessor.GPU_IMCLT_ALL_kernel,
+		cuLaunchKernel(this.gpuTileProcessor.GPU_IMCLT_ALL_kernel, // failed with LWIR
 				GridFullWarps[0],    GridFullWarps[1],   GridFullWarps[2],   // Grid dimension
 				ThreadsFullWarps[0], ThreadsFullWarps[1],ThreadsFullWarps[2],// Block dimension
 				0, null,                   // Shared memory size and stream (shared - only dynamic, static is in code)
@@ -2231,7 +2231,8 @@ public class GpuQuad{ // quad camera description
 		copyD2H.dstPitch =        width_in_bytes;
 
 		copyD2H.WidthInBytes =    width_in_bytes;
-		copyD2H.Height =          3 * height; // /2;
+//		copyD2H.Height =          3 * height; // /2;
+		copyD2H.Height =          num_colors * height; // /2;
 
 		cuMemcpy2D(copyD2H); // run copy
 

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

@@ -2260,12 +2260,34 @@ public class QuadCLT extends QuadCLTCPU {
 
 		int out_width =  gpuQuad.getImageWidth()  + gpuQuad.getDttSize();
 		int out_height = gpuQuad.getImageHeight() + gpuQuad.getDttSize();
+		if (isLwir() && colorProcParameters.lwir_autorange) {
+			double rel_low =  colorProcParameters.lwir_low;
+			double rel_high = colorProcParameters.lwir_high;
+			if (!Double.isNaN(getLwirOffset())) {
+				rel_low -=  getLwirOffset();
+				rel_high -= getLwirOffset();
+			}
+			double [] cold_hot =  autorange(
+					iclt_fimg,                         // iclt_data, // double [][][] iclt_data, //  [iQuad][ncol][i] - normally only [][2][] is non-null
+					rel_low,                           // double hard_cold,// matches data, DC (this.lwir_offset)  subtracted
+					rel_high,                          // double hard_hot,   // matches data, DC (this.lwir_offset)  subtracted
+					colorProcParameters.lwir_too_cold, // double too_cold, // pixels per image
+					colorProcParameters.lwir_too_hot,  // double too_hot,  // pixels per image
+					1024); // int num_bins)
+			if (cold_hot != null) {
+				if (!Double.isNaN(getLwirOffset())) {
+					cold_hot[0] += getLwirOffset();
+					cold_hot[1] += getLwirOffset();
+				}
+			}
+			setColdHot(cold_hot); // will be used for shifted images and for texture tiles
+		}
 		
 		/* Prepare 4-channel images*/
 		ImagePlus [] imps_RGB = new ImagePlus[iclt_fimg.length];
 		for (int ncam = 0; ncam < iclt_fimg.length; ncam++) {
             String title=String.format("%s%s-%02d",image_name, sAux(), ncam);
-			imps_RGB[ncam] = linearStackToColor( // probably no need to separate and process the second half with quadCLT_aux
+			imps_RGB[ncam] = linearStackToColor( // probably no need to separate and process the second half with quadCLT_aux (!)
 					clt_parameters,
 					colorProcParameters,
 					rgbParameters,
@@ -2284,22 +2306,22 @@ public class QuadCLT extends QuadCLTCPU {
 		
 		if (clt_parameters.gen_chn_img || only4slice) { // save and show 4-slice image
 			// combine to a sliced color image
-			// assuming total number of images to be multiple of 4
-			//			  int [] slice_seq = {0,1,3,2}; //clockwise
-//			int [] slice_seq = new int[gpuQuad.getNumCams()]; //results.length];
-			int [] slice_seq = new int[getNumSensors()]; //results.length];
+			int [] slice_seq = {0,1,3,2}; //clockwise
+			if (imps_RGB.length > 4) {
+				slice_seq = new int [imps_RGB.length];
 				for (int i = 0; i < slice_seq.length; i++) {
-				slice_seq[i] = i ^ ((i >> 1) & 1); // 0,1,3,2,4,5,7,6, ...
+					slice_seq[i] = i;
+				}
 			}
 			int width = imps_RGB[0].getWidth();
 			int height = imps_RGB[0].getHeight();
 			ImageStack array_stack=new ImageStack(width,height);
 			for (int i = 0; i<slice_seq.length; i++){
-///				if (imps_RGB[slice_seq[i]] != null) {
+				///				if (imps_RGB[slice_seq[i]] != null) {
 				array_stack.addSlice("port_"+slice_seq[i], imps_RGB[slice_seq[i]].getProcessor().getPixels());
-///				} else {
-///					array_stack.addSlice("port_"+slice_seq[i], results[slice_seq[i]].getProcessor().getPixels());
-///				}
+				///				} else {
+				///					array_stack.addSlice("port_"+slice_seq[i], results[slice_seq[i]].getProcessor().getPixels());
+				///				}
 			}
 			ImagePlus imp_stack = new ImagePlus(image_name+sAux()+"-SHIFTED-D"+clt_parameters.disparity, array_stack);
 			imp_stack.getProcessor().resetMinAndMax();

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

@@ -5829,7 +5829,7 @@ public class QuadCLTCPU {
 								  debugLevel);
 						  for (int ii = 0; ii < clt_set.length; ii++) clt[chn*4+ii] = clt_set[ii];
 					  }
-/*
+					  /*
 					  if (debugLevel > 0){
 						  sdfa_instance.showArrays(clt,
 								  tilesX*image_dtt.transform_size,
@@ -5863,7 +5863,7 @@ public class QuadCLTCPU {
 
 
 			  // Use iclt_data here for LWIR autorange
-//			  if (colorProcParameters.isLwir() && colorProcParameters.lwir_autorange) {
+			  //			  if (colorProcParameters.isLwir() && colorProcParameters.lwir_autorange) {
 			  if (isLwir() && colorProcParameters.lwir_autorange) {
 				  double rel_low =  colorProcParameters.lwir_low;
 				  double rel_high = colorProcParameters.lwir_high;
@@ -5924,7 +5924,7 @@ public class QuadCLTCPU {
 					  if (imps_RGB[slice_seq[i]] != null) {
 						  array_stack.addSlice("port_"+slice_seq[i], imps_RGB[slice_seq[i]].getProcessor().getPixels());
 					  } else { // not used in lwir
-///						  array_stack.addSlice("port_"+slice_seq[i], results[slice_seq[i]].getProcessor().getPixels());
+						  ///						  array_stack.addSlice("port_"+slice_seq[i], results[slice_seq[i]].getProcessor().getPixels());
 					  }
 				  }
 				  ImagePlus imp_stack = new ImagePlus(image_name+sAux()+"-SHIFTED-D"+clt_parameters.disparity, array_stack);
@@ -6633,6 +6633,21 @@ public class QuadCLTCPU {
 		  }
 		  return hist;
 	  }
+	  public int [] getLwirHistogram( // USED in lwir
+			  float [] data,
+			  double    hard_cold,
+			  double    hard_hot,
+			  int       num_bins) {
+		  int [] hist = new int [num_bins];
+		  double k = num_bins / (hard_hot - hard_cold);
+		  for (double d:data) {
+			  int bin = (int) ((d - hard_cold)*k);
+			  if (bin < 0) bin = 0;
+			  else if (bin >= num_bins) bin = (num_bins -1);
+			  hist[bin]++;
+		  }
+		  return hist;
+	  }
 	  public int [] addHist( // USED in lwir
 			  int [] this_hist,
 			  int [] other_hist) {
@@ -6723,6 +6738,49 @@ public class QuadCLTCPU {
 		  return abs_lim;
 	  }
 
+	  public double [] autorange( // USED in lwir
+			  float [][][] iclt_data, //  [iQuad][ncol][i] - normally only [][2][] is non-null
+			  double hard_cold,// matches data, DC (this.lwir_offset)  subtracted
+			  double hard_hot, // matches data, DC (this.lwir_offset)  subtracted
+			  double too_cold, // pixels per image
+			  double too_hot,  // pixels per image
+			  int num_bins) {
+		  int ncol;
+		  for (ncol = 0; ncol < iclt_data[0].length; ncol++) {
+			  if (iclt_data[0][ncol] != null) break;
+		  }
+		  too_cold *= iclt_data.length;
+		  too_hot *= iclt_data.length;
+		  int [] hist = null;
+		  for (int iQuad = 0; iQuad < iclt_data.length; iQuad++) {
+			  int [] this_hist = getLwirHistogram(
+					  iclt_data[iQuad][ncol], // double [] data,
+					  hard_cold,
+					  hard_hot,
+					  num_bins);
+			  if (hist == null) {
+				  hist = this_hist;
+			  } else {
+				  addHist(
+						  hist,
+						  this_hist);
+			  }
+		  }
+		  double [] rel_lim = {
+				  getMarginFromHist(
+						  hist, // histogram
+						  too_cold, // double cumul_val, // cummulative number of items to be ignored
+						  false), // boolean high_marg)
+				  getMarginFromHist(
+						  hist, // histogram
+						  too_hot, // double cumul_val, // cummulative number of items to be ignored
+						  true)}; // boolean high_marg)
+		  double [] abs_lim = {
+				  rel_lim[0] * (hard_hot - hard_cold) + hard_cold,
+				  rel_lim[1] * (hard_hot - hard_cold) + hard_cold,
+		  };
+		  return abs_lim;
+	  }
 
 
 	  // float
@@ -6747,9 +6805,14 @@ public class QuadCLTCPU {
 		  // convert to ImageStack of 3 slices
 		  String [] sliceNames = {"red", "blue", "green"};
 		  int green_index = 2;
-		  float [][] rbg_in = {iclt_data[0],iclt_data[1],iclt_data[2]};
+		  float [][] rbg_in;
+		  if (iclt_data.length >= 3) {
+			  rbg_in = new float [][] {iclt_data[0],iclt_data[1],iclt_data[2]}; // RBG or LWIR CPU
+		  } else {
+			  rbg_in = new float [][] {iclt_data[0],iclt_data[0],iclt_data[0]}; // after LWIR/GPU
+			  green_index = 0;
+		  }
 		  float []   alpha = null; // (0..1.0)
-//		  float [][] rgb_in = {iclt_data[0],iclt_data[1],iclt_data[2]};
 		  if (iclt_data.length > 3) alpha = iclt_data[3];
 		  if (isLwir()) {
 			  String [] rgb_titles =  {"red","green","blue"};