From 12e4266950691cfc67068017aaf6ed792bcf74f7 Mon Sep 17 00:00:00 2001 From: Yaossg Date: Sun, 11 Oct 2026 12:27:26 +0800 Subject: [PATCH] feat: add CUDA benchmark CLI and shared search algorithms --- .gitignore | 52 +++++- bench/.gitignore | 2 + bench/Makefile | 21 +++ bench/README.md | 115 +++++++++++++ bench/results/local-r512-n5.json | 1 + bench/results/local-r64-n5.json | 1 + bench/slime_bench.cu | 201 ++++++++++++++++++++++ bench/test_cli.py | 65 +++++++ bench/test_math.cu | 25 +++ include/README.md | 8 + include/slime/solver_core.hpp | 286 +++++++++++++++++++++++++++++++ include/slime/upper_interval.hpp | 71 ++++++++ 12 files changed, 846 insertions(+), 2 deletions(-) create mode 100644 bench/.gitignore create mode 100644 bench/Makefile create mode 100644 bench/README.md create mode 100644 bench/results/local-r512-n5.json create mode 100644 bench/results/local-r64-n5.json create mode 100644 bench/slime_bench.cu create mode 100644 bench/test_cli.py create mode 100644 bench/test_math.cu create mode 100644 include/README.md create mode 100644 include/slime/solver_core.hpp create mode 100644 include/slime/upper_interval.hpp diff --git a/.gitignore b/.gitignore index e9f9e20..1505ee2 100644 --- a/.gitignore +++ b/.gitignore @@ -1,2 +1,50 @@ -slime -slime-cuda +# Local build output (source assembly snapshots are intentionally retained). +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 diff --git a/bench/.gitignore b/bench/.gitignore new file mode 100644 index 0000000..0f28468 --- /dev/null +++ b/bench/.gitignore @@ -0,0 +1,2 @@ +/bin/ +/__pycache__/ diff --git a/bench/Makefile b/bench/Makefile new file mode 100644 index 0000000..92c43cd --- /dev/null +++ b/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 diff --git a/bench/README.md b/bench/README.md new file mode 100644 index 0000000..50867cd --- /dev/null +++ b/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 上执行。 diff --git a/bench/results/local-r512-n5.json b/bench/results/local-r512-n5.json new file mode 100644 index 0000000..d1b25fb --- /dev/null +++ b/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} diff --git a/bench/results/local-r64-n5.json b/bench/results/local-r64-n5.json new file mode 100644 index 0000000..88036e1 --- /dev/null +++ b/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} diff --git a/bench/slime_bench.cu b/bench/slime_bench.cu new file mode 100644 index 0000000..4ec676a --- /dev/null +++ b/bench/slime_bench.cu @@ -0,0 +1,201 @@ +#include "../include/slime/upper_interval.hpp" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +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()||nINT32_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=(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=n)return;const auto job=jobs[k]; + for(int i=0;i>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>4,lo=(1U<=(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);} +} +templatestruct 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 sample;}; +struct Shared{std::atomic next{0},done{0};std::atomic finished{0};int ready=0,workerCount=0;std::condition_variable cv;std::atomic 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 dcells(size_t(c.positionBatch)*count);Bufferdcoords(c.positionBatch);Bufferdroots(rootsCapacity),djobs(c.jobBatch); + Bufferprep(size_t(c.jobBatch)*count*2);Buffertotal(1);Buffermatches(c.maxResults); + std::map> cache;std::deque fifo;size_t cachedRoots=0; + result.setupMs=std::chrono::duration(Clock::now()-setup).count();{ + std::unique_lock 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(c.positionBatch,end-first);std::vector coords;std::vector keys; + std::map> lists;std::vector missing;std::vector missingKeys; + for(uint32_t i=0;isecond;++result.cacheHits;}else{lists[key]={};missing.push_back(p);missingKeys.push_back(key);} + } + if(!missing.empty()){ + auto start=Clock::now();std::vector cells(missing.size()*count);for(size_t i=0;i>>(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 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(Clock::now()-start).count(); + } + // Construct concrete positions only for classes with roots. + std::vector active;std::vector cells;std::vector jobs;jobs.reserve(c.jobBatch); + for(uint32_t i=0;i>>(dcells.p,djobs.p,jobs.size(),count,prep.p); + searchKernel<<>>(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(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(Clock::now()-start).count();jobs.clear();}; + for(uint32_t i=0;i(result.end-result.start).count(); +}catch(const std::exception& e){shared.cancel=true;std::lock_guard 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<<'\\'<=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)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(nhi)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;iINT32_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< devices;if(c.devices=="all"){for(int i=0;i=uint64_t(available))throw std::runtime_error("CUDA device unavailable: "+v);devices.push_back(id);} + } + if(devices.empty()||std::set(devices.begin(),devices.end()).size()!=devices.size())throw std::runtime_error("device list must be nonempty and unique"); + report<<",\"position_batch\":"< timings;uint64_t expected=0; + for(int round=0;round metrics(devices.size());std::vector workers;auto start=Clock::now(); + for(size_t i=0;i lock(shared.mutex);while(shared.finished(Clock::now()-start).count();auto first=metrics[0].start,last=metrics[0].end; + uint64_t visited=0,jobs=0,matches=0;std::vector 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()(last-first).count();timings.push_back(scan); + if(round)report<<",";report<<"{\"round\":"<sample.size()?"true":"false")<<",\"sample\":["; + for(size_t i=0;i>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() diff --git a/bench/test_math.cu b/bench/test_math.cu new file mode 100644 index 0000000..16960e0 --- /dev/null +++ b/bench/test_math.cu @@ -0,0 +1,25 @@ +#define main cli_entry_for_tests +#include "slime_bench.cu" +#undef main +#include +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> actual; + for(uint64_t i=0;i cells(count);terms(coord,n,cells.data());uint32_t root=rng()&262143,base=(rng()&131071)&~240U; + std::vector prepared(count*2); + for(int i=0;i>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< +#include + +#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(static_cast(static_cast(value))); +} + +SLIME_HD SLIME_INLINE uint64_t coordinateTerm(int32_t chunkX, int32_t chunkZ) { + const uint32_t x = static_cast(chunkX); + const uint32_t z = static_cast(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(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(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(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( + 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(kUpperSize - v0), + static_cast(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 +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( + 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(kUpperSize - v0), + static_cast(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 diff --git a/include/slime/upper_interval.hpp b/include/slime/upper_interval.hpp new file mode 100644 index 0000000..cc294a7 --- /dev/null +++ b/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 +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