Minecraft 1.19 Panorama seedcracking code [by andrew#0858]

//Minecraft 1.19 Panorama seedcracking code for CUDA
//created by Andrew (andrew#0858) from Minecraft@Home
//https://github.com/Gaider10

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>

#define QPC_FREQUENCY 10000000

__device__ __host__ uint64_t qpc_tick_to_nano_time(uint64_t qpc_tick) {
    return (uint64_t)((double)qpc_tick / (double)QPC_FREQUENCY * (double)1000000000);
}

#define DEF_BASE_RANDOM_NEXT_LONG(random_type, random_name) \
__device__ __host__ int64_t random_name##_next_long(random_type *rng) { \
    return ((int64_t)random_name##_next(rng, 32) << 32) + (int64_t)random_name##_next(rng, 32); \
}

#define RANDOM_MASK 281474976710655LLU
#define RANDOM_MULTIPLIER 25214903917LLU
#define RANDOM_ADDEND 11LLU

__device__ __host__ int32_t random_next(uint64_t *random, int bits) {
    *random = (*random * RANDOM_MULTIPLIER + RANDOM_ADDEND) & RANDOM_MASK;
    return *random >> 48 - bits;
}

DEF_BASE_RANDOM_NEXT_LONG(uint64_t, random)

__device__ __host__ uint64_t nano_time_to_world_seed(uint64_t nano_time) {
    uint64_t random = (nano_time ^ RANDOM_MULTIPLIER) & RANDOM_MASK;
    return random_next_long(&random);
}

#define XRSR_MIX1 0xbf58476d1ce4e5b9ULL
#define XRSR_MIX2 0x94d049bb133111ebULL
#define XRSR_MIX1_INVERSE 0x96de1b173f119089ULL
#define XRSR_MIX2_INVERSE 0x319642b2d24d8ec3ULL
#define XRSR_SILVER_RATIO 0x6a09e667f3bcc909ULL
#define XRSR_GOLDEN_RATIO 0x9e3779b97f4a7c15ULL

typedef struct {
    uint64_t lo;
    uint64_t hi;
} Xrsr;

__device__ __host__ uint64_t mix64(uint64_t a) {
    a = (a ^ a >> 30) * XRSR_MIX1;
    a = (a ^ a >> 27) * XRSR_MIX2;
    return a ^ a >> 31;
}

__device__ __host__ uint64_t fix64(uint64_t a) {
    a = (a ^ a >> 31 ^ a >> 62) * XRSR_MIX2_INVERSE;
    a = (a ^ a >> 27 ^ a >> 54) * XRSR_MIX1_INVERSE;
    return a ^ a >> 30 ^ a >> 60;
}

__device__ __host__ void xrsr_seed(Xrsr *xrsr, uint64_t seed) {
    seed ^= XRSR_SILVER_RATIO;
    xrsr->lo = mix64(seed);
    xrsr->hi = mix64(seed + XRSR_GOLDEN_RATIO);
}

__device__ __host__ uint64_t rol64(uint64_t a, int bits) {
    return (a << bits) | (a >> 64 - bits);
}

__device__ __host__ uint64_t xrsr_next_internal(Xrsr *xrsr) {
    uint64_t l = xrsr->lo;
    uint64_t m = xrsr->hi;
    uint64_t r = rol64(l + m, 17) + l;
    m ^= l;
    xrsr->lo = rol64(l, 49) ^ m ^ m << 21;
    xrsr->hi = rol64(m, 28);
    return r;
}

__device__ __host__ int32_t xrsr_next(Xrsr *xrsr, int bits) {
    return (int32_t)(xrsr_next_internal(xrsr) >> (64 - bits));
}

DEF_BASE_RANDOM_NEXT_LONG(Xrsr, xrsr)

__device__ __host__ void xrsr_set_decoration_seed(Xrsr *xrsr, uint64_t world_seed, int32_t chunk_x, int32_t chunk_z, int32_t index, int32_t step) {
    xrsr_seed(xrsr, world_seed);
    int64_t a = xrsr_next_long(xrsr) | 1LL;
    int64_t b = xrsr_next_long(xrsr) | 1LL;
    xrsr_seed(xrsr, ((a * (chunk_x * 16) + b * (chunk_z * 16)) ^ world_seed) + index + step * 10000);
}

__device__ __host__ uint32_t xrsr_next_int_pow2(Xrsr *xrsr, int32_t bits) {
    return (uint32_t)xrsr_next(xrsr, 31) >> (31 - bits);
}

__constant__ uint16_t mask_lilypads_inside[16];
// __constant__ uint32_t mask_lilypads_all[32];
__device__ uint64_t found_buffer[1024];
__device__ uint32_t found_count;

__global__
void filter_first(uint64_t start) {
    uint32_t index = blockIdx.x * blockDim.x + threadIdx.x;

    uint64_t qpc_tick = start | index;
    // uint64_t nano_time = qpc_tick_to_nano_time(qpc_tick);
    // uint64_t world_seed = nano_time_to_world_seed(nano_time);
    uint64_t world_seed = (uint64_t)(int64_t)(int32_t)(uint32_t)qpc_tick;
    Xrsr xrsr;
    xrsr_set_decoration_seed(&xrsr, world_seed, 238, 34, 55, 9);

    // if (qpc_tick == 6527241171) {
    //     printf("QPC MATCHING!!!\n");
    // } else {
    //     return;
    // }

    int matching_count = 0;

    for (int count = 0; count < 4; count++) {
        uint32_t patch_x = xrsr_next_int_pow2(&xrsr, 4);
        uint32_t patch_z = xrsr_next_int_pow2(&xrsr, 4);

        bool patch_on_lilypad = (mask_lilypads_inside[patch_z] >> patch_x) & 1;

        for (int attempt = 0; attempt < 10; attempt++) {
            uint32_t x = patch_x + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);
            uint32_t y = xrsr_next_int_pow2(&xrsr, 2) - xrsr_next_int_pow2(&xrsr, 2);
            uint32_t z = patch_z + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);

            // bool on_lilypad = (mask_lilypads_all[z + 8] >> (x + 8)) & 1;
            bool on_lilypad = x < 16 && z < 16 && (mask_lilypads_inside[z] >> x) & 1;

            if (on_lilypad && (y == 0 || (patch_on_lilypad && y == (uint32_t)-1))) {
                matching_count += 1;
            }
        }
    }

    // if (world_seed == ((uint64_t)-4026714142520107754LL)) {
    //     printf("MATCHING COUNT ON SEED: %i\n", matching_count);
    // }

    if (matching_count < 8) return;

    // if (world_seed == ((uint64_t)-4026714142520107754LL)) {
    //     printf("DIDN'T FILTER OUT THE SEED!!!\n");
    // }

    found_buffer[atomicAdd(&found_count, 1)] = world_seed;
}

