gistfile1.c

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <OpenCL/opencl.h>

const char *KernelSource = "\n"                                        \
    "#define DELTA 0x9e3779b9 \n"                                       \
    "#define ENC_ROUND(sum)"                                            \
    "{ "                                                                \
    "  v.s0 += ((v.s1 << 4) + key.s0) ^ (v.s1 + sum) ^ ((v.s1 >> 5) + key.s1); " \
    "  v.s1 += ((v.s0 << 4) + key.s2) ^ (v.s0 + sum) ^ ((v.s0 >> 5) + key.s3); " \
    "} \n"                                                              \
    "__kernel void encrypt(__global uint* input, \n"         \
    "                      const uint4 key, \n"              \
    "                      const uint blocks) \n"            \
    "{ \n"                                                   \
    "  int i; int id = get_global_id(0); \n"                 \
    "  if (id > blocks) return; \n"                          \
    "  uint2 v = {input[id * 2], input[(id * 2) + 1]}; \n"   \
    "  // no unroll pragma? \n"                              \
    "  ENC_ROUND(DELTA); \n"                                 \
    "  ENC_ROUND(DELTA * 2); ENC_ROUND(DELTA * 3); \n"       \
    "  ENC_ROUND(DELTA * 4); ENC_ROUND(DELTA * 5); \n"       \
    "  ENC_ROUND(DELTA * 6); ENC_ROUND(DELTA * 7); \n"       \
    "  ENC_ROUND(DELTA * 8); ENC_ROUND(DELTA * 9); \n"       \
    "  ENC_ROUND(DELTA * 10); ENC_ROUND(DELTA * 11); \n"     \
    "  ENC_ROUND(DELTA * 12); ENC_ROUND(DELTA * 13); \n"     \
    "  ENC_ROUND(DELTA * 14); ENC_ROUND(DELTA * 15); \n"     \
    "  ENC_ROUND(DELTA * 16); ENC_ROUND(DELTA * 17); \n"     \
    "  ENC_ROUND(DELTA * 18); ENC_ROUND(DELTA * 19); \n"     \
    "  ENC_ROUND(DELTA * 20); ENC_ROUND(DELTA * 21); \n"     \
    "  ENC_ROUND(DELTA * 22); ENC_ROUND(DELTA * 23); \n"     \
    "  ENC_ROUND(DELTA * 24); ENC_ROUND(DELTA * 25); \n"     \
    "  ENC_ROUND(DELTA * 26); ENC_ROUND(DELTA * 27); \n"     \
    "  ENC_ROUND(DELTA * 28); ENC_ROUND(DELTA * 29); \n"     \
    "  ENC_ROUND(DELTA * 30); ENC_ROUND(DELTA * 31); \n"     \
    "  ENC_ROUND(DELTA * 32); \n"                            \
    "  input[id * 2] = v.s0; input[(id * 2) + 1] = v.s1; \n" \
    "} \n"                                                   \
    "#define DEC_SUM 0xC6EF3720 \n"                                     \
    "#define DEC_ROUND(sum)"                                           \
    "{ "                                                                \
    "  v.s1 -= ((v.s0 << 4) + key.s2) ^ (v.s0 + sum) ^ ((v.s0 >> 5) + key.s3); " \
    "  v.s0 -= ((v.s1 << 4) + key.s0) ^ (v.s1 + sum) ^ ((v.s1 >> 5) + key.s1); " \
    "} \n"                                                              \
    "__kernel void decrypt(__global uint* input, \n"                    \
    "                      const uint4 key, \n"                         \
    "                      const uint blocks) \n"                       \
    "{ \n"                                                              \
    "  int i; \n"
    "  int id = get_global_id(0); \n"                                   \
    "  if (id > blocks) return; \n"                                     \
    "  uint2 v = {input[id * 2], input[(id * 2) + 1]}; \n"              \
    "  DEC_ROUND(DEC_SUM); \n"                                          \
    "  DEC_ROUND(DEC_SUM - DELTA * 1); DEC_ROUND(DEC_SUM - DELTA * 2); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 3); DEC_ROUND(DEC_SUM - DELTA * 4); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 5); DEC_ROUND(DEC_SUM - DELTA * 6); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 7); DEC_ROUND(DEC_SUM - DELTA * 8); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 9); DEC_ROUND(DEC_SUM - DELTA * 10); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 11); DEC_ROUND(DEC_SUM - DELTA * 12); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 13); DEC_ROUND(DEC_SUM - DELTA * 14); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 15); DEC_ROUND(DEC_SUM - DELTA * 16); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 17); DEC_ROUND(DEC_SUM - DELTA * 18); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 19); DEC_ROUND(DEC_SUM - DELTA * 20); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 21); DEC_ROUND(DEC_SUM - DELTA * 22); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 23); DEC_ROUND(DEC_SUM - DELTA * 24); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 25); DEC_ROUND(DEC_SUM - DELTA * 26); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 27); DEC_ROUND(DEC_SUM - DELTA * 28); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 29); DEC_ROUND(DEC_SUM - DELTA * 30); " \
    "  DEC_ROUND(DEC_SUM - DELTA * 31); "                               \
    "  input[id * 2] = v.s0; input[(id * 2) + 1] = v.s1; \n"            \
    "} \n"                                                              \
    "\n";

