CLAUDE: tp_proc_set_const — upload LPF coefficients to named __constant__ symbols (native module)

Mirrors GpuQuad.setLpfRbg/setLpfCorr (cuModuleGetGlobal + cuMemcpyHtoD) against the native module's
lpf_data / lpf_corr / lpf_rb_corr symbols. Fixes the first JNA-mode NPE (setLpfRbg hit the null
gpuTileProcessor.module).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

jna/tp_jna.cpp

@@ -657,6 +657,20 @@ int tp_proc_set_tasks(TpProc* p, const float* ftasks, int ntiles, int total_floa
   if(p->ftasks) cudaFree(p->ftasks); cudaMalloc((void**)&p->ftasks,(size_t)total_floats*sizeof(float));
   p->ntiles=ntiles; return cudaMemcpy(p->ftasks,ftasks,(size_t)total_floats*sizeof(float),cudaMemcpyHostToDevice)==cudaSuccess?0:-2; }
 
+// Upload to a named __constant__ symbol in the module (LPF filters: lpf_data / lpf_corr / lpf_rb_corr).
+// Mirrors GpuQuad.setLpfRbg/setLpfCorr (cuModuleGetGlobal + cuMemcpyHtoD), against the native module.
+int tp_proc_set_const(TpProc* p, const char* name, const float* data, int n){
+  if(!p){ seterr("tp_proc_set_const: null"); return -1; }
+  cuCtxSetCurrent(p->mod->ctx);
+  CUdeviceptr dptr=0; size_t sz=0;
+  CUresult cr=cuModuleGetGlobal(&dptr,&sz,p->mod->mod,name);
+  if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es); seterr("cuModuleGetGlobal(%s) -> %d (%s)",name,cr,es); return -2; }
+  size_t want=(size_t)n*sizeof(float); size_t cpy = (want<sz)?want:sz;
+  cr=cuMemcpyHtoD(dptr,data,cpy);
+  if(cr!=CUDA_SUCCESS){ const char* es; cuGetErrorString(cr,&es); seterr("HtoD const %s -> %d (%s)",name,cr,es); return -3; }
+  return 0;
+}
+
 // geometry: calc_reverse_distortions + rot_derivs + (uniform_grid) calculate_tiles_offsets
 int tp_proc_exec_geometry(TpProc* p, int uniform_grid){
   if(!p)return -1; cuCtxSetCurrent(p->mod->ctx); CUresult cr; const char* es;