#define cudaCheckError(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) {
   if (code != cudaSuccess) {
      fprintf(stderr, "Cuda Error: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

constexpr uint16_t reverse16(uint16_t v) {
    v = (v & 0xFF00) >> 8 | (v & 0x00FF) << 8;
    v = (v & 0xF0F0) >> 4 | (v & 0x0F0F) << 4;
    v = (v & 0xCCCC) >> 2 | (v & 0x3333) << 2;
    v = (v & 0xAAAA) >> 1 | (v & 0x5555) << 1;
    return v;
}

// look down facing north
const uint16_t mask_lilypads_inside_host[16] = {
    reverse16(0b0000000000000000),
    reverse16(0b0000000000001000),
    reverse16(0b0000001000000000),
    reverse16(0b0000000000000100),
    reverse16(0b0000001010000000),
    reverse16(0b0000000000000000),
    reverse16(0b0000000000100000),
    reverse16(0b0000000000010000),
    reverse16(0b0100011000000000),
    reverse16(0b0000000000000000),
    reverse16(0b0000000000000000),
    reverse16(0b0000000000000000),
    reverse16(0b0000100000000000),
    reverse16(0b1000100000000000),
    reverse16(0b0000000000000000),
    reverse16(0b0000000000000000),
};
// const uint32_t mask_lilypads_all_host[32] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0b00000000000111111111100000000000, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };

void init_constants() {
    cudaCheckError( cudaMemcpyToSymbol(mask_lilypads_inside, mask_lilypads_inside_host, sizeof(mask_lilypads_inside_host)) );
    // cudaCheckError( cudaMemcpyToSymbol(mask_lilypads_all, mask_lilypads_all_host, sizeof(mask_lilypads_all_host)) );
}

constexpr uint32_t reverse32(uint32_t v) {
    v = (v & 0xFFFF0000) >> 16 | (v & 0x0000FFFF) << 16;
    v = (v & 0xFF00FF00) >>  8 | (v & 0x00FF00FF) <<  8;
    v = (v & 0xF0F0F0F0) >>  4 | (v & 0x0F0F0F0F) <<  4;
    v = (v & 0xCCCCCCCC) >>  2 | (v & 0x33333333) <<  2;
    v = (v & 0xAAAAAAAA) >>  1 | (v & 0x55555555) <<  1;
    return v;
}

const uint32_t mask_lilypads_combined_host[16] = {
    reverse32(0b00000000000000000000000000000000),
    reverse32(0b00000000000000000000000000001000),
    reverse32(0b00000000000000000000001000000000),
    reverse32(0b00000000000000000000000000000100),
    reverse32(0b00000000000000010000001010000000),
    reverse32(0b00000000000011000000000000000000),
    reverse32(0b00000100000100000000000000100000),
    reverse32(0b00000000010000000000000000010000),
    reverse32(0b00000000000000000100011000000000),
    reverse32(0b00000000000000000000000000000000),
    reverse32(0b00000000000001000000000000000000),
    reverse32(0b00000000000000000000000000000000),
    reverse32(0b00000000000000100000100000000000),
    reverse32(0b00000000000000001000100000000000),
    reverse32(0b00000000000000000000000000000000),
    reverse32(0b00000000000001000000000000000000),
};

bool test_2_host(uint64_t world_seed) {
    int matching_count = 0;

    Xrsr xrsr;

    uint32_t generated[16] = { 0 };

    xrsr_set_decoration_seed(&xrsr, world_seed, 237, 34, 55, 9);
    for (int count = 0; count < 4; count++) {
        uint32_t patch_x = xrsr_next_int_pow2(&xrsr, 4);
        uint32_t patch_z = xrsr_next_int_pow2(&xrsr, 4);

        bool patch_on_lilypad = (mask_lilypads_combined_host[patch_z] >> patch_x) & 1;

        for (int attempt = 0; attempt < 10; attempt++) {
            uint32_t x = patch_x + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);
            uint32_t y = xrsr_next_int_pow2(&xrsr, 2) - xrsr_next_int_pow2(&xrsr, 2);
            uint32_t z = patch_z + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);

            if (x < 32 && z < 16) {
                if (!((generated[z] >> x) & 1)) {
                    bool on_lilypad = (mask_lilypads_combined_host[z] >> x) & 1;
                    if (on_lilypad && (y == 0 || (patch_on_lilypad && y == (uint32_t)-1))) {
                        matching_count += 1;
                        generated[z] |= 1 << x;
                    }
                }
            }
        }
    }

    xrsr_set_decoration_seed(&xrsr, world_seed, 238, 34, 55, 9);
    for (int count = 0; count < 4; count++) {
        uint32_t patch_x = xrsr_next_int_pow2(&xrsr, 4) + 16;
        uint32_t patch_z = xrsr_next_int_pow2(&xrsr, 4);

        bool patch_on_lilypad = (mask_lilypads_combined_host[patch_z] >> patch_x) & 1;

        for (int attempt = 0; attempt < 10; attempt++) {
            uint32_t x = patch_x + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);
            uint32_t y = xrsr_next_int_pow2(&xrsr, 2) - xrsr_next_int_pow2(&xrsr, 2);
            uint32_t z = patch_z + xrsr_next_int_pow2(&xrsr, 3) - xrsr_next_int_pow2(&xrsr, 3);

            if (x < 32 && z < 16) {
                if (!((generated[z] >> x) & 1)) {
                    bool on_lilypad = (mask_lilypads_combined_host[z] >> x) & 1;
                    if (on_lilypad && (y == 0 || (patch_on_lilypad && y == (uint32_t)-1))) {
                        matching_count += 1;
                        generated[z] |= 1 << x;
                    }
                }
            }
        }
    }

    return matching_count >= 14;
}

