2
\$\begingroup\$

I have a cryptographic function with two 24 bit keys.

I have two blocks of input and two blocks of output, and want to brute force the keys using CUDA.

Overview:
The function is composed to two back-to-back 3 round substitution-permutation-networks with block size of 24 bits and sbox size of 4 bits.

Due to the low block size, the substitution and permutation steps can be coalesced into a single lookup table

The encrypt function can be split into two sections, the first of which only uses the first key.
So partition the threads such that each thread handles one value for the first key, and iterates over the second key.
This minimizes the amount of redundant computation.

The problem is that I get a ~3x speedup using CUDA vs a single cpu core. How do I get the fabled 100x CUDA speedup?

I've included "easy" input that can be found almost instantly, and "hard" input that would take days/weeks to find, given the current performance

//cuda
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

//io
#include <stdio.h>
#include <stdlib.h>

//timing
#include <ctime>


__device__ __host__ int encrypt(int block, int key, const int * lookup);
__global__ void brute_Kernel(int block, int block2, int target, int target2,const int * lookup, int * solution);
int permute(int block, const int *pbox);
int substitute(int block, const int sboxes[][16]);
int cudaBrute(const int * lookup, int block, int block2, int target, int target2);
int cpuBrute(const int * lookup, int block, int block2, int target, int target2);


int main()
{
    const unsigned size = (1 << 24) * sizeof(int);
    int * lookup = (int *)malloc(size);

    const int sboxes[6][16] = { { 15, 1, 7, 0, 9, 6, 2, 14, 11, 8, 5, 3, 12, 13, 4, 10 }, { 3, 7, 8, 9, 11, 0, 15, 13, 4, 1, 10, 2, 14, 6, 12, 5 }, { 4, 12, 9, 8, 5, 13, 11, 7, 6, 3, 10, 14, 15, 1, 2, 0 }, { 2, 4, 10, 5, 7, 13, 1, 15, 0, 11, 3, 12, 14, 9, 8, 6 }, { 3, 8, 0, 2, 13, 14, 5, 11, 9, 1, 7, 12, 4, 6, 10, 15 }, { 14, 12, 7, 0, 11, 4, 13, 15, 10, 3, 8, 9, 2, 6, 1, 5 } };
    const int pbox[] = { 13, 3, 15, 23, 6, 5, 22, 21, 19, 1, 18, 17, 20, 10, 7, 8, 12, 2, 16, 9, 14, 0, 11, 4 };

    for (int i = 0; i < (1 << 24); ++i)
    {
        lookup[i] = permute(substitute(i, sboxes), pbox);
    }

    //timing
    std::clock_t start;
    double duration;
    start = std::clock();

    //easy mode; both keys are < 0xfff
    //printf("[+] brute returned: %d\n", cudaBrute(lookup, 0x313233, 0x343536, 0x0cc3fb, 0x409d0d));

    //hard mode; I have no idea what the keys are
    printf("[+] Brute returned: %d\n", cudaBrute(lookup, 0x313233, 0x343536, 0x0cc3fb, 0x409d0e));

    duration = ( std::clock() - start ) / (double) CLOCKS_PER_SEC;
    printf("[+] time: %f\n", duration);

    free(lookup);

}

int cpuBrute(const int * lookup, int block, int block2, int target, int target2)
{
    for(int key = 0; key <= 0xffffff; ++key)
    {
        int first = encrypt(block, key, lookup);

        //should be up to 0xffffff, but it'd take months to complete if it was
        for(int key2 = 0; key2 <= 0xfff; ++key2)
        {
            if (encrypt(first, key2, lookup) == target && encrypt(encrypt(block2, key, lookup), key2, lookup) == target2)
            {
                printf("[+] key: 0x%06x key: 0x%06x\n", key, key2);
                return 1;
            }
        }
    }
    return 0;
}


int cudaBrute(const int * lookup, int block, int block2, int target, int target2)
{
    int * dev_lookup;
    int * dev_solution;
    cudaError_t cudaStatus;

    int host_solution;
    const unsigned size = (1 << 24) * sizeof(int);

    // two dimension, spanning 12 bits each
    dim3 DimGrid(256, 256, 1);
    dim3 DimBlock(16, 16, 1);

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);

    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_solution, sizeof(int));
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_lookup, size);
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_lookup, lookup, size, cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemset(dev_solution, 0x00, 4);
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaMemset failed!");
        goto Error;
    }

    // Launch the kernel
    brute_Kernel <<<DimGrid, DimBlock >>>(block, block2, target, target2, dev_lookup, dev_solution);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "Kernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(&host_solution, dev_solution, sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaFree(dev_lookup);
    cudaFree(dev_solution);
    return host_solution != 0;

