feat: add CUDA benchmark CLI and shared search algorithms

.gitignore

@@ -1,2 +1,50 @@
-slime
+# Local build output (source assembly snapshots are intentionally retained).
-slime-cuda
+bin/
+build/
+cmake-build-*/
+*.o
+*.obj
+*.a
+*.so
+*.dll
+*.exe
+*.d
+*.ptx
+*.cubin
+*.fatbin
+/slime
+/slime-cuda
+/sliming
+/codex/slime_cpu
+/codex/slime_cuda
+
+# Python environments and caches.
+__pycache__/
+*.py[cod]
+.venv/
+venv/
+.pytest_cache/
+.mypy_cache/
+.ruff_cache/
+
+# Temporary logs and scratch files, not benchmark JSON or research fixtures.
+*.log
+*.stderr
+*.tmp
+*.swp
+*.swo
+*~
+/exp/*.smt2
+/artifacts/
+
+# Machine-local configuration and credentials.
+.env
+.env.*
+!.env.example
+/.aws/
+/.codex/
+/.agents/
+/.idea/
+/.vscode/
+.DS_Store
+Thumbs.db

bench/.gitignore

@@ -0,0 +1,2 @@
+/bin/
+/__pycache__/

bench/Makefile

@@ -0,0 +1,21 @@
+NVCC ?= nvcc
+CUDA_ARCH ?= 86
+NVCCFLAGS ?= -O3 -std=c++17
+TARGET := bin/slime-bench
+HEADERS := ../include/slime/solver_core.hpp ../include/slime/upper_interval.hpp
+
+.PHONY: all clean test test-gpu
+all: $(TARGET)
+$(TARGET): slime_bench.cu $(HEADERS) Makefile
+	mkdir -p bin
+	$(NVCC) $(NVCCFLAGS) -Xcompiler -pthread -gencode arch=compute_$(CUDA_ARCH),code=sm_$(CUDA_ARCH) -gencode arch=compute_$(CUDA_ARCH),code=compute_$(CUDA_ARCH) slime_bench.cu -o $@
+bin/test-math: test_math.cu slime_bench.cu $(HEADERS) Makefile
+	mkdir -p bin
+	$(NVCC) $(NVCCFLAGS) -Xcompiler -pthread test_math.cu -o $@
+test: $(TARGET) bin/test-math
+	bin/test-math
+	python3 test_cli.py
+test-gpu: $(TARGET)
+	python3 test_cli.py --gpu
+clean:
+	rm -f $(TARGET) bin/test-math

bench/README.md

@@ -0,0 +1,115 @@
+# slime-bench：普通分支全种子搜索基准
+
+指定中心区块、曼哈顿半径和实心 `n×n` 图案，搜索范围内每个锚点的全部有效种子，输出 JSON 结果及计时。支持 NVIDIA CUDA 单卡、多卡。
+
+## 本地编译与运行
+
+需要 Linux/WSL、CUDA Toolkit（nvcc，推荐 12 或以上）、支持 C++17 的主机编译器。运行时需要 NVIDIA 驱动。项目根目录执行：
+
+```sh
+make -C bench
+mkdir -p artifacts
+./bench/bin/slime-bench --list-devices
+./bench/bin/slime-bench --x 0 --z 0 --radius 64 --size 5 \
+  --devices 0 --repeat 3 --output artifacts/local-r64-n5.json
+```
+
+默认编译 `sm_86` 并附带 PTX，适合本机 RTX 3060 Laptop。换架构时用 `-B` 强制重建；仅改变 Make 变量不会使已有文件自动过期。
+
+## H100 服务器
+
+将仓库内的 `bench/` 和 `include/` 一起复制到服务器，保持相对路径；CLI 依赖 `include/slime/` 中的共用头文件。
+
+```sh
+make -C bench -B CUDA_ARCH=90
+mkdir -p artifacts
+./bench/bin/slime-bench --list-devices
+
+# 先测单张 H100，和本地使用相同范围、尺寸及批大小。
+./bench/bin/slime-bench --x 0 --z 0 --radius 64 --size 5 \
+  --devices 0 --repeat 3 --output artifacts/h100-one-r64-n5.json
+
+# 再测全部可见 GPU；八卡机器即可使用八张卡。
+./bench/bin/slime-bench --x 0 --z 0 --radius 64 --size 5 \
+  --devices all --repeat 3 --output artifacts/h100-all-r64-n5.json
+```
+
+也可以指定 `--devices 0,1,2,3,4,5,6,7`。编号遵循 `CUDA_VISIBLE_DEVICES` 的可见编号。重复/不存在的设备编号会报错。多卡以动态位置批次分配任务，完整覆盖范围一次，各卡独立缓存类根，结果总数合并；所有 GPU 初始化完成后一起开始计时。实际八卡性能尚未在本机验证。
+
+半径 64 包含 8,321 个锚点，适合先验证环境。正式比较可以在两台机器上使用下面同一套较长参数（半径 512，共 525,313 个锚点）：
+
+```sh
+./bench/bin/slime-bench --x 0 --z 0 --radius 512 --size 5 \
+  --position-batch 128 --devices all --repeat 3 --output artifacts/benchmark-r512-n5.json
+```
+
+若八卡上仍耗时太短，可同时增大半径。`--repeat 3` 是完整扫描三遍，不是只跑三次 kernel。默认无额外隐藏预热；第一轮包含首次 kernel/JIT 开销，报告保留每一轮数据。
+
+## 范围与搜索口径
+
+```text
+abs(anchor_x - x) + abs(anchor_z - z) <= radius
+锚点数量 = 1 + 2*radius*(radius+1)
+```
+
+- `x,z,radius` 的单位都是区块，不是方块。
+- 锚点是图案最小 X、最小 Z 的角；图案沿 `+X,+Z` 延伸。
+- 半径限制锚点，图案允许超出菱形范围。
+- 不自动按世界边界裁剪；锚点和整块图案须在 int32 坐标范围内，否则报错。
+- `--size` 支持 1～6，默认 5。小尺寸匹配种子极多，全种子计数也可能耗时很久，不能用 1×1 推断 5×5 性能。
+- 遍历所有有效低 48 位种子的**普通候选分支**。沿用低位筛选、16 候选组 mask、区间求交、精确 Java 验证，**不包含拒绝重抽补扫**。
+- 返回的种子是 `[0,2^48)` 内的有效代表；不重复枚举高 16 位等价的 64 位世界种子。
+
+这是范围基准 CLI，不提供全世界长跑的断点续跑。Ctrl+C 可终止，未完成扫描不会输出成功报告。
+
+## 参数及输出
+
+```sh
+./bench/bin/slime-bench --help
+
+# 不使用 GPU，只检查输入范围
+./bench/bin/slime-bench --x 51069 --z 419124 --radius 2 --size 5 --dry-run
+
+# 已知位置验证：应该得到 24 个种子
+./bench/bin/slime-bench --x 51069 --z 419124 --radius 0 --size 5 \
+  --max-results 100 --repeat 3 --output artifacts/known.json
+```
+
+本文示例均从项目根目录运行，输出前请先执行 `mkdir -p artifacts`。日常跑分写入已忽略的 `artifacts/`；`results/` 只保存需要长期引用的基准样例。
+
+JSON 输出到 stdout，状态/约每 5 秒进度输出到 stderr；`--output` 同时保存 JSON，已有文件会报错以免覆盖。默认只计数；`--max-results N` 最多保留 N 个示例三元组 `[seed48,chunk_x,chunk_z]`，**不提前停止扫描**。示例顺序受 GPU 调度影响，不能按示例子集或其顺序对比性能/正确性。`sample_truncated` 表示有命中没有保存在示例中，`matches` 始终是完整计数；整数计数溢出会失败。
+
+主要计时字段：
+
+| 字段 | 含义 |
+|---|---|
+| `rounds[].elapsed_ms` | 本轮启动工作线程、设备初始化/分配、完整扫描及资源清理的墙钟时间 |
+| `rounds[].scan_wall_ms` | 所有 GPU 就绪后，从最早开始扫描到最后完成的墙钟时间，含本轮分组、根生成、传输、展开、搜索、结果回传 |
+| `median_scan_wall_ms` | 各轮 `scan_wall_ms` 的中位数，适合本地与服务器比较 |
+| `rounds[].anchors_per_second` | 实际锚点数 / 扫描墙钟秒数 |
+| `gpus[].root_stage_ms` | 该 GPU 工作线程的代表类构造、根筛选、回传、整理及缓存耗时 |
+| `gpus[].search_stage_ms` | 该线程的根任务上传、公共项准备、成组搜索和命中回传耗时 |
+| `gpus[].scan_ms` | 该线程完整扫描耗时；还包含以上两阶段外的调度、具体位置构造等工作 |
+
+多卡阶段时间相互重叠，不要把八卡 `scan_ms` 相加当作完成时间。比较相同的中心、半径、尺寸和批大小；位置根分布差异很大，单个已知命中点不适合代表全世界平均吞吐。
+
+可调批参数：
+
+- `--position-batch 32`：每个动态任务最多 32 个锚点，支持 1～256。
+- `--job-batch 2048`：每次搜索最多 2048 个低位根任务，支持 1～65536。
+- `--cache-classes 4096`：每卡缓存的坐标类数量，0 禁用；另限制最多约 1,048,576 个缓存根，FIFO 淘汰。每轮清空缓存，避免重复测量只跑缓存命中。
+
+按坐标类归一化并只构造有根类的具体位置。此 CLI 使用有限缓存、按曼哈顿行顺序调度；不是此前设想的全世界按类遍历计划。小半径下通常没有重复类，所以测到的根生成费用是真实费用，不会按假设将其省略。半径范围的实测吞吐不能直接替代此前约 50 天的全世界类复用搜索阶段外推。
+
+## 检查
+
+```sh
+make -C bench test       # 无需 GPU：范围、输入、可变尺寸数学对照
+make -C bench test-gpu   # 需要 GPU：实际搜索与文件输出验证
+```
+
+验证内容包括：半径 0～64 的完整菱形枚举，较大半径抽样，1～6 中多个尺寸的组 mask 对照，已知 24 个命中，小根批分拆一致性，4×4 的全部 422 个示例经独立 Java 算法验证，6×6 的已知种子子集对照，以及半径 2 批量结果与全部 13 个位置分别搜索的完整列表一致。
+
+本机实测：RTX 3060 Laptop，中心 `(0,0)`，半径 512，5×5，`--position-batch 128 --repeat 3`。525,313 个锚点，三轮均有 368,935 个低位根任务、0 个命中，扫描墙钟中位数约 **3.990 秒**。原始 JSON 在 [results/local-r512-n5.json](results/local-r512-n5.json)。另有半径 64 的快速环境验证记录；其批大小为默认 32，与上述正式样例不同。
+
+H100 的 `sm_90` 原生代码和 PTX 已在本机交叉编译通过，尚未在 H100 或八张物理 GPU 上执行。

bench/results/local-r512-n5.json

@@ -0,0 +1 @@
+{"schema_version":1,"mode":"ordinary_all_effective_seeds","rejection_tail_included":false,"center":[0,0],"radius_chunks":512,"square_size":5,"anchors":525313,"effective_seed_bits":48,"dry_run":false,"position_batch":128,"job_batch":2048,"cache_classes":4096,"cuda_runtime_version":12000,"rounds":[{"round":1,"elapsed_ms":4145.453479,"scan_wall_ms":3973.099405,"anchors_per_second":132217.431897,"root_jobs":368935,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":525313,"root_classes_generated":525313,"cache_hits":0,"root_jobs":368935,"matches":0,"setup_ms":170.953854,"scan_ms":3973.099405,"root_stage_ms":3429.569304,"search_stage_ms":369.153727}]},{"round":2,"elapsed_ms":3991.578504,"scan_wall_ms":3989.615391,"anchors_per_second":131670.085589,"root_jobs":368935,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":525313,"root_classes_generated":525313,"cache_hits":0,"root_jobs":368935,"matches":0,"setup_ms":0.879198,"scan_ms":3989.615391,"root_stage_ms":3441.271225,"search_stage_ms":371.001978}]},{"round":3,"elapsed_ms":4101.32096,"scan_wall_ms":4099.380455,"anchors_per_second":128144.485677,"root_jobs":368935,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":525313,"root_classes_generated":525313,"cache_hits":0,"root_jobs":368935,"matches":0,"setup_ms":0.86543,"scan_ms":4099.380455,"root_stage_ms":3543.591974,"search_stage_ms":374.196725}]}],"median_scan_wall_ms":3989.615391}

bench/results/local-r64-n5.json

@@ -0,0 +1 @@
+{"schema_version":1,"mode":"ordinary_all_effective_seeds","rejection_tail_included":false,"center":[0,0],"radius_chunks":64,"square_size":5,"anchors":8321,"effective_seed_bits":48,"dry_run":false,"position_batch":32,"job_batch":2048,"cache_classes":4096,"cuda_runtime_version":12000,"rounds":[{"round":1,"elapsed_ms":284.341397,"scan_wall_ms":95.570636,"anchors_per_second":87066.4918459,"root_jobs":6835,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":8321,"root_classes_generated":8321,"cache_hits":0,"root_jobs":6835,"matches":0,"setup_ms":187.64489,"scan_ms":95.570636,"root_stage_ms":81.438496,"search_stage_ms":9.997073}]},{"round":2,"elapsed_ms":95.692767,"scan_wall_ms":94.367954,"anchors_per_second":88176.1196179,"root_jobs":6835,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":8321,"root_classes_generated":8321,"cache_hits":0,"root_jobs":6835,"matches":0,"setup_ms":0.577818,"scan_ms":94.367954,"root_stage_ms":80.619006,"search_stage_ms":9.56154}]},{"round":3,"elapsed_ms":88.980653,"scan_wall_ms":87.740238,"anchors_per_second":94836.7612133,"root_jobs":6835,"matches":0,"sample_truncated":false,"sample":[],"gpus":[{"device":0,"name":"NVIDIA GeForce RTX 3060 Laptop GPU","anchors":8321,"root_classes_generated":8321,"cache_hits":0,"root_jobs":6835,"matches":0,"setup_ms":0.466152,"scan_ms":87.740238,"root_stage_ms":74.091536,"search_stage_ms":9.67283}]}],"median_scan_wall_ms":94.367954}

bench/slime_bench.cu

@@ -0,0 +1,201 @@
+#include "../include/slime/upper_interval.hpp"
+#include <cuda_runtime.h>
+#include <algorithm>
+#include <atomic>
+#include <chrono>
+#include <condition_variable>
+#include <cmath>
+#include <cstdint>
+#include <deque>
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+#include <map>
+#include <mutex>
+#include <set>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+#include <thread>
+#include <vector>
+namespace sr=slime_research;
+namespace fs=full_search_experiment;
+using Clock=std::chrono::steady_clock;
+#define CUDA(call) do{cudaError_t status_=(call);if(status_!=cudaSuccess)throw std::runtime_error(std::string(#call)+": "+cudaGetErrorString(status_));}while(0)
+struct Coord{int32_t x,z;};
+struct Job{uint32_t position,root;};
+struct Match{uint64_t seed;int32_t x,z;};
+struct Config{int32_t x=0,z=0;uint64_t radius=0;int size=5,positionBatch=32,jobBatch=2048,repeats=1;uint64_t maxResults=0;size_t cacheClasses=4096;std::string devices="0",output;bool dry=false,list=false;};
+uint64_t number(const std::string& s){size_t used=0;if(s.empty()||s[0]=='-')throw std::runtime_error("expected nonnegative integer: "+s);auto n=std::stoull(s,&used);if(used!=s.size())throw std::runtime_error("invalid integer: "+s);return n;}
+int32_t signed32(const std::string&s){size_t used=0;auto n=std::stoll(s,&used);if(used!=s.size()||n<INT32_MIN||n>INT32_MAX)throw std::runtime_error("invalid int32 coordinate");return int32_t(n);}
+uint64_t anchors(uint64_t r){return 1+2*r*(r+1);}
+uint64_t isqrt(uint64_t v){uint64_t r=std::sqrt(double(v));while((r+1)*(r+1)<=v)++r;while(r*r>v)--r;return r;}
+Coord coordinate(const Config& c,uint64_t index){
+    uint64_t top=(c.radius+1)*(c.radius+1);int64_t dx,dz;
+    if(index<top){uint64_t row=isqrt(index);dz=int64_t(row)-c.radius;dx=int64_t(index-row*row)-int64_t(row);}
+    else{uint64_t back=anchors(c.radius)-1-index,row=isqrt(back);dz=int64_t(c.radius)-row;dx=int64_t(row)-int64_t(back-row*row);}
+    return {int32_t(int64_t(c.x)+dx),int32_t(int64_t(c.z)+dz)};
+}
+uint64_t classKey(Coord p){return (uint64_t(uint32_t(p.x)&65535)<<17)|(uint32_t(p.z)&131071);}
+void terms(Coord p,int n,sr::CellTerm* out){uint64_t anchor=sr::coordinateTerm(p.x,p.z);for(int z=0;z<n;++z)for(int x=0;x<n;++x)out[z*n+x]={(sr::coordinateTerm(p.x+x,p.z+z)-anchor)&sr::kMask48,sr::kSlime};}
+__global__ void rootsKernel(const sr::CellTerm* cells,uint32_t positions,int count,Job* out,unsigned long long* total){
+    uint64_t t=uint64_t(blockIdx.x)*blockDim.x+threadIdx.x;if(t>=(uint64_t(positions)<<18))return;
+    uint32_t pos=t>>18,root=t&262143;const auto* p=cells+uint64_t(pos)*count;
+    for(int i=0;i<count;++i)if((((root+uint32_t(p[i].term))^uint32_t(sr::kCombinedXor))*uint32_t(sr::kJavaMultiplier)+11U)&(1U<<17))return;
+    out[atomicAdd(total,1ULL)]={pos,root};
+}
+__global__ void prepareKernel(const sr::CellTerm* cells,const Job* jobs,uint32_t n,int count,uint32_t* prepared){
+    uint32_t k=blockIdx.x*blockDim.x+threadIdx.x;if(k>=n)return;const auto job=jobs[k];
+    for(int i=0;i<count;++i){uint64_t a=job.root+cells[uint64_t(job.position)*count+i].term;
+        prepared[(uint64_t(k)*count+i)*2]=a>>18;
+        prepared[(uint64_t(k)*count+i)*2+1]=((((a^sr::kCombinedXor)&262143)*sr::kJavaMultiplier+11)>>17);
+    }
+}
+__host__ __device__ __forceinline__ uint32_t q(const uint32_t*p,int i,uint32_t m){return (p[2*i+1]+2U*uint32_t(sr::kJavaMultiplier)*((m+p[2*i])^uint32_t(sr::kCombinedXor>>18)))&0x7fffffffU;}
+__host__ __device__ uint32_t groupMask(const uint32_t*p,uint32_t base,int count){
+    uint32_t mask=65535,q0=q(p,0,base),qe=q(p,0,base+240);
+    for(int i=1;i<count&&mask;++i){uint32_t cut=(256-((base+p[2*i])&255)+15)>>4,lo=(1U<<cut)-1;
+        uint32_t r0=((q(p,i,base)-q0)&0x7fffffffU)%10,r1=((q(p,i,base+240)-qe)&0x7fffffffU)%10;
+        mask&=((r0==0||r0==8)?lo:0U)|((r1==0||r1==8)?(65535U^lo):0U);
+    }
+    return mask;
+}
+struct Emit{
+    Coord p;Match* out;uint64_t capacity;unsigned long long* total;
+    __device__ void operator()(uint64_t seed){uint64_t slot=atomicAdd(total,1ULL);if(slot<capacity)out[slot]={(seed-sr::coordinateTerm(p.x,p.z))&sr::kMask48,p.x,p.z};}
+};
+__global__ void searchKernel(const sr::CellTerm*cells,const Coord*coords,const Job*jobs,uint32_t n,int count,const uint32_t*prepared,Match*out,uint64_t capacity,unsigned long long*total){
+    uint64_t index=uint64_t(blockIdx.x)*blockDim.x+threadIdx.x;if(index>=(uint64_t(n)<<13))return;
+    uint32_t g=index&8191,k=index>>13,base=((g>>4)<<8)+(g&15);const uint32_t* p=prepared+uint64_t(k)*count*2;
+    uint32_t mask=groupMask(p,base,count);
+    Job job=jobs[k];Emit emit{coords[job.position],out,capacity,total};
+    while(mask){uint32_t j=__ffs(mask)-1;mask&=mask-1;fs::solveIntervals(job.root+(uint64_t(base+16*j)<<18),cells+uint64_t(job.position)*count,count,emit);}
+}
+template<class T>struct Buffer{T* p=nullptr;explicit Buffer(size_t n){if(n)CUDA(cudaMalloc(&p,n*sizeof(T)));}~Buffer(){if(p)cudaFree(p);}Buffer(const Buffer&)=delete;};
+struct Metrics{int device;std::string name;uint64_t positions=0,classes=0,cacheHits=0,jobs=0,matches=0;double setupMs=0,scanMs=0,rootMs=0,searchMs=0;Clock::time_point start,end;std::vector<Match> sample;};
+struct Shared{std::atomic<uint64_t> next{0},done{0};std::atomic<int> finished{0};int ready=0,workerCount=0;std::condition_variable cv;std::atomic<bool> cancel{false};std::mutex mutex;std::string error;};
+void worker(const Config& c,int device,Shared& shared,Metrics& result){try{
+    auto setup=Clock::now();CUDA(cudaSetDevice(device));cudaDeviceProp prop{};CUDA(cudaGetDeviceProperties(&prop,device));result.device=device;result.name=prop.name;
+    int count=c.size*c.size;size_t rootsCapacity=size_t(c.positionBatch)*(1U<<17);
+    Buffer<sr::CellTerm> dcells(size_t(c.positionBatch)*count);Buffer<Coord>dcoords(c.positionBatch);Buffer<Job>droots(rootsCapacity),djobs(c.jobBatch);
+    Buffer<uint32_t>prep(size_t(c.jobBatch)*count*2);Buffer<unsigned long long>total(1);Buffer<Match>matches(c.maxResults);
+    std::map<uint64_t,std::vector<uint32_t>> cache;std::deque<uint64_t> fifo;size_t cachedRoots=0;
+    result.setupMs=std::chrono::duration<double,std::milli>(Clock::now()-setup).count();{
+        std::unique_lock<std::mutex> lock(shared.mutex);++shared.ready;shared.cv.notify_all();
+        shared.cv.wait(lock,[&](){return shared.ready==shared.workerCount||shared.cancel.load();});
+        if(shared.cancel)return;
+    }
+    result.start=Clock::now();
+    const uint64_t end=anchors(c.radius);
+    while(!shared.cancel){uint64_t first=shared.next.fetch_add(c.positionBatch);if(first>=end)break;
+        uint32_t num=std::min<uint64_t>(c.positionBatch,end-first);std::vector<Coord> coords;std::vector<uint64_t> keys;
+        std::map<uint64_t,std::vector<uint32_t>> lists;std::vector<Coord> missing;std::vector<uint64_t> missingKeys;
+        for(uint32_t i=0;i<num;++i){Coord p=coordinate(c,first+i);coords.push_back(p);uint64_t key=classKey(p);keys.push_back(key);
+            if(lists.count(key)){++result.cacheHits;continue;}auto it=cache.find(key);
+            if(it!=cache.end()){lists[key]=it->second;++result.cacheHits;}else{lists[key]={};missing.push_back(p);missingKeys.push_back(key);}
+        }
+        if(!missing.empty()){
+            auto start=Clock::now();std::vector<sr::CellTerm> cells(missing.size()*count);for(size_t i=0;i<missing.size();++i)terms(missing[i],c.size,cells.data()+i*count);
+            CUDA(cudaMemcpy(dcells.p,cells.data(),cells.size()*sizeof(sr::CellTerm),cudaMemcpyHostToDevice));CUDA(cudaMemset(total.p,0,sizeof(*total.p)));
+            rootsKernel<<<unsigned((uint64_t(missing.size())<<18)/256),256>>>(dcells.p,missing.size(),count,droots.p,total.p);CUDA(cudaGetLastError());
+            unsigned long long found;CUDA(cudaMemcpy(&found,total.p,sizeof(found),cudaMemcpyDeviceToHost));if(found>rootsCapacity)throw std::runtime_error("root capacity invariant failed");
+            std::vector<Job> roots(found);if(found)CUDA(cudaMemcpy(roots.data(),droots.p,found*sizeof(Job),cudaMemcpyDeviceToHost));
+            for(auto job:roots)lists[missingKeys[job.position]].push_back(job.root);
+            result.classes+=missing.size();
+            for(auto key:missingKeys){auto& list=lists[key];std::sort(list.begin(),list.end());if(c.cacheClasses){
+                while(!fifo.empty()&&(cache.size()>=c.cacheClasses||cachedRoots+list.size()>(1U<<20))){auto old=fifo.front();fifo.pop_front();cachedRoots-=cache.at(old).size();cache.erase(old);}
+                cache[key]=list;cachedRoots+=list.size();fifo.push_back(key);
+            }}
+            result.rootMs+=std::chrono::duration<double,std::milli>(Clock::now()-start).count();
+        }
+        // Construct concrete positions only for classes with roots.
+        std::vector<Coord> active;std::vector<sr::CellTerm> cells;std::vector<Job> jobs;jobs.reserve(c.jobBatch);
+        for(uint32_t i=0;i<num;++i)if(!lists.at(keys[i]).empty())active.push_back(coords[i]);
+        if(!active.empty()){
+            cells.resize(active.size()*count);for(size_t i=0;i<active.size();++i)terms(active[i],c.size,cells.data()+i*count);
+            CUDA(cudaMemcpy(dcells.p,cells.data(),cells.size()*sizeof(sr::CellTerm),cudaMemcpyHostToDevice));CUDA(cudaMemcpy(dcoords.p,active.data(),active.size()*sizeof(Coord),cudaMemcpyHostToDevice));
+            auto search=[&](){if(jobs.empty()||shared.cancel)return;auto start=Clock::now();
+                CUDA(cudaMemcpy(djobs.p,jobs.data(),jobs.size()*sizeof(Job),cudaMemcpyHostToDevice));CUDA(cudaMemset(total.p,0,sizeof(*total.p)));
+                prepareKernel<<<unsigned((jobs.size()+127)/128),128>>>(dcells.p,djobs.p,jobs.size(),count,prep.p);
+                searchKernel<<<unsigned((uint64_t(jobs.size())<<13)/256),256>>>(dcells.p,dcoords.p,djobs.p,jobs.size(),count,prep.p,matches.p,c.maxResults,total.p);CUDA(cudaGetLastError());
+                unsigned long long found;CUDA(cudaMemcpy(&found,total.p,sizeof(found),cudaMemcpyDeviceToHost));if(found>UINT64_MAX-result.matches)throw std::runtime_error("match count exceeds uint64");result.matches+=found;result.jobs+=jobs.size();
+                size_t take=std::min<uint64_t>(found,c.maxResults-result.sample.size());if(take){size_t old=result.sample.size();result.sample.resize(old+take);CUDA(cudaMemcpy(result.sample.data()+old,matches.p,take*sizeof(Match),cudaMemcpyDeviceToHost));}
+                result.searchMs+=std::chrono::duration<double,std::milli>(Clock::now()-start).count();jobs.clear();};
+            for(uint32_t i=0;i<active.size()&&!shared.cancel;++i)for(uint32_t root:lists.at(classKey(active[i]))){jobs.push_back({i,root});if(jobs.size()==size_t(c.jobBatch))search();}
+            search();
+        }
+        result.positions+=num;shared.done.fetch_add(num);
+    }
+    result.end=Clock::now();result.scanMs=std::chrono::duration<double,std::milli>(result.end-result.start).count();
+}catch(const std::exception& e){shared.cancel=true;std::lock_guard<std::mutex> lock(shared.mutex);shared.error="GPU "+std::to_string(device)+": "+e.what();shared.cv.notify_all();}}
+std::string escape(const std::string&s){std::ostringstream o;for(unsigned char ch:s){if(ch=='"'||ch=='\\')o<<'\\'<<char(ch);else if(ch<32)o<<"\\u"<<std::hex<<std::setw(4)<<std::setfill('0')<<int(ch)<<std::dec;else o<<char(ch);}return o.str();}
+void help(){std::cout<<R"(slime-bench --radius R [options]
+Search all effective 48-bit seeds for solid n x n slime chunk squares.
+Ordinary candidate branch only; rejection-tail supplement is NOT included.
+
+  --x X --z Z          Center chunk coordinates (default 0,0)
+  --radius R           Manhattan radius in chunks (required, >=0)
+  --size N             Square side, 1..6 (default 5)
+  --devices 0,1,...     CUDA device indices, or all (default 0)
+  --repeat N           Complete passes; fresh root cache each pass (default 1)
+  --position-batch N   Anchors per scheduled batch, 1..256 (default 32)
+  --job-batch N        Low18-root jobs per GPU launch, 1..65536 (default 2048)
+  --cache-classes N    Per-GPU bounded FIFO root cache entries (default 4096)
+  --max-results N      Save up to N example matches; count ALL (default 0)
+  --output FILE        Also save the complete JSON report (file must not exist)
+  --dry-run            Print scope JSON without using CUDA
+  --list-devices       List visible CUDA devices
+  --help               Show this help
+
+Scope: abs(anchor_x-X)+abs(anchor_z-Z)<=R, including boundary.
+Anchor is the minimum-x/minimum-z corner; the square may extend outside
+this diamond. Coordinates are not clipped to Minecraft's world border.
+Small squares can have enormous numbers of matching seeds and take very long.
+)";}
+int main(int argc,char**argv){try{
+    Config c;bool radiusSet=false;for(int i=1;i<argc;++i){std::string a=argv[i];if(a=="--help"||a=="-h"){help();return 0;}if(a=="--dry-run"){c.dry=true;continue;}if(a=="--list-devices"){c.list=true;continue;}
+        if(i+1>=argc)throw std::runtime_error("missing value for "+a);std::string v=argv[++i];
+        auto bounded=[&](uint64_t lo,uint64_t hi){auto n=number(v);if(n<lo||n>hi)throw std::runtime_error("out of range: "+a);return n;};
+        if(a=="--x")c.x=signed32(v);else if(a=="--z")c.z=signed32(v);else if(a=="--radius"){c.radius=bounded(0,INT32_MAX);radiusSet=true;}
+        else if(a=="--size")c.size=bounded(1,6);else if(a=="--devices")c.devices=v;else if(a=="--repeat")c.repeats=bounded(1,1000);
+        else if(a=="--position-batch")c.positionBatch=bounded(1,256);else if(a=="--job-batch")c.jobBatch=bounded(1,65536);
+        else if(a=="--cache-classes")c.cacheClasses=bounded(0,1000000);else if(a=="--max-results")c.maxResults=bounded(0,1000000);
+        else if(a=="--output")c.output=v;else throw std::runtime_error("unknown option: "+a);
+    }
+    if(c.list){int n;CUDA(cudaGetDeviceCount(&n));for(int i=0;i<n;++i){cudaDeviceProp p{};CUDA(cudaGetDeviceProperties(&p,i));std::cout<<i<<"\t"<<p.name<<"\tsm_"<<p.major<<p.minor<<"\t"<<p.totalGlobalMem/(1ULL<<20)<<" MiB\n";}return 0;}
+    if(!radiusSet)throw std::runtime_error("--radius is required");
+    for(int64_t center:{int64_t(c.x),int64_t(c.z)})if(center-int64_t(c.radius)<INT32_MIN||center+int64_t(c.radius)+c.size-1>INT32_MAX)throw std::runtime_error("coordinates including square must fit int32");
+    if(!c.output.empty()&&std::ifstream(c.output).good())throw std::runtime_error("output already exists: "+c.output);
+    std::ostringstream report;report<<std::setprecision(12)<<"{\"schema_version\":1,\"mode\":\"ordinary_all_effective_seeds\",\"rejection_tail_included\":false,\"center\":["<<c.x<<","<<c.z<<"],\"radius_chunks\":"<<c.radius<<",\"square_size\":"<<c.size<<",\"anchors\":"<<anchors(c.radius)<<",\"effective_seed_bits\":48,\"dry_run\":"<<(c.dry?"true":"false");
+    if(c.dry){auto first=coordinate(c,0),last=coordinate(c,anchors(c.radius)-1);report<<",\"first_anchor\":["<<first.x<<","<<first.z<<"],\"last_anchor\":["<<last.x<<","<<last.z<<"]}";}
+    else{
+        int available;CUDA(cudaGetDeviceCount(&available));std::vector<int> devices;if(c.devices=="all"){for(int i=0;i<available;++i)devices.push_back(i);}else{
+            std::stringstream ss(c.devices);std::string v;while(std::getline(ss,v,',')){auto id=number(v);if(id>=uint64_t(available))throw std::runtime_error("CUDA device unavailable: "+v);devices.push_back(id);}
+        }
+        if(devices.empty()||std::set<int>(devices.begin(),devices.end()).size()!=devices.size())throw std::runtime_error("device list must be nonempty and unique");
+        report<<",\"position_batch\":"<<c.positionBatch<<",\"job_batch\":"<<c.jobBatch<<",\"cache_classes\":"<<c.cacheClasses<<",\"cuda_runtime_version\":"<<CUDART_VERSION<<",\"rounds\":[";
+        std::vector<double> timings;uint64_t expected=0;
+        for(int round=0;round<c.repeats;++round){Shared shared;shared.workerCount=devices.size();std::vector<Metrics> metrics(devices.size());std::vector<std::thread> workers;auto start=Clock::now();
+            for(size_t i=0;i<devices.size();++i)workers.emplace_back([&,i](){worker(c,devices[i],shared,metrics[i]);++shared.finished;shared.cv.notify_all();});
+            {std::unique_lock<std::mutex> lock(shared.mutex);while(shared.finished<int(devices.size())){
+                if(!shared.cv.wait_for(lock,std::chrono::seconds(5),[&](){return shared.finished==int(devices.size());}))
+                    std::cerr<<"progress: "<<shared.done<<"/"<<anchors(c.radius)<<" anchors (round "<<round+1<<")\n";
+            }}
+            for(auto&t:workers)t.join();
+            if(!shared.error.empty())throw std::runtime_error(shared.error);
+            double elapsed=std::chrono::duration<double,std::milli>(Clock::now()-start).count();auto first=metrics[0].start,last=metrics[0].end;
+            uint64_t visited=0,jobs=0,matches=0;std::vector<Match> sample;
+            for(auto&m:metrics){first=std::min(first,m.start);last=std::max(last,m.end);visited+=m.positions;jobs+=m.jobs;if(m.matches>UINT64_MAX-matches)throw std::runtime_error("total match count exceeds uint64");matches+=m.matches;for(auto hit:m.sample)if(sample.size()<c.maxResults)sample.push_back(hit);}
+            if(visited!=anchors(c.radius))throw std::runtime_error("incomplete anchor coverage");if(round&&matches!=expected)throw std::runtime_error("repeat match count mismatch");expected=matches;
+            double scan=std::chrono::duration<double,std::milli>(last-first).count();timings.push_back(scan);
+            if(round)report<<",";report<<"{\"round\":"<<round+1<<",\"elapsed_ms\":"<<elapsed<<",\"scan_wall_ms\":"<<scan<<",\"anchors_per_second\":"<<visited/(scan/1000)<<",\"root_jobs\":"<<jobs<<",\"matches\":"<<matches<<",\"sample_truncated\":"<<(matches>sample.size()?"true":"false")<<",\"sample\":[";
+            for(size_t i=0;i<sample.size();++i){if(i)report<<",";auto m=sample[i];report<<"["<<m.seed<<","<<m.x<<","<<m.z<<"]";}report<<"],\"gpus\":[";
+            for(size_t i=0;i<metrics.size();++i){if(i)report<<",";auto&m=metrics[i];report<<"{\"device\":"<<m.device<<",\"name\":\""<<escape(m.name)<<"\",\"anchors\":"<<m.positions<<",\"root_classes_generated\":"<<m.classes<<",\"cache_hits\":"<<m.cacheHits<<",\"root_jobs\":"<<m.jobs<<",\"matches\":"<<m.matches<<",\"setup_ms\":"<<m.setupMs<<",\"scan_ms\":"<<m.scanMs<<",\"root_stage_ms\":"<<m.rootMs<<",\"search_stage_ms\":"<<m.searchMs<<"}";}report<<"]}";
+            std::cerr<<"round "<<round+1<<"/"<<c.repeats<<": "<<visited<<" anchors, "<<matches<<" matches, scan "<<scan<<" ms, elapsed "<<elapsed<<" ms\n";
+        }
+        std::sort(timings.begin(),timings.end());double median=timings[timings.size()/2];if(timings.size()%2==0)median=(median+timings[timings.size()/2-1])/2;
+        report<<"],\"median_scan_wall_ms\":"<<median<<"}";
+    }
+    if(!c.output.empty()){std::ofstream file(c.output);if(!file)throw std::runtime_error("cannot write report: "+c.output);file<<report.str()<<"\n";file.close();if(!file)throw std::runtime_error("failed writing report");}
+    std::cout<<report.str()<<"\n";return 0;
+}catch(const std::exception&e){std::cerr<<"error: "<<e.what()<<"\n";return 1;}}

bench/test_cli.py

@@ -0,0 +1,65 @@
+#!/usr/bin/env python3
+import argparse,json,subprocess,tempfile
+from pathlib import Path
+HERE=Path(__file__).resolve().parent
+BIN=HERE/'bin/slime-bench'
+def run(*args,ok=True):
+    p=subprocess.run([str(BIN),*map(str,args)],capture_output=True,text=True,timeout=180)
+    if ok:
+        assert p.returncode==0,p.stderr
+        return json.loads(p.stdout)
+    assert p.returncode!=0,p.stdout
+
+def term(x,z):
+    def i32(v):return ((v+2**31)%2**32)-2**31
+    return i32(x*x*4987142)+i32(x*5947611)+i32(z*z)*4392871+i32(z*389711)
+def slime(seed,x,z):
+    s=((seed+term(x,z))^987234911^0x5DEECE66D)&((1<<48)-1)
+    while True:
+        s=(s*0x5DEECE66D+11)&((1<<48)-1);q=s>>17
+        if q<2147483640:return q%10==0
+
+def main():
+    ap=argparse.ArgumentParser();ap.add_argument('--gpu',action='store_true');args=ap.parse_args()
+    for radius in [0,1,2,31,1000000]:
+        r=run('--x',-123,'--z',456,'--radius',radius,'--size',6,'--dry-run')
+        assert r['anchors']==1+2*radius*(radius+1)
+        assert r['first_anchor']==[-123,456-radius] and r['last_anchor']==[-123,456+radius]
+    for bad in [[],['--radius','-1'],['--radius','1','--size','0'],['--radius','1','--size','7'],['--radius','abc'],['--radius','0','--x','2147483647'],['--radius','1','--repeat','0'],['--radius','1','--position-batch','0']]:run(*bad,'--dry-run',ok=False)
+    summary={'cli_scope_and_errors':True}
+    if args.gpu:
+        common=['--x',51069,'--z',419124]
+        reference=run(*common,'--radius',0,'--size',5,'--repeat',2,'--max-results',100)
+        expected={61692151016558,15959335054446}
+        assert all(r['matches']==24 and r['root_jobs']==84 for r in reference['rounds'])
+        samples=[set(map(tuple,r['sample'])) for r in reference['rounds']]
+        assert samples[0]==samples[1] and expected <= {s for s,x,z in samples[0]}
+        sliced=run(*common,'--radius',0,'--size',5,'--job-batch',17,'--max-results',100,'--cache-classes',0)
+        assert set(map(tuple,sliced['rounds'][0]['sample']))==samples[0]
+        partial=run(*common,'--radius',0,'--size',5,'--max-results',3)
+        assert partial['rounds'][0]['matches']==24 and len(partial['rounds'][0]['sample'])==3 and partial['rounds'][0]['sample_truncated']
+        for n in [4,6]:
+            r=run(*common,'--radius',0,'--size',n,'--max-results',1000)
+            for seed,x,z in r['rounds'][0]['sample']:assert all(slime(seed,x+dx,z+dz) for dx in range(n) for dz in range(n))
+            if n>=5:
+                wanted={t for t in samples[0] if all(slime(t[0],t[1]+dx,t[2]+dz) for dx in range(n) for dz in range(n))}
+                assert set(map(tuple,r['rounds'][0]['sample']))==wanted
+            summary[f'size_{n}']=r['rounds'][0]['matches']
+        diamond=run(*common,'--radius',2,'--size',5,'--max-results',1000,'--position-batch',3)
+        split=[]
+        for dz in range(-2,3):
+            for dx in range(-2,3):
+                if abs(dx)+abs(dz)>2:continue
+                r=run('--x',51069+dx,'--z',419124+dz,'--radius',0,'--size',5,'--max-results',1000)
+                split.extend(map(tuple,r['rounds'][0]['sample']))
+        assert sorted(split)==sorted(map(tuple,diamond['rounds'][0]['sample']))
+        assert diamond['anchors']==13
+        run('--radius',0,'--devices','0,0',ok=False)
+        with tempfile.TemporaryDirectory() as tmp:
+            out=Path(tmp)/'report.json'
+            r=run(*common,'--radius',0,'--output',out)
+            assert json.loads(out.read_text())==r
+            run(*common,'--radius',0,'--output',out,ok=False)
+        summary['gpu_known_matches']=24;summary['diamond_vs_individual_anchors']=True
+    print(json.dumps(summary,indent=2))
+if __name__=='__main__':main()

bench/test_math.cu

@@ -0,0 +1,25 @@
+#define main cli_entry_for_tests
+#include "slime_bench.cu"
+#undef main
+#include <random>
+void check(bool ok,const char*msg){if(!ok)throw std::runtime_error(msg);}
+int main(){
+    uint64_t points=0,groups=0;std::mt19937_64 rng(123456);
+    for(int r=0;r<=64;++r){Config c;c.x=-123;c.z=98;c.radius=r;std::set<std::pair<int32_t,int32_t>> actual;
+        for(uint64_t i=0;i<anchors(r);++i){auto p=coordinate(c,i);actual.emplace(p.x,p.z);}
+        check(actual.size()==anchors(r),"duplicate anchor");
+        for(int dz=-r;dz<=r;++dz)for(int dx=-r;dx<=r;++dx)check(bool(actual.count({c.x+dx,c.z+dz}))==(std::abs(dx)+std::abs(dz)<=r),"diamond coverage mismatch");points+=actual.size();
+    }
+    Config large;large.radius=1000000;
+    for(int t=0;t<100000;++t){auto p=coordinate(large,rng()%anchors(large.radius));check(std::abs(int64_t(p.x))+std::abs(int64_t(p.z))<=int64_t(large.radius),"large diamond mismatch");}
+    for(int n:{1,2,3,4,5,6})for(int t=0;t<512;++t){
+        Coord coord{int32_t(rng()%3750000)-1875000,int32_t(rng()%3750000)-1875000};int count=n*n;
+        std::vector<sr::CellTerm> cells(count);terms(coord,n,cells.data());uint32_t root=rng()&262143,base=(rng()&131071)&~240U;
+        std::vector<uint32_t> prepared(count*2);
+        for(int i=0;i<count;++i){uint64_t a=root+cells[i].term;prepared[2*i]=a>>18;prepared[2*i+1]=((((a^sr::kCombinedXor)&262143)*sr::kJavaMultiplier+11)>>17);
+            for(int j=0;j<16;++j)check(q(prepared.data(),i,base+16*j)==(sr::nextStateFromTerm(root+(uint64_t(base+16*j)<<18),cells[i].term)>>17),"prepared q mismatch");}
+        uint32_t expected=0;for(int j=0;j<16;++j)if(fs::differenceFilter(root+(uint64_t(base+16*j)<<18),cells.data(),count))expected|=1U<<j;
+        check(groupMask(prepared.data(),base,count)==expected,"variable-size group mask mismatch");++groups;
+    }
+    std::cout<<"{\"diamond_points\":"<<points<<",\"variable_size_groups\":"<<groups<<",\"passed\":true}\n";
+}

include/README.md

@@ -0,0 +1,8 @@
+# 共用算法头文件
+
+- `slime/solver_core.hpp`：Java 整数回绕、48 位随机数、精确史莱姆判定、低位约束和原高位求解器。
+- `slime/upper_interval.hpp`：输出差必要条件与高位区间求交，用于实心图案的普通候选分支。
+
+CLI 直接依赖这里。`research/slime_solver_core.hpp` 和原实验目录的 `interval_solver.hpp` 仅转发包含，以便已有研究程序、教学代码和历史复现命令继续工作。
+
+本次只迁移实现位置，保留现有命名空间及算法语义。精确验证支持 Java 重试，不代表普通候选生成已经包含拒绝尾部补扫。

include/slime/solver_core.hpp

@@ -0,0 +1,286 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+
+#if defined(__CUDACC__)
+#define SLIME_HD __host__ __device__
+#define SLIME_INLINE __forceinline__
+#else
+#define SLIME_HD
+#define SLIME_INLINE inline
+#endif
+
+namespace slime_research {
+
+constexpr uint64_t kJavaMultiplier = 0x5DEECE66DULL;
+constexpr uint64_t kJavaAddend = 0xBULL;
+constexpr uint64_t kMask48 = (1ULL << 48) - 1ULL;
+constexpr uint64_t kSlimeXor = 987234911ULL;
+constexpr uint64_t kCombinedXor = kSlimeXor ^ kJavaMultiplier;
+constexpr int kLowerBits = 35;
+constexpr int kUpperBits = 13;
+constexpr uint32_t kUpperSize = 1U << kUpperBits;
+constexpr uint32_t kUpperMask = kUpperSize - 1U;
+constexpr uint32_t kHalfDeltaInverse = 4965U;
+constexpr int kMaxPatternCells = 64;
+constexpr uint32_t kX2Coeff = 4987142U;
+constexpr uint32_t kXCoeff = 5947611U;
+constexpr uint32_t kZ2Coeff = 4392871U;
+constexpr uint32_t kZCoeff = 389711U;
+
+enum Requirement : uint8_t {
+    kNonSlime = 0,
+    kSlime = 1,
+};
+
+struct CellTerm {
+    uint64_t term;
+    uint8_t requirement;
+};
+
+struct UpperEvent {
+    uint16_t threshold;
+    uint8_t before;
+};
+
+SLIME_HD SLIME_INLINE uint64_t javaIntAsLongBits(uint32_t value) {
+    return static_cast<uint64_t>(static_cast<int64_t>(static_cast<int32_t>(value)));
+}
+
+SLIME_HD SLIME_INLINE uint64_t coordinateTerm(int32_t chunkX, int32_t chunkZ) {
+    const uint32_t x = static_cast<uint32_t>(chunkX);
+    const uint32_t z = static_cast<uint32_t>(chunkZ);
+    return javaIntAsLongBits(x * x * kX2Coeff) +
+           javaIntAsLongBits(x * kXCoeff) +
+           javaIntAsLongBits(z * z) * kZ2Coeff +
+           javaIntAsLongBits(z * kZCoeff);
+}
+
+SLIME_HD SLIME_INLINE uint64_t coordinateXTerm(int32_t chunkX) {
+    const uint32_t x = static_cast<uint32_t>(chunkX);
+    return javaIntAsLongBits(x * x * kX2Coeff) +
+           javaIntAsLongBits(x * kXCoeff);
+}
+
+SLIME_HD SLIME_INLINE uint64_t coordinateZTerm(int32_t chunkZ) {
+    const uint32_t z = static_cast<uint32_t>(chunkZ);
+    return javaIntAsLongBits(z * z) * kZ2Coeff +
+           javaIntAsLongBits(z * kZCoeff);
+}
+
+SLIME_HD SLIME_INLINE uint64_t nextStateFromTerm(uint64_t seed48, uint64_t term) {
+    uint64_t state = ((seed48 + term) ^ kCombinedXor) & kMask48;
+    return (state * kJavaMultiplier + kJavaAddend) & kMask48;
+}
+
+SLIME_HD SLIME_INLINE bool exactIsSlimeFromTerm(uint64_t seed48, uint64_t term) {
+    uint64_t state = ((seed48 + term) ^ kCombinedXor) & kMask48;
+    for (;;) {
+        state = (state * kJavaMultiplier + kJavaAddend) & kMask48;
+        const uint32_t bits = static_cast<uint32_t>(state >> 17);
+        const uint32_t value = bits % 10U;
+        if (bits - value + 9U <= 0x7FFFFFFFU) {
+            return value == 0U;
+        }
+    }
+}
+
+SLIME_HD SLIME_INLINE bool verifyPattern(
+    uint64_t seed48,
+    const CellTerm* cells,
+    int cellCount
+) {
+    for (int i = 0; i < cellCount; ++i) {
+        const bool slime = exactIsSlimeFromTerm(seed48, cells[i].term);
+        if (slime != (cells[i].requirement == kSlime)) {
+            return false;
+        }
+    }
+    return true;
+}
+
+SLIME_HD SLIME_INLINE bool partialPatternPossible(
+    uint64_t seedPrefix,
+    int knownBits,
+    const CellTerm* cells,
+    int cellCount
+) {
+    const uint64_t mask = (1ULL << knownBits) - 1ULL;
+    seedPrefix &= mask;
+    for (int i = 0; i < cellCount; ++i) {
+        uint64_t state = (((seedPrefix + cells[i].term) ^ kCombinedXor) *
+                          kJavaMultiplier + kJavaAddend) & mask;
+        const uint64_t knownOutput = state >> 17;
+        if (cells[i].requirement == kSlime && (knownOutput & 1ULL) != 0ULL) {
+            return false;
+        }
+    }
+    return true;
+}
+
+SLIME_HD SLIME_INLINE uint32_t firstWithResidue(
+    uint32_t start,
+    uint32_t residue
+) {
+    return start + ((residue + 5U - (start % 5U)) % 5U);
+}
+
+SLIME_HD SLIME_INLINE bool upper13HasCandidate(
+    uint64_t lower35,
+    const CellTerm* cells,
+    int cellCount
+) {
+    UpperEvent events[kMaxPatternCells];
+    uint16_t residueCounts[5] = {0, 0, 0, 0, 0};
+    int eventCount = 0;
+    lower35 &= (1ULL << kLowerBits) - 1ULL;
+    for (int i = 0; i < cellCount; ++i) {
+        if (cells[i].requirement != kSlime) continue;
+        const uint32_t output = static_cast<uint32_t>(
+            nextStateFromTerm(lower35, cells[i].term) >> 17);
+        const uint32_t low17 = output & ((1U << 17) - 1U);
+        if ((low17 & 1U) != 0U) return false;
+        const uint32_t upper14 = output >> 17;
+        const uint32_t parity = upper14 & 1U;
+        const uint32_t v0 = (upper14 - parity) >> 1;
+        const uint32_t wanted = (3U * ((5U - ((low17 >> 1) % 5U) +
+                                              5U - parity) % 5U)) % 5U;
+        const uint32_t before = (wanted + 5U - (v0 % 5U)) % 5U;
+        const UpperEvent event{static_cast<uint16_t>(kUpperSize - v0),
+                               static_cast<uint8_t>(before)};
+        int insert = eventCount;
+        while (insert > 0 && events[insert - 1].threshold > event.threshold) {
+            events[insert] = events[insert - 1];
+            --insert;
+        }
+        events[insert] = event;
+        ++eventCount;
+        ++residueCounts[before];
+    }
+    if (eventCount == 0) return true;
+    uint32_t start = 0;
+    int index = 0;
+    for (;;) {
+        const uint32_t stop = index < eventCount ? events[index].threshold : kUpperSize;
+        for (uint32_t residue = 0; residue < 5; ++residue) {
+            if (residueCounts[residue] == eventCount &&
+                firstWithResidue(start, residue) < stop) {
+                return true;
+            }
+        }
+        if (index == eventCount) return false;
+        const uint32_t threshold = events[index].threshold;
+        while (index < eventCount && events[index].threshold == threshold) {
+            const uint32_t before = events[index].before;
+            --residueCounts[before];
+            ++residueCounts[(before + 2U) % 5U];
+            ++index;
+        }
+        start = threshold;
+    }
+}
+
+template <typename Emit>
+SLIME_HD SLIME_INLINE uint32_t solveUpper13(
+    uint64_t lower35,
+    const CellTerm* cells,
+    int cellCount,
+    Emit& emit
+) {
+    UpperEvent events[kMaxPatternCells];
+    uint16_t residueCounts[5] = {0, 0, 0, 0, 0};
+    int eventCount = 0;
+    lower35 &= (1ULL << kLowerBits) - 1ULL;
+
+    for (int i = 0; i < cellCount; ++i) {
+        if (cells[i].requirement != kSlime) {
+            continue;
+        }
+
+        const uint32_t output = static_cast<uint32_t>(
+            nextStateFromTerm(lower35, cells[i].term) >> 17
+        );
+        const uint32_t low17 = output & ((1U << 17) - 1U);
+        if ((low17 & 1U) != 0U) {
+            return 0;
+        }
+
+        const uint32_t upper14 = output >> 17;
+        const uint32_t parity = upper14 & 1U;
+        const uint32_t v0 = (upper14 - parity) >> 1;
+        const uint32_t wanted = (3U * ((5U - ((low17 >> 1) % 5U) +
+                                              5U - parity) % 5U)) % 5U;
+        const uint32_t before = (wanted + 5U - (v0 % 5U)) % 5U;
+        UpperEvent event{
+            static_cast<uint16_t>(kUpperSize - v0),
+            static_cast<uint8_t>(before),
+        };
+
+        int insert = eventCount;
+        while (insert > 0 && events[insert - 1].threshold > event.threshold) {
+            events[insert] = events[insert - 1];
+            --insert;
+        }
+        events[insert] = event;
+        ++eventCount;
+        ++residueCounts[before];
+    }
+
+    uint32_t emitted = 0;
+    if (eventCount == 0) {
+        for (uint32_t n = 0; n < kUpperSize; ++n) {
+            const uint32_t m = (kHalfDeltaInverse * n) & kUpperMask;
+            const uint64_t seed = lower35 + (uint64_t(m) << kLowerBits);
+            if (verifyPattern(seed, cells, cellCount)) {
+                emit(seed);
+                ++emitted;
+            }
+        }
+        return emitted;
+    }
+
+    uint32_t start = 0;
+    int index = 0;
+    for (;;) {
+        const uint32_t stop = index < eventCount ? events[index].threshold : kUpperSize;
+        int commonResidue = -1;
+        for (int residue = 0; residue < 5; ++residue) {
+            if (residueCounts[residue] == eventCount) {
+                commonResidue = residue;
+                break;
+            }
+        }
+
+        if (commonResidue >= 0 && start < stop) {
+            for (uint32_t n = firstWithResidue(start, uint32_t(commonResidue));
+                 n < stop;
+                 n += 5U) {
+                const uint32_t m = (kHalfDeltaInverse * n) & kUpperMask;
+                const uint64_t seed = lower35 + (uint64_t(m) << kLowerBits);
+                if (verifyPattern(seed, cells, cellCount)) {
+                    emit(seed);
+                    ++emitted;
+                }
+            }
+        }
+
+        if (index == eventCount) {
+            break;
+        }
+        const uint32_t threshold = events[index].threshold;
+        while (index < eventCount && events[index].threshold == threshold) {
+            const uint32_t before = events[index].before;
+            --residueCounts[before];
+            ++residueCounts[(before + 2U) % 5U];
+            ++index;
+        }
+        start = threshold;
+    }
+    return emitted;
+}
+
+}  // namespace slime_research
+
+#undef SLIME_HD
+#undef SLIME_INLINE

include/slime/upper_interval.hpp

@@ -0,0 +1,71 @@
+#pragma once
+
+#include "solver_core.hpp"
+
+#if defined(__CUDACC__)
+#define FS_HD __host__ __device__
+#define FS_INLINE __forceinline__
+#else
+#define FS_HD
+#define FS_INLINE inline
+#endif
+
+namespace full_search_experiment {
+namespace sr = slime_research;
+
+// Solid patterns only. Each surviving residue has a single interval.
+// Keeping just the anchor's two residues makes intersection constant-space.
+template<class Emit>
+FS_HD FS_INLINE void solveIntervals(uint64_t low, const sr::CellTerm* cells,
+                                    int count, Emit& emit) {
+    low &= (1ULL << 35) - 1;
+    uint32_t q = uint32_t(sr::nextStateFromTerm(low, cells[0].term) >> 17);
+    if (q & 1U) return;
+    uint32_t b = q >> 18;
+    uint32_t r = q & ((1U << 18) - 1);
+    uint32_t threshold = 8192 - b;
+    uint32_t residue[2] = {(r % 5 + 5 - b % 5) % 5, 0};
+    residue[1] = (residue[0] + 2) % 5;
+    uint32_t lo[2] = {0, threshold}, hi[2] = {threshold, 8192};
+    unsigned active = 0;
+    for (int k = 0; k < 2; ++k)
+        if (sr::firstWithResidue(lo[k], residue[k]) < hi[k]) active |= 1U << k;
+    for (int i = 1; i < count && active; ++i) {
+        q = uint32_t(sr::nextStateFromTerm(low, cells[i].term) >> 17);
+        if (q & 1U) return;
+        b = q >> 18;
+        r = q & ((1U << 18) - 1);
+        threshold = 8192 - b;
+        const uint32_t before = (r % 5 + 5 - b % 5) % 5;
+        const uint32_t after = (before + 2) % 5;
+        for (int k = 0; k < 2; ++k) if (active & (1U << k)) {
+            if (residue[k] == before) hi[k] = hi[k] < threshold ? hi[k] : threshold;
+            else if (residue[k] == after) lo[k] = lo[k] > threshold ? lo[k] : threshold;
+            else { active &= ~(1U << k); continue; }
+            if (sr::firstWithResidue(lo[k], residue[k]) >= hi[k]) active &= ~(1U << k);
+        }
+    }
+    for (int k = 0; k < 2; ++k) if (active & (1U << k)) {
+        for (uint32_t t = sr::firstWithResidue(lo[k], residue[k]); t < hi[k]; t += 5) {
+            const uint64_t seed = low + (uint64_t((4965U * t) & 8191U) << 35);
+            if (sr::verifyPattern(seed, cells, count)) emit(seed);
+        }
+    }
+}
+
+// A cheap necessary condition before either exact high-bit solver.
+// Adding the common high-bit shift can wrap q_i-q_0 once: delta mod 10 is 0 or 8.
+FS_HD FS_INLINE bool differenceFilter(uint64_t low, const sr::CellTerm* cells, int count) {
+    const uint32_t q0 = uint32_t(sr::nextStateFromTerm(low, cells[0].term) >> 17);
+    for (int i = 1; i < count; ++i) {
+        const uint32_t qi = uint32_t(sr::nextStateFromTerm(low, cells[i].term) >> 17);
+        const uint32_t delta = (qi - q0) & 0x7fffffffU;
+        const uint32_t r = delta % 10;
+        if (r != 0 && r != 8) return false;
+    }
+    return true;
+}
+}
+
+#undef FS_HD
+#undef FS_INLINE