uint32_t found_count_host = 0;
uint64_t found_buffer_host[1024];

#define COUNT(BITS) (1LLU << (BITS))

#define BLOCK_BITS 20
#define BLOCK_COUNT COUNT(BLOCK_BITS)
#define THREAD_BITS 8
#define THREAD_COUNT COUNT(THREAD_BITS)
#define RUN_BITS (BLOCK_BITS + THREAD_BITS)
#define RUN_COUNT COUNT(RUN_BITS)

int main() {
    init_constants();

    FILE *out = fopen("out.txt", "w");

    uint64_t start = 0;
    const uint64_t end = COUNT(42);
    while (start < end) {
        clock_t clock_start = clock();

        found_count_host = 0;
        cudaMemcpyToSymbol(found_count, &found_count_host, sizeof(found_count_host));

        filter_first<<<BLOCK_COUNT, THREAD_COUNT>>>(start);
        cudaCheckError( cudaGetLastError() );
        cudaCheckError( cudaDeviceSynchronize() );

        cudaMemcpyFromSymbol(&found_count_host, found_count, sizeof(found_count_host));
        cudaMemcpyFromSymbol(&found_buffer_host, found_buffer, sizeof(uint64_t) * found_count_host);

        start += RUN_COUNT;

        clock_t clock_end = clock();

        // printf("Matching test 1 count: %i\n", found_count_host);
        for (uint32_t i = 0; i < found_count_host; i++) {
            uint64_t world_seed = found_buffer_host[i];
            if (test_2_host(world_seed)) {
                printf("Matching test 2 seed: %lli\n", (int64_t)world_seed);
                fprintf(out, "%lli\n", (int64_t)world_seed);
            }
        }

        if (start % COUNT(32) == 0) {
            double progress = (double)start / end;
            double elapsed = (double)(clock_end - clock_start) / CLOCKS_PER_SEC;
            printf("%fs %f%% %fM/s\n", elapsed, progress * 100, RUN_COUNT / elapsed / 1000000);
        }
    }

    return 0;
}