Error:
    cudaFree(dev_lookup);
    cudaFree(dev_solution);
    return -1;
}


int permute(int block, const int *pbox)
{
    int output = 0;
    for (int i = 0; i < 24; ++i)
        output |= ((block >> pbox[i]) & 1) << i;
    return output;
}

int substitute(int block, const int sboxes[][16])
{
    int output = 0;
    for (int i = 0; i < 6; ++i)
        output |= sboxes[i][(block >> 4 * i) & 0xf] << 4 * i;
    return output;
}

/*
SPN network with 3 rounds
    block size = 24 bits
    substitution box size = 4 bits
key derivation function = identity
due to small block size, substition and permutation steps have been
    coalesced into a single lookup table
*/
__device__ __host__ int encrypt(int block, int key, const int * lookup)
{
    for (int i = 0; i < 3; ++i)
    {
        block ^= key;
        block = lookup[block];
    }
    return block ^ key;
}


__global__ void brute_Kernel(int block, int block2, int target, int target2, const int * lookup, int * solution)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;
    int key = idx + (idy << 12);
    // 12 LSB from idx, 12 MSB from idy -> 24 bit key

    int first = encrypt(block, key, lookup);
    int first_2 = encrypt(block2, key, lookup);

    /*
    to reduce repetition, have each cuda thread iterate over ALL of key2
    this way, you don't have to repeat the first encryption 0xffffff times
    */
    for (int key2 = 0; key2 <= 0xffffff; ++key2)
    {
        if ((encrypt(first, key2, lookup) == target) && (encrypt(first_2, key2, lookup) == target2))
        {
            printf("[+] key: 0x%06x key2: 0x%06x\n", key, key2);
            *solution = 1;
            /*
            "gracefully" crash the kernel lol
            This kills all threads once I find a match
            It'll screw up future cudaMemcpy, so I need to printf from the kernel
            */
            asm("trap;");
        }
    }
}

BTW, an easier way to do this problem is via a meet-in-the-middle attack, which takes ~2^25 encryptions.
I'm doing this to learn CUDA

\$\endgroup\$
  • \$\begingroup\$ What is in "cuda_runtime.h" and "device_launch_parameters.h"? Additionally, you're probably not even going to get a 3x speedup with a fair comparison - your cpu brute force one should be threaded. \$\endgroup\$ – Dannnno May 28 '16 at 5:59
3
\$\begingroup\$

I'm going to start with simple style and organizational things, because they're important and will help make the rest of the code easier to review.

The first thing you should do is move your cuda stuff into .cu and .cuh files. This lets you keep the actual cuda things separate, and also lets you use nice C++11 (and newer) features in your other files.

This will probably end up complicating the build process a bit. Assuming you use a Makefile to compile, you'll probably have something like this (your flags/variables will vary)

GENCODE_FLAGS   := -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35
CUDA_PATH       ?= /usr/local/cuda
CUDA_LIB_PATH   ?= $(CUDA_PATH)/lib/
LD_FLAGS        += -L$(CUDA_LIB_PATH) -lcudart -lcufft -lcublas
CC_FLAGS      += -m64
NVCC_FLAGS    += -m64

CC              := g++
NVCCSCC         := $(CC)

CUDA_INC_PATH   ?= $(CUDA_PATH)/include
CUDA_BIN_PATH   ?= $(CUDA_PATH)/bin
NVCC            ?= $(CUDA_BIN_PATH)/nvcc

CC_FLAGS        += -O3 -std=c++11 -Wall -Wextra -fopenmp
NVCC_FLAGS      += -O3 -std=c++11

CC_INCLUDE      += -I$(CUDA_INC_PATH)

Encryptor_cuda.o: Encryptor_cuda.cu Encryptor_cuda.cuh
    $(NVCC) -ccbin=$(NVCCSCC) $(NVCC_FLAGS) $(GENCODE_FLAGS) -I$(CUDA_INC_PATH) $(NVCC_INCLUDE) -o $@ -c $<

Encryptor: Encryptor.cc Encryptor_cuda.o
    $(CC) $< -o $@ $(CC_INCLUDE) $(CC_FLAGS) Encryptor_cuda.o $(LD_FLAGS)

Next, if you actually want to time things you probably want to use std::chrono::high_resolution_clock assuming you have C++11 available. You'll also want to repeat the tests a few times - even if you have a dedicated machine that doesn't have anything else on there, noise and black magic can cause unusual timings. My setups normally look like this:

