Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- //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;
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement