CLAUDE: step-3 green-ctx carve - TP_GREEN_SM env (count|skip+count|skip+rem) + run_cases --green

Module context optionally built from a green-ctx SM carve (probe2k pattern);
disjoint partitions via one granularity split + group merge (remainder re-split
is CUDA 915). Gates: scene_dp CDP chain + render_pipe tol-0 PASS in 8/16/8+16/
8+rem/24+rem carves; no-carve suite unchanged; sanitizer 0.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

jna/tp_jna.cpp

@@ -66,7 +66,67 @@ static const char* KERNELS[] = {
   "lma_normal_products"}; // deterministic double normal-equation products // By Codex on 07/14/2026
 static const int N_KERNELS = sizeof(KERNELS)/sizeof(KERNELS[0]);
 
-struct TpModule { CUcontext ctx; CUmodule mod; int nfun; };
+struct TpModule { CUcontext ctx; CUmodule mod; int nfun; CUgreenCtx gctx; }; // gctx: item-4 step-3 SM carve (null = plain ctx) // By Claude on 07/19/2026
+
+// item-4 step-3 (GPU partitioning): optional green-context SM carve for the module's
+// context. TP_GREEN_SM = "<count>" | "<skip>+<count>" | "<skip>+rem"; unset/"0" = plain
+// cuCtxCreate (bit-for-bit the old path). The split is deterministic, so "<skip>+..."
+// yields partitions DISJOINT from a "<count>"-carved module in the same process
+// (e.g. pose "8" + render "8+16"). Everything the module does afterwards (allocs,
+// streams incl. the R1a render stream, CDP link, launches) inherits the carve because
+// every API entry re-binds m->ctx. By Claude on 07/19/2026.
+static CUresult make_tp_context(CUdevice dev, CUcontext* out_ctx, CUgreenCtx* out_g){
+  *out_ctx=nullptr; *out_g=nullptr;
+  const char* spec = getenv("TP_GREEN_SM");
+  if(!spec || !*spec || !strcmp(spec,"0")) return cuCtxCreate(out_ctx,0,dev);
+  int skip=0, count=0, use_rem=0;
+  const char* plus = strchr(spec,'+');
+  if(plus){ skip=atoi(spec); if(!strcmp(plus+1,"rem")) use_rem=1; else count=atoi(plus+1); }
+  else count=atoi(spec);
+  if(skip<0 || (!use_rem && count<=0)){
+    seterr("TP_GREEN_SM='%s' invalid (want <count> | <skip>+<count> | <skip>+rem)",spec);
+    return CUDA_ERROR_INVALID_VALUE; }
+  CUresult cr;
+  CUdevResource res={}; if((cr=cuDeviceGetDevResource(dev,&res,CU_DEV_RESOURCE_TYPE_SM))!=CUDA_SUCCESS) return cr;
+  CUdevResourceDesc desc; unsigned got=0;
+  if(skip==0 && !use_rem){ // plain "<count>": one split, take the first group
+    CUdevResource grp={}, rem={}; unsigned n=1;
+    if((cr=cuDevSmResourceSplitByCount(&grp,&n,&res,&rem,0,(unsigned)count))!=CUDA_SUCCESS) return cr;
+    if((cr=cuDevResourceGenerateDesc(&desc,&grp,1))!=CUDA_SUCCESS) return cr;
+    got=grp.sm.smCount;
+  } else {
+    // Disjoint partitions: a split's REMAINDER cannot be split again (CUDA 915),
+    // so split ONCE into granularity groups (minCount=1 rounds up to the card's
+    // co-scheduling granularity, 8 on sm_120/36-SM) and MERGE disjoint slices
+    // into one descriptor (merge verified: 8+8 -> 16-SM green ctx). Accumulate
+    // until >= target, so off-granularity requests overshoot (actual printed).
+    CUdevResource parts[65]; unsigned np=0;
+    CUdevResource grpv[64]={}, remg={}; unsigned ng=64;
+    if((cr=cuDevSmResourceSplitByCount(grpv,&ng,&res,&remg,0,1))!=CUDA_SUCCESS) return cr;
+    if(remg.sm.smCount) grpv[ng++]=remg; // treat the tail remainder as a final (bigger) group
+    unsigned i=0, acc=0;
+    for(; i<ng && acc<(unsigned)skip; i++) acc+=grpv[i].sm.smCount;
+    if(acc<(unsigned)skip){ seterr("TP_GREEN_SM='%s': only %u SMs available to skip",spec,acc); return CUDA_ERROR_INVALID_VALUE; }
+    for(unsigned j=i; j<ng && (use_rem || got<(unsigned)count); j++){ parts[np++]=grpv[j]; got+=grpv[j].sm.smCount; }
+    if(!np || (!use_rem && got<(unsigned)count)){
+      seterr("TP_GREEN_SM='%s': only %u SMs left after skipping %u",spec,got,acc); return CUDA_ERROR_INVALID_VALUE; }
+    if((cr=cuDevResourceGenerateDesc(&desc,parts,np))!=CUDA_SUCCESS) return cr;
+  }
+  CUgreenCtx g;
+  if((cr=cuGreenCtxCreate(&g,desc,dev,CU_GREEN_CTX_DEFAULT_STREAM))!=CUDA_SUCCESS) return cr;
+  if((cr=cuCtxFromGreenCtx(out_ctx,g))!=CUDA_SUCCESS){ cuGreenCtxDestroy(g); return cr; }
+  // unlike cuCtxCreate, cuCtxFromGreenCtx does NOT bind the context // By Claude on 07/19/2026
+  if((cr=cuCtxSetCurrent(*out_ctx))!=CUDA_SUCCESS){ cuGreenCtxDestroy(g); *out_ctx=nullptr; return cr; }
+  *out_g=g;
+  printf("tp_jna: GREEN CTX carve TP_GREEN_SM=%s -> %u SMs%s\n", spec, got,
+         skip>0?" (disjoint: carved after skip)":"");
+  fflush(stdout);
+  return CUDA_SUCCESS;
+}
+// A context converted from a green ctx must NOT be cuCtxDestroy'ed. By Claude on 07/19/2026.
+static void destroy_tp_context(CUcontext ctx, CUgreenCtx g){
+  if(g) cuGreenCtxDestroy(g); else if(ctx) cuCtxDestroy(ctx);
+}
 
 // Stage-1 geometry buffer sizes (floats). Must match GpuQuad.java alloc:
 //   gpu_geometry_correction = GeometryCorrection.arrayLength(16) (== sizeof(struct gc)/4 == 165),
