/* CUDA running-key corpus attack for the Paradigm Kryptos-CTF PK8 (and PK9/PK10).
 *
 * PK8's author hint (Dan Robinson): "the algorithm is simple. The key has quite a lot of
 * entropy, but some structure."  That is the signature of a RUNNING-KEY Vigenere: a simple
 * position-wise add whose key is a long English text (high entropy, but English structure).
 * The break is to try every offset of a large corpus as the key and score the decryption -
 * embarrassingly parallel, one GPU thread per corpus offset.
 *
 * This is a GPU-pod deliverable: it needs an NVIDIA GPU + CudaLink (WL) or nvcc.  The Mac
 * dev box has no NVIDIA GPU, so run this on a pod.  Two ways to run:
 *   (a) via WL CudaLink: see cuda_runkey.wls (CUDAFunctionLoad on this kernel).
 *   (b) standalone: nvcc -O3 cuda_runkey.cu -o runkey  (add a host main to load arrays).
 *
 * Index space: all letters are 0..25 ints in a caller-chosen alphabet (plain or KRYPTOS).
 * The quadgram table is a 26^4 float array of log-probs (English, relabelled to that space).
 * op: 0 = plain=ct-key, 1 = plain=key-ct (Beaufort), 2 = plain=ct+key.
 *
 * Corpus scale this is built for: 10^8-10^9 letters (all of Project Gutenberg, the CIA
 * reading room, Sanborn/Scheidt interviews, etc.).  At ~10^9 offsets x n quadgram lookups a
 * single modern GPU sweeps a book-length corpus in well under a second; a CPU cannot. */

extern "C" __global__ void runKeyCorpus(
        const int* __restrict__ ct, int n,
        const int* __restrict__ corpus, long clen,
        const float* __restrict__ QG, int op,
        float* __restrict__ scoreOut)      /* length clen-n+1, one score per offset */
{
    long off = blockIdx.x * (long)blockDim.x + threadIdx.x;
    if (off + n > clen) return;
    float s = 0.0f;
    int p0, p1, p2, p3;
    /* prime the first three plaintext letters */
    #define PLAIN(i) ( op==0 ? ((ct[i]-corpus[off+i]+26)%26) \
                     : op==1 ? ((corpus[off+i]-ct[i]+26)%26) \
                     :         ((ct[i]+corpus[off+i])%26) )
    p0 = PLAIN(0); p1 = PLAIN(1); p2 = PLAIN(2);
    for (int i = 3; i < n; i++) {
        p3 = PLAIN(i);
        s += QG[((p0*26 + p1)*26 + p2)*26 + p3];
        p0 = p1; p1 = p2; p2 = p3;
    }
    #undef PLAIN
    scoreOut[off] = s;
}

/* Host-side reduction (argmax over scoreOut) is trivial and done on the CPU or with a second
 * kernel / thrust::max_element.  The best offset's corpus slice is PK8's running key. */
