523 lines
19 KiB
Text
523 lines
19 KiB
Text
#include <cuda_runtime.h>
|
|
|
|
#include <algorithm>
|
|
#include <chrono>
|
|
#include <climits>
|
|
#include <cstdint>
|
|
#include <cstdlib>
|
|
#include <iostream>
|
|
#include <stdexcept>
|
|
#include <string>
|
|
#include <vector>
|
|
|
|
namespace {
|
|
|
|
constexpr uint64_t kJavaMultiplier = 0x5DEECE66DULL;
|
|
constexpr uint64_t kJavaAddend = 0xBULL;
|
|
constexpr uint64_t kJavaMask = (1ULL << 48) - 1;
|
|
constexpr uint64_t kSlimeXor = 987234911ULL;
|
|
constexpr int32_t kTileW = 16;
|
|
constexpr int32_t kTileH = 16;
|
|
constexpr int32_t kMaskFastPathMaxN = 49;
|
|
|
|
struct Args {
|
|
int64_t startSeed = 0;
|
|
int64_t endSeed = 1'000'000;
|
|
int32_t n = 3;
|
|
int32_t required = -1;
|
|
int32_t radius = 100;
|
|
int64_t seedBatch = 0;
|
|
int32_t maxResults = 32;
|
|
bool allLocations = false;
|
|
};
|
|
|
|
struct Hit {
|
|
int64_t seed;
|
|
int32_t x;
|
|
int32_t z;
|
|
int32_t count;
|
|
};
|
|
|
|
void checkCuda(cudaError_t error, const char* expr, const char* file, int line) {
|
|
if (error == cudaSuccess) {
|
|
return;
|
|
}
|
|
std::cerr << file << ':' << line << ": CUDA call failed: " << expr
|
|
<< ": " << cudaGetErrorString(error) << '\n';
|
|
std::exit(1);
|
|
}
|
|
|
|
#define CUDA_CHECK(expr) checkCuda((expr), #expr, __FILE__, __LINE__)
|
|
|
|
[[noreturn]] void usage(int exitCode = 1) {
|
|
std::cerr
|
|
<< "Usage: slime_cuda [args...]\n"
|
|
<< " -s <startSeed> First world seed, inclusive (default: 0)\n"
|
|
<< " -e <endSeed> Last world seed, exclusive (default: 1000000)\n"
|
|
<< " -n <chunkSize> N for the NxN window (default: 3)\n"
|
|
<< " -d <numChunks> Required slime chunks in the window (default: N*N)\n"
|
|
<< " -r <radius> Search start positions from -r..r on both axes (default: 100)\n"
|
|
<< " --batch <seeds> Seeds per kernel launch (default: max(total/100, 1000000))\n"
|
|
<< " --max-results <n> Print at most n hits, while still counting all hits (default: 32)\n"
|
|
<< " --all-locations Count every matching location, not just the first per seed\n"
|
|
<< " --help Show this help\n";
|
|
std::exit(exitCode);
|
|
}
|
|
|
|
const char* requireValue(int argc, char** argv, int& i) {
|
|
if (i + 1 >= argc) {
|
|
usage();
|
|
}
|
|
return argv[++i];
|
|
}
|
|
|
|
Args parseArgs(int argc, char** argv) {
|
|
Args args;
|
|
for (int i = 1; i < argc; ++i) {
|
|
const std::string arg = argv[i];
|
|
if (arg == "--help") {
|
|
usage(0);
|
|
} else if (arg == "-s") {
|
|
args.startSeed = std::stoll(requireValue(argc, argv, i));
|
|
} else if (arg == "-e") {
|
|
args.endSeed = std::stoll(requireValue(argc, argv, i));
|
|
} else if (arg == "-n") {
|
|
args.n = std::stoi(requireValue(argc, argv, i));
|
|
} else if (arg == "-d") {
|
|
args.required = std::stoi(requireValue(argc, argv, i));
|
|
} else if (arg == "-r") {
|
|
args.radius = std::stoi(requireValue(argc, argv, i));
|
|
} else if (arg == "--batch") {
|
|
args.seedBatch = std::stoll(requireValue(argc, argv, i));
|
|
} else if (arg == "--max-results") {
|
|
args.maxResults = std::stoi(requireValue(argc, argv, i));
|
|
} else if (arg == "--all-locations") {
|
|
args.allLocations = true;
|
|
} else {
|
|
usage();
|
|
}
|
|
}
|
|
|
|
const int64_t cells = int64_t(args.n) * args.n;
|
|
if (args.required < 0) {
|
|
args.required = static_cast<int32_t>(cells);
|
|
}
|
|
if (args.startSeed >= args.endSeed || args.n <= 0 || args.radius < 0 ||
|
|
args.seedBatch < 0 || args.maxResults < 0 || cells > INT32_MAX ||
|
|
args.required <= 0 || args.required > cells) {
|
|
usage();
|
|
}
|
|
return args;
|
|
}
|
|
|
|
__host__ __device__ __forceinline__ uint64_t javaIntAsLongBits(uint32_t value) {
|
|
return static_cast<uint64_t>(static_cast<int64_t>(static_cast<int32_t>(value)));
|
|
}
|
|
|
|
__host__ __device__ __forceinline__ uint64_t coordXTerm(int32_t chunkX) {
|
|
const uint32_t x = static_cast<uint32_t>(chunkX);
|
|
return javaIntAsLongBits(x * x * 4987142u) + javaIntAsLongBits(x * 5947611u);
|
|
}
|
|
|
|
__host__ __device__ __forceinline__ uint64_t coordZTerm(int32_t chunkZ) {
|
|
const uint32_t z = static_cast<uint32_t>(chunkZ);
|
|
return javaIntAsLongBits(z * z) * 4392871ULL + javaIntAsLongBits(z * 389711u);
|
|
}
|
|
|
|
__host__ __device__ __forceinline__ uint64_t coordTerm(int32_t chunkX, int32_t chunkZ) {
|
|
return coordXTerm(chunkX) + coordZTerm(chunkZ);
|
|
}
|
|
|
|
__device__ __forceinline__ bool isSlimeFromTerm(uint64_t worldSeedBits, uint64_t term) {
|
|
uint64_t randomSeed = ((worldSeedBits + term) ^ kSlimeXor ^ kJavaMultiplier) & kJavaMask;
|
|
|
|
for (;;) {
|
|
randomSeed = (randomSeed * kJavaMultiplier + kJavaAddend) & kJavaMask;
|
|
const uint32_t bits = static_cast<uint32_t>(randomSeed >> 17);
|
|
const uint32_t value = bits % 10u;
|
|
if (bits - value + 9u <= 0x7FFFFFFFu) {
|
|
return value == 0;
|
|
}
|
|
}
|
|
}
|
|
|
|
__device__ __forceinline__ bool isNextInt10ZeroFast(uint32_t bits) {
|
|
const uint32_t product = bits * 0xCCCCCCCDu;
|
|
const uint32_t rotated = (product >> 1) | (product << 31);
|
|
return rotated <= 429496729u;
|
|
}
|
|
|
|
__device__ __forceinline__ bool isSlimeFromSplitTerms(uint64_t worldSeedBits, uint64_t xTerm, uint64_t zTerm) {
|
|
uint64_t randomSeed = ((worldSeedBits + xTerm + zTerm) ^ kSlimeXor ^ kJavaMultiplier) & kJavaMask;
|
|
randomSeed = (randomSeed * kJavaMultiplier + kJavaAddend) & kJavaMask;
|
|
return isNextInt10ZeroFast(static_cast<uint32_t>(randomSeed >> 17));
|
|
}
|
|
|
|
__device__ __forceinline__ bool isSlimeChunkInline(uint64_t worldSeedBits, int32_t chunkX, int32_t chunkZ) {
|
|
return isSlimeFromSplitTerms(worldSeedBits, coordXTerm(chunkX), coordZTerm(chunkZ));
|
|
}
|
|
|
|
std::vector<uint64_t> makeCoordTerms(int32_t radius, int32_t n) {
|
|
const int32_t window = radius * 2 + 1;
|
|
const int32_t chunkWidth = window + n - 1;
|
|
std::vector<uint64_t> terms(size_t(chunkWidth) * chunkWidth);
|
|
for (int32_t z = 0; z < chunkWidth; ++z) {
|
|
for (int32_t x = 0; x < chunkWidth; ++x) {
|
|
terms[size_t(z) * chunkWidth + x] = coordTerm(x - radius, z - radius);
|
|
}
|
|
}
|
|
return terms;
|
|
}
|
|
|
|
__device__ void recordHit(
|
|
int32_t onePerSeed,
|
|
int32_t* seedFound,
|
|
int64_t seedOffset,
|
|
unsigned long long* totalHits,
|
|
Hit* hits,
|
|
int32_t maxHits,
|
|
int64_t seed,
|
|
int32_t x,
|
|
int32_t z,
|
|
int32_t count
|
|
) {
|
|
if (onePerSeed != 0) {
|
|
if (atomicCAS(seedFound + seedOffset, 0, 1) != 0) {
|
|
return;
|
|
}
|
|
}
|
|
|
|
const unsigned long long index = atomicAdd(totalHits, 1ULL);
|
|
if (index < static_cast<unsigned long long>(maxHits)) {
|
|
hits[index] = {seed, x, z, count};
|
|
}
|
|
}
|
|
|
|
__global__ void searchTilesKernel(
|
|
int64_t startSeed,
|
|
int64_t seedCount,
|
|
const uint64_t* terms,
|
|
int32_t chunkWidth,
|
|
int32_t window,
|
|
int32_t n,
|
|
int32_t required,
|
|
int32_t radius,
|
|
int32_t tilesPerSeed,
|
|
int32_t tilesX,
|
|
int32_t onePerSeed,
|
|
int32_t* seedFound,
|
|
unsigned long long* totalHits,
|
|
Hit* hits,
|
|
int32_t maxHits
|
|
) {
|
|
extern __shared__ uint8_t tile[];
|
|
__shared__ unsigned long long horizontalRows[kTileH + kMaskFastPathMaxN - 1];
|
|
|
|
const int64_t linearBlock = blockIdx.x;
|
|
const int64_t seedOffset = linearBlock / tilesPerSeed;
|
|
if (seedOffset >= seedCount) {
|
|
return;
|
|
}
|
|
if (onePerSeed != 0 && seedFound[seedOffset] != 0) {
|
|
return;
|
|
}
|
|
|
|
const int32_t tileId = static_cast<int32_t>(linearBlock - seedOffset * tilesPerSeed);
|
|
const int32_t tileBaseX = (tileId % tilesX) * kTileW;
|
|
const int32_t tileBaseZ = (tileId / tilesX) * kTileH;
|
|
const int32_t actualW = (window - tileBaseX) < kTileW ? (window - tileBaseX) : kTileW;
|
|
const int32_t actualH = (window - tileBaseZ) < kTileH ? (window - tileBaseZ) : kTileH;
|
|
const int32_t sharedW = actualW + n - 1;
|
|
const int32_t sharedH = actualH + n - 1;
|
|
const int32_t sharedCells = sharedW * sharedH;
|
|
const int64_t seed = startSeed + seedOffset;
|
|
const uint64_t seedBits = static_cast<uint64_t>(seed);
|
|
|
|
const int32_t localThread = threadIdx.y * blockDim.x + threadIdx.x;
|
|
const int32_t threadCount = blockDim.x * blockDim.y;
|
|
for (int32_t index = localThread; index < sharedCells; index += threadCount) {
|
|
const int32_t localX = index % sharedW;
|
|
const int32_t localZ = index / sharedW;
|
|
const size_t termIndex = size_t(tileBaseZ + localZ) * chunkWidth + tileBaseX + localX;
|
|
tile[index] = isSlimeFromTerm(seedBits, terms[termIndex]) ? 1 : 0;
|
|
}
|
|
__syncthreads();
|
|
|
|
if (required == n * n && n <= kMaskFastPathMaxN) {
|
|
for (int32_t row = localThread; row < sharedH; row += threadCount) {
|
|
unsigned long long rowMask = 0;
|
|
for (int32_t x = 0; x < sharedW; ++x) {
|
|
rowMask |= static_cast<unsigned long long>(tile[row * sharedW + x]) << x;
|
|
}
|
|
unsigned long long mask = rowMask;
|
|
for (int32_t shift = 1; shift < n; ++shift) {
|
|
mask &= rowMask >> shift;
|
|
}
|
|
horizontalRows[row] = mask;
|
|
}
|
|
__syncthreads();
|
|
|
|
const int32_t localX = threadIdx.x;
|
|
const int32_t localZ = threadIdx.y;
|
|
if (localX < actualW && localZ < actualH) {
|
|
unsigned long long mask = ~0ULL;
|
|
for (int32_t dz = 0; dz < n; ++dz) {
|
|
mask &= horizontalRows[localZ + dz];
|
|
}
|
|
if (((mask >> localX) & 1ULL) != 0) {
|
|
recordHit(
|
|
onePerSeed,
|
|
seedFound,
|
|
seedOffset,
|
|
totalHits,
|
|
hits,
|
|
maxHits,
|
|
seed,
|
|
tileBaseX + localX - radius,
|
|
tileBaseZ + localZ - radius,
|
|
n * n
|
|
);
|
|
}
|
|
}
|
|
return;
|
|
}
|
|
|
|
const int32_t localX = threadIdx.x;
|
|
const int32_t localZ = threadIdx.y;
|
|
if (localX >= actualW || localZ >= actualH) {
|
|
return;
|
|
}
|
|
if (onePerSeed != 0 && seedFound[seedOffset] != 0) {
|
|
return;
|
|
}
|
|
|
|
int32_t count = 0;
|
|
int32_t remaining = n * n;
|
|
for (int32_t dz = 0; dz < n; ++dz) {
|
|
const int32_t rowBase = (localZ + dz) * sharedW + localX;
|
|
for (int32_t dx = 0; dx < n; ++dx) {
|
|
count += tile[rowBase + dx];
|
|
--remaining;
|
|
if (count + remaining < required) {
|
|
return;
|
|
}
|
|
}
|
|
}
|
|
|
|
if (count >= required) {
|
|
recordHit(
|
|
onePerSeed,
|
|
seedFound,
|
|
seedOffset,
|
|
totalHits,
|
|
hits,
|
|
maxHits,
|
|
seed,
|
|
tileBaseX + localX - radius,
|
|
tileBaseZ + localZ - radius,
|
|
count
|
|
);
|
|
}
|
|
}
|
|
|
|
__global__ void searchSeedsFullKernel(
|
|
int64_t startSeed,
|
|
int64_t seedCount,
|
|
int32_t window,
|
|
int32_t n,
|
|
int32_t radius,
|
|
unsigned long long* totalHits,
|
|
Hit* hits,
|
|
int32_t maxHits
|
|
) {
|
|
const int64_t seedOffset = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
|
|
if (seedOffset >= seedCount) {
|
|
return;
|
|
}
|
|
|
|
const int64_t seed = startSeed + seedOffset;
|
|
const uint64_t seedBits = static_cast<uint64_t>(seed);
|
|
|
|
for (int32_t x = 0; x < window; ++x) {
|
|
const int32_t chunkX = x - radius;
|
|
const uint64_t firstXTerm = coordXTerm(chunkX);
|
|
for (int32_t z = 0; z < window; ++z) {
|
|
const int32_t chunkZ = z - radius;
|
|
if (!isSlimeFromSplitTerms(seedBits, firstXTerm, coordZTerm(chunkZ))) {
|
|
continue;
|
|
}
|
|
|
|
bool found = true;
|
|
for (int32_t dx = 0; dx < n && found; ++dx) {
|
|
const uint64_t xTerm = dx == 0 ? firstXTerm : coordXTerm(chunkX + dx);
|
|
const int32_t firstDz = dx == 0 ? 1 : 0;
|
|
for (int32_t dz = firstDz; dz < n; ++dz) {
|
|
if (!isSlimeFromSplitTerms(seedBits, xTerm, coordZTerm(chunkZ + dz))) {
|
|
found = false;
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
if (found) {
|
|
const unsigned long long index = atomicAdd(totalHits, 1ULL);
|
|
if (index < static_cast<unsigned long long>(maxHits)) {
|
|
hits[index] = {seed, chunkX, chunkZ, n * n};
|
|
}
|
|
return;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
void printHits(const std::vector<Hit>& hits) {
|
|
for (const Hit& hit : hits) {
|
|
std::cout << "hit seed=" << hit.seed
|
|
<< " x=" << hit.x
|
|
<< " z=" << hit.z
|
|
<< " count=" << hit.count << '\n';
|
|
}
|
|
}
|
|
|
|
} // namespace
|
|
|
|
int main(int argc, char** argv) {
|
|
try {
|
|
const Args args = parseArgs(argc, argv);
|
|
const int32_t window = args.radius * 2 + 1;
|
|
const int32_t chunkWidth = window + args.n - 1;
|
|
const int32_t tilesX = (window + kTileW - 1) / kTileW;
|
|
const int32_t tilesZ = (window + kTileH - 1) / kTileH;
|
|
const int32_t tilesPerSeed = tilesX * tilesZ;
|
|
const size_t dynamicSharedBytes = size_t(kTileW + args.n - 1) * (kTileH + args.n - 1);
|
|
const size_t staticSharedBytes = size_t(kTileH + kMaskFastPathMaxN - 1) * sizeof(unsigned long long);
|
|
const bool useSeedKernel = !args.allLocations && args.required == args.n * args.n;
|
|
const int32_t seedThreadsPerBlock = 64;
|
|
|
|
int device = 0;
|
|
cudaDeviceProp prop{};
|
|
CUDA_CHECK(cudaGetDevice(&device));
|
|
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
|
|
if (!useSeedKernel && dynamicSharedBytes + staticSharedBytes > static_cast<size_t>(prop.sharedMemPerBlock)) {
|
|
throw std::runtime_error("requested N needs more shared memory per block than this GPU allows");
|
|
}
|
|
|
|
uint64_t* deviceTerms = nullptr;
|
|
if (!useSeedKernel) {
|
|
std::vector<uint64_t> hostTerms = makeCoordTerms(args.radius, args.n);
|
|
CUDA_CHECK(cudaMalloc(&deviceTerms, hostTerms.size() * sizeof(uint64_t)));
|
|
CUDA_CHECK(cudaMemcpy(deviceTerms, hostTerms.data(), hostTerms.size() * sizeof(uint64_t), cudaMemcpyHostToDevice));
|
|
}
|
|
|
|
const int64_t totalSeeds = args.endSeed - args.startSeed;
|
|
const int64_t requestedBatch = args.seedBatch > 0
|
|
? args.seedBatch
|
|
: std::max(totalSeeds / 100, int64_t(1'000'000));
|
|
const int64_t maxSeedsPerLaunch = useSeedKernel
|
|
? int64_t(INT_MAX) * seedThreadsPerBlock
|
|
: std::max<int64_t>(1, INT_MAX / std::max<int32_t>(1, tilesPerSeed));
|
|
const int64_t maxBatchSeeds = std::min({requestedBatch, maxSeedsPerLaunch, totalSeeds});
|
|
|
|
int32_t* deviceSeedFound = nullptr;
|
|
if (!args.allLocations && !useSeedKernel) {
|
|
CUDA_CHECK(cudaMalloc(&deviceSeedFound, size_t(maxBatchSeeds) * sizeof(int32_t)));
|
|
}
|
|
|
|
unsigned long long* deviceHitCount = nullptr;
|
|
CUDA_CHECK(cudaMalloc(&deviceHitCount, sizeof(unsigned long long)));
|
|
|
|
Hit* deviceHits = nullptr;
|
|
if (args.maxResults > 0) {
|
|
CUDA_CHECK(cudaMalloc(&deviceHits, size_t(args.maxResults) * sizeof(Hit)));
|
|
}
|
|
|
|
std::cout << "CUDA search on " << prop.name << ": seeds [" << args.startSeed << ", "
|
|
<< args.endSeed << "), N=" << args.n << ", required=" << args.required
|
|
<< ", radius=" << args.radius << ", tile=" << kTileW << "x" << kTileH
|
|
<< ", mode=" << (useSeedKernel ? "seed-early-exit" : "tile-shared")
|
|
<< (args.allLocations ? ", all locations\n" : ", first location per seed\n");
|
|
|
|
const auto started = std::chrono::steady_clock::now();
|
|
uint64_t totalHits = 0;
|
|
int32_t printed = 0;
|
|
|
|
for (int64_t seed = args.startSeed; seed < args.endSeed;) {
|
|
const int64_t remainingSeeds = args.endSeed - seed;
|
|
const int64_t batchSeeds = std::min(maxBatchSeeds, remainingSeeds);
|
|
const int32_t remainingPrints = std::max<int32_t>(0, args.maxResults - printed);
|
|
|
|
CUDA_CHECK(cudaMemset(deviceHitCount, 0, sizeof(unsigned long long)));
|
|
if (!args.allLocations && !useSeedKernel) {
|
|
CUDA_CHECK(cudaMemset(deviceSeedFound, 0, size_t(batchSeeds) * sizeof(int32_t)));
|
|
}
|
|
|
|
if (useSeedKernel) {
|
|
const int64_t blocks = (batchSeeds + seedThreadsPerBlock - 1) / seedThreadsPerBlock;
|
|
searchSeedsFullKernel<<<static_cast<unsigned int>(blocks), seedThreadsPerBlock>>>(
|
|
seed,
|
|
batchSeeds,
|
|
window,
|
|
args.n,
|
|
args.radius,
|
|
deviceHitCount,
|
|
deviceHits,
|
|
remainingPrints
|
|
);
|
|
} else {
|
|
const int64_t blocks = batchSeeds * tilesPerSeed;
|
|
searchTilesKernel<<<static_cast<unsigned int>(blocks), dim3(kTileW, kTileH), dynamicSharedBytes>>>(
|
|
seed,
|
|
batchSeeds,
|
|
deviceTerms,
|
|
chunkWidth,
|
|
window,
|
|
args.n,
|
|
args.required,
|
|
args.radius,
|
|
tilesPerSeed,
|
|
tilesX,
|
|
args.allLocations ? 0 : 1,
|
|
deviceSeedFound,
|
|
deviceHitCount,
|
|
deviceHits,
|
|
remainingPrints
|
|
);
|
|
}
|
|
CUDA_CHECK(cudaGetLastError());
|
|
CUDA_CHECK(cudaDeviceSynchronize());
|
|
|
|
unsigned long long batchHits = 0;
|
|
CUDA_CHECK(cudaMemcpy(&batchHits, deviceHitCount, sizeof(batchHits), cudaMemcpyDeviceToHost));
|
|
totalHits += batchHits;
|
|
|
|
const int32_t copyHits = static_cast<int32_t>(
|
|
std::min<unsigned long long>(batchHits, static_cast<unsigned long long>(remainingPrints))
|
|
);
|
|
if (copyHits > 0) {
|
|
std::vector<Hit> hostHits(copyHits);
|
|
CUDA_CHECK(cudaMemcpy(hostHits.data(), deviceHits, size_t(copyHits) * sizeof(Hit), cudaMemcpyDeviceToHost));
|
|
printHits(hostHits);
|
|
printed += copyHits;
|
|
}
|
|
|
|
seed += batchSeeds;
|
|
std::cerr << "\rchecked " << (seed - args.startSeed) << "/" << totalSeeds << " seeds" << std::flush;
|
|
}
|
|
std::cerr << '\n';
|
|
|
|
const auto finished = std::chrono::steady_clock::now();
|
|
const std::chrono::duration<double> elapsed = finished - started;
|
|
std::cout << "checked " << totalSeeds << " seeds, hits=" << totalHits
|
|
<< ", elapsed=" << elapsed.count() << "s\n";
|
|
|
|
CUDA_CHECK(cudaFree(deviceHits));
|
|
CUDA_CHECK(cudaFree(deviceHitCount));
|
|
CUDA_CHECK(cudaFree(deviceSeedFound));
|
|
if (deviceTerms != nullptr) {
|
|
CUDA_CHECK(cudaFree(deviceTerms));
|
|
}
|
|
return 0;
|
|
} catch (const std::exception& ex) {
|
|
std::cerr << "error: " << ex.what() << '\n';
|
|
return 1;
|
|
}
|
|
}
|