@@ -116,7 +176,9 @@ TpModule* tp_create_module(const char* srcdir, const char* devrt){
   FAILCU(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev));
   FAILCU(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev));
   int cap = major*10+minor;
-  CUcontext ctx; FAILCU(cuCtxCreate(&ctx,0,dev));
+  CUcontext ctx; CUgreenCtx gctx; // plain ctx, or TP_GREEN_SM carve // By Claude on 07/19/2026
+  cr = make_tp_context(dev,&ctx,&gctx);
+  if(cr!=CUDA_SUCCESS){ if(g_err[0]==0){ cuGetErrorString(cr,&es); seterr("make_tp_context (TP_GREEN_SM) -> CUDA %d (%s)",cr,es); } return nullptr; }
 
   int narch=0; nvrtcGetNumSupportedArchs(&narch);
   std::vector<int> archs(narch>0?narch:1); nvrtcGetSupportedArchs(archs.data());
@@ -125,39 +187,39 @@ TpModule* tp_create_module(const char* srcdir, const char* devrt){
 
   std::string src = TP_DEFINES;
   for(const char* fn : SRC_FILES){ std::string c=readFile(std::string(srcdir)+"/"+fn);
-    if(c.empty()){ seterr("missing/empty kernel source: %s/%s", srcdir, fn); cuCtxDestroy(ctx); return nullptr; }
+    if(c.empty()){ seterr("missing/empty kernel source: %s/%s", srcdir, fn); destroy_tp_context(ctx,gctx); return nullptr; }
     src += "\n// ==== "; src += fn; src += " ====\n"; src += c; }
 
   nvrtcProgram prog; nvrtcResult nr = nvrtcCreateProgram(&prog, src.c_str(), "tileproc.cu", 0, nullptr, nullptr);
-  if(nr!=NVRTC_SUCCESS){ seterr("nvrtcCreateProgram -> %s", nvrtcGetErrorString(nr)); cuCtxDestroy(ctx); return nullptr; }
+  if(nr!=NVRTC_SUCCESS){ seterr("nvrtcCreateProgram -> %s", nvrtcGetErrorString(nr)); destroy_tp_context(ctx,gctx); return nullptr; }
   std::string archopt = "--gpu-architecture=compute_" + std::to_string(cap);
   const char* opts[] = { archopt.c_str(), "--extensible-whole-program" };
   nr = nvrtcCompileProgram(prog, 2, opts);
   if(nr!=NVRTC_SUCCESS){ size_t ls=0; nvrtcGetProgramLogSize(prog,&ls); std::vector<char> lg(ls?ls:1); nvrtcGetProgramLog(prog,lg.data());
-    seterr("nvrtcCompileProgram -> %s\n%s", nvrtcGetErrorString(nr), lg.data()); nvrtcDestroyProgram(&prog); cuCtxDestroy(ctx); return nullptr; }
+    seterr("nvrtcCompileProgram -> %s\n%s", nvrtcGetErrorString(nr), lg.data()); nvrtcDestroyProgram(&prog); destroy_tp_context(ctx,gctx); return nullptr; }
   size_t ptxsz=0; nvrtcGetPTXSize(prog,&ptxsz); std::vector<char> ptx(ptxsz); nvrtcGetPTX(prog,ptx.data()); nvrtcDestroyProgram(&prog);
 
   CUlinkState ls;
   FAILCU(cuLinkCreate(0,nullptr,nullptr,&ls));
