/perf/kseta

#include <sys/types.h>
#include <stdio.h>
#include <string.h>

#include <CL/cl.h>

#define BLOCK_SIZE 16

#ifndef PPT
# define PPT 1
#endif 

#ifndef SHMEM
# define SHMEM 1
#endif

static cl_int err;
static cl_platform_id platform = 0;
static cl_device_id device = 0;

static cl_context_properties props[3] = { CL_CONTEXT_PLATFORM, 0, 0 };
static cl_context ctx = 0;
static cl_command_queue queue = 0;

static cl_mem dev_res, dev_a, dev_b;

static cl_program app;
static cl_kernel kernel;

static size_t matrix_size;

int exercise_required_alignment = PPT * BLOCK_SIZE;

int exercise_init(const char *name, size_t size) {
    int ret = 0;
    char param[1024];

    FILE *f;
    size_t len;
    char *source;

    err = clGetPlatformIDs(1, &platform, NULL);
    if (err != CL_SUCCESS) {
        printf( "clGetPlatformIDs() failed with %d\n", err );
        return 1;
    }

    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    if (err != CL_SUCCESS) {
        printf( "clGetDeviceIDs() failed with %d\n", err );
        return 1;
    }

    err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(param), param, NULL);
    if (!err) {
	printf("Using device: %s\n", param);
    }


    props[1] = (cl_context_properties)platform;
    ctx = clCreateContext(props, 1, &device, NULL, NULL, &err);
    if (err != CL_SUCCESS) {
        printf( "clCreateContext() failed with %d\n", err );
        return 1;
    }

    queue = clCreateCommandQueue(ctx, device, CL_QUEUE_PROFILING_ENABLE, &err);
    if (err != CL_SUCCESS) {
        printf( "clCreateCommandQueue() failed with %d\n", err );
        clReleaseContext(ctx);
        return 1;
    }

    dev_a = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size * size * sizeof(float), NULL, &err);
    if (err != CL_SUCCESS) {
	printf("clCreateBuffer() failed with %d\n", err);
	return 1;
    }

    dev_b = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size * size * sizeof(float), NULL, &err);
    if (err != CL_SUCCESS) {
	printf("clCreateBuffer() failed with %d\n", err);
	return 1;
    }

    dev_res = clCreateBuffer(ctx, CL_MEM_READ_WRITE, size * size * sizeof(float), NULL, &err);
    if (err != CL_SUCCESS) {
	printf("clCreateBuffer() failed with %d\n", err);
	return 1;
    }

    char *res = calloc(size * size, sizeof(float));
    if (res) {
	err = clEnqueueWriteBuffer(queue, dev_res, CL_TRUE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
	    printf("clEnqueueWriteBuffer() failed with %d\n", err);
	    return -1;
	}
	free(res);

        err = clFinish(queue);
	if (err != CL_SUCCESS) {
	    printf("clFinish() failed with %d\n", err);
	    return -1;
	}
    }

    sprintf(param, "%s.cl", name);
    f = fopen(param, "r");
    if (!f) {
	printf("Can't open file with OpenCL kernels\n");
	return 1;
    }

    fseek(f, 0, SEEK_END);
    len = ftell(f);
    fseek(f, 0, SEEK_SET);
    
    source = (char*)malloc(len + 128);
    if (!source) {
	printf("Can't allocate memory for OpenCL source\n");
	return 1;
    }
    
    sprintf(source, "#define BLOCK_SIZE %u\n#define PPT %u\n\n", BLOCK_SIZE, PPT);
    if (fread(source + strlen(source), 1, len, f) != len) {
	printf("Can't read OpenCL source\n");
	return 1;
    }
    fclose(f);

    len = strlen(source);

    app = clCreateProgramWithSource(ctx, 1, (const char**)&source, &len, &err);
    if (err != CL_SUCCESS) {
	printf("clCreateProgramWithSource() failed with %d\n", err);
	return 1;
    }
    
    err = clBuildProgram(app, 1, &device,  "", NULL, NULL);
    if (err != CL_SUCCESS) {
	printf("clBuildProgram() failed with %d\n", err);
	return 1;
    }
    
    free(source);
    
    kernel = clCreateKernel(app, "multiply", &err);
    if (err != CL_SUCCESS) {
	printf("clCreateKernel() failed with %d\n", err);
	return 1;
    }

    clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_res);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_a);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_b);
    matrix_size = size;
    clSetKernelArg(kernel, 3, sizeof(unsigned long), &matrix_size);
//    printf("%lu %lu %lu\n", PPT, PPT, 2 * PPT * PPT * BLOCK_SIZE * BLOCK_SIZE * sizeof(float) / 1024);
    err = clSetKernelArg(kernel, 4, 2 * SHMEM * BLOCK_SIZE * BLOCK_SIZE * sizeof(float), NULL);

    return 0;
}

int exercise_allocate(float **res, float **a, float **b, size_t size) {
    return 0;
}


void exercise_free() {
    clReleaseKernel(kernel);
    clReleaseProgram(app);
    
    clReleaseCommandQueue(queue);
    clReleaseContext(ctx);

    clReleaseMemObject(dev_res);
    clReleaseMemObject(dev_b);
    clReleaseMemObject(dev_a);
}

size_t exercise(float *res, float *a, float *b, size_t size, int iterations) {
    int i;
    
    cl_event event = NULL;
    size_t runtime = 0;
    cl_ulong start, end;

    size_t local_size[] = {BLOCK_SIZE, BLOCK_SIZE};
    size_t global_size[] = {size/PPT, size/PPT};

    for (i = 0; i < iterations; i++) {
	err = clEnqueueWriteBuffer(queue, dev_a, CL_TRUE, 0, size * size * sizeof(float), a, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
	    printf("clEnqueueWriteBuffer() failed with %d\n", err);
	    return -11;
	}

	err = clEnqueueWriteBuffer(queue, dev_b, CL_TRUE, 0, size * size * sizeof(float), b, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
	    printf("clEnqueueWriteBuffer() failed with %d\n", err);
	    return -1;
	}

	err = clEnqueueNDRangeKernel(queue, kernel, 2, 0, global_size, local_size, 0, NULL, &event);
	if (err != CL_SUCCESS) {
	    printf("clEnqueueNDRangeKernel() failed with %d\n", err);
	    return -1;
	}

	err = clEnqueueReadBuffer(queue, dev_res, CL_TRUE, 0, size * size * sizeof(float), res, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
	    printf("clEnqueueReadBuffer() failed with %d\n", err);
	    return -1;
	}

	err = clFinish(queue);
	if (err != CL_SUCCESS) {
	    printf("clFinish() failed with %d\n", err);
	    return -1;
	}
	
	err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
	if (err != CL_SUCCESS) {
	    printf("clGetEventProfilingInfo() failed with %d\n", err);
	    return -1;
	}
	
	err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
	if (err != CL_SUCCESS) {
	    printf("clGetEventProfilingInfo() failed with %d\n", err);
	    return -1;
	}
	
	runtime += end - start;
    }

    return runtime;
}