CODEX: Add CUDA pose Jacobian kernel

Co-authored-by: Codex <codex@elphel.com>

jna/test_pose_fx_jacobian_jna.cu

+// test_pose_fx_jacobian_jna.cu - launch/ABI smoke gate for the primitive-float
+// pose getFxDerivs CUDA kernel. The real numerical oracle is
+// imagej-elphel/IntersceneLmaFloat on captured ERS camera state; this synthetic
+// identity-camera case verifies the JNA buffer contract, validity mask, fx and
+// both calculateJacobian modes without requiring captured scene files.
+//
+// Build after jna/build_lib.sh:
+//   cd tile_processor_gpu && /usr/local/cuda-12.8/bin/nvcc -std=c++17 -O2 \
+//     jna/test_pose_fx_jacobian_jna.cu -o tests_bin/test_pose_fx_jacobian_jna \
+//     -L jna -ltileproc -Xlinker -rpath -Xlinker $PWD/jna
+//
+// Created Jul 15, 2026 by Codex for Elphel.
+
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <vector>
+
+extern "C" {
+  void* tp_create_module(const char*, const char*);
+  const char* tp_last_error();
+  void* tp_proc_create(void*);
+  int   tp_proc_setup(void*, int,int,int,int,int,int);
+  int   tp_proc_exec_pose_fx_jacobian(
+      void*, const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,const float*,int,const float*,int,
+      const float*,const float*,const unsigned char*,int,int,
+      float*,float*,unsigned char*);
+  void  tp_proc_destroy(void*);
+  void  tp_destroy_module(void*);
+}
+
+int main(int argc, char** argv)
+{
+  const char* srcdir=(argc>1)?argv[1]:"src";
+  const char* devrt=(argc>2)?argv[2]:
+      "/usr/local/cuda-12.8/targets/x86_64-linux/lib/libcudadevrt.a";
+  void* module=tp_create_module(srcdir,devrt);
+  if(!module){ std::fprintf(stderr,"tp_create_module: %s\n",tp_last_error()); return EXIT_FAILURE; }
+  void* proc=tp_proc_create(module);
+  if(!proc || tp_proc_setup(proc,1,1,16,16,0,0)){
+    std::fprintf(stderr,"tp_proc setup: %s\n",tp_last_error());
+    if(proc) tp_proc_destroy(proc); tp_destroy_module(module); return EXIT_FAILURE;
+  }
+
+  const int width=64,height=48,rbr_len=256,ers_len=height*14,num_tiles=3;
+  const float meta[8]={
+      (float)width,(float)height,4.0f,5.0f,2.85f,1.0f,0.0f,0.0004f};
+  const float radial[7]={0,0,0,0,0,0,0};
+  std::vector<float> rbr(rbr_len,1.0f),ers(ers_len,0.0f);
+  for(int line=0;line<height;line++) ers[(size_t)line*14+6]=1.0f;
+  const float pose[12]={0,0,0, 0,0,0, 0,0,0, 0,0,0};
+  const float centers[3*num_tiles]={20.0f,10.0f,5.0f, 40.0f,30.0f,0.0f, 12.0f,16.0f,3.0f};
+  const unsigned char selection[num_tiles]={1,1,0};
+  float fx[2*num_tiles],jt[3*2*num_tiles]; unsigned char valid[num_tiles];
+  int bad=0;
+  int rc=tp_proc_exec_pose_fx_jacobian(
+      proc,meta,radial,rbr.data(),rbr_len,ers.data(),ers_len,
+      meta,radial,rbr.data(),rbr_len,ers.data(),ers_len,
+      pose,centers,selection,num_tiles,1,fx,jt,valid);
+  if(rc){ std::fprintf(stderr,"pose jacobian rc=%d: %s\n",rc,tp_last_error()); bad++; }
+  if(!rc){
+    for(int tile=0;tile<2;tile++){
+      if(!valid[tile]){ std::printf("tile %d unexpectedly invalid\n",tile); bad++; continue; }
+      for(int component=0;component<2;component++){
+        const float error=std::fabs(fx[2*tile+component]-centers[3*tile+component]);
+        if(error>1.0e-4f){ std::printf("tile %d fx[%d] error %g\n",tile,component,error); bad++; }
+      }
+      for(int par=0;par<3;par++) for(int component=0;component<2;component++){
+        const float value=jt[(size_t)par*2*num_tiles+2*tile+component];
+        if(!std::isfinite(value)){ std::printf("tile %d J[%d,%d] non-finite\n",tile,par,component); bad++; }
+      }
+    }
+    if(valid[2] || !std::isnan(fx[4]) || !std::isnan(fx[5])){
+      std::printf("unselected tile was not masked\n"); bad++;
+    }
+  }
+
+  // Candidate-only mode passes a null Jacobian and must preserve fx/valid.
+  rc=tp_proc_exec_pose_fx_jacobian(
+      proc,meta,radial,rbr.data(),rbr_len,ers.data(),ers_len,
+      meta,radial,rbr.data(),rbr_len,ers.data(),ers_len,
+      pose,centers,selection,num_tiles,0,fx,nullptr,valid);
+  if(rc){ std::fprintf(stderr,"pose fx-only rc=%d: %s\n",rc,tp_last_error()); bad++; }
+  else if(!valid[0]||!valid[1]||valid[2]){ std::printf("fx-only validity mismatch\n"); bad++; }
+
+  std::printf("NVRTC pose_fx_jacobian: identity fx + mask + finite J -> %s\n",bad?"FAIL":"PASS");
+  tp_proc_destroy(proc); tp_destroy_module(module);
+  return bad?EXIT_FAILURE:EXIT_SUCCESS;
+}

src/tp_lma.h

 /*
  * tp_lma.h
  *
- * Small normal-equation product kernel for the pose-chain LMA. Host policy
- * chooses the active parameters/tiles and prepares the eigen-transformed
- * Jacobian, weights and weighted residual; the GPU returns only
+ * Pose projection/Jacobian and normal-equation kernels for the pose-chain LMA.
+ * The first kernel is the primitive-float CUDA counterpart of
+ * IntersceneLmaFloat.getFxDerivs(): common pose matrices are prepared once per
+ * block and one thread processes each selected tile. The older product kernel
+ * consumes a host-prepared eigen-transformed Jacobian, weights and residual.
  *
  *   H = J^T W J    and    b = J^T W (y - f(x)).
  *
@@ -21,6 +23,40 @@
 #include <cuda_runtime.h>
 #endif
 
+#define POSE_NUM_PARAMS       3
+#define POSE_NUM_COMPONENTS   2
+#define POSE_RADIAL_FLOATS    7
+#define POSE_ERS_LINE_STRIDE 14
+
+typedef struct {
+	int   width;
+	int   height;
+	int   rbr_len;
+	float focal_length;
+	float pixel_size;
+	float distortion_radius;
+	float disparity_radius;
+	float line_time;
+	float step_r;
+} PoseCameraMeta;
+
+extern "C" __global__ void pose_fx_jacobian(
+		int                    num_tiles,
+		int                    calculate_jacobian,
+		const PoseCameraMeta * cameras,       // reference, scene
+		const float          * ref_radial,     // [7]
+		const float          * ref_rbr,
+		const float          * ref_ers,        // [height][14]
+		const float          * scene_radial,   // [7]
+		const float          * scene_rbr,
+		const float          * scene_ers,      // [height][14]
+		const float          * pose_vectors,   // ref xyz/atr, scene xyz/atr: 4x3
+		const float          * centers,        // packed xyz triples
+		const unsigned char  * selection,
+		float                * fx,             // [2*num_tiles]
+		float                * jt,             // parameter-major [3][2*num_tiles], nullable
+		unsigned char        * valid);
+
 extern "C" __global__ void lma_normal_products(
 		int            num_params,
 		int            num_values,