template <typename Functor>
void runTest(const some_type& inputs, Functor function, some_type* output, size_t numTrials, double* elapsedTime) {

    *elapsedTime = std::numeric_limits<double>::max();

    for (size_t i = 0; i < numTrials; ++i) {
        std::chrono::high_resolution_clock::time_point tic = std::chrono::high_resolution_clock::now();
        function(inputs, output, .....);
        std::chrono::high_resolution_clock::time_point toc = std::chrono::high_resolution_clock::now();
        const double thisTrialElapsedTime = duration_cast<duration<double>>(toc - tic).count();
        *elapsedTime = std::min(thisTrialElapsedTime, *elapsedTime);

    }
}

This then makes it really easy to have a much higher degree of control over the timing of the test. If you don't want the minimum you can do average, or something else. YMMV.

Now I'm going to look at your actual code. From the get-go, your cpuBrute is fighting a losing battle because it isn't a fair comparison - your CPU has vector registers and multiple cores (probably) so you should take advantage of them.

Your encrypt function seems pretty straightforward. The one thing might be looking into reordering the lookup table somehow to improve caching behavior, but nothing is jumping out at me there. Judicious prefetching might help there.

Otherwise, you should use your good old friend OpenMP here. I've made a few changes - instead of printing the keys, it writes them to an output variable. The caller can then handle things.

int cpuBrute(const int * lookup, int block, int block2, int target, int target2, int* key1out, int* key2out)
{
    bool done = false

    #pragma omp parallel for
    for(int key = 0; key <= 0xffffff; ++key)
    {
        if (!done) {
            int first = encrypt(block, key, lookup);
            int second = encrypt(block2, key, lookup);

            //should be up to 0xffffff, but it'd take months to complete if it was
            for(int key2 = 0; key2 <= 0xfff && !done; ++key2)
            {
                if (encrypt(first, key2, lookup) == target && encrypt(encrypt(block2, key, lookup), key2, lookup) == target2) {
                    #pragma omp critical
                    {
                        *key1out = key;
                        *key2out = key2;
                        done = true;
                    }
                }
            }
        }
    }

    return done;
}

The call would then look like this

int key1result, key2result;
if (cpuBrute(lookup, block, block2, target, &key1result, &key2result)) {
    printf("[+] key: 0x%06x key: 0x06x\n", key1result, key2result);
}

The biggest problem I see is that you have a conditional in your kernel - this is going to cause a lot of context switches and make your warps swap around a lot. I also wouldn't print from the kernel - as above, it should write to some output variable, then copy that to the CPU, then display based on that. You want to use a mask or something to avoid the conditional.

Given that the lookups into your lookup table don't seem to be in any particular order you're going to have a ton of cache misses (GPUs like synchronous memroy access patterns). Again, reordering that or improving that should help a lot. Alternately, loading it into shared memory should make the problem moot.

Lastly, the reason I suspect that you aren't going to see much of a speedup is that your problem size is too small. If you make each computation expensive, and give the GPU a million of them, that should beat the pants off the CPU version.

I know I didn't give a ton of specific suggestions for the GPU version, but the issue imo is really a design issue, and a problem-choice issue.

\$\endgroup\$
  • \$\begingroup\$ Thanks for the CPU improvements! I'm not sure if it'd be possible to reorder the lookup table, as it's used with "random" inputs several times per encrypt. What do you mean by the problem bring too small? I gave the GPU a million (0xffffff) computations, each of which is expensive (call encrypt 0xffffff times with different keys). Can you clarify what you mean? Finally, I'm not sure that it's possible to remove the conditional from the kernel, as it's iterating over a key space until it finds a match. \$\endgroup\$ – robertkin May 28 '16 at 7:44
  • \$\begingroup\$ @robertkin Sorry - I had just gotten off a 12-hour plane ride, and was on a layover for another 6 hour one and was clearly not thinking straight. On closer inspection, the problem size seems reasonable (although I'm not sure that you're applying the GPU to it in the best way). If I get a chance later I'll edit and maybe expand my answer to not suck in that regard \$\endgroup\$ – Dannnno May 28 '16 at 18:57
  • \$\begingroup\$ Thanks for your time! This is my first time applying cuda, so I'm sure my approach leaves much to be desired. Would loading the global lookup table into each block's shared memory, or possibly having multiple threads handle the first 24 bit key, be an improvement? \$\endgroup\$ – robertkin May 28 '16 at 21:41
  • \$\begingroup\$ Having the lookup table in shared memory will help - it essentially removes the need for coalesced memory accesses. I really don't know if multiple threads handling the first 24 bit key would be helpful - I doubt it tbh, since its ideal if each thread handles some other part of the work. \$\endgroup\$ – Dannnno May 29 '16 at 1:16
  • \$\begingroup\$ That makes sense. Are there any other speedups, other than copying the table every block? \$\endgroup\$ – robertkin May 29 '16 at 1:18

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Not the answer you're looking for? Browse other questions tagged or ask your own question.