From a1e43bb679c66f2ca9d531d0e376e3ffdd90b414 Mon Sep 17 00:00:00 2001 From: Yaossg Date: Sun, 11 Oct 2026 12:33:56 +0800 Subject: [PATCH] chore: remove obsolete codex implementations --- .gitignore | 2 - Makefile | 1 - README.md | 1 - codex/Makefile | 18 -- codex/README.md | 9 - codex/bench_wsl.sh | 310 -------------------------- codex/slime_cpu.cpp | 403 ---------------------------------- codex/slime_cuda.cu | 523 -------------------------------------------- 8 files changed, 1267 deletions(-) delete mode 100644 codex/Makefile delete mode 100644 codex/README.md delete mode 100644 codex/bench_wsl.sh delete mode 100644 codex/slime_cpu.cpp delete mode 100644 codex/slime_cuda.cu diff --git a/.gitignore b/.gitignore index 1505ee2..630f9d7 100644 --- a/.gitignore +++ b/.gitignore @@ -15,8 +15,6 @@ cmake-build-*/ /slime /slime-cuda /sliming -/codex/slime_cpu -/codex/slime_cuda # Python environments and caches. __pycache__/ diff --git a/Makefile b/Makefile index 88178b9..69cb2bb 100644 --- a/Makefile +++ b/Makefile @@ -27,7 +27,6 @@ clean: rm -f slime slime-cuda sliming $(MAKE) -C bench clean $(MAKE) -C research clean - $(MAKE) -C codex clean $(MAKE) -C show clean $(MAKE) -C tools clean help: diff --git a/README.md b/README.md index 59ec4e6..df125c7 100644 --- a/README.md +++ b/README.md @@ -19,7 +19,6 @@ H100 编译：`make -C bench -B CUDA_ARCH=90`。需要 CUDA Toolkit 和 C++17 | [`research/`](research/README.md) | 求解原型、算法实验、正确性验证和测量报告 | | [`research/experiments/`](research/experiments/README.md) | 最近的完整覆盖、成组筛选和调度实验 | | [`show/`](show/README.md) | 分步骤教学代码、博客/视频稿及 Manim 工程 | -| [`codex/`](codex/README.md) | 早期 CPU/CUDA 对照实现及脚本 | | [`exp/`](exp/README.md) | 独立 SMT 逆向搜索探索 | | [`tools/`](tools/README.md) | 单区块检查小工具与历史汇编参考 | | `slime.cpp`、`slime.cu` | 最初的暴力扫描实现，保留原入口 | diff --git a/codex/Makefile b/codex/Makefile deleted file mode 100644 index 7f71318..0000000 --- a/codex/Makefile +++ /dev/null @@ -1,18 +0,0 @@ -CXX ?= g++ -NVCC ?= nvcc - -CXXFLAGS ?= -O3 -std=c++17 -march=native -pthread -NVCCFLAGS ?= -O3 -std=c++17 - -.PHONY: all clean - -all: slime_cpu slime_cuda - -slime_cpu: slime_cpu.cpp - $(CXX) $(CXXFLAGS) -o $@ $< - -slime_cuda: slime_cuda.cu - $(NVCC) $(NVCCFLAGS) -o $@ $< - -clean: - rm -f slime_cpu slime_cuda slime_cpu.exe slime_cuda.exe diff --git a/codex/README.md b/codex/README.md deleted file mode 100644 index 8a32da5..0000000 --- a/codex/README.md +++ /dev/null @@ -1,9 +0,0 @@ -# 早期 CPU/CUDA 对照实现 - -这里保留早期按种子扫描的 CPU、CUDA 实现及 WSL 对照脚本。它不是当前全有效种子范围基准 CLI 的入口。 - -- 当前 CLI：[bench/README.md](../bench/README.md) -- 早期实现构建：`make -C codex` -- 对照脚本参数：`bash codex/bench_wsl.sh --help` - -目录名和源码路径保留，避免破坏既有脚本。编译输出忽略；脚本默认把临时构建和运行结果放在 `/tmp` 下。 diff --git a/codex/bench_wsl.sh b/codex/bench_wsl.sh deleted file mode 100644 index 4fa812f..0000000 --- a/codex/bench_wsl.sh +++ /dev/null @@ -1,310 +0,0 @@ -#!/usr/bin/env bash -set -euo pipefail - -SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)" -ROOT_DIR="$(cd -- "$SCRIPT_DIR/.." && pwd)" -BUILD_DIR="${SLIME_BENCH_BUILD_DIR:-/tmp/slime-bench-build}" -OUT_DIR="${SLIME_BENCH_OUT_DIR:-/tmp/slime-bench-results}" - -CXX="${CXX:-g++}" -NVCC="${NVCC:-nvcc}" -CXXFLAGS="${CXXFLAGS:--O3 -std=c++17 -march=native -pthread -fwrapv}" -NVCCFLAGS="${NVCCFLAGS:--O3 -std=c++17 -Xcompiler=-fwrapv}" - -RUNS=3 -WARMUPS=1 -DO_CPU=1 -DO_CUDA=1 -KEEP_OUTPUT=0 -SEARCH_ARGS=() - -usage() { - cat <<'EOF' -Usage: bash codex/bench_wsl.sh [bench options] [-- common slime args] - -Bench options: - --runs Timed runs per binary (default: 3) - --warmups Untimed warmup runs per binary (default: 1) - --skip-cpu Do not build/run CPU binaries - --skip-cuda Do not build/run CUDA binaries - --build-dir Build directory (default: /tmp/slime-bench-build) - --out-dir Output directory (default: /tmp/slime-bench-results) - --keep-output Keep per-run stdout/stderr logs - --help Show this help - -Common slime args are passed to all binaries. Use only options supported by both -the original programs and the codex programs: -s, -e, -n, -d, -r. - -Default common args: - -s 0 -e 100000 -n 3 -r 60 - -Examples: - bash codex/bench_wsl.sh - bash codex/bench_wsl.sh --runs 5 -- -s 0 -e 1000000 -n 3 -r 100 - bash codex/bench_wsl.sh --skip-cuda -- -s 0 -e 200000 -n 4 -r 80 - -Notes: - * CPU comparison runs codex/slime_cpu with --all-locations to match the - original CPU program, which scans every matching location. - * CUDA comparison runs codex/slime_cuda in first-hit-per-seed mode to match - the original CUDA program. - * The original CPU program ignores non-zero -s while splitting thread ranges, - so use -s 0 for fair CPU result counts. -EOF -} - -die() { - echo "error: $*" >&2 - exit 1 -} - -need_value() { - local name="$1" - local value="${2:-}" - [[ -n "$value" ]] || die "$name needs a value" -} - -while [[ $# -gt 0 ]]; do - case "$1" in - --runs) - need_value "$1" "${2:-}" - RUNS="$2" - shift 2 - ;; - --warmups) - need_value "$1" "${2:-}" - WARMUPS="$2" - shift 2 - ;; - --skip-cpu) - DO_CPU=0 - shift - ;; - --skip-cuda) - DO_CUDA=0 - shift - ;; - --build-dir) - need_value "$1" "${2:-}" - BUILD_DIR="$2" - shift 2 - ;; - --out-dir) - need_value "$1" "${2:-}" - OUT_DIR="$2" - shift 2 - ;; - --keep-output) - KEEP_OUTPUT=1 - shift - ;; - --help) - usage - exit 0 - ;; - --) - shift - SEARCH_ARGS=("$@") - break - ;; - *) - SEARCH_ARGS+=("$1") - shift - ;; - esac -done - -[[ "$RUNS" =~ ^[0-9]+$ && "$RUNS" -gt 0 ]] || die "--runs must be a positive integer" -[[ "$WARMUPS" =~ ^[0-9]+$ ]] || die "--warmups must be a non-negative integer" - -if [[ "${#SEARCH_ARGS[@]}" -eq 0 ]]; then - SEARCH_ARGS=(-s 0 -e 10000000 -n 3 -r 60) -fi - -START_SEED_ARG=0 -for ((i = 0; i < ${#SEARCH_ARGS[@]}; ++i)); do - if [[ "${SEARCH_ARGS[$i]}" == "-s" && $((i + 1)) -lt ${#SEARCH_ARGS[@]} ]]; then - START_SEED_ARG="${SEARCH_ARGS[$((i + 1))]}" - fi -done - -mkdir -p "$BUILD_DIR" "$OUT_DIR" - -read -r -a CXXFLAGS_ARR <<< "$CXXFLAGS" -read -r -a NVCCFLAGS_ARR <<< "$NVCCFLAGS" - -build_cpu() { - command -v "$CXX" >/dev/null 2>&1 || die "cannot find C++ compiler '$CXX'" - - echo "Building CPU binaries with: $CXX $CXXFLAGS" - "$CXX" "${CXXFLAGS_ARR[@]}" "$ROOT_DIR/slime.cpp" -o "$BUILD_DIR/user_cpu" - "$CXX" "${CXXFLAGS_ARR[@]}" "$ROOT_DIR/codex/slime_cpu.cpp" -o "$BUILD_DIR/codex_cpu" -} - -build_cuda() { - if ! command -v "$NVCC" >/dev/null 2>&1; then - echo "Skipping CUDA: cannot find '$NVCC'" >&2 - DO_CUDA=0 - return - fi - - echo "Building CUDA binaries with: $NVCC $NVCCFLAGS" - "$NVCC" "${NVCCFLAGS_ARR[@]}" "$ROOT_DIR/slime.cu" -o "$BUILD_DIR/user_cuda" - "$NVCC" "${NVCCFLAGS_ARR[@]}" "$ROOT_DIR/codex/slime_cuda.cu" -o "$BUILD_DIR/codex_cuda" -} - -extract_count() { - local file="$1" - local count - count="$(sed -n 's/.*hits=\([0-9][0-9]*\).*/\1/p' "$file" | tail -n 1)" - if [[ -z "$count" ]]; then - count="$(sed -n 's/.*found \([0-9][0-9]*\) seeds.*/\1/p' "$file" | tail -n 1)" - fi - if [[ -z "$count" ]]; then - count="$(grep -Ec '^(hit|found seed)' "$file" || true)" - fi - echo "${count:-unknown}" -} - -run_once() { - local name="$1" - shift - - local stdout_file="$OUT_DIR/$name.stdout" - local stderr_file="$OUT_DIR/$name.stderr" - local time_file="$OUT_DIR/$name.time" - - set +e - /usr/bin/time -f '%e' -o "$time_file" "$@" >"$stdout_file" 2>"$stderr_file" - local status=$? - set -e - - if [[ "$status" -ne 0 ]]; then - echo "Command failed: $*" >&2 - echo "--- stderr ---" >&2 - tail -n 80 "$stderr_file" >&2 || true - exit "$status" - fi - - local elapsed - elapsed="$(tr -d '\r\n' < "$time_file")" - local count - count="$(extract_count "$stdout_file")" - echo "$elapsed $count" -} - -mean() { - awk ' - NF { sum += $1; n += 1 } - END { - if (n == 0) print "nan"; - else printf "%.6f", sum / n; - } - ' -} - -speedup() { - local base="$1" - local candidate="$2" - awk -v base="$base" -v candidate="$candidate" ' - BEGIN { - if (candidate > 0) printf "%.2fx", base / candidate; - else printf "n/a"; - } - ' -} - -bench_binary() { - local label="$1" - shift - local -a command=("$@") - - echo - echo "== $label ==" - echo "Command: ${command[*]}" - - for ((i = 1; i <= WARMUPS; ++i)); do - echo "Warmup $i/$WARMUPS..." - run_once "${label// /_}.warmup.$i" "${command[@]}" >/dev/null - done - - local times_file="$OUT_DIR/${label// /_}.times" - : > "$times_file" - local last_count="unknown" - for ((i = 1; i <= RUNS; ++i)); do - local result elapsed count - result="$(run_once "${label// /_}.run.$i" "${command[@]}")" - elapsed="${result%% *}" - count="${result#* }" - last_count="$count" - echo "$elapsed" >> "$times_file" - printf 'Run %d/%d: %ss, reported=%s\n' "$i" "$RUNS" "$elapsed" "$count" - done - - local avg - avg="$(mean < "$times_file")" - printf '%s\t%s\t%s\n' "$label" "$avg" "$last_count" >> "$OUT_DIR/summary.tsv" -} - -print_summary() { - echo - echo "Summary" - printf '%-16s %12s %12s %12s\n' "binary" "avg_s" "reported" "speedup" - - local user_cpu_avg="" codex_cpu_avg="" user_cuda_avg="" codex_cuda_avg="" - while IFS=$'\t' read -r label avg count; do - case "$label" in - user_cpu) user_cpu_avg="$avg" ;; - codex_cpu) codex_cpu_avg="$avg" ;; - user_cuda) user_cuda_avg="$avg" ;; - codex_cuda) codex_cuda_avg="$avg" ;; - esac - done < "$OUT_DIR/summary.tsv" - - while IFS=$'\t' read -r label avg count; do - local base="" sp="baseline" - case "$label" in - codex_cpu) base="$user_cpu_avg" ;; - codex_cuda) base="$user_cuda_avg" ;; - esac - if [[ -n "$base" ]]; then - sp="$(speedup "$base" "$avg")" - fi - printf '%-16s %12s %12s %12s\n' "$label" "$avg" "$count" "$sp" - done < "$OUT_DIR/summary.tsv" -} - -echo "Root: $ROOT_DIR" -echo "Build dir: $BUILD_DIR" -echo "Output dir: $OUT_DIR" -echo "Args: ${SEARCH_ARGS[*]}" -echo "Runs: $RUNS timed, $WARMUPS warmup" - -if [[ "$DO_CPU" -eq 1 && "$START_SEED_ARG" != "0" ]]; then - echo "Warning: original slime.cpp ignores non-zero -s in its CPU thread split; CPU counts will not be apples-to-apples." >&2 -fi - -: > "$OUT_DIR/summary.tsv" - -if [[ "$DO_CPU" -eq 1 ]]; then - build_cpu - bench_binary user_cpu "$BUILD_DIR/user_cpu" "${SEARCH_ARGS[@]}" - bench_binary codex_cpu "$BUILD_DIR/codex_cpu" "${SEARCH_ARGS[@]}" --all-locations -fi - -if [[ "$DO_CUDA" -eq 1 ]]; then - build_cuda - if [[ "$DO_CUDA" -eq 1 ]]; then - bench_binary user_cuda "$BUILD_DIR/user_cuda" "${SEARCH_ARGS[@]}" - bench_binary codex_cuda "$BUILD_DIR/codex_cuda" "${SEARCH_ARGS[@]}" - fi -fi - -print_summary - -if [[ "$KEEP_OUTPUT" -eq 0 ]]; then - find "$OUT_DIR" -type f \( -name '*.stdout' -o -name '*.stderr' -o -name '*.time' -o -name '*.times' \) -delete -else - echo - echo "Kept logs in $OUT_DIR" -fi diff --git a/codex/slime_cpu.cpp b/codex/slime_cpu.cpp deleted file mode 100644 index effe6de..0000000 --- a/codex/slime_cpu.cpp +++ /dev/null @@ -1,403 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#if defined(_MSC_VER) -#include -#endif - -namespace { - -constexpr uint64_t kJavaMultiplier = 0x5DEECE66DULL; -constexpr uint64_t kJavaAddend = 0xBULL; -constexpr uint64_t kJavaMask = (1ULL << 48) - 1; -constexpr uint64_t kSlimeXor = 987234911ULL; - -struct Args { - int64_t startSeed = 0; - int64_t endSeed = 1'000'000; - int32_t n = 3; - int32_t required = -1; - int32_t radius = 100; - int32_t threads = std::max(1u, std::thread::hardware_concurrency()); - int64_t seedBatch = 256; - int32_t maxResults = 32; - bool allLocations = false; -}; - -struct Hit { - int64_t seed; - int32_t x; - int32_t z; - int32_t count; -}; - -[[noreturn]] void usage(int exitCode = 1) { - std::cerr - << "Usage: slime_cpu [args...]\n" - << " -s First world seed, inclusive (default: 0)\n" - << " -e Last world seed, exclusive (default: 1000000)\n" - << " -n N for the NxN window (default: 3)\n" - << " -d Required slime chunks in the window (default: N*N)\n" - << " -r Search start positions from -r..r on both axes (default: 100)\n" - << " -p CPU worker threads (default: hardware concurrency)\n" - << " --batch Seeds fetched per worker batch (default: 256)\n" - << " --max-results 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 == "-p") { - args.threads = 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(cells); - } - if (args.startSeed >= args.endSeed || args.n <= 0 || args.radius < 0 || - args.threads <= 0 || args.seedBatch <= 0 || args.maxResults < 0 || - cells > INT32_MAX || args.required <= 0 || args.required > cells) { - usage(); - } - return args; -} - -uint64_t javaIntAsLongBits(uint32_t value) { - return (value & 0x80000000u) ? (0xFFFFFFFF00000000ULL | value) : uint64_t(value); -} - -uint64_t coordTerm(int32_t chunkX, int32_t chunkZ) { - const uint32_t x = static_cast(chunkX); - const uint32_t z = static_cast(chunkZ); - - uint64_t term = 0; - term += javaIntAsLongBits(x * x * 4987142u); - term += javaIntAsLongBits(x * 5947611u); - term += javaIntAsLongBits(z * z) * 4392871ULL; - term += javaIntAsLongBits(z * 389711u); - return term; -} - -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(randomSeed >> 17); - const uint32_t value = bits % 10u; - if (bits - value + 9u <= 0x7FFFFFFFu) { - return value == 0; - } - } -} - -int trailingZeroes(uint64_t value) { -#if defined(_MSC_VER) - unsigned long index = 0; - _BitScanForward64(&index, value); - return static_cast(index); -#else - return __builtin_ctzll(value); -#endif -} - -uint64_t lowBitsMask(int32_t bits) { - const int32_t rem = bits & 63; - return rem == 0 ? ~0ULL : ((1ULL << rem) - 1ULL); -} - -struct SearchSpace { - int32_t n; - int32_t radius; - int32_t window; - int32_t chunkWidth; - int32_t chunkRows; - int32_t chunkWords; - int32_t windowWords; - uint64_t lastWindowMask; - std::vector terms; - - SearchSpace(int32_t radiusIn, int32_t nIn) - : n(nIn), - radius(radiusIn), - window(radiusIn * 2 + 1), - chunkWidth(window + nIn - 1), - chunkRows(window + nIn - 1), - chunkWords((chunkWidth + 63) / 64), - windowWords((window + 63) / 64), - lastWindowMask(lowBitsMask(window)), - terms(size_t(chunkWidth) * chunkRows) { - for (int32_t z = 0; z < chunkRows; ++z) { - for (int32_t x = 0; x < chunkWidth; ++x) { - terms[size_t(z) * chunkWidth + x] = coordTerm(x - radius, z - radius); - } - } - } -}; - -struct Workspace { - std::vector rows; - std::vector horizontal; - std::vector rowValues; - std::vector rowRing; - std::vector vertical; - - explicit Workspace(const SearchSpace& space) - : rows(size_t(space.chunkRows) * space.chunkWords), - horizontal(size_t(space.chunkRows) * space.chunkWords), - rowValues(space.chunkWidth), - rowRing(size_t(space.n) * space.window), - vertical(space.window) {} -}; - -struct ResultSink { - std::mutex outputMutex; - std::atomic hits{0}; - int32_t maxResults; - - explicit ResultSink(int32_t maxResultsIn) : maxResults(maxResultsIn) {} - - void record(const Hit& hit) { - const uint64_t index = hits.fetch_add(1, std::memory_order_relaxed); - if (index >= static_cast(maxResults)) { - return; - } - - std::lock_guard lock(outputMutex); - std::cout << "hit seed=" << hit.seed - << " x=" << hit.x - << " z=" << hit.z - << " count=" << hit.count << '\n'; - } -}; - -uint64_t shiftedWord(const uint64_t* row, int32_t words, int32_t word, int32_t shift) { - const int32_t whole = shift >> 6; - const int32_t bits = shift & 63; - const int32_t src = word + whole; - const uint64_t lo = src < words ? row[src] : 0ULL; - if (bits == 0) { - return lo; - } - const uint64_t hi = src + 1 < words ? row[src + 1] : 0ULL; - return (lo >> bits) | (hi << (64 - bits)); -} - -uint64_t findFullSquare( - int64_t seed, - const SearchSpace& space, - Workspace& workspace, - ResultSink& sink, - bool allLocations -) { - const uint64_t seedBits = static_cast(seed); - std::fill(workspace.rows.begin(), workspace.rows.end(), 0ULL); - - for (int32_t z = 0; z < space.chunkRows; ++z) { - uint64_t* row = workspace.rows.data() + size_t(z) * space.chunkWords; - const uint64_t* termRow = space.terms.data() + size_t(z) * space.chunkWidth; - for (int32_t x = 0; x < space.chunkWidth; ++x) { - if (isSlimeFromTerm(seedBits, termRow[x])) { - row[x >> 6] |= 1ULL << (x & 63); - } - } - } - - for (int32_t z = 0; z < space.chunkRows; ++z) { - const uint64_t* row = workspace.rows.data() + size_t(z) * space.chunkWords; - uint64_t* out = workspace.horizontal.data() + size_t(z) * space.chunkWords; - for (int32_t word = 0; word < space.chunkWords; ++word) { - uint64_t mask = row[word]; - for (int32_t shift = 1; shift < space.n; ++shift) { - mask &= shiftedWord(row, space.chunkWords, word, shift); - } - out[word] = mask; - } - } - - uint64_t found = 0; - for (int32_t z = 0; z < space.window; ++z) { - for (int32_t word = 0; word < space.windowWords; ++word) { - uint64_t mask = ~0ULL; - for (int32_t dz = 0; dz < space.n; ++dz) { - mask &= workspace.horizontal[(size_t(z + dz) * space.chunkWords) + word]; - } - if (word == space.windowWords - 1) { - mask &= space.lastWindowMask; - } - - while (mask != 0) { - const int32_t bit = trailingZeroes(mask); - const int32_t x = word * 64 + bit; - sink.record({seed, x - space.radius, z - space.radius, space.n * space.n}); - ++found; - if (!allLocations) { - return found; - } - mask &= mask - 1; - } - } - } - - return found; -} - -uint64_t findAtLeastCount( - int64_t seed, - int32_t required, - const SearchSpace& space, - Workspace& workspace, - ResultSink& sink, - bool allLocations -) { - const uint64_t seedBits = static_cast(seed); - std::fill(workspace.vertical.begin(), workspace.vertical.end(), 0); - - uint64_t found = 0; - for (int32_t z = 0; z < space.chunkRows; ++z) { - const uint64_t* termRow = space.terms.data() + size_t(z) * space.chunkWidth; - for (int32_t x = 0; x < space.chunkWidth; ++x) { - workspace.rowValues[x] = isSlimeFromTerm(seedBits, termRow[x]) ? 1 : 0; - } - - int32_t* ringRow = workspace.rowRing.data() + size_t(z % space.n) * space.window; - if (z >= space.n) { - for (int32_t x = 0; x < space.window; ++x) { - workspace.vertical[x] -= ringRow[x]; - } - } - - int32_t rowSum = 0; - for (int32_t x = 0; x < space.n; ++x) { - rowSum += workspace.rowValues[x]; - } - for (int32_t x = 0; x < space.window; ++x) { - if (x != 0) { - rowSum += workspace.rowValues[x + space.n - 1]; - rowSum -= workspace.rowValues[x - 1]; - } - ringRow[x] = rowSum; - workspace.vertical[x] += rowSum; - } - - if (z + 1 < space.n) { - continue; - } - - const int32_t startZ = z - space.n + 1; - for (int32_t x = 0; x < space.window; ++x) { - const int32_t count = workspace.vertical[x]; - if (count >= required) { - sink.record({seed, x - space.radius, startZ - space.radius, count}); - ++found; - if (!allLocations) { - return found; - } - } - } - } - - return found; -} - -void worker( - const Args& args, - const SearchSpace& space, - std::atomic& nextSeed, - ResultSink& sink -) { - Workspace workspace(space); - for (;;) { - const int64_t begin = nextSeed.fetch_add(args.seedBatch, std::memory_order_relaxed); - if (begin >= args.endSeed) { - return; - } - const int64_t end = std::min(args.endSeed, begin + args.seedBatch); - for (int64_t seed = begin; seed < end; ++seed) { - if (args.required == args.n * args.n) { - findFullSquare(seed, space, workspace, sink, args.allLocations); - } else { - findAtLeastCount(seed, args.required, space, workspace, sink, args.allLocations); - } - } - } -} - -} // namespace - -int main(int argc, char** argv) { - try { - const Args args = parseArgs(argc, argv); - const SearchSpace space(args.radius, args.n); - ResultSink sink(args.maxResults); - std::atomic nextSeed(args.startSeed); - - std::cout << "CPU search: seeds [" << args.startSeed << ", " << args.endSeed - << "), N=" << args.n << ", required=" << args.required - << ", radius=" << args.radius << ", threads=" << args.threads - << (args.allLocations ? ", all locations\n" : ", first location per seed\n"); - - const auto started = std::chrono::steady_clock::now(); - std::vector workers; - workers.reserve(args.threads); - for (int32_t i = 0; i < args.threads; ++i) { - workers.emplace_back(worker, std::cref(args), std::cref(space), std::ref(nextSeed), std::ref(sink)); - } - for (std::thread& thread : workers) { - thread.join(); - } - - const auto finished = std::chrono::steady_clock::now(); - const std::chrono::duration elapsed = finished - started; - std::cout << "checked " << (args.endSeed - args.startSeed) - << " seeds, hits=" << sink.hits.load(std::memory_order_relaxed) - << ", elapsed=" << elapsed.count() << "s\n"; - return 0; - } catch (const std::exception& ex) { - std::cerr << "error: " << ex.what() << '\n'; - return 1; - } -} diff --git a/codex/slime_cuda.cu b/codex/slime_cuda.cu deleted file mode 100644 index 63dc6bc..0000000 --- a/codex/slime_cuda.cu +++ /dev/null @@ -1,523 +0,0 @@ -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -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 First world seed, inclusive (default: 0)\n" - << " -e Last world seed, exclusive (default: 1000000)\n" - << " -n N for the NxN window (default: 3)\n" - << " -d Required slime chunks in the window (default: N*N)\n" - << " -r Search start positions from -r..r on both axes (default: 100)\n" - << " --batch Seeds per kernel launch (default: max(total/100, 1000000))\n" - << " --max-results 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(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(static_cast(static_cast(value))); -} - -__host__ __device__ __forceinline__ uint64_t coordXTerm(int32_t chunkX) { - const uint32_t x = static_cast(chunkX); - return javaIntAsLongBits(x * x * 4987142u) + javaIntAsLongBits(x * 5947611u); -} - -__host__ __device__ __forceinline__ uint64_t coordZTerm(int32_t chunkZ) { - const uint32_t z = static_cast(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(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(randomSeed >> 17)); -} - -__device__ __forceinline__ bool isSlimeChunkInline(uint64_t worldSeedBits, int32_t chunkX, int32_t chunkZ) { - return isSlimeFromSplitTerms(worldSeedBits, coordXTerm(chunkX), coordZTerm(chunkZ)); -} - -std::vector makeCoordTerms(int32_t radius, int32_t n) { - const int32_t window = radius * 2 + 1; - const int32_t chunkWidth = window + n - 1; - std::vector 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(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(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(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(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(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(maxHits)) { - hits[index] = {seed, chunkX, chunkZ, n * n}; - } - return; - } - } - } -} - -void printHits(const std::vector& 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(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 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(1, INT_MAX / std::max(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(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<<(blocks), seedThreadsPerBlock>>>( - seed, - batchSeeds, - window, - args.n, - args.radius, - deviceHitCount, - deviceHits, - remainingPrints - ); - } else { - const int64_t blocks = batchSeeds * tilesPerSeed; - searchTilesKernel<<(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( - std::min(batchHits, static_cast(remainingPrints)) - ); - if (copyHits > 0) { - std::vector 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 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; - } -}