int main() {
    cl_device_id device_id;
    cl_context context;
    cl_command_queue commands;
    cl_program program;
    cl_kernel enc_kernel, dec_kernel;
    cl_int err;

    cl_mem data_buf;

    const unsigned int BLOCKS = pow(2, 24);

    size_t global = BLOCKS;
    
    cl_uint *data = calloc(BLOCKS * 2, sizeof(cl_uint));
    cl_uint4 key = {1, 2, 3, 4};

    err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS) {
        printf("Error getting device ID.\n");
        exit(1);
    }

    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (!context) {
        printf("Error creating context.\n");
        exit(1);
    }

    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands) {
        printf("Error creating command queue.\n");
        exit(1);
    }

    program = clCreateProgramWithSource(context, 1,
                                        (const char **)&KernelSource,
                                        NULL, &err);
    if (!program) {
        printf("Error creating program.\n");
        exit(1);
    }

    err = clBuildProgram(program, 0, NULL, "-Werror", NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t len;
        char buffer[2048];
        
        printf("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,
                              sizeof(buffer), buffer, &len);
        printf("%s\n", buffer);
        exit(1);
    }

    enc_kernel = clCreateKernel(program, "encrypt", &err);
    if (!enc_kernel || err != CL_SUCCESS) {
        printf("Error creating kernel.\n");
        exit(1);
    }

    dec_kernel = clCreateKernel(program, "decrypt", &err);
    if (!dec_kernel || err != CL_SUCCESS) {
        printf("Error creating decrypt kernel.\n");
        exit(1);
    }

    data_buf = clCreateBuffer(context, CL_MEM_READ_WRITE,
                              sizeof(cl_uint) * BLOCKS * 2, NULL, NULL);
    if (!data_buf) {
        printf("Error creating buffer.\n");
        exit(1);
    }

    err = clEnqueueWriteBuffer(commands, data_buf, CL_TRUE, 0,
                               sizeof(cl_uint) * BLOCKS * 2, data,
                               0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("Error writing data buffer.\n");
        exit(1);
    }

    err = 0;
    err = clSetKernelArg(enc_kernel, 0, sizeof(cl_mem), &data_buf);
    err |= clSetKernelArg(enc_kernel, 1, sizeof(cl_uint4), &key);
    err |= clSetKernelArg(enc_kernel, 2, sizeof(unsigned int), &BLOCKS);
    if (err != CL_SUCCESS) {
        printf("Error setting args.\n");
        exit(1);
    }
    
    err = clEnqueueNDRangeKernel(commands, enc_kernel, 1, NULL, &global,
                                 NULL, 0, NULL, NULL);
    if (err) {
        printf("Error executing encrypt kernel: %d.\n", err);
        exit(1);
    }

    clEnqueueBarrier(commands);

    err = 0;
    err = clSetKernelArg(dec_kernel, 0, sizeof(cl_mem), &data_buf);
    err |= clSetKernelArg(dec_kernel, 1, sizeof(cl_uint4), &key);
    err |= clSetKernelArg(dec_kernel, 2, sizeof(unsigned int), &BLOCKS);
    if (err != CL_SUCCESS) {
        printf("Error setting args.\n");
        exit(1);
    }
    
    err = clEnqueueNDRangeKernel(commands, dec_kernel, 1, NULL, &global,
                                 NULL, 0, NULL, NULL);
    if (err) {
        printf("Error executing decrypt kernel: %d\n", err);
        exit(1);
    }

    clFinish(commands);

    err = clEnqueueReadBuffer(commands, data_buf, CL_TRUE, 0,
                              sizeof(cl_uint) * BLOCKS * 2,
                              data, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("Error reading buffer.\n");
        exit(1);
    }

    printf("%u %u\n", data[0], data[1]);
    
    clReleaseMemObject(data_buf);
    clReleaseProgram(program);
    clReleaseKernel(enc_kernel);
    clReleaseKernel(dec_kernel);
    clReleaseCommandQueue(commands);
    clReleaseContext(context);

    free(data);
    
    return 0;
}