-  if((cr=cuLinkAddFile(ls,CU_JIT_INPUT_LIBRARY,devrt,0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkAddFile(%s) -> %d (%s)",devrt,cr,es); cuLinkDestroy(ls); cuCtxDestroy(ctx); return nullptr; }
-  if((cr=cuLinkAddData(ls,CU_JIT_INPUT_PTX,ptx.data(),ptxsz,"tileproc.ptx",0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkAddData -> %d (%s)",cr,es); cuLinkDestroy(ls); cuCtxDestroy(ctx); return nullptr; }
+  if((cr=cuLinkAddFile(ls,CU_JIT_INPUT_LIBRARY,devrt,0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkAddFile(%s) -> %d (%s)",devrt,cr,es); cuLinkDestroy(ls); destroy_tp_context(ctx,gctx); return nullptr; }
+  if((cr=cuLinkAddData(ls,CU_JIT_INPUT_PTX,ptx.data(),ptxsz,"tileproc.ptx",0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkAddData -> %d (%s)",cr,es); cuLinkDestroy(ls); destroy_tp_context(ctx,gctx); return nullptr; }
   void* cubin=nullptr; size_t cubinsz=0;
-  if((cr=cuLinkComplete(ls,&cubin,&cubinsz))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkComplete -> %d (%s)",cr,es); cuLinkDestroy(ls); cuCtxDestroy(ctx); return nullptr; }
+  if((cr=cuLinkComplete(ls,&cubin,&cubinsz))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuLinkComplete -> %d (%s)",cr,es); cuLinkDestroy(ls); destroy_tp_context(ctx,gctx); return nullptr; }
   CUmodule mod;
-  if((cr=cuModuleLoadDataEx(&mod,cubin,0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuModuleLoadDataEx -> %d (%s)",cr,es); cuLinkDestroy(ls); cuCtxDestroy(ctx); return nullptr; }
+  if((cr=cuModuleLoadDataEx(&mod,cubin,0,nullptr,nullptr))!=CUDA_SUCCESS){ cuGetErrorString(cr,&es); seterr("cuModuleLoadDataEx -> %d (%s)",cr,es); cuLinkDestroy(ls); destroy_tp_context(ctx,gctx); return nullptr; }
   cuLinkDestroy(ls);
 
   int found=0;
   for(int i=0;i<N_KERNELS;i++){ CUfunction f; if(cuModuleGetFunction(&f,mod,KERNELS[i])==CUDA_SUCCESS) found++;
     else { if(g_err[0]==0) seterr("missing kernel %s (and maybe others)", KERNELS[i]); } }
-  TpModule* m = new TpModule{ctx,mod,found};
+  TpModule* m = new TpModule{ctx,mod,found,gctx}; // By Claude on 07/19/2026
   return m;
   #undef FAILCU
 }
 
 int  tp_module_num_functions(TpModule* m){ return m ? m->nfun : -1; }
 const char* tp_last_error(){ return g_err; }
-void tp_destroy_module(TpModule* m){ if(m){ if(m->mod) cuModuleUnload(m->mod); if(m->ctx) cuCtxDestroy(m->ctx); delete m; } }
+void tp_destroy_module(TpModule* m){ if(m){ cuCtxSetCurrent(m->ctx); if(m->mod) cuModuleUnload(m->mod); destroy_tp_context(m->ctx,m->gctx); delete m; } } // green-aware // By Claude on 07/19/2026
 
 // ---- Stage 1: instance + geometry path (calcReverseDistortionTable, calc_rot_deriv) ----
 // Mirrors GpuQuad.execCalcReverseDistortions / execRotDerivs (driver-API cuLaunchKernel),

run_cases.sh

@@ -17,6 +17,8 @@
 #   ./run_cases.sh -l <scene.list> -v v013-MB2x ...   # explicit overrides
 #   ./run_cases.sh --model <dir> ...      # pin ${MODEL} directly
 #   ./run_cases.sh --list                 # show cases + resolved paths, run nothing
+#   ./run_cases.sh --green 8 ...          # run inside a green-ctx SM carve (exports
+#                                         # TP_GREEN_SM: <count>|<skip>+<count>|<skip>+rem)
 #
 # By Claude on 07/13/2026, from Andrey's design.
 set -u
@@ -34,7 +36,8 @@ while [ $# -gt 0 ]; do
         -v)      version="$2"; shift 2;;
         --model) model="$2"; shift 2;;
         --list)  show_only=1; shift;;
-        -h|--help) sed -n '2,21p' "$0"; exit 0;;
+        --green) export TP_GREEN_SM="$2"; shift 2;;   # item-4 step-3 carve // By Claude on 07/19/2026
+        -h|--help) sed -n '2,23p' "$0"; exit 0;;
         *)       want+=("$1"); shift;;
     esac
 done