Yaossg/slime
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Support custom slime chunk patterns in CUDA finder

Browse source
This commit is contained in:
Yaossg 2026-10-16 23:51:12 +08:00
parent dd5fab9540
commit 7952207735
10 changed files with 2302 additions and 34 deletions
Download patch file Download diff file Expand all files Collapse all files

80
bench/slime_bench.cu
View file

@ -25,7 +25,25 @@ using Clock=std::chrono::steady_clock;
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;};
struct PatternCell{int x,z;uint8_t requirement;};
struct Pattern{int width=0,height=0,positive=0;std::vector<std::string> rows;std::vector<PatternCell> cells;};
Pattern parsePattern(std::string text){
std::replace(text.begin(),text.end(),'/','\n');
Pattern p;std::stringstream input(text);std::string row;
while(std::getline(input,row)){if(!row.empty()&&row.back()=='\r')row.pop_back();if(row.empty())throw std::runtime_error("empty pattern row");p.rows.push_back(row);}
if(p.rows.empty())throw std::runtime_error("empty pattern");
p.width=p.rows[0].size();p.height=p.rows.size();
if(p.width>64||p.height>64)throw std::runtime_error("pattern dimensions must be 1..64");
for(auto& r:p.rows)if(r.size()!=size_t(p.width))throw std::runtime_error("pattern rows must have equal width");
for(char wanted:{'#','.'})for(int z=0;z<p.height;++z)for(int x=0;x<p.width;++x){char ch=p.rows[z][x];
if(ch!='#'&&ch!='.'&&ch!='?')throw std::runtime_error("pattern characters must be #, . or ?");
if(ch==wanted)p.cells.push_back({x,z,uint8_t(ch=='#'?sr::kSlime:sr::kNonSlime)});
if(wanted=='#'&&ch=='#')++p.positive;
}
if(!p.positive)throw std::runtime_error("pattern requires at least one # cell");
return p;
}
struct Config{Pattern pattern;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);}
@ -38,15 +56,20 @@ Coord coordinate(const Config& c,uint64_t index){
}
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){
void terms(Coord p,const Pattern& pattern,sr::CellTerm* out){
// The grouped filter needs its first positive term to be zero.
auto first=pattern.cells.front();uint64_t anchor=sr::coordinateTerm(p.x+first.x,p.z+first.z);size_t i=0;
for(auto cell:pattern.cells)out[i++]={(sr::coordinateTerm(p.x+cell.x,p.z+cell.z)-anchor)&sr::kMask48,cell.requirement};
}
__global__ void rootsKernel(const sr::CellTerm* cells,uint32_t positions,int count,int positive,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;
for(int i=0;i<positive;++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){
__global__ void prepareKernel(const sr::CellTerm* cells,const Job* jobs,uint32_t n,int count,int positive,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;
for(int i=0;i<positive;++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);
}
@ -61,22 +84,22 @@ __host__ __device__ uint32_t groupMask(const uint32_t*p,uint32_t base,int count)
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};}
Coord p;uint64_t originTerm;const sr::CellTerm* negative;int negativeCount;Match* out;uint64_t capacity;unsigned long long* total;
__device__ void operator()(uint64_t seed){if(!sr::verifyPattern(seed,negative,negativeCount))return;uint64_t slot=atomicAdd(total,1ULL);if(slot<capacity)out[slot]={(seed-originTerm)&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){
__global__ void searchKernel(const sr::CellTerm*cells,const Coord*coords,const Job*jobs,uint32_t n,int count,int positive,int originX,int originZ,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);}
uint32_t mask=groupMask(p,base,positive);
Job job=jobs[k];Coord origin=coords[job.position];Emit emit{origin,sr::coordinateTerm(origin.x+originX,origin.z+originZ),cells+uint64_t(job.position)*count+positive,count-positive,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,positive,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);
int count=c.pattern.cells.size(),positive=c.pattern.positive;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;
@ -95,9 +118,9 @@ void worker(const Config& c,int device,Shared& shared,Metrics& result){try{
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);
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.pattern,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());
rootsKernel<<<unsigned((uint64_t(missing.size())<<18)/256),256>>>(dcells.p,missing.size(),count,positive,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);
@ -112,12 +135,12 @@ void worker(const Config& c,int device,Shared& shared,Metrics& result){try{
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);
cells.resize(active.size()*count);for(size_t i=0;i<active.size();++i)terms(active[i],c.pattern,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());
prepareKernel<<<unsigned((jobs.size()+127)/128),128>>>(dcells.p,djobs.p,jobs.size(),count,positive,prep.p);
searchKernel<<<unsigned((uint64_t(jobs.size())<<13)/256),256>>>(dcells.p,dcoords.p,djobs.p,jobs.size(),count,positive,c.pattern.cells.front().x,c.pattern.cells.front().z,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();};
@ -130,12 +153,15 @@ void worker(const Config& c,int device,Shared& shared,Metrics& result){try{
}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.
Search effective 48-bit seeds for slime chunk patterns.
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)
--size N Solid square side, 1..6 (default 5)
--pattern TEXT Rows separated by /; # slime, . non-slime, ? ignored
--pattern-file FILE Same pattern, one row per line (dimensions 1..64)
Requires at least one #; exclusive with --size
--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)
@ -153,20 +179,28 @@ 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;}
Config c;bool radiusSet=false,sizeSet=false;std::string patternText;bool patternSet=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=="--size"){c.size=bounded(1,6);sizeSet=true;}
else if(a=="--pattern"||a=="--pattern-file"){if(patternSet)throw std::runtime_error("specify only one pattern");patternSet=true;
if(a=="--pattern")patternText=v;else{std::ifstream file(v);if(!file)throw std::runtime_error("cannot read pattern: "+v);std::ostringstream text;text<<file.rdbuf();patternText=text.str();}}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(patternSet&&sizeSet)throw std::runtime_error("--size and pattern options are mutually exclusive");
if(!patternSet){for(int z=0;z<c.size;++z){if(z)patternText+='/';patternText+=std::string(c.size,'#');}}
c.pattern=parsePattern(patternText);
for(auto axis:{std::make_pair(c.x,c.pattern.width),std::make_pair(c.z,c.pattern.height)})
if(int64_t(axis.first)-int64_t(c.radius)<INT32_MIN||int64_t(axis.first)+int64_t(c.radius)+axis.second-1>INT32_MAX)throw std::runtime_error("coordinates including pattern 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");
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\":"<<(patternSet?"null":std::to_string(c.size))<<",\"anchors\":"<<anchors(c.radius)<<",\"effective_seed_bits\":48,\"dry_run\":"<<(c.dry?"true":"false");
report<<",\"pattern_width\":"<<c.pattern.width<<",\"pattern_height\":"<<c.pattern.height<<",\"slime_cells\":"<<c.pattern.positive<<",\"non_slime_cells\":"<<c.pattern.cells.size()-c.pattern.positive<<",\"pattern\":[";
for(size_t i=0;i<c.pattern.rows.size();++i){if(i)report<<",";report<<"\""<<c.pattern.rows[i]<<"\"";}report<